From d7c05d19c647ad723eb6c2c8c0125585d0f96bbf Mon Sep 17 00:00:00 2001
From: blozano-tt <blozano@tenstorrent.com>
Date: Wed, 4 Dec 2024 01:02:10 +0000
Subject: [PATCH] Add get_num_nocs() to Hal

---
 tt_metal/impl/debug/watcher_device_reader.cpp | 3 ++-
 tt_metal/impl/debug/watcher_server.cpp        | 3 +--
 tt_metal/llrt/blackhole/bh_hal.cpp            | 2 ++
 tt_metal/llrt/grayskull/gs_hal.cpp            | 2 ++
 tt_metal/llrt/hal.hpp                         | 3 +++
 tt_metal/llrt/wormhole/wh_hal.cpp             | 2 ++
 6 files changed, 12 insertions(+), 3 deletions(-)

diff --git a/tt_metal/impl/debug/watcher_device_reader.cpp b/tt_metal/impl/debug/watcher_device_reader.cpp
index 002a4a9303e..365b20f957d 100644
--- 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
index eb17467b6df..8b92836a1fb 100644
--- 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
index 56bee5d6af7..1096da8ec65 100644
--- 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
index 94fcc312530..5477beeec65 100644
--- 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
index 51c0619f60a..f7da19e2f97 100644
--- 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
index 32b5e454b79..4f20cfb9993 100644
--- 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