Remove "noc_parameters.h" from device.cpp (#16582)

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);
         }
     }

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;

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;

tt_metal/llrt/hal.hpp

@@ -153,6 +153,8 @@ class Hal {
     std::vector<uint32_t> dram_sizes_;
     std::vector<uint32_t> 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 <typename IndexType, typename SizeType, typename CoordType>
     auto noc_coordinate(IndexType noc_index, SizeType noc_size, CoordType coord) const

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;