From 23ce00c668d7bff14c98657d1245f58a65fe0cef Mon Sep 17 00:00:00 2001 From: asaigal Date: Tue, 10 Dec 2024 22:45:10 +0000 Subject: [PATCH] #0: Resolve BH failure after Virtual Coord Changes --- .../debug_tools/watcher/test_noc_sanitize.cpp | 2 +- .../perf_microbenchmark/dispatch/common.h | 6 +-- tt_metal/impl/device/device.cpp | 43 +++++++++++++------ tt_metal/impl/device/device.hpp | 5 ++- 4 files changed, 38 insertions(+), 18 deletions(-) diff --git a/tests/tt_metal/tt_metal/debug_tools/watcher/test_noc_sanitize.cpp b/tests/tt_metal/tt_metal/debug_tools/watcher/test_noc_sanitize.cpp index 0ac4f6ce267..dc3624789cf 100644 --- a/tests/tt_metal/tt_metal/debug_tools/watcher/test_noc_sanitize.cpp +++ b/tests/tt_metal/tt_metal/debug_tools/watcher/test_noc_sanitize.cpp @@ -130,7 +130,7 @@ void RunTestOnCore(WatcherFixture* fixture, Device* device, CoreCoord &core, boo // We should be able to find the expected watcher error in the log as well. string expected; int noc = (use_ncrisc) ? 1 : 0; - CoreCoord target_core = device->virtual_noc_coordinate(noc, input_buf_noc_xy); + CoreCoord target_core = device->virtual_noc0_coordinate(noc, input_buf_noc_xy); string risc_name = (is_eth_core) ? "erisc" : "brisc"; if (use_ncrisc) { risc_name = "ncrisc"; 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 08661f7d616..00868b341a8 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h @@ -534,9 +534,9 @@ void configure_kernel_variant( NOC my_noc_index, NOC upstream_noc_index, NOC downstream_noc_index) { - auto my_virtual_noc_coords = device->virtual_noc_coordinate(my_noc_index, phys_my_core); - auto upstream_virtual_noc_coords = device->virtual_noc_coordinate(upstream_noc_index, phys_upstream_core); - auto downstream_virtual_noc_coords = device->virtual_noc_coordinate(downstream_noc_index, phys_downstream_core); + auto my_virtual_noc_coords = device->virtual_noc0_coordinate(my_noc_index, phys_my_core); + auto upstream_virtual_noc_coords = device->virtual_noc0_coordinate(upstream_noc_index, phys_upstream_core); + auto downstream_virtual_noc_coords = device->virtual_noc0_coordinate(downstream_noc_index, phys_downstream_core); std::map defines = { {"DISPATCH_KERNEL", "1"}, diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index e2aaf7fec52..5c2875b4885 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -837,10 +837,10 @@ void Device::configure_kernel_variant( is_active_eth_core ? hal.get_programmable_core_type_index(HalProgrammableCoreType::ACTIVE_ETH) : hal.get_programmable_core_type_index(HalProgrammableCoreType::IDLE_ETH); - auto my_virtual_noc_coords = this->virtual_noc_coordinate(my_noc_index, kernel_virtual_core); - auto upstream_virtual_noc_coords = this->virtual_noc_coordinate(upstream_noc_index, upstream_virtual_core); - auto downstream_virtual_noc_coords = this->virtual_noc_coordinate(downstream_noc_index, downstream_virtual_core); - auto downstream_slave_virtual_noc_coords = this->virtual_noc_coordinate(downstream_noc_index, downstream_slave_virtual_core); + auto my_virtual_noc_coords = this->virtual_noc0_coordinate(my_noc_index, kernel_virtual_core); + auto upstream_virtual_noc_coords = this->virtual_noc0_coordinate(upstream_noc_index, upstream_virtual_core); + auto downstream_virtual_noc_coords = this->virtual_noc0_coordinate(downstream_noc_index, downstream_virtual_core); + auto downstream_slave_virtual_noc_coords = this->virtual_noc0_coordinate(downstream_noc_index, downstream_slave_virtual_core); std::map defines = { {"DISPATCH_KERNEL", "1"}, @@ -3142,18 +3142,37 @@ CoreType Device::core_type_from_virtual_core(const CoreCoord &virtual_coord) con } -CoreCoord Device::virtual_noc_coordinate(uint8_t noc_index, CoreCoord coord) const { +CoreCoord Device::virtual_noc0_coordinate(uint8_t noc_index, CoreCoord coord) const { if (coord.x >= this->grid_size().x || coord.y >= this->grid_size().y) { // Coordinate already in virtual space: NOC0 and NOC1 are the same return coord; } else { const auto& grid_size = this->grid_size(); - // Coordinate in Physical Space. Convert to Virtual. - CoreCoord phys_coord = { + // Coordinate in Physical NOC0 Space. Convert to Virtual. + coord = this->virtual_core_from_physical_core(coord, this->core_type_from_physical_core(coord)); + // Derive virtual coord in noc_index space. + CoreCoord virtual_coord = { + hal.noc_coordinate(noc_index, grid_size.x, coord.x), + hal.noc_coordinate(noc_index, grid_size.y, coord.y) + }; + return virtual_coord; + } +} + +CoreCoord Device::virtual_noc_coordinate(uint8_t noc_index, CoreCoord coord) const { + if (coord.x >= this->grid_size().x || coord.y >= this->grid_size().y) { + // Coordinate already in virtual space: NOC0 and NOC1 are the same + return coord; + } else { + const auto& grid_size = this->grid_size(); + // Coordinate passed in can be NOC0 or NOC1. The noc_index corresponds to + // the system this coordinate belongs to. + // Use this to convert to NOC0 coordinates and then derive Virtual Coords from it. + CoreCoord physical_coord = { hal.noc_coordinate(noc_index, grid_size.x, coord.x), hal.noc_coordinate(noc_index, grid_size.y, coord.y) }; - return this->virtual_core_from_physical_core(phys_coord, this->core_type_from_physical_core(phys_coord)); + return this->virtual_core_from_physical_core(physical_coord, this->core_type_from_physical_core(physical_coord)); } } @@ -3198,7 +3217,7 @@ CoreCoord Device::logical_core_from_ethernet_core(const CoreCoord ðernet_core } uint32_t Device::get_noc_unicast_encoding(uint8_t noc_index, const CoreCoord& core) const { - auto virtual_noc_coord = this->virtual_noc_coordinate(noc_index, core); + auto virtual_noc_coord = this->virtual_noc0_coordinate(noc_index, core); return tt::tt_metal::hal.noc_xy_encoding( virtual_noc_coord.x, virtual_noc_coord.y @@ -3206,8 +3225,8 @@ uint32_t Device::get_noc_unicast_encoding(uint8_t noc_index, const CoreCoord& co } uint32_t Device::get_noc_multicast_encoding(uint8_t noc_index, const CoreRange& cores) const { - auto virtual_noc_start = this->virtual_noc_coordinate(noc_index, cores.start_coord); - auto virtual_noc_end = this->virtual_noc_coordinate(noc_index, cores.end_coord); + auto virtual_noc_start = this->virtual_noc0_coordinate(noc_index, cores.start_coord); + auto virtual_noc_end = this->virtual_noc0_coordinate(noc_index, cores.end_coord); // NOC 1 mcasts from bottom left to top right, so we need to reverse the coords if (noc_index == 0) { @@ -3649,7 +3668,7 @@ void Device::generate_device_bank_to_noc_tables() l1_bank_to_noc_xy_.reserve(tt::tt_metal::hal.get_num_nocs() * l1_noc_coord_per_bank.size()); for (unsigned int noc = 0; noc < tt::tt_metal::hal.get_num_nocs(); noc++) { for (unsigned int bank_id = 0; bank_id < l1_noc_coord_per_bank.size(); bank_id++) { - auto l1_noc_coords = this->virtual_noc_coordinate(noc, l1_noc_coord_per_bank[bank_id]); + auto l1_noc_coords = this->virtual_noc0_coordinate(noc, l1_noc_coord_per_bank[bank_id]); uint16_t noc_x = l1_noc_coords.x; uint16_t noc_y = l1_noc_coords.y; uint16_t xy = ((noc_y << NOC_ADDR_NODE_ID_BITS) | noc_x) << NOC_COORD_REG_OFFSET; diff --git a/tt_metal/impl/device/device.hpp b/tt_metal/impl/device/device.hpp index f2ef56a3a83..a8cdb1f23b0 100644 --- a/tt_metal/impl/device/device.hpp +++ b/tt_metal/impl/device/device.hpp @@ -116,9 +116,10 @@ class Device { CoreCoord dram_grid_size() const; CoreType core_type_from_virtual_core(const CoreCoord& virtual_coord) const; - + // Given a Virtual coordinate in noc_index space, get the equivalent coordinate in Virtual NOC0 space CoreCoord virtual_noc_coordinate(uint8_t noc_index, CoreCoord coord) const; - + // Given a coordinate in Virtual NOC0 Space, get the equivalent coordinate in Virtual noc_index space + CoreCoord virtual_noc0_coordinate(uint8_t noc_index, CoreCoord coord) const; std::vector worker_cores_from_logical_cores(const std::vector &logical_cores) const; std::vector ethernet_cores_from_logical_cores(const std::vector &logical_cores) const; std::vector get_optimal_dram_bank_to_logical_worker_assignment();