Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

#14634: Remove usage of ARCH_NAME sp constants MEM_L1_SIZE #14878

Open
wants to merge 7 commits into
base: main
Choose a base branch
from
Open

#14634: Remove usage of ARCH_NAME sp constants MEM_L1_SIZE #14878

Show file tree
Hide file tree
Changes from 6 commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
e3a1de2
#14634: Remove usage of ARCH_NAME sp constants MEM_L1_SIZE
blozano-tt Nov 8, 2024
a3343f6
Update sanitize_noc.h
blozano-tt Nov 8, 2024
ac56615
Merge branch 'main' into blozano/rm_arch_name_constants
blozano-tt Nov 14, 2024
7861dd2
Put MEM_L1/ETH_BASE and MEM_L1/ETH_SIZE behind Hal
blozano-tt Nov 14, 2024
42f2129
Update to use Hal
blozano-tt Nov 14, 2024
1717136
Merge branch 'main' into blozano/rm_arch_name_constants
blozano-tt Nov 15, 2024
c863e17
Update tt_metal/llrt/hal.hpp
blozano-tt Nov 15, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 7 additions & 3 deletions tt_metal/impl/debug/sanitize_noc_host.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,15 +2,19 @@
//
// SPDX-License-Identifier: Apache-2.0

#pragma once

#include <cstdint>

#include "noc/noc_parameters.h"
#include "noc/noc_overlay_parameters.h"

#pragma once
#include "llrt/hal.hpp"

namespace tt {

// Host MMIO reads/writes don't have alignment restrictions, so no need to check alignment here.
#define DEBUG_VALID_L1_ADDR(a, l) (((a) >= MEM_L1_BASE) && ((a) + (l) <= MEM_L1_BASE + MEM_L1_SIZE))
#define DEBUG_VALID_L1_ADDR(a, l) (((a) >= HAL_MEM_L1_BASE) && ((a) + (l) <= HAL_MEM_L1_BASE + HAL_MEM_L1_SIZE))

// what's the size of the NOC<n> address space? using 0x1000 for now
#define DEBUG_VALID_REG_ADDR(a) \
Expand All @@ -22,7 +26,7 @@ namespace tt {
#define DEBUG_VALID_WORKER_ADDR(a, l) (DEBUG_VALID_L1_ADDR(a, l) || (DEBUG_VALID_REG_ADDR(a) && (l) == 4))
#define DEBUG_VALID_DRAM_ADDR(a, l, b, e) (((a) >= b) && ((a) + (l) <= e))

#define DEBUG_VALID_ETH_ADDR(a, l) (((a) >= MEM_ETH_BASE) && ((a) + (l) <= MEM_ETH_BASE + MEM_ETH_SIZE))
#define DEBUG_VALID_ETH_ADDR(a, l) (((a) >= HAL_MEM_ETH_BASE) && ((a) + (l) <= HAL_MEM_ETH_BASE + HAL_MEM_ETH_SIZE))

static bool coord_found_p(std::vector<CoreCoord>coords, CoreCoord core) {
for (CoreCoord item : coords) {
Expand Down
2 changes: 1 addition & 1 deletion tt_metal/impl/device/device.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1567,7 +1567,7 @@ void Device::update_workers_build_settings(std::vector<std::vector<std::tuple<tt
uint32_t pcie_alignment = hal.get_alignment(HalMemType::HOST);
uint32_t scratch_db_base = (prefetch_d_settings.cb_start_address + prefetch_d_settings.cb_size_bytes + pcie_alignment - 1) & (~(pcie_alignment - 1));
uint32_t scratch_db_size = dispatch_constants::get(dispatch_core_type).scratch_db_size();
const uint32_t l1_size = dispatch_core_type == CoreType::WORKER ? MEM_L1_SIZE : MEM_ETH_SIZE;
const uint32_t l1_size = dispatch_core_type == CoreType::WORKER ? HAL_MEM_L1_SIZE : HAL_MEM_ETH_SIZE;
uint32_t dispatch_s_buffer_base;
uint32_t dispatch_buffer_base = dispatch_constants::get(dispatch_core_type).dispatch_buffer_base();
if (dispatch_core_type == CoreType::WORKER) {
Expand Down
3 changes: 2 additions & 1 deletion tt_metal/impl/dispatch/command_queue_interface.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,6 +15,7 @@
#include "tt_metal/llrt/hal.hpp"
#include "tt_metal/llrt/llrt.hpp"

// FIXME: Don't do this in header files
using namespace tt::tt_metal;

namespace tt::tt_metal {
Expand Down Expand Up @@ -212,7 +213,7 @@ struct dispatch_constants {
uint32_t prefetch_dispatch_unreserved_base = device_cq_addrs_[tt::utils::underlying_type<CommandQueueDeviceAddrType>(CommandQueueDeviceAddrType::UNRESERVED)];
cmddat_q_base_ = prefetch_dispatch_unreserved_base + ((prefetch_q_size_ + pcie_alignment - 1) / pcie_alignment * pcie_alignment);
scratch_db_base_ = cmddat_q_base_ + ((cmddat_q_size_ + pcie_alignment - 1) / pcie_alignment * pcie_alignment);
const uint32_t l1_size = core_type == CoreType::WORKER ? MEM_L1_SIZE : MEM_ETH_SIZE;
const uint32_t l1_size = core_type == CoreType::WORKER ? HAL_MEM_L1_SIZE : HAL_MEM_ETH_SIZE;
TT_ASSERT(scratch_db_base_ + scratch_db_size_ < l1_size);
dispatch_buffer_base_ = ((prefetch_dispatch_unreserved_base - 1) | ((1 << DISPATCH_BUFFER_LOG_PAGE_SIZE) - 1)) + 1;
dispatch_buffer_block_size_pages_ =
Expand Down
2 changes: 2 additions & 0 deletions tt_metal/llrt/blackhole/bh_hal_active_eth.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@ HalCoreInfoType create_active_eth_mem_map() {
std::vector<DeviceAddr> mem_map_bases;

mem_map_bases.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = 0x0; // Anything better to use?
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch);
Expand All @@ -46,6 +47,7 @@ HalCoreInfoType create_active_eth_mem_map() {

std::vector<uint32_t> mem_map_sizes;
mem_map_sizes.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = eth_l1_mem::address_map::MAX_SIZE; // Anything better to use?
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t);
Expand Down
2 changes: 2 additions & 0 deletions tt_metal/llrt/blackhole/bh_hal_idle_eth.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@ HalCoreInfoType create_idle_eth_mem_map() {
std::vector<DeviceAddr> mem_map_bases;

mem_map_bases.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_ETH_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = GET_IERISC_MAILBOX_ADDRESS_HOST(launch);
Expand All @@ -46,6 +47,7 @@ HalCoreInfoType create_idle_eth_mem_map() {

std::vector<uint32_t> mem_map_sizes;
mem_map_sizes.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_ETH_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t);
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t);
Expand Down
2 changes: 2 additions & 0 deletions tt_metal/llrt/blackhole/bh_hal_tensix.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@ HalCoreInfoType create_tensix_mem_map() {
std::vector<DeviceAddr> mem_map_bases;

mem_map_bases.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_L1_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = GET_MAILBOX_ADDRESS_HOST(launch);
Expand All @@ -44,6 +45,7 @@ HalCoreInfoType create_tensix_mem_map() {

std::vector<uint32_t> mem_map_sizes;
mem_map_sizes.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_L1_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t);
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t);
Expand Down
2 changes: 2 additions & 0 deletions tt_metal/llrt/grayskull/gs_hal.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,7 @@ void Hal::initialize_gs() {
std::vector<DeviceAddr> mem_map_bases;

mem_map_bases.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_L1_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = GET_MAILBOX_ADDRESS_HOST(launch);
Expand All @@ -54,6 +55,7 @@ void Hal::initialize_gs() {

std::vector<uint32_t> mem_map_sizes;
mem_map_sizes.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_L1_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t);
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t);
Expand Down
9 changes: 8 additions & 1 deletion tt_metal/llrt/hal.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,7 @@ enum class HalProcessorClassType : uint8_t {
};

enum class HalL1MemAddrType : uint8_t {
BASE,
BARRIER,
MAILBOX,
LAUNCH,
Expand Down Expand Up @@ -137,7 +138,7 @@ class Hal {

void initialize(tt::ARCH arch);

tt::ARCH get_arch() {return arch_;}
tt::ARCH get_arch() const {return arch_;}

template <typename IndexType, typename SizeType, typename CoordType>
auto noc_coordinate(IndexType noc_index, SizeType noc_size, CoordType coord) const -> decltype(noc_size - 1 - coord) {
Expand Down Expand Up @@ -263,3 +264,9 @@ extern Hal hal;

} // namespace tt_metal
} // namespace tt


#define HAL_MEM_L1_BASE tt::tt_metal::hal.get_dev_size(tt::tt_metal::HalProgrammableCoreType::TENSIX, tt::tt_metal::HalL1MemAddrType::BASE)
#define HAL_MEM_L1_SIZE tt::tt_metal::hal.get_dev_addr(tt::tt_metal::HalProgrammableCoreType::TENSIX, tt::tt_metal::HalL1MemAddrType::BASE)
#define HAL_MEM_ETH_BASE tt::tt_metal::hal.get_dev_size(tt::tt_metal::HalProgrammableCoreType::IDLE_ETH, tt::tt_metal::HalL1MemAddrType::BASE)
#define HAL_MEM_ETH_SIZE tt::tt_metal::hal.get_dev_addr(tt::tt_metal::HalProgrammableCoreType::IDLE_ETH, tt::tt_metal::HalL1MemAddrType::BASE)
blozano-tt marked this conversation as resolved.
Show resolved Hide resolved
2 changes: 2 additions & 0 deletions tt_metal/llrt/wormhole/wh_hal_active_eth.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,6 +31,7 @@ HalCoreInfoType create_active_eth_mem_map() {
std::vector<DeviceAddr> mem_map_bases;

mem_map_bases.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = 0x0; // Anything better to use?
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch);
Expand All @@ -46,6 +47,7 @@ HalCoreInfoType create_active_eth_mem_map() {

std::vector<uint32_t> mem_map_sizes;
mem_map_sizes.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = eth_l1_mem::address_map::MAX_SIZE; // Anything better to use?
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t);
Expand Down
2 changes: 2 additions & 0 deletions tt_metal/llrt/wormhole/wh_hal_idle_eth.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,6 +32,7 @@ HalCoreInfoType create_idle_eth_mem_map() {
std::vector<DeviceAddr> mem_map_bases;

mem_map_bases.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_ETH_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = GET_IERISC_MAILBOX_ADDRESS_HOST(launch);
Expand All @@ -46,6 +47,7 @@ HalCoreInfoType create_idle_eth_mem_map() {

std::vector<uint32_t> mem_map_sizes;
mem_map_sizes.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_ETH_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t);
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t);
Expand Down
2 changes: 2 additions & 0 deletions tt_metal/llrt/wormhole/wh_hal_tensix.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@ HalCoreInfoType create_tensix_mem_map() {
std::vector<DeviceAddr> mem_map_bases;

mem_map_bases.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_L1_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_BASE;
mem_map_bases[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = GET_MAILBOX_ADDRESS_HOST(launch);
Expand All @@ -43,6 +44,7 @@ HalCoreInfoType create_tensix_mem_map() {

std::vector<uint32_t> mem_map_sizes;
mem_map_sizes.resize(utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::COUNT));
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BASE)] = MEM_L1_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t);
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_SIZE;
mem_map_sizes[utils::underlying_type<HalL1MemAddrType>(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t);
Expand Down
Loading