#18332: Remove unused args in BN

tenstorrent · Feb 26, 2025 · c1b2b8f · c1b2b8f
1 parent 39d0f66
commit c1b2b8f
Show file tree

Hide file tree

Showing 2 changed files with 6 additions and 9 deletions.
diff --git a/ttnn/cpp/ttnn/operations/normalization/batch_norm/device/batch_norm_program_factory.cpp b/ttnn/cpp/ttnn/operations/normalization/batch_norm/device/batch_norm_program_factory.cpp
@@ -54,8 +54,8 @@ void set_or_update_runtime_arguments(
         tt::tt_metal::split_work_to_cores(compute_with_storage_grid_size, num_output_tiles, row_major);
 
     auto cores = grid_to_cores(num_cores_total, num_cores_x, num_cores_y, row_major);
-    constexpr size_t num_reader_args = 11;
-    constexpr size_t num_writer_args = 14;
+    constexpr size_t num_reader_args = 9;
+    constexpr size_t num_writer_args = 12;
     constexpr size_t num_kernel_args = 3;
     for (uint32_t i = 0, start_tile_id = 0; i < num_cores_total; i++) {
         const auto& core = cores[i];
@@ -87,9 +87,7 @@ void set_or_update_runtime_arguments(
             aHt * aWt * aC * (aN > 1),
             aHt * aWt * (aC > 1),
             cN,
-            cC,
-            cHt,
-            cWt};
+            cC};
         handle_args(program, reader_kernel_id, core, reader_runtime_args);
 
         const auto weight_addr = weight_has_value ? weight_tensor->buffer()->address() : 0;
@@ -106,9 +104,7 @@ void set_or_update_runtime_arguments(
             bHt * bWt * bC * (bN > 1),
             bHt * bWt * (bC > 1),
             cN,
-            cC,
-            cHt,
-            cWt};
+            cC};
         handle_args(program, writer_kernel_id, core, writer_runtime_args);
 
         auto counter = start_tile_id % cHtWt;

diff --git a/...cpp/ttnn/operations/normalization/batch_norm/device/kernels/compute/batch_norm_kernel.cpp b/...cpp/ttnn/operations/normalization/batch_norm/device/kernels/compute/batch_norm_kernel.cpp
@@ -163,7 +163,8 @@ void MAIN {
     uint32_t complete_iterations = (num_tiles + tile_start) / tile_freq;
     uint32_t remaining_iterations = (num_tiles + tile_start) % tile_freq;
 
-    cb_wait_front(cb_eps, 1);
+    constexpr uint32_t onetile = 1;
+    cb_wait_front(cb_eps, onetile);
 
     for (uint32_t i = 0; i < complete_iterations; ++i, tile_start = 0) {
         batchnorm_bcast_tiles(