Add get_num_nocs() to Hal

tenstorrent · Dec 4, 2024 · d7c05d1 · d7c05d1
1 parent 3ef4b98
commit d7c05d1
Show file tree

Hide file tree

Showing 6 changed files with 12 additions and 3 deletions.
diff --git a/tt_metal/impl/debug/watcher_device_reader.cpp b/tt_metal/impl/debug/watcher_device_reader.cpp
@@ -356,7 +356,8 @@ void WatcherDeviceReader::DumpCore(CoreDescriptor& logical_core, bool is_active_
             DumpL1Status(core, &mbox_data->launch[launch_msg_read_ptr]);
         }
         if (!tt::llrt::OptionsG.watcher_noc_sanitize_disabled()) {
-            for (uint32_t noc = 0; noc < NUM_NOCS; noc++) {
+            const auto NUM_NOCS_ = tt::tt_metal::hal.get_num_nocs();
+            for (uint32_t noc = 0; noc < NUM_NOCS_; noc++) {
                 DumpNocSanitizeStatus(core, core_str, mbox_data, noc);
             }
         }

diff --git a/tt_metal/impl/debug/watcher_server.cpp b/tt_metal/impl/debug/watcher_server.cpp
@@ -18,8 +18,6 @@
 #include "dev_msgs.h"
 #include "llrt/llrt.hpp"
 #include "llrt/rtoptions.hpp"
-#include "noc/noc_overlay_parameters.h"
-#include "noc/noc_parameters.h"
 #include "debug/ring_buffer.h"
 #include "watcher_device_reader.hpp"
 
@@ -220,6 +218,7 @@ void watcher_init(Device* device) {
     }
 
     // Initialize debug sanity L1/NOC addresses to sentinel "all ok"
+    const auto NUM_NOCS = tt::tt_metal::hal.get_num_nocs();
     for (int i = 0; i < NUM_NOCS; i++) {
         data->sanitize_noc[i].noc_addr = watcher::DEBUG_SANITIZE_NOC_SENTINEL_OK_64;
         data->sanitize_noc[i].l1_addr = watcher::DEBUG_SANITIZE_NOC_SENTINEL_OK_32;

diff --git a/tt_metal/llrt/blackhole/bh_hal.cpp b/tt_metal/llrt/blackhole/bh_hal.cpp
@@ -76,6 +76,8 @@ void Hal::initialize_bh() {
     this->noc_multicast_encoding_func_ = [](uint32_t x_start, uint32_t y_start, uint32_t x_end, uint32_t y_end) {
         return NOC_MULTICAST_ENCODING(x_start, y_start, x_end, y_end);
     };
+
+    num_nocs_ = NUM_NOCS;
 }
 
 }  // namespace tt_metal

diff --git a/tt_metal/llrt/grayskull/gs_hal.cpp b/tt_metal/llrt/grayskull/gs_hal.cpp
@@ -158,6 +158,8 @@ void Hal::initialize_gs() {
     this->noc_multicast_encoding_func_ = [](uint32_t x_start, uint32_t y_start, uint32_t x_end, uint32_t y_end) {
         return NOC_MULTICAST_ENCODING(x_start, y_start, x_end, y_end);
     };
+
+    num_nocs_ = NUM_NOCS;
 }
 
 }  // namespace tt_metal

diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp
@@ -147,6 +147,7 @@ class Hal {
     std::vector<DeviceAddr> dram_bases_;
     std::vector<uint32_t> dram_sizes_;
     std::vector<uint32_t> mem_alignments_;
+    uint32_t num_nocs_;
 
     void initialize_gs();
     void initialize_wh();
@@ -163,6 +164,8 @@ class Hal {
 
     tt::ARCH get_arch() const { return arch_; }
 
+    uint32_t get_num_nocs() const { return num_nocs_; }
+
     template <typename IndexType, typename SizeType, typename CoordType>
     auto noc_coordinate(IndexType noc_index, SizeType noc_size, CoordType coord) const
         -> decltype(noc_size - 1 - coord) {

diff --git a/tt_metal/llrt/wormhole/wh_hal.cpp b/tt_metal/llrt/wormhole/wh_hal.cpp
@@ -77,6 +77,8 @@ void Hal::initialize_wh() {
     this->noc_multicast_encoding_func_ = [](uint32_t x_start, uint32_t y_start, uint32_t x_end, uint32_t y_end) {
         return NOC_MULTICAST_ENCODING(x_start, y_start, x_end, y_end);
     };
+
+    num_nocs_ = NUM_NOCS;
 }
 
 }  // namespace tt_metal