Move relocate_dev_addr behind Hal (#15312)

tt_metal/llrt/blackhole/bh_hal.cpp

@@ -6,6 +6,7 @@
 #include <cstdint>
 
 #include "core_config.h" // ProgrammableCoreType
+#include "dev_mem_map.h"
 #include "noc/noc_parameters.h"
 
 #include "hal.hpp"
@@ -43,6 +44,20 @@ void Hal::initialize_bh() {
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::L1)] = L1_ALIGNMENT;
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::DRAM)] = DRAM_ALIGNMENT;
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::HOST)] = PCIE_ALIGNMENT;
+
+    this->relocate_func_ = [](uint64_t addr, uint64_t local_init_addr) {
+        if ((addr & MEM_LOCAL_BASE) == MEM_LOCAL_BASE) {
+            // Move addresses in the local memory range to l1 (copied by kernel)
+            return (addr & ~MEM_LOCAL_BASE) + local_init_addr;
+        }
+
+        // Note: Blackhole does not have IRAM
+
+        // No relocation needed
+        return addr;
+    };
+
+
 }
 
 }  // namespace tt_metal

tt_metal/llrt/grayskull/gs_hal.cpp

@@ -129,6 +129,20 @@ void Hal::initialize_gs() {
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::L1)] = L1_ALIGNMENT;
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::DRAM)] = DRAM_ALIGNMENT;
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::HOST)] = PCIE_ALIGNMENT;
+
+    this->relocate_func_ = [](uint64_t addr, uint64_t local_init_addr) {
+        if ((addr & MEM_LOCAL_BASE) == MEM_LOCAL_BASE) {
+            // Move addresses in the local memory range to l1 (copied by kernel)
+            return (addr & ~MEM_LOCAL_BASE) + local_init_addr;
+        }
+        else if ((addr & MEM_NCRISC_IRAM_BASE) == MEM_NCRISC_IRAM_BASE) {
+            // Move addresses in the NCRISC memory range to l1 (copied by kernel)
+            return (addr & ~MEM_NCRISC_IRAM_BASE) + MEM_NCRISC_INIT_IRAM_L1_BASE;
+        }
+
+        // No relocation needed
+        return addr;
+    };
 }
 
 }  // namespace tt_metal

tt_metal/llrt/hal.hpp

@@ -10,6 +10,7 @@
 //
 
 #include <cstdint>
+#include <functional>
 #include <variant>
 #include <vector>
 #include <memory>
@@ -143,6 +144,10 @@ 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)>;
+
   private:
     tt::ARCH arch_;
     std::vector<HalCoreInfoType> core_info_;
@@ -154,6 +159,9 @@ class Hal {
     void initialize_wh();
     void initialize_bh();
 
+    // Functions where implementation varies by architecture
+    RelocateFunc relocate_func_;
+
   public:
     Hal();
 
@@ -195,6 +203,11 @@ class Hal {
     T get_base_firmware_addr(uint32_t programmable_core_type_index, uint32_t processor_class_idx, uint32_t processor_type_idx) const;
     template <typename T = DeviceAddr>
     T get_binary_local_init_addr(uint32_t programmable_core_type_index, uint32_t processor_class_idx, uint32_t processor_type_idx) const;
+
+    uint64_t relocate_dev_addr(uint64_t addr, uint64_t local_init_addr = 0) {
+        return relocate_func_(addr, local_init_addr);
+    }
+
 };
 
 inline uint32_t Hal::get_programmable_core_type_count() const {

tt_metal/llrt/llrt.cpp

@@ -138,7 +138,7 @@ ll_api::memory read_mem_from_core(chip_id_t chip, const CoreCoord &core, const l
 
     ll_api::memory read_mem;
     read_mem.fill_from_mem_template(mem, [&](std::vector<uint32_t>::iterator mem_ptr, uint64_t addr, uint32_t len) {
-        uint64_t relo_addr = relocate_dev_addr(addr, local_init_addr);
+        uint64_t relo_addr = tt::tt_metal::hal.relocate_dev_addr(addr, local_init_addr);
         tt::Cluster::instance().read_core(&*mem_ptr, len * sizeof(uint32_t), tt_cxy_pair(chip, core), relo_addr);
     });
     return read_mem;
@@ -185,7 +185,7 @@ bool test_load_write_read_risc_binary(
 
     log_debug(tt::LogLLRuntime, "hex_vec size = {}, size_in_bytes = {}", mem.size(), mem.size()*sizeof(uint32_t));
     mem.process_spans([&](std::vector<uint32_t>::const_iterator mem_ptr, uint64_t addr, uint32_t len_words) {
-        uint64_t relo_addr = relocate_dev_addr(addr, local_init_addr);
+        uint64_t relo_addr = tt::tt_metal::hal.relocate_dev_addr(addr, local_init_addr);
 
         tt::Cluster::instance().write_core(&*mem_ptr, len_words * sizeof(uint32_t), tt_cxy_pair(chip_id, core), relo_addr);
     });

tt_metal/llrt/llrt.hpp

@@ -120,24 +120,6 @@ void wait_until_cores_done(
 
 }  // namespace internal_
 
-inline uint64_t relocate_dev_addr(uint64_t addr, uint64_t local_init_addr = 0) {
-    uint64_t relo_addr;
-    if ((addr & MEM_LOCAL_BASE) == MEM_LOCAL_BASE) {
-        // Move addresses in the local memory range to l1 (copied by kernel)
-        relo_addr = (addr & ~MEM_LOCAL_BASE) + local_init_addr;
-    }
-#ifdef NCRISC_HAS_IRAM
-    else if ((addr & MEM_NCRISC_IRAM_BASE) == MEM_NCRISC_IRAM_BASE) {
-        // Move addresses in the trisc memory range to l1 (copied by kernel)
-        relo_addr = (addr & ~MEM_NCRISC_IRAM_BASE) + MEM_NCRISC_INIT_IRAM_L1_BASE;
-    }
-#endif
-    else {
-        relo_addr = addr;
-    }
-    return relo_addr;
-}
-
 }  // namespace llrt
 
 }  // namespace tt

tt_metal/llrt/wormhole/wh_hal.cpp

@@ -6,6 +6,7 @@
 #include <cstdint>
 
 #include "core_config.h" // ProgrammableCoreType
+#include "dev_mem_map.h" // MEM_LOCAL_BASE
 #include "noc/noc_parameters.h"
 
 #include "hal.hpp"
@@ -43,6 +44,21 @@ void Hal::initialize_wh() {
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::L1)] = L1_ALIGNMENT;
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::DRAM)] = DRAM_ALIGNMENT;
     this->mem_alignments_[static_cast<std::size_t>(HalMemType::HOST)] = PCIE_ALIGNMENT;
+
+    this->relocate_func_ = [](uint64_t addr, uint64_t local_init_addr) {
+        if ((addr & MEM_LOCAL_BASE) == MEM_LOCAL_BASE) {
+            // Move addresses in the local memory range to l1 (copied by kernel)
+            return (addr & ~MEM_LOCAL_BASE) + local_init_addr;
+        }
+        else if ((addr & MEM_NCRISC_IRAM_BASE) == MEM_NCRISC_IRAM_BASE) {
+            // Move addresses in the NCRISC memory range to l1 (copied by kernel)
+            return (addr & ~MEM_NCRISC_IRAM_BASE) + MEM_NCRISC_INIT_IRAM_L1_BASE;
+        }
+
+        // No relocation needed
+        return addr;
+    };
+
 }
 
 }  // namespace tt_metal