diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h index 3140eec26a9..e4790b61218 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h +++ b/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 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); diff --git a/tests/tt_metal/tt_metal/unit_tests_common/watcher/test_noc_sanitize.cpp b/tests/tt_metal/tt_metal/unit_tests_common/watcher/test_noc_sanitize.cpp index ecb5bc1afcb..416ffece9bd 100644 --- a/tests/tt_metal/tt_metal/unit_tests_common/watcher/test_noc_sanitize.cpp +++ b/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) diff --git a/tt_metal/hostdevcommon/common_runtime_address_map.h b/tt_metal/hostdevcommon/common_runtime_address_map.h index 4d49751e22b..5a3b10267be 100644 --- a/tt_metal/hostdevcommon/common_runtime_address_map.h +++ b/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); diff --git a/tt_metal/impl/debug/watcher_device_reader.cpp b/tt_metal/impl/debug/watcher_device_reader.cpp index f3c7f8529be..5ea724d4825 100644 --- a/tt_metal/impl/debug/watcher_device_reader.cpp +++ b/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 { diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index fba7276a1c0..ecaf333924f 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -24,8 +24,6 @@ #include "tt_metal/tools/profiler/tt_metal_tracy.hpp" #include "llrt/hal.hpp" -#include "tt_metal/hostdevcommon/common_runtime_address_map.h" // NOC_0_X - namespace tt { namespace tt_metal { @@ -769,15 +767,15 @@ void Device::configure_kernel_variant( std::map 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) { @@ -3099,8 +3097,8 @@ std::vector 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) ); } @@ -3110,17 +3108,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) ); } } diff --git a/tt_metal/jit_build/genfiles.cpp b/tt_metal/jit_build/genfiles.cpp index de1b29d8257..d8b8194484f 100644 --- a/tt_metal/jit_build/genfiles.cpp +++ b/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" @@ -588,8 +587,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; } @@ -609,8 +608,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; } @@ -667,19 +666,4 @@ void jit_build_genfiles_bank_to_noc_coord_descriptor( file_stream_siec.close(); } -static string generate_noc_core_xy_range_define(const std::vector& cores) { - stringstream ss; - - string end_of_line = " \\\n ( \\"; - for (const auto& core : cores) { - ss << end_of_line << endl; - ss << " ((x) == NOC_0_X(noc_idx, noc_size_x, (uint32_t)" << core.x - << ") && (y) == NOC_0_Y(noc_idx, noc_size_y, (uint32_t)" << core.y << "))"; - end_of_line = " || \\"; - } - ss << ")" << endl; - - return ss.str(); -} - } // namespace tt::tt_metal diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp index 1ba7a104e84..d8a698ebb12 100644 --- a/tt_metal/llrt/hal.hpp +++ b/tt_metal/llrt/hal.hpp @@ -137,6 +137,11 @@ class Hal { tt::ARCH get_arch() {return arch_;} + template + 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;