#0: Cleanup 1: Reference virtual coords correctly and remove debug code

tt_metal/hw/inc/blackhole/noc/noc_parameters.h

@@ -5,6 +5,7 @@
 #ifndef _NOC_PARAMETERS_H_
 #define _NOC_PARAMETERS_H_
 
+// Coordinate Virtualization is not currently supported on BH (requires syseng support for updating FW).
 #define VIRTUAL_TENSIX_START_X 0
 #define VIRTUAL_TENSIX_START_Y 0
 #define COORDINATE_VIRTUALIZATION_ENABLED 0

tt_metal/hw/inc/dataflow_api.h

@@ -2002,6 +2002,7 @@ void noc_async_read_barrier_with_trid(uint32_t trid, uint8_t noc = noc_index) {
 template<bool DRAM>
 FORCE_INLINE
 uint64_t get_noc_addr_from_bank_id(uint32_t bank_id, uint32_t bank_address_offset, uint8_t noc = noc_index) {
+    // Use addrgen tables to convert bank_ids to physical NOC coordinates
     uint64_t noc_addr = 0;
     if constexpr (DRAM) {
         noc_addr = dram_bank_to_noc_xy[noc_index][bank_id];

tt_metal/hw/inc/debug/sanitize_noc.h

@@ -16,7 +16,7 @@
 
 // NOC logging enabled independently of watcher, need to include it here because it hooks into DEBUG_SANITIZE_NOC_*
 #include "noc_logging.h"
-#include "dprint.h"
+
 #if (                                                                                          \
     defined(COMPILE_FOR_BRISC) || defined(COMPILE_FOR_NCRISC) || defined(COMPILE_FOR_ERISC) || \
     defined(COMPILE_FOR_IDLE_ERISC)) &&                                                        \

tt_metal/hw/inc/grayskull/noc/noc_parameters.h

@@ -13,6 +13,7 @@
 #define NOC_Y_SIZE 1
 #endif
 
+// Coordinate Virtualization is not supported on GS (feature does not exist in NOC Hardware).
 #define VIRTUAL_TENSIX_START_X 0
 #define VIRTUAL_TENSIX_START_Y 0
 #define COORDINATE_VIRTUALIZATION_ENABLED 0

tt_metal/hw/inc/risc_common.h

@@ -21,6 +21,7 @@
 #define NOC_Y(y) NOC_0_Y(noc_index, noc_size_y, (y))
 #define DYNAMIC_NOC_X(noc, x) NOC_0_X(noc, noc_size_x, (x))
 #define DYNAMIC_NOC_Y(noc, y) NOC_0_Y(noc, noc_size_y, (y))
+
 #define TILE_WORD_2_BIT ((256 + 64 + 32) >> 4)
 #define TILE_WORD_4_BIT ((512 + 64 + 32) >> 4)
 #define TILE_WORD_8_BIT ((32 * 32 * 1 + 64 + 32) >> 4)

tt_metal/hw/inc/wormhole/noc/noc_parameters.h

@@ -13,6 +13,8 @@
 #define NOC_Y_SIZE 1
 #endif
 
+// Coordinate Virtualization is fully supported by WH NOC Hardware and Firmware.
+// Tensix cores start at coorddinate <x = 18, y = 18> in Virtual Space.
 #define VIRTUAL_TENSIX_START_X 18
 #define VIRTUAL_TENSIX_START_Y 18
 #define COORDINATE_VIRTUALIZATION_ENABLED 1

tt_metal/impl/buffers/buffer.cpp

@@ -436,8 +436,15 @@ DeviceAddr Buffer::page_address(uint32_t bank_id, uint32_t page_index) const {
 DeviceAddr Buffer::bank_local_page_address(uint32_t bank_id, uint32_t page_index) const {
     uint32_t num_banks = allocator::num_banks(*this->allocator_, this->buffer_type_);
     TT_FATAL(bank_id < num_banks, "Invalid Bank ID: {} exceeds total numbers of banks ({})!", bank_id, num_banks);
-    int pages_offset_within_bank = (int)page_index / num_banks;
-    auto offset = (round_up(this->page_size(), this->alignment()) * pages_offset_within_bank);
+    uint32_t offset;
+    if (is_sharded(this->buffer_layout())) {
+        auto shard_spec = this->shard_spec();
+        uint32_t pages_offset_within_bank = page_index % shard_spec.size();
+        offset = (round_up(this->page_size(), this->alignment()) * pages_offset_within_bank);
+    } else {
+        uint32_t pages_offset_within_bank = page_index / num_banks;
+        offset = (round_up(this->page_size(), this->alignment()) * pages_offset_within_bank);
+    }
     return this->address() + offset;
 }
 
@@ -457,14 +464,6 @@ DeviceAddr Buffer::aligned_size_per_bank() const {
     return tt::tt_metal::detail::SizeBytesPerBank(this->size_, this->page_size_, num_banks, this->alignment());
 }
 
-DeviceAddr Buffer::sharded_bank_local_page_address(uint32_t bank_id, uint32_t page_index) const {
-    TT_FATAL(is_sharded(this->buffer_layout()), "Buffer not sharded");
-    auto shard_spec = this->shard_spec();
-    uint32_t pages_offset_within_bank = page_index % shard_spec.size();
-    auto offset = (round_up(this->page_size(), this->alignment()) * pages_offset_within_bank);
-    return this->address() + offset;
-}
-
 DeviceAddr Buffer::sharded_page_address(uint32_t bank_id, uint32_t page_index) const {
     TT_FATAL(is_sharded(this->buffer_layout()), "Buffer not sharded");
     auto shard_spec = this->shard_spec();

tt_metal/impl/buffers/buffer.hpp

@@ -224,7 +224,6 @@ class Buffer final {
     DeviceAddr page_address(uint32_t bank_id, uint32_t page_index) const;
 
     DeviceAddr bank_local_page_address(uint32_t bank_id, uint32_t page_index) const;
-    DeviceAddr sharded_bank_local_page_address(uint32_t bank_id, uint32_t page_index) const;
     uint32_t alignment() const;
     DeviceAddr aligned_page_size() const;
     DeviceAddr aligned_size() const;

tt_metal/impl/debug/dprint_server.cpp

@@ -701,7 +701,6 @@ void DebugPrintServerContext::DetachDevice(Device* device) {
         outstanding_prints = false;
         for (auto& logical_core : device_to_core_range_.at(device)) {
             CoreCoord phys_core = device->virtual_core_from_logical_core(logical_core.coord, logical_core.type);
-            ;
             for (int risc_id = 0; risc_id < GetNumRiscs(logical_core); risc_id++) {
                 if (risc_mask & (1 << risc_id)) {
                     // No need to check if risc is not dprint-enabled.
@@ -764,7 +763,6 @@ void DebugPrintServerContext::DetachDevice(Device* device) {
     CoreDescriptorSet all_cores = GetAllCores(device);
     for (auto& logical_core : all_cores) {
         CoreCoord phys_core = device->virtual_core_from_logical_core(logical_core.coord, logical_core.type);
-        ;
         for (int hart_index = 0; hart_index < GetNumRiscs(logical_core); hart_index++) {
             WriteInitMagic(device, phys_core, hart_index, false);
         }
@@ -794,7 +792,6 @@ bool DebugPrintServerContext::PeekOneHartNonBlocking(
     Device* device, const CoreDescriptor& logical_core, int hart_id, bool new_data_this_iter) {
     // If init magic isn't cleared for this risc, then dprint isn't enabled on it, don't read it.
     CoreCoord phys_core = device->virtual_core_from_logical_core(logical_core.coord, logical_core.type);
-    ;
     if (!CheckInitMagicCleared(device, phys_core, hart_id)) {
         return false;
     }