Fix simulator setup

- HACK for passing BH pcie coordinates during simulation mode - Correct compute grid size for BH simulator
tenstorrent · Dec 3, 2024 · 2c31478 · 2c31478
1 parent 6a524ab
commit 2c31478
Show file tree

Hide file tree

Showing 3 changed files with 15 additions and 9 deletions.
diff --git a/tt_metal/common/core_descriptor.cpp b/tt_metal/common/core_descriptor.cpp
@@ -121,7 +121,7 @@ const core_descriptor_t& get_core_descriptor_config(
         }
         dispatch_cores.push_back(coord);
     }
-    TT_ASSERT(dispatch_cores.size(), "Dispatch cores size must be positive");
+    TT_ASSERT(dispatch_cores.size() || std::getenv("TT_METAL_SIMULATOR_EN"), "Dispatch cores size must be positive");
 
     config_by_num_cqs[num_hw_cqs] = core_descriptor_t{
         .compute_grid_size = compute_grid_size,

diff --git a/tt_metal/core_descriptors/blackhole_simulation_1x2_arch.yaml b/tt_metal/core_descriptors/blackhole_simulation_1x2_arch.yaml
@@ -9,8 +9,8 @@ blackhole:
   col:
     1:
       compute_with_storage_grid_range: # Logical only start and end [x, y]
-        start: [0, 1]
-        end: [1, 1]
+        start: [0, 0]
+        end: [1, 0]
 
       storage_cores: # Relative to grid of tensix cores
         []
@@ -22,8 +22,8 @@ blackhole:
         "tensix"
     2:
       compute_with_storage_grid_range: # Logical only start and end [x, y]
-        start: [0, 1]
-        end: [1, 1]
+        start: [0, 0]
+        end: [1, 0]
 
       storage_cores: # Relative to grid of tensix cores
         []

diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp
@@ -321,10 +321,16 @@ void Device::initialize_device_kernel_defines()
     const metal_SocDescriptor& soc_d = tt::Cluster::instance().get_soc_desc(this->id());
     auto pcie_cores = soc_d.get_pcie_cores();
     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(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)));
+
+    // Workaround for Simulator integration as they use a 2x2 grid which would underflow PCIE_NOC1*
+    CoreCoord pcie_core = pcie_cores.empty() ? grid_size : pcie_cores[0];
+    auto pcie_noc1_x = pcie_cores.empty() ? 14 : tt::tt_metal::hal.noc_coordinate(NOC::NOC_1, grid_size.x, pcie_cores[0].x);
+    auto pcie_noc1_y = pcie_cores.empty() ? 11 : tt::tt_metal::hal.noc_coordinate(NOC::NOC_1, grid_size.x, pcie_cores[0].y);
+
+    this->device_kernel_defines_.emplace("PCIE_NOC_X", std::to_string(pcie_core.x));
+    this->device_kernel_defines_.emplace("PCIE_NOC_Y", std::to_string(pcie_core.y));
+    this->device_kernel_defines_.emplace("PCIE_NOC1_X", std::to_string(pcie_noc1_x));
+    this->device_kernel_defines_.emplace("PCIE_NOC1_Y", std::to_string(pcie_noc1_x));
 }
 
 void Device::initialize_build() {