diff --git a/tt_metal/impl/debug/sanitize_noc_host.hpp b/tt_metal/impl/debug/sanitize_noc_host.hpp
index 2fb0049a404..2c575ba5dca 100644
--- a/tt_metal/impl/debug/sanitize_noc_host.hpp
+++ b/tt_metal/impl/debug/sanitize_noc_host.hpp
@@ -16,13 +16,7 @@ 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) >= HAL_MEM_L1_BASE) && ((a) + (l) <= HAL_MEM_L1_BASE + HAL_MEM_L1_SIZE))
 
-// what's the size of the NOC<n> address space?  using 0x1000 for now
-#define DEBUG_VALID_REG_ADDR(a)                                                        \
-    ((((a) >= NOC_OVERLAY_START_ADDR) &&                                               \
-      ((a) < NOC_OVERLAY_START_ADDR + NOC_STREAM_REG_SPACE_SIZE * NOC_NUM_STREAMS)) || \
-     (((a) >= NOC0_REGS_START_ADDR) && ((a) < NOC0_REGS_START_ADDR + 0x1000)) ||       \
-     (((a) >= NOC1_REGS_START_ADDR) && ((a) < NOC1_REGS_START_ADDR + 0x1000)) ||       \
-     ((a) == RISCV_DEBUG_REG_SOFT_RESET_0))
+#define DEBUG_VALID_REG_ADDR(a) tt::tt_metal::hal.valid_reg_addr(a)
 #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))
 
diff --git a/tt_metal/llrt/blackhole/bh_hal.cpp b/tt_metal/llrt/blackhole/bh_hal.cpp
index c583176d239..2fab117e8c7 100644
--- a/tt_metal/llrt/blackhole/bh_hal.cpp
+++ b/tt_metal/llrt/blackhole/bh_hal.cpp
@@ -8,6 +8,8 @@
 #include "core_config.h"  // ProgrammableCoreType
 #include "dev_mem_map.h"
 #include "noc/noc_parameters.h"
+#include "noc/noc_overlay_parameters.h"
+#include "tensix.h"
 
 #include "hal.hpp"
 #include "blackhole/bh_hal.hpp"
@@ -60,6 +62,15 @@ void Hal::initialize_bh() {
         // No relocation needed
         return addr;
     };
+
+    this->valid_reg_addr_func_ = [](uint32_t addr) {
+        return (
+            ((addr >= NOC_OVERLAY_START_ADDR) &&
+             (addr < NOC_OVERLAY_START_ADDR + NOC_STREAM_REG_SPACE_SIZE * NOC_NUM_STREAMS)) ||
+            ((addr >= NOC0_REGS_START_ADDR) && (addr < NOC0_REGS_START_ADDR + 0x1000)) ||
+            ((addr >= NOC1_REGS_START_ADDR) && (addr < NOC1_REGS_START_ADDR + 0x1000)) ||
+            (addr == RISCV_DEBUG_REG_SOFT_RESET_0));
+    };
 }
 
 }  // namespace tt_metal
diff --git a/tt_metal/llrt/grayskull/gs_hal.cpp b/tt_metal/llrt/grayskull/gs_hal.cpp
index 8bde1590ab2..800355e106e 100644
--- a/tt_metal/llrt/grayskull/gs_hal.cpp
+++ b/tt_metal/llrt/grayskull/gs_hal.cpp
@@ -12,6 +12,8 @@
 #include "dev_mem_map.h"
 #include "dev_msgs.h"
 #include "noc/noc_parameters.h"
+#include "noc/noc_overlay_parameters.h"
+#include "tensix.h"
 
 #include "hal.hpp"
 
@@ -142,6 +144,15 @@ void Hal::initialize_gs() {
         // No relocation needed
         return addr;
     };
+
+    this->valid_reg_addr_func_ = [](uint32_t addr) {
+        return (
+            ((addr >= NOC_OVERLAY_START_ADDR) &&
+             (addr < NOC_OVERLAY_START_ADDR + NOC_STREAM_REG_SPACE_SIZE * NOC_NUM_STREAMS)) ||
+            ((addr >= NOC0_REGS_START_ADDR) && (addr < NOC0_REGS_START_ADDR + 0x1000)) ||
+            ((addr >= NOC1_REGS_START_ADDR) && (addr < NOC1_REGS_START_ADDR + 0x1000)) ||
+            (addr == RISCV_DEBUG_REG_SOFT_RESET_0));
+    };
 }
 
 }  // namespace tt_metal
diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp
index 52a1d3c65bf..38fe0a6da86 100644
--- a/tt_metal/llrt/hal.hpp
+++ b/tt_metal/llrt/hal.hpp
@@ -137,6 +137,7 @@ inline T HalCoreInfoType::get_binary_local_init_addr(uint32_t processor_class_id
 class Hal {
 public:
     using RelocateFunc = std::function<uint64_t(uint64_t, uint64_t)>;
+    using ValidRegAddrFunc = std::function<bool(uint32_t)>;
 
 private:
     tt::ARCH arch_;
@@ -151,6 +152,7 @@ class Hal {
 
     // Functions where implementation varies by architecture
     RelocateFunc relocate_func_;
+    ValidRegAddrFunc valid_reg_addr_func_;
 
 public:
     Hal();
@@ -199,6 +201,8 @@ class Hal {
     uint64_t relocate_dev_addr(uint64_t addr, uint64_t local_init_addr = 0) {
         return relocate_func_(addr, local_init_addr);
     }
+
+    uint32_t valid_reg_addr(uint32_t addr) { return valid_reg_addr_func_(addr); }
 };
 
 inline uint32_t Hal::get_programmable_core_type_count() const { return core_info_.size(); }
diff --git a/tt_metal/llrt/wormhole/wh_hal.cpp b/tt_metal/llrt/wormhole/wh_hal.cpp
index c00f95d1448..4c4fed3aeb5 100644
--- a/tt_metal/llrt/wormhole/wh_hal.cpp
+++ b/tt_metal/llrt/wormhole/wh_hal.cpp
@@ -8,6 +8,8 @@
 #include "core_config.h"  // ProgrammableCoreType
 #include "dev_mem_map.h"  // MEM_LOCAL_BASE
 #include "noc/noc_parameters.h"
+#include "noc/noc_overlay_parameters.h"
+#include "tensix.h"
 
 #include "hal.hpp"
 #include "wormhole/wh_hal.hpp"
@@ -61,6 +63,15 @@ void Hal::initialize_wh() {
         // No relocation needed
         return addr;
     };
+
+    this->valid_reg_addr_func_ = [](uint32_t addr) {
+        return (
+            ((addr >= NOC_OVERLAY_START_ADDR) &&
+             (addr < NOC_OVERLAY_START_ADDR + NOC_STREAM_REG_SPACE_SIZE * NOC_NUM_STREAMS)) ||
+            ((addr >= NOC0_REGS_START_ADDR) && (addr < NOC0_REGS_START_ADDR + 0x1000)) ||
+            ((addr >= NOC1_REGS_START_ADDR) && (addr < NOC1_REGS_START_ADDR + 0x1000)) ||
+            (addr == RISCV_DEBUG_REG_SOFT_RESET_0));
+    };
 }
 
 }  // namespace tt_metal