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 4bbf0e94f73..5c44063b71b 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 @@ -47,12 +47,12 @@ void RunTestOnCore(WatcherFixture* fixture, IDevice* device, CoreCoord &core, bo uint32_t single_tile_size = 2 * 1024; uint32_t num_tiles = 50; uint32_t l1_buffer_size = single_tile_size * num_tiles; - uint32_t l1_buffer_addr = 400 * 1024; + uint32_t l1_buffer_addr = hal.get_dev_addr(HalProgrammableCoreType::TENSIX, HalL1MemAddrType::UNRESERVED); // For ethernet core, need to have smaller buffer at a different address if (is_eth_core) { l1_buffer_size = 1024; - l1_buffer_addr = 200 * 1024; + l1_buffer_addr = hal.get_dev_addr(HalProgrammableCoreType::ACTIVE_ETH, HalL1MemAddrType::UNRESERVED); } tt_metal::InterleavedBufferConfig l1_config{ @@ -103,8 +103,8 @@ void RunTestOnCore(WatcherFixture* fixture, IDevice* device, CoreCoord &core, bo // depending on the flags passed in. switch(feature) { case SanitizeAddress: - output_buf_noc_xy.x = 16; - output_buf_noc_xy.y = 16; + output_buf_noc_xy.x = 26; + output_buf_noc_xy.y = 18; break; case SanitizeAlignmentL1Write: output_l1_buffer_addr++; // This is illegal because reading DRAM->L1 needs DRAM alignment diff --git a/tt_metal/api/tt-metalium/dev_msgs.h b/tt_metal/api/tt-metalium/dev_msgs.h index 7a84a260776..4fde76aff8a 100644 --- a/tt_metal/api/tt-metalium/dev_msgs.h +++ b/tt_metal/api/tt-metalium/dev_msgs.h @@ -324,7 +324,9 @@ struct core_info_msg_t { volatile uint8_t virtual_harvested_y[MAX_HARVESTED_ROWS]; volatile uint8_t noc_size_x; volatile uint8_t noc_size_y; - volatile uint8_t pad[27]; + volatile uint8_t worker_grid_size_x; + volatile uint8_t worker_grid_size_y; + volatile uint8_t pad[25]; }; constexpr uint32_t launch_msg_buffer_num_entries = 4; diff --git a/tt_metal/hw/inc/debug/sanitize_noc.h b/tt_metal/hw/inc/debug/sanitize_noc.h index 85a0997ae96..1d3b9aa202f 100644 --- a/tt_metal/hw/inc/debug/sanitize_noc.h +++ b/tt_metal/hw/inc/debug/sanitize_noc.h @@ -111,11 +111,18 @@ AddressableCoreType get_core_type(uint8_t noc_id, uint8_t x, uint8_t y, bool& is } } if constexpr (COORDINATE_VIRTUALIZATION_ENABLED) { - // Check if NOC endpoint is valid in the Tensix Virtual Coordinate Space. + // Check if NOC endpoint is valid in the Tensix Virtual Coordinate Space. Use worker grid size instead of noc + // size because virtual coords are continuous. if (x >= NOC_0_X(noc_id, core_info->noc_size_x, (uint32_t)VIRTUAL_TENSIX_START_X) && - x <= NOC_0_X(noc_id, core_info->noc_size_x, (uint32_t)VIRTUAL_TENSIX_START_X + core_info->noc_size_x - 1) && + x <= NOC_0_X( + noc_id, + core_info->noc_size_x, + (uint32_t)VIRTUAL_TENSIX_START_X + core_info->worker_grid_size_x - 1) && y >= NOC_0_Y(noc_id, core_info->noc_size_y, (uint32_t)VIRTUAL_TENSIX_START_Y) && - y <= NOC_0_Y(noc_id, core_info->noc_size_y, (uint32_t)VIRTUAL_TENSIX_START_Y + core_info->noc_size_y - 1)) { + y <= NOC_0_Y( + noc_id, + core_info->noc_size_y, + (uint32_t)VIRTUAL_TENSIX_START_Y + core_info->worker_grid_size_y - 1)) { is_virtual_coord = true; return AddressableCoreType::TENSIX; } diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index 3de32389ce5..e360bf508ef 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -749,6 +749,8 @@ void Device::initialize_and_launch_firmware() { core_info->noc_size_x = soc_d.grid_size.x; core_info->noc_size_y = soc_d.grid_size.y; + core_info->worker_grid_size_x = this->logical_grid_size().x; // Grid size as virtual coords see it (workers only) + core_info->worker_grid_size_y = this->logical_grid_size().y; // Download to worker cores log_debug("Initializing firmware");