diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index 742a53b8976..0f8ae984ccc 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -28,9 +28,6 @@ #include "tt_metal/impl/sub_device/sub_device_types.hpp" #include "tt_metal/tt_stl/span.hpp" #include "tt_metal/types.hpp" -#include "noc/noc_parameters.h" - -// FIXME: ARCH_NAME specific #include "impl/dispatch/topology.hpp" namespace tt { @@ -1716,7 +1713,8 @@ void Device::generate_device_bank_to_noc_tables() for (unsigned int bank_id = 0; bank_id < dram_noc_coord_per_bank.size(); bank_id++) { uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.x, dram_noc_coord_per_bank[bank_id].x); uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.y, dram_noc_coord_per_bank[bank_id].y); - uint16_t xy = ((noc_y << NOC_ADDR_NODE_ID_BITS) | noc_x) << NOC_COORD_REG_OFFSET; + uint16_t xy = ((noc_y << tt::tt_metal::hal.get_noc_addr_node_id_bits()) | noc_x) + << tt::tt_metal::hal.get_noc_coord_reg_offset(); dram_bank_to_noc_xy_.push_back(xy); } } @@ -1728,7 +1726,8 @@ void Device::generate_device_bank_to_noc_tables() auto l1_noc_coords = this->virtual_noc0_coordinate(noc, l1_noc_coord_per_bank[bank_id]); uint16_t noc_x = l1_noc_coords.x; uint16_t noc_y = l1_noc_coords.y; - uint16_t xy = ((noc_y << NOC_ADDR_NODE_ID_BITS) | noc_x) << NOC_COORD_REG_OFFSET; + uint16_t xy = ((noc_y << tt::tt_metal::hal.get_noc_addr_node_id_bits()) | noc_x) + << tt::tt_metal::hal.get_noc_coord_reg_offset(); l1_bank_to_noc_xy_.push_back(xy); } } diff --git a/tt_metal/llrt/blackhole/bh_hal.cpp b/tt_metal/llrt/blackhole/bh_hal.cpp index ec46fa6efc6..a4a415907b7 100644 --- a/tt_metal/llrt/blackhole/bh_hal.cpp +++ b/tt_metal/llrt/blackhole/bh_hal.cpp @@ -93,6 +93,8 @@ void Hal::initialize_bh() { }; this->num_nocs_ = NUM_NOCS; + this->noc_addr_node_id_bits_ = NOC_ADDR_NODE_ID_BITS; + this->noc_coord_reg_offset_ = NOC_COORD_REG_OFFSET; this->coordinate_virtualization_enabled_ = COORDINATE_VIRTUALIZATION_ENABLED; this->virtual_worker_start_x_ = VIRTUAL_TENSIX_START_X; this->virtual_worker_start_y_ = VIRTUAL_TENSIX_START_Y; diff --git a/tt_metal/llrt/grayskull/gs_hal.cpp b/tt_metal/llrt/grayskull/gs_hal.cpp index 7a2d51791cb..27fda5760b1 100644 --- a/tt_metal/llrt/grayskull/gs_hal.cpp +++ b/tt_metal/llrt/grayskull/gs_hal.cpp @@ -176,6 +176,8 @@ void Hal::initialize_gs() { }; this->num_nocs_ = NUM_NOCS; + this->noc_addr_node_id_bits_ = NOC_ADDR_NODE_ID_BITS; + this->noc_coord_reg_offset_ = NOC_COORD_REG_OFFSET; this->coordinate_virtualization_enabled_ = COORDINATE_VIRTUALIZATION_ENABLED; this->virtual_worker_start_x_ = VIRTUAL_TENSIX_START_X; this->virtual_worker_start_y_ = VIRTUAL_TENSIX_START_Y; diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp index ea06aad8fb4..0457f69593c 100644 --- a/tt_metal/llrt/hal.hpp +++ b/tt_metal/llrt/hal.hpp @@ -153,6 +153,8 @@ class Hal { std::vector dram_sizes_; std::vector mem_alignments_; uint32_t num_nocs_; + uint32_t noc_addr_node_id_bits_; + uint32_t noc_coord_reg_offset_; bool coordinate_virtualization_enabled_; uint32_t virtual_worker_start_x_; uint32_t virtual_worker_start_y_; @@ -174,6 +176,8 @@ class Hal { tt::ARCH get_arch() const { return arch_; } uint32_t get_num_nocs() const { return num_nocs_; } + uint32_t get_noc_addr_node_id_bits() const { return noc_addr_node_id_bits_; } + uint32_t get_noc_coord_reg_offset() const { return noc_coord_reg_offset_; } template auto noc_coordinate(IndexType noc_index, SizeType noc_size, CoordType coord) const diff --git a/tt_metal/llrt/wormhole/wh_hal.cpp b/tt_metal/llrt/wormhole/wh_hal.cpp index 29e4c619339..3b141c03f0a 100644 --- a/tt_metal/llrt/wormhole/wh_hal.cpp +++ b/tt_metal/llrt/wormhole/wh_hal.cpp @@ -94,6 +94,8 @@ void Hal::initialize_wh() { }; this->num_nocs_ = NUM_NOCS; + this->noc_addr_node_id_bits_ = NOC_ADDR_NODE_ID_BITS; + this->noc_coord_reg_offset_ = NOC_COORD_REG_OFFSET; this->coordinate_virtualization_enabled_ = COORDINATE_VIRTUALIZATION_ENABLED; this->virtual_worker_start_x_ = VIRTUAL_TENSIX_START_X; this->virtual_worker_start_y_ = VIRTUAL_TENSIX_START_Y;