#9117: Remove hardcoding of l1_bank_size in core desc yamls to allow …

…l1 bank size to be wtorker l1 size - reserved region
tenstorrent · Jun 17, 2024 · 16ed7bc · 16ed7bc
1 parent e87f5db
commit 16ed7bc
Show file tree

Hide file tree

Showing 11 changed files with 26 additions and 68 deletions.
diff --git a/tests/tt_metal/tt_metal/unit_tests/allocator/test_l1_banking_allocator.cpp b/tests/tt_metal/tt_metal/unit_tests/allocator/test_l1_banking_allocator.cpp
@@ -18,7 +18,7 @@ TEST_F(BasicFixture, TestL1BuffersAllocatedTopDown) {
     size_t total_size_bytes = 0;
 
     const metal_SocDescriptor &soc_desc = tt::Cluster::instance().get_soc_desc(device->id());
-    const uint32_t interleaved_l1_bank_size = tt::get_storage_core_bank_size(device->id(), device->num_hw_cqs());
+    const uint32_t interleaved_l1_bank_size = tt::get_l1_bank_size(device->id(), device->num_hw_cqs());
     uint64_t alloc_limit = interleaved_l1_bank_size - STORAGE_ONLY_UNRESERVED_BASE;
 
     std::vector<std::unique_ptr<Buffer>> buffers;
@@ -45,7 +45,7 @@ TEST_F(BasicFixture, TestL1BuffersDoNotGrowBeyondBankSize) {
     tt::tt_metal::Device *device = tt::tt_metal::CreateDevice(0, 1, 0);
 
     const metal_SocDescriptor &soc_desc = tt::Cluster::instance().get_soc_desc(device->id());
-    const uint32_t interleaved_l1_bank_size = tt::get_storage_core_bank_size(device->id(), device->num_hw_cqs());
+    const uint32_t interleaved_l1_bank_size = tt::get_l1_bank_size(device->id(), device->num_hw_cqs());    
     uint64_t alloc_limit = interleaved_l1_bank_size - STORAGE_ONLY_UNRESERVED_BASE;
 
     tt::tt_metal::InterleavedBufferConfig l1_config{

diff --git a/tt_metal/common/core_descriptor.cpp b/tt_metal/common/core_descriptor.cpp
@@ -37,7 +37,6 @@ const core_descriptor_t &get_core_descriptor_config(chip_id_t device_id, const u
     YAML::Node desc_yaml = core_descriptor_yaml[product_name][std::to_string(num_hw_cqs)];
 
     // Parse the yaml into core_descriptor_t
-    uint32_t storage_core_bank_size = desc_yaml["l1_bank_size"].as<uint32_t>();
     std::vector<RelativeCoreCoord> storage_cores;
     for (const auto& core_node : desc_yaml["storage_cores"]) {
         RelativeCoreCoord coord = {};
@@ -49,6 +48,17 @@ const core_descriptor_t &get_core_descriptor_config(chip_id_t device_id, const u
         }
         storage_cores.push_back(coord);
     }
+    std::optional<uint32_t> storage_core_bank_size = std::nullopt;
+    if (not storage_cores.empty()) {
+        try {
+            storage_core_bank_size = desc_yaml["storage_core_bank_size"].as<uint32_t>();
+        } catch (std::runtime_error &ex) {
+            TT_THROW(
+                "Core descriptor yaml for {} needs to specify storage_core_bank_size since there are {} storage cores!",
+                get_string_lowercase(arch),
+                storage_cores.size());
+        }
+    }
 
     auto compute_with_storage_start = desc_yaml["compute_with_storage_grid_range"]["start"];
     auto compute_with_storage_end = desc_yaml["compute_with_storage_grid_range"]["end"];

diff --git a/tt_metal/common/core_descriptor.hpp b/tt_metal/common/core_descriptor.hpp
@@ -13,7 +13,7 @@ struct core_descriptor_t {
     CoreCoord compute_grid_size;
     std::vector<RelativeCoreCoord> relative_compute_cores;
     std::vector<RelativeCoreCoord> relative_storage_cores;
-    uint32_t storage_core_bank_size;
+    std::optional<uint32_t> storage_core_bank_size = std::nullopt;
     std::vector<RelativeCoreCoord> relative_dispatch_cores;
     CoreType dispatch_core_type;
 };
@@ -76,9 +76,14 @@ inline const std::string get_product_name(tt::ARCH arch, uint32_t num_harvested_
 
 const core_descriptor_t &get_core_descriptor_config(chip_id_t device_id, const uint8_t num_hw_cqs);
 
-inline uint32_t get_storage_core_bank_size(chip_id_t device_id, const uint8_t num_hw_cqs) {
+inline uint32_t get_l1_bank_size(chip_id_t device_id, const uint8_t num_hw_cqs) {
     const core_descriptor_t &core_desc = get_core_descriptor_config(device_id, num_hw_cqs);
-    return core_desc.storage_core_bank_size;
+    const metal_SocDescriptor &soc_desc = tt::Cluster::instance().get_soc_desc(device_id);
+    uint32_t l1_bank_size = core_desc.storage_core_bank_size.has_value()
+                                ? core_desc.storage_core_bank_size.value()
+                                : (soc_desc.worker_l1_size - L1_UNRESERVED_BASE);
+    TT_FATAL(l1_bank_size % L1_ALIGNMENT == 0, "L1 bank size must be {} B aligned", L1_ALIGNMENT);
+    return l1_bank_size;
 }
 
 inline const std::vector<CoreCoord> &get_logical_storage_cores(chip_id_t device_id, const uint8_t num_hw_cqs) {

diff --git a/tt_metal/core_descriptors/blackhole_140_arch.yaml b/tt_metal/core_descriptors/blackhole_140_arch.yaml
@@ -7,9 +7,6 @@
 
 blackhole:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [13, 8]
@@ -24,9 +21,6 @@ blackhole:
       "tensix"
 
   2:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [13, 8]

diff --git a/tt_metal/core_descriptors/blackhole_versim_1x1_arch.yaml b/tt_metal/core_descriptors/blackhole_versim_1x1_arch.yaml
@@ -7,9 +7,6 @@
 
 blackhole:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range:
       start: [0, 0]
       end: [0, 0]

diff --git a/tt_metal/core_descriptors/grayskull_120_arch.yaml b/tt_metal/core_descriptors/grayskull_120_arch.yaml
@@ -7,7 +7,7 @@
 
 E150:
   1:
-    l1_bank_size:
+    storage_core_bank_size:
       524288
 
     compute_with_storage_grid_range: # Logical only start and end [x, y]
@@ -20,7 +20,7 @@ E150:
     dispatch_cores:
       [[0, -1], [6, -1]]
   2:
-    l1_bank_size:
+    storage_core_bank_size:
       524288
 
     compute_with_storage_grid_range: # Logical only start and end [x, y]
@@ -35,7 +35,7 @@ E150:
 
 E75:
   1:
-    l1_bank_size:
+    storage_core_bank_size:
       1048576
 
     compute_with_storage_grid_range: # Logical only start and end [x, y]
@@ -48,7 +48,7 @@ E75:
     dispatch_cores:
       [[11, 0], [11, 4]]
   2:
-    l1_bank_size:
+    storage_core_bank_size:
       1048576
 
     compute_with_storage_grid_range: # Logical only start and end [x, y]

diff --git a/tt_metal/core_descriptors/grayskull_versim_1x1_arch.yaml b/tt_metal/core_descriptors/grayskull_versim_1x1_arch.yaml
@@ -7,9 +7,6 @@
 
 E150:
   1:
-    l1_bank_size:
-      524288
-
     compute_with_storage_grid_range:
       start: [0, 0]
       end: [0, 0]

diff --git a/tt_metal/core_descriptors/wormhole_b0_80_arch.yaml b/tt_metal/core_descriptors/wormhole_b0_80_arch.yaml
@@ -7,9 +7,6 @@
 
 galaxy:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -24,9 +21,6 @@ galaxy:
       "tensix"
 
   2:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -42,9 +36,6 @@ galaxy:
 
 nebula_x1:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -68,9 +59,6 @@ nebula_x1:
       "tensix"
 
   2:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -86,9 +74,6 @@ nebula_x1:
 
 nebula_x2:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 6]
@@ -103,9 +88,6 @@ nebula_x2:
       "tensix"
 
   2:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 6]

diff --git a/tt_metal/core_descriptors/wormhole_b0_80_arch_eth_dispatch.yaml b/tt_metal/core_descriptors/wormhole_b0_80_arch_eth_dispatch.yaml
@@ -7,9 +7,6 @@
 
 galaxy:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -24,9 +21,6 @@ galaxy:
       "ethernet"
 
   2:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -42,9 +36,6 @@ galaxy:
 
 nebula_x1:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -59,9 +50,6 @@ nebula_x1:
       "ethernet"
 
   2:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -77,9 +65,6 @@ nebula_x1:
 
 nebula_x2:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]
@@ -94,9 +79,6 @@ nebula_x2:
       "ethernet"
 
   2:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range: # Logical only start and end [x, y]
       start: [0, 0]
       end: [7, 7]

diff --git a/tt_metal/core_descriptors/wormhole_b0_versim_1x1_arch.yaml b/tt_metal/core_descriptors/wormhole_b0_versim_1x1_arch.yaml
@@ -7,9 +7,6 @@
 
 galaxy:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range:
       start: [0, 0]
       end: [0, 0]
@@ -22,9 +19,6 @@ galaxy:
 
 nebula_x1:
   1:
-    l1_bank_size:
-      1376256
-
     compute_with_storage_grid_range:
       start: [0, 0]
       end: [0, 0]
@@ -37,9 +31,6 @@ nebula_x1:
 
 nebula_x2:
   1:
-    l1_bank_size:
-      749568
-
     compute_with_storage_grid_range:
       start: [0, 0]
       end: [0, 0]

diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp
@@ -64,7 +64,7 @@ void Device::initialize_allocator(size_t l1_small_size, size_t trace_region_size
          .dram_bank_offsets = {},
          .worker_grid_size = this->logical_grid_size(),
          .worker_l1_size = static_cast<size_t>(soc_desc.worker_l1_size),
-         .l1_bank_size = static_cast<size_t>(get_storage_core_bank_size(this->id_, this->num_hw_cqs_)),
+         .l1_bank_size = static_cast<size_t>(get_l1_bank_size(this->id_, this->num_hw_cqs_)),
          .l1_small_size = l1_small_size,
          .trace_region_size = trace_region_size,
          .core_type_from_noc_coord_table = {},  // Populated later