Move NOC_0_X/Y behind Hal (#14920)

The NOC_0_X and NOC_0_Y macros which were used for determining noc
coordinates have been collapsed into a single Hal API. The new Hal API
can later be specialized to behave differently.

CMakeLists.txt

@@ -203,6 +203,9 @@ target_link_libraries(
         numa
 )
 
+if(NOT DEFINED ENV{ARCH_NAME})
+    message(FATAL_ERROR "Please set ARCH_NAME to grayskull, wormhole_b0, or blackhole")
+endif(NOT DEFINED ENV{ARCH_NAME})
 string(TOUPPER "$ENV{ARCH_NAME}" ARCH_NAME_DEF)
 add_compile_definitions(ARCH_${ARCH_NAME_DEF})
 add_compile_options(

tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h

@@ -13,7 +13,7 @@
 #include "tt_metal/impl/dispatch/cq_commands.hpp"
 #include "noc/noc_parameters.h"
 
-#include "tt_metal/hostdevcommon/common_runtime_address_map.h" // NOC_0_X
+#include "tt_metal/llrt/hal.hpp"
 
 extern bool debug_g;
 extern bool use_coherent_data_g;
@@ -490,15 +490,15 @@ void configure_kernel_variant(
     const auto& grid_size = device->grid_size();
 
     std::map<string, string> defines = {
-        {"MY_NOC_X", std::to_string(NOC_0_X(my_noc_index, grid_size.x, phys_my_core.x))},
-        {"MY_NOC_Y", std::to_string(NOC_0_Y(my_noc_index, grid_size.y, phys_my_core.y))},
+        {"MY_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.x, phys_my_core.x))},
+        {"MY_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.y, phys_my_core.y))},
         {"UPSTREAM_NOC_INDEX", std::to_string(upstream_noc_index)},
-        {"UPSTREAM_NOC_X", std::to_string(NOC_0_X(upstream_noc_index, grid_size.x, phys_upstream_core.x))},
-        {"UPSTREAM_NOC_Y", std::to_string(NOC_0_Y(upstream_noc_index, grid_size.y, phys_upstream_core.y))},
-        {"DOWNSTREAM_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, phys_downstream_core.x))},
-        {"DOWNSTREAM_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, phys_downstream_core.y))},
-        {"DOWNSTREAM_SLAVE_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, 0xff))},
-        {"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, 0xff))}, // todo, add testing with dispatch_s once it processes more than go signals
+        {"UPSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.x, phys_upstream_core.x))},
+        {"UPSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.y, phys_upstream_core.y))},
+        {"DOWNSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, phys_downstream_core.x))},
+        {"DOWNSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, phys_downstream_core.y))},
+        {"DOWNSTREAM_SLAVE_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, 0xff))},
+        {"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, 0xff))}, // todo, add testing with dispatch_s once it processes more than go signals
         {"FD_CORE_TYPE", std::to_string(0)}, // todo, support dispatch on eth
     };
     compile_args.push_back(is_dram_variant);

tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_prefetcher.cpp

@@ -16,8 +16,6 @@
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp"
 
-#include "tt_metal/hostdevcommon/common_runtime_address_map.h" // NOC_0_X
-
 #include "llrt/hal.hpp"
 
 #define CQ_PREFETCH_CMD_BARE_MIN_SIZE tt::tt_metal::hal.get_alignment(tt::tt_metal::HalMemType::HOST)

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_2ep.cpp

@@ -6,7 +6,6 @@
 #include "tt_metal/detail/tt_metal.hpp"
 #include "tt_metal/llrt/rtoptions.hpp"
 #include "tt_metal/impl/dispatch/cq_commands.hpp"
-#include "tt_metal/hostdevcommon/common_runtime_address_map.h"
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
 #include "tt_metal/impl/device/device.hpp"

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_4ep.cpp

@@ -6,7 +6,6 @@
 #include "tt_metal/detail/tt_metal.hpp"
 #include "tt_metal/llrt/rtoptions.hpp"
 #include "tt_metal/impl/dispatch/cq_commands.hpp"
-#include "tt_metal/hostdevcommon/common_runtime_address_map.h"
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
 #include "tt_metal/impl/device/device.hpp"

tests/tt_metal/tt_metal/unit_tests_common/watcher/test_noc_sanitize.cpp

@@ -8,7 +8,10 @@
 #include "tt_metal/detail/tt_metal.hpp"
 #include "tt_metal/host_api.hpp"
 #include "common/bfloat16.hpp"
-#include "hostdevcommon/common_runtime_address_map.h"
+
+// Do we really want to expose Hal like this?
+// This looks like an API level test
+#include "llrt/hal.hpp"
 
 //////////////////////////////////////////////////////////////////////////////////////////
 // A test for checking watcher NOC sanitization.
@@ -155,8 +158,8 @@ void RunTestOnCore(WatcherFixture* fixture, Device* device, CoreCoord &core, boo
             const metal_SocDescriptor& soc_d = tt::Cluster::instance().get_soc_desc(device->id());
             int noc = (use_ncrisc) ? 1 : 0;
             CoreCoord target_phys_core = {
-                NOC_0_X(noc, soc_d.grid_size.x, input_dram_noc_xy.x),
-                NOC_0_Y(noc, soc_d.grid_size.y, input_dram_noc_xy.y)
+                tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.x, input_dram_noc_xy.x),
+                tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.y, input_dram_noc_xy.y)
             };
             string risc_name = (is_eth_core) ? "erisc" : "brisc";
             if (use_ncrisc)

tests/tt_metal/tt_metal/unit_tests_fast_dispatch/common/command_queue_fixture.hpp

@@ -11,7 +11,6 @@
 #include "common/core_coord.hpp"
 #include "common/env_lib.hpp"
 #include "gtest/gtest.h"
-#include "hostdevcommon/common_runtime_address_map.h"
 #include "hostdevcommon/common_values.hpp"
 #include "impl/buffers/circular_buffer_types.hpp"
 #include "impl/device/device.hpp"

tt_metal/hostdevcommon/common_runtime_address_map.h

@@ -17,8 +17,4 @@
 constexpr static std::uint32_t L1_KERNEL_CONFIG_BASE = MEM_MAP_END;
 constexpr static std::uint32_t L1_KERNEL_CONFIG_SIZE = 69 * 1024;
 
-// Helper functions to convert NoC coordinates to NoC-0 coordinates, used in metal as "physical" coordinates.
-#define NOC_0_X(noc_index, noc_size_x, x) (noc_index == 0 ? (x) : (noc_size_x-1-(x)))
-#define NOC_0_Y(noc_index, noc_size_y, y) (noc_index == 0 ? (y) : (noc_size_y-1-(y)))
-
 static_assert(L1_KERNEL_CONFIG_BASE % L1_ALIGNMENT == 0);

tt_metal/impl/debug/watcher_device_reader.cpp

@@ -16,7 +16,6 @@
 // FIXME: Avoid dependence on ARCH_NAME specific includes
 #include "dev_mem_map.h" // for MEM_BRISC_STAC...
 #include "eth_l1_address_map.h" // for address_map
-#include "hostdevcommon/common_runtime_address_map.h" // for NOC_0_X, NOC_0_Y
 #include "hw/inc/dev_msgs.h"
 
 #include "third_party/umd/device/tt_arch_types.h"
@@ -79,7 +78,7 @@ static string get_noc_target_str(Device *device, CoreDescriptor &core, int noc,
         // Get the physical coord from the noc coord
         const metal_SocDescriptor &soc_d = tt::Cluster::instance().get_soc_desc(device->id());
         CoreCoord phys_core = {
-            NOC_0_X(noc, soc_d.grid_size.x, noc_coord.x), NOC_0_Y(noc, soc_d.grid_size.y, noc_coord.y)};
+            tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.x, noc_coord.x), tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.y, noc_coord.y)};
 
         CoreType core_type;
         try {

tt_metal/impl/device/device.cpp

@@ -29,8 +29,6 @@
 #include "tt_metal/impl/sub_device/sub_device_types.hpp"
 #include "tt_metal/tt_stl/span.hpp"
 
-#include "tt_metal/hostdevcommon/common_runtime_address_map.h" // NOC_0_X
-
 namespace tt {
 
 namespace tt_metal {
@@ -323,8 +321,8 @@ void Device::initialize_device_kernel_defines()
     auto grid_size = this->grid_size();
     this->device_kernel_defines_.emplace("PCIE_NOC_X", std::to_string(pcie_cores[0].x));
     this->device_kernel_defines_.emplace("PCIE_NOC_Y", std::to_string(pcie_cores[0].y));
-    this->device_kernel_defines_.emplace("PCIE_NOC1_X", std::to_string(NOC_0_X(NOC::NOC_1, grid_size.x, pcie_cores[0].x)));
-    this->device_kernel_defines_.emplace("PCIE_NOC1_Y", std::to_string(NOC_0_X(NOC::NOC_1, grid_size.x, pcie_cores[0].y)));
+    this->device_kernel_defines_.emplace("PCIE_NOC1_X", std::to_string(tt::tt_metal::hal.noc_coordinate(NOC::NOC_1, grid_size.x, pcie_cores[0].x)));
+    this->device_kernel_defines_.emplace("PCIE_NOC1_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(NOC::NOC_1, grid_size.x, pcie_cores[0].y)));
 }
 
 void Device::initialize_build() {
@@ -797,15 +795,15 @@ void Device::configure_kernel_variant(
 
     std::map<string, string> defines = {
         {"DISPATCH_KERNEL", "1"},
-        {"MY_NOC_X", std::to_string(NOC_0_X(my_noc_index, grid_size.x, kernel_physical_core.x))},
-        {"MY_NOC_Y", std::to_string(NOC_0_Y(my_noc_index, grid_size.y, kernel_physical_core.y))},
+        {"MY_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.x, kernel_physical_core.x))},
+        {"MY_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.y, kernel_physical_core.y))},
         {"UPSTREAM_NOC_INDEX", std::to_string(upstream_noc_index)},
-        {"UPSTREAM_NOC_X", std::to_string(NOC_0_X(upstream_noc_index, grid_size.x, upstream_physical_core.x))},
-        {"UPSTREAM_NOC_Y", std::to_string(NOC_0_Y(upstream_noc_index, grid_size.y, upstream_physical_core.y))},
-        {"DOWNSTREAM_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, downstream_physical_core.x))},
-        {"DOWNSTREAM_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, downstream_physical_core.y))},
-        {"DOWNSTREAM_SLAVE_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, downstream_slave_physical_core.x))},
-        {"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, downstream_slave_physical_core.y))},
+        {"UPSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.x, upstream_physical_core.x))},
+        {"UPSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.y, upstream_physical_core.y))},
+        {"DOWNSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, downstream_physical_core.x))},
+        {"DOWNSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, downstream_physical_core.y))},
+        {"DOWNSTREAM_SLAVE_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, downstream_slave_physical_core.x))},
+        {"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, downstream_slave_physical_core.y))},
         {"FD_CORE_TYPE", std::to_string(programmable_core_type_index)},
     };
     if (force_watcher_no_inline) {
@@ -3134,8 +3132,8 @@ std::vector<CoreCoord> Device::ethernet_cores_from_logical_cores(const std::vect
 uint32_t Device::get_noc_unicast_encoding(uint8_t noc_index, const CoreCoord& physical_core) const {
     const auto& grid_size = this->grid_size();
     return NOC_XY_ENCODING(
-        NOC_0_X(noc_index, grid_size.x, physical_core.x),
-        NOC_0_Y(noc_index, grid_size.y, physical_core.y)
+        tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_core.x),
+        tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_core.y)
     );
 }
 
@@ -3145,17 +3143,17 @@ uint32_t Device::get_noc_multicast_encoding(uint8_t noc_index, const CoreRange&
     // NOC 1 mcasts from bottom left to top right, so we need to reverse the coords
     if (noc_index == 0) {
         return NOC_MULTICAST_ENCODING(
-            NOC_0_X(noc_index, grid_size.x, physical_cores.start_coord.x),
-            NOC_0_Y(noc_index, grid_size.y, physical_cores.start_coord.y),
-            NOC_0_X(noc_index, grid_size.x, physical_cores.end_coord.x),
-            NOC_0_Y(noc_index, grid_size.y, physical_cores.end_coord.y)
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.start_coord.x),
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.start_coord.y),
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.end_coord.x),
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.end_coord.y)
         );
     } else {
         return NOC_MULTICAST_ENCODING(
-            NOC_0_X(noc_index, grid_size.x, physical_cores.end_coord.x),
-            NOC_0_Y(noc_index, grid_size.y, physical_cores.end_coord.y),
-            NOC_0_X(noc_index, grid_size.x, physical_cores.start_coord.x),
-            NOC_0_Y(noc_index, grid_size.y, physical_cores.start_coord.y)
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.end_coord.x),
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.end_coord.y),
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.start_coord.x),
+            tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.start_coord.y)
         );
     }
 }

tt_metal/impl/program/program.cpp

@@ -1153,7 +1153,8 @@ uint32_t detail::Program_::finalize_rt_args(uint32_t programmable_core_type_inde
 
     // TODO: this is asserted here as the leveling above can break the limits enforced by the API
     // Once we use a ring buffer, memory space will be dynamic and this assert won't matter
-    TT_FATAL(offset <= L1_KERNEL_CONFIG_SIZE, "offset {} cannot exceed config size {}", offset, L1_KERNEL_CONFIG_SIZE);
+    std::uint32_t l1_kernel_config_size = tt::tt_metal::hal.get_dev_size(tt::tt_metal::HalProgrammableCoreType::TENSIX, tt::tt_metal::HalL1MemAddrType::KERNEL_CONFIG);
+    TT_FATAL(offset <= l1_kernel_config_size, "offset {} cannot exceed config size {}", offset, l1_kernel_config_size);
 
     return max_unique_rta_size + total_crta_size;
 }

tt_metal/jit_build/genfiles.cpp

@@ -11,7 +11,6 @@
 
 #include "common/tt_backend_api_types.hpp"
 #include "common/utils.hpp"
-#include "hostdevcommon/common_runtime_address_map.h" // NOC_0_X
 #include "hostdevcommon/common_values.hpp"
 #include "jit_build/build.hpp"
 #include "jit_build/settings.hpp"
@@ -589,8 +588,8 @@ std::string generate_bank_to_noc_coord_descriptor_string(
         ss << "    {"
            << "\t// noc=" << noc << endl;
         for (unsigned int bank_id = 0; bank_id < dram_bank_map.size(); bank_id++) {
-            uint16_t noc_x = NOC_0_X(noc, grid_size.x, dram_bank_map[bank_id].x);
-            uint16_t noc_y = NOC_0_Y(noc, grid_size.y, dram_bank_map[bank_id].y);
+            uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, grid_size.x, dram_bank_map[bank_id].x);
+            uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, grid_size.y, dram_bank_map[bank_id].y);
             ss << "        (((" << noc_y << " << NOC_ADDR_NODE_ID_BITS) | " << noc_x << ") << NOC_COORD_REG_OFFSET),"
                << "\t// NOC_X=" << noc_x << " NOC_Y=" << noc_y << endl;
         }
@@ -610,8 +609,8 @@ std::string generate_bank_to_noc_coord_descriptor_string(
         ss << "    {"
            << "\t// noc=" << noc << endl;
         for (unsigned int bank_id = 0; bank_id < l1_bank_map.size(); bank_id++) {
-            uint16_t noc_x = NOC_0_X(noc, grid_size.x, l1_bank_map[bank_id].x);
-            uint16_t noc_y = NOC_0_Y(noc, grid_size.y, l1_bank_map[bank_id].y);
+            uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, grid_size.x, l1_bank_map[bank_id].x);
+            uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, grid_size.y, l1_bank_map[bank_id].y);
             ss << "        (((" << noc_y << " << NOC_ADDR_NODE_ID_BITS) | " << noc_x << ") << NOC_COORD_REG_OFFSET),"
                << "\t// NOC_X=" << noc_x << " NOC_Y=" << noc_y << endl;
         }

tt_metal/llrt/hal.hpp

@@ -139,6 +139,11 @@ class Hal {
 
     tt::ARCH get_arch() {return arch_;}
 
+    template <typename IndexType, typename SizeType, typename CoordType>
+    auto noc_coordinate(IndexType noc_index, SizeType noc_size, CoordType coord) const -> decltype(noc_size - 1 - coord) {
+        return noc_index == 0 ? coord : (noc_size - 1 - coord);
+    }
+
     uint32_t get_programmable_core_type_count() const;
     HalProgrammableCoreType get_programmable_core_type(uint32_t core_type_index) const;
     uint32_t get_programmable_core_type_index(HalProgrammableCoreType programmable_core_type_index) const;