From e3a1de2881f51341d4f3828b92066b4cd48426ec Mon Sep 17 00:00:00 2001 From: blozano-tt Date: Fri, 8 Nov 2024 06:29:16 +0000 Subject: [PATCH 1/5] #14634: Remove usage of ARCH_NAME sp constants MEM_L1_SIZE --- tt_metal/hw/inc/debug/sanitize_noc.h | 7 +++++ tt_metal/impl/debug/sanitize_noc_host.hpp | 27 +++++++++++++++++-- tt_metal/impl/device/device.cpp | 2 +- .../impl/dispatch/command_queue_interface.hpp | 19 ++++++++++++- tt_metal/llrt/hal.hpp | 2 +- 5 files changed, 52 insertions(+), 5 deletions(-) diff --git a/tt_metal/hw/inc/debug/sanitize_noc.h b/tt_metal/hw/inc/debug/sanitize_noc.h index 35f005b9801..3a5f0a86c34 100644 --- a/tt_metal/hw/inc/debug/sanitize_noc.h +++ b/tt_metal/hw/inc/debug/sanitize_noc.h @@ -14,6 +14,8 @@ // #pragma once +#include + // NOC logging enabled independently of watcher, need to include it here because it hooks into DEBUG_SANITIZE_NOC_* #include "noc_logging.h" @@ -29,6 +31,11 @@ #include "noc_overlay_parameters.h" #include "noc_parameters.h" +#undef MEM_L1_BASE +#undef MEM_ETH_BASE +constexpr std::uint32_t MEM_L1_BASE = 0x0; +constexpr std::uint32_t MEM_ETH_BASE = 0x0; + // A couple defines for specifying read/write and multi/unicast #define DEBUG_SANITIZE_NOC_READ true #define DEBUG_SANITIZE_NOC_WRITE false diff --git a/tt_metal/impl/debug/sanitize_noc_host.hpp b/tt_metal/impl/debug/sanitize_noc_host.hpp index 58029fe339d..f6deb4f8203 100644 --- a/tt_metal/impl/debug/sanitize_noc_host.hpp +++ b/tt_metal/impl/debug/sanitize_noc_host.hpp @@ -2,15 +2,38 @@ // // SPDX-License-Identifier: Apache-2.0 +#include + #include "noc/noc_parameters.h" #include "noc/noc_overlay_parameters.h" #pragma once +#undef MEM_L1_BASE +#undef MEM_ETH_BASE +constexpr std::uint32_t MEM_L1_BASE = 0x0; +constexpr std::uint32_t MEM_ETH_BASE = 0x0; + +namespace { + inline std::uint32_t compute_size(const tt::tt_metal::Hal& hal, tt::tt_metal::HalProgrammableCoreType core_type) { + return hal.get_dev_size(core_type, tt::tt_metal::HalL1MemAddrType::UNRESERVED) + + hal.get_dev_addr(core_type, tt::tt_metal::HalL1MemAddrType::UNRESERVED); + } + inline std::uint32_t mem_l1_size(const tt::tt_metal::Hal& hal) { + return compute_size(hal, tt::tt_metal::HalProgrammableCoreType::TENSIX); + } + inline std::uint32_t mem_eth_size(const tt::tt_metal::Hal& hal) { + if (hal.get_arch() == tt::ARCH::GRAYSKULL) { + return 0; + } + return compute_size(hal, tt::tt_metal::HalProgrammableCoreType::IDLE_ETH); + } +} + 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) >= MEM_L1_BASE) && ((a) + (l) <= MEM_L1_BASE + mem_l1_size(tt_metal::hal))) // what's the size of the NOC address space? using 0x1000 for now #define DEBUG_VALID_REG_ADDR(a) \ @@ -22,7 +45,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) >= MEM_ETH_BASE) && ((a) + (l) <= MEM_ETH_BASE + mem_eth_size(tt_metal::hal))) static bool coord_found_p(std::vectorcoords, CoreCoord core) { for (CoreCoord item : coords) { diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index ee0d65ba059..a37d9772429 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -1510,7 +1510,7 @@ void Device::update_workers_build_settings(std::vector +namespace { + inline std::uint32_t compute_size(const tt::tt_metal::Hal& hal, tt::tt_metal::HalProgrammableCoreType core_type) { + return hal.get_dev_size(core_type, tt::tt_metal::HalL1MemAddrType::UNRESERVED) + + hal.get_dev_addr(core_type, tt::tt_metal::HalL1MemAddrType::UNRESERVED); + } + inline std::uint32_t mem_l1_size(const tt::tt_metal::Hal& hal) { + return compute_size(hal, tt::tt_metal::HalProgrammableCoreType::TENSIX); + } + inline std::uint32_t mem_eth_size(const tt::tt_metal::Hal& hal) { + if (hal.get_arch() == tt::ARCH::GRAYSKULL) { + return 0; + } + return compute_size(hal, tt::tt_metal::HalProgrammableCoreType::IDLE_ETH); + } +} + +// FIXME: Don't do this in header files using namespace tt::tt_metal; // todo consider moving these to dispatch_addr_map @@ -192,7 +209,7 @@ struct dispatch_constants { uint32_t prefetch_dispatch_unreserved_base = device_cq_addrs_[tt::utils::underlying_type(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 ? mem_l1_size(hal) : mem_eth_size(hal); 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_ = diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp index 875252ca8da..9dca730049f 100644 --- a/tt_metal/llrt/hal.hpp +++ b/tt_metal/llrt/hal.hpp @@ -134,7 +134,7 @@ class Hal { void initialize(tt::ARCH arch); - tt::ARCH get_arch() {return arch_;} + tt::ARCH get_arch() const {return arch_;} uint32_t get_programmable_core_type_count() const; HalProgrammableCoreType get_programmable_core_type(uint32_t core_type_index) const; From a3343f6b2cd3fc747b43e59ad555558f37cef934 Mon Sep 17 00:00:00 2001 From: Bryan Wilder Field Lozano Date: Fri, 8 Nov 2024 01:44:30 -0800 Subject: [PATCH 2/5] Update sanitize_noc.h --- tt_metal/hw/inc/debug/sanitize_noc.h | 7 ------- 1 file changed, 7 deletions(-) diff --git a/tt_metal/hw/inc/debug/sanitize_noc.h b/tt_metal/hw/inc/debug/sanitize_noc.h index 3a5f0a86c34..35f005b9801 100644 --- a/tt_metal/hw/inc/debug/sanitize_noc.h +++ b/tt_metal/hw/inc/debug/sanitize_noc.h @@ -14,8 +14,6 @@ // #pragma once -#include - // NOC logging enabled independently of watcher, need to include it here because it hooks into DEBUG_SANITIZE_NOC_* #include "noc_logging.h" @@ -31,11 +29,6 @@ #include "noc_overlay_parameters.h" #include "noc_parameters.h" -#undef MEM_L1_BASE -#undef MEM_ETH_BASE -constexpr std::uint32_t MEM_L1_BASE = 0x0; -constexpr std::uint32_t MEM_ETH_BASE = 0x0; - // A couple defines for specifying read/write and multi/unicast #define DEBUG_SANITIZE_NOC_READ true #define DEBUG_SANITIZE_NOC_WRITE false From 7861dd20dc8b028d66c1de286d07c6a0f7e8858c Mon Sep 17 00:00:00 2001 From: blozano-tt Date: Thu, 14 Nov 2024 18:28:38 +0000 Subject: [PATCH 3/5] Put MEM_L1/ETH_BASE and MEM_L1/ETH_SIZE behind Hal --- tt_metal/llrt/blackhole/bh_hal_active_eth.cpp | 2 ++ tt_metal/llrt/blackhole/bh_hal_idle_eth.cpp | 2 ++ tt_metal/llrt/blackhole/bh_hal_tensix.cpp | 2 ++ tt_metal/llrt/grayskull/gs_hal.cpp | 2 ++ tt_metal/llrt/hal.hpp | 1 + tt_metal/llrt/wormhole/wh_hal_active_eth.cpp | 2 ++ tt_metal/llrt/wormhole/wh_hal_idle_eth.cpp | 2 ++ tt_metal/llrt/wormhole/wh_hal_tensix.cpp | 2 ++ 8 files changed, 15 insertions(+) diff --git a/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp b/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp index 6695b4d3c39..68772aa77ff 100644 --- a/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp +++ b/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp @@ -32,6 +32,7 @@ HalCoreInfoType create_active_eth_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BASE)] = 0x0; // Anything better to use? mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch); @@ -47,6 +48,7 @@ HalCoreInfoType create_active_eth_mem_map() { std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BASE)] = eth_l1_mem::address_map::MAX_SIZE; // Anything better to use? mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); diff --git a/tt_metal/llrt/blackhole/bh_hal_idle_eth.cpp b/tt_metal/llrt/blackhole/bh_hal_idle_eth.cpp index 693bc25f243..54eaa3390bd 100644 --- a/tt_metal/llrt/blackhole/bh_hal_idle_eth.cpp +++ b/tt_metal/llrt/blackhole/bh_hal_idle_eth.cpp @@ -33,6 +33,7 @@ HalCoreInfoType create_idle_eth_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_ETH_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_IERISC_MAILBOX_ADDRESS_HOST(launch); @@ -47,6 +48,7 @@ HalCoreInfoType create_idle_eth_mem_map() { std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_ETH_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); diff --git a/tt_metal/llrt/blackhole/bh_hal_tensix.cpp b/tt_metal/llrt/blackhole/bh_hal_tensix.cpp index f8ff8589d7c..cffa511aa5b 100644 --- a/tt_metal/llrt/blackhole/bh_hal_tensix.cpp +++ b/tt_metal/llrt/blackhole/bh_hal_tensix.cpp @@ -30,6 +30,7 @@ HalCoreInfoType create_tensix_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_L1_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_MAILBOX_ADDRESS_HOST(launch); @@ -45,6 +46,7 @@ HalCoreInfoType create_tensix_mem_map() { std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_L1_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); diff --git a/tt_metal/llrt/grayskull/gs_hal.cpp b/tt_metal/llrt/grayskull/gs_hal.cpp index 5e6530b5226..baff0a3b780 100644 --- a/tt_metal/llrt/grayskull/gs_hal.cpp +++ b/tt_metal/llrt/grayskull/gs_hal.cpp @@ -41,6 +41,7 @@ void Hal::initialize_gs() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_L1_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_MAILBOX_ADDRESS_HOST(launch); @@ -55,6 +56,7 @@ void Hal::initialize_gs() { std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_L1_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp index 35378eee358..e6543c5cf09 100644 --- a/tt_metal/llrt/hal.hpp +++ b/tt_metal/llrt/hal.hpp @@ -41,6 +41,7 @@ enum class HalProcessorClassType : uint8_t { }; enum class HalL1MemAddrType : uint8_t { + BASE, BARRIER, MAILBOX, LAUNCH, diff --git a/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp b/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp index 5bd82ef3c9e..1296b7e123f 100644 --- a/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp +++ b/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp @@ -32,6 +32,7 @@ HalCoreInfoType create_active_eth_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BASE)] = 0x0; // Anything better to use? mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch); @@ -47,6 +48,7 @@ HalCoreInfoType create_active_eth_mem_map() { std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BASE)] = eth_l1_mem::address_map::MAX_SIZE; // Anything better to use? mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); diff --git a/tt_metal/llrt/wormhole/wh_hal_idle_eth.cpp b/tt_metal/llrt/wormhole/wh_hal_idle_eth.cpp index b4d091c0999..3402f1f8294 100644 --- a/tt_metal/llrt/wormhole/wh_hal_idle_eth.cpp +++ b/tt_metal/llrt/wormhole/wh_hal_idle_eth.cpp @@ -33,6 +33,7 @@ HalCoreInfoType create_idle_eth_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_ETH_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_IERISC_MAILBOX_ADDRESS_HOST(launch); @@ -47,6 +48,7 @@ HalCoreInfoType create_idle_eth_mem_map() { std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_ETH_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_IERISC_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); diff --git a/tt_metal/llrt/wormhole/wh_hal_tensix.cpp b/tt_metal/llrt/wormhole/wh_hal_tensix.cpp index 265c9ca0266..7e37718d9e4 100644 --- a/tt_metal/llrt/wormhole/wh_hal_tensix.cpp +++ b/tt_metal/llrt/wormhole/wh_hal_tensix.cpp @@ -30,6 +30,7 @@ HalCoreInfoType create_tensix_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_L1_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_MAILBOX_ADDRESS_HOST(launch); @@ -44,6 +45,7 @@ HalCoreInfoType create_tensix_mem_map() { std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BASE)] = MEM_L1_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = MEM_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); From 42f21295115ba533d78ef0c14446fb195c1dd0e8 Mon Sep 17 00:00:00 2001 From: blozano-tt Date: Thu, 14 Nov 2024 19:05:50 +0000 Subject: [PATCH 4/5] Update to use Hal --- tt_metal/impl/debug/sanitize_noc_host.hpp | 32 +++---------------- tt_metal/impl/device/device.cpp | 2 +- .../impl/dispatch/command_queue_interface.hpp | 18 +---------- tt_metal/llrt/hal.hpp | 6 ++++ 4 files changed, 13 insertions(+), 45 deletions(-) diff --git a/tt_metal/impl/debug/sanitize_noc_host.hpp b/tt_metal/impl/debug/sanitize_noc_host.hpp index d61085c5de7..967d69ea7e6 100644 --- a/tt_metal/impl/debug/sanitize_noc_host.hpp +++ b/tt_metal/impl/debug/sanitize_noc_host.hpp @@ -2,41 +2,19 @@ // // SPDX-License-Identifier: Apache-2.0 +#pragma once + #include #include "noc/noc_parameters.h" #include "noc/noc_overlay_parameters.h" -// FIXME: ARCH_NAME specific include -#include "dev_mem_map.h" // MEM_[L1/ETH]_BASE - -#pragma once - -#undef MEM_L1_BASE -#undef MEM_ETH_BASE -constexpr std::uint32_t MEM_L1_BASE = 0x0; -constexpr std::uint32_t MEM_ETH_BASE = 0x0; - -namespace { - inline std::uint32_t compute_size(const tt::tt_metal::Hal& hal, tt::tt_metal::HalProgrammableCoreType core_type) { - return hal.get_dev_size(core_type, tt::tt_metal::HalL1MemAddrType::UNRESERVED) + - hal.get_dev_addr(core_type, tt::tt_metal::HalL1MemAddrType::UNRESERVED); - } - inline std::uint32_t mem_l1_size(const tt::tt_metal::Hal& hal) { - return compute_size(hal, tt::tt_metal::HalProgrammableCoreType::TENSIX); - } - inline std::uint32_t mem_eth_size(const tt::tt_metal::Hal& hal) { - if (hal.get_arch() == tt::ARCH::GRAYSKULL) { - return 0; - } - return compute_size(hal, tt::tt_metal::HalProgrammableCoreType::IDLE_ETH); - } -} +#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(tt_metal::hal))) +#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 address space? using 0x1000 for now #define DEBUG_VALID_REG_ADDR(a) \ @@ -48,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(tt_metal::hal))) +#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::vectorcoords, CoreCoord core) { for (CoreCoord item : coords) { diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index 0b1b375522d..f81fc2fbe6c 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -1567,7 +1567,7 @@ void Device::update_workers_build_settings(std::vector(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(hal) : mem_eth_size(hal); + 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_ = diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp index e6543c5cf09..4fbbe44fb24 100644 --- a/tt_metal/llrt/hal.hpp +++ b/tt_metal/llrt/hal.hpp @@ -264,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) From c863e17b409fda5632f5b98ae717e76366e9522c Mon Sep 17 00:00:00 2001 From: Bryan Wilder Field Lozano Date: Thu, 14 Nov 2024 21:04:26 -0800 Subject: [PATCH 5/5] Update tt_metal/llrt/hal.hpp --- tt_metal/llrt/hal.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp index 4fbbe44fb24..7dfcdff349d 100644 --- a/tt_metal/llrt/hal.hpp +++ b/tt_metal/llrt/hal.hpp @@ -266,7 +266,7 @@ extern Hal hal; } // 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) +#define HAL_MEM_L1_BASE tt::tt_metal::hal.get_dev_addr(tt::tt_metal::HalProgrammableCoreType::TENSIX, tt::tt_metal::HalL1MemAddrType::BASE) +#define HAL_MEM_L1_SIZE tt::tt_metal::hal.get_dev_size(tt::tt_metal::HalProgrammableCoreType::TENSIX, tt::tt_metal::HalL1MemAddrType::BASE) +#define HAL_MEM_ETH_BASE tt::tt_metal::hal.get_dev_addr(tt::tt_metal::HalProgrammableCoreType::IDLE_ETH, tt::tt_metal::HalL1MemAddrType::BASE) +#define HAL_MEM_ETH_SIZE tt::tt_metal::hal.get_dev_size(tt::tt_metal::HalProgrammableCoreType::IDLE_ETH, tt::tt_metal::HalL1MemAddrType::BASE)