#0: fix BH brisc fw logic

tenstorrent · Feb 28, 2025 · e6e65eb · e6e65eb
1 parent a4422ec
commit e6e65eb
Show file tree

Hide file tree

Showing 2 changed files with 20 additions and 27 deletions.
diff --git a/tests/tt_metal/tt_metal/noc/test_dynamic_noc.cpp b/tests/tt_metal/tt_metal/noc/test_dynamic_noc.cpp
@@ -83,7 +83,7 @@ void build_and_run_program(
         DataMovementConfig{
             .processor = DataMovementProcessor::RISCV_0,
             .noc = NOC::RISCV_0_default,
-            .noc_mode = tt_metal::NOC_MODE::DM_DEDICATED_NOC,
+            .noc_mode = mix_noc_mode ? tt_metal::NOC_MODE::DM_DEDICATED_NOC : tt_metal::NOC_MODE::DM_DYNAMIC_NOC,
             .compile_args = compile_args});
 
     auto ncrisc_kernel2 = CreateKernel(
@@ -94,7 +94,7 @@ void build_and_run_program(
         DataMovementConfig{
             .processor = DataMovementProcessor::RISCV_1,
             .noc = NOC::RISCV_1_default,
-            .noc_mode = tt_metal::NOC_MODE::DM_DEDICATED_NOC,
+            .noc_mode = mix_noc_mode ? tt_metal::NOC_MODE::DM_DEDICATED_NOC : tt_metal::NOC_MODE::DM_DYNAMIC_NOC,
             .compile_args = compile_args});
 
     for (int core_idx_y = 0; core_idx_y < worker_grid_size.y; core_idx_y++) {

diff --git a/tt_metal/hw/firmware/src/brisc.cc b/tt_metal/hw/firmware/src/brisc.cc
@@ -456,6 +456,10 @@ int main() {
                     31 /*wrap*/,
                     false /*linked*/,
                     post_atomic_increments /*posted*/);
+#if defined(ARCH_BLACKHOLE)
+                // flush for BH since this is non-posted, which could cause counter mismatch
+                while (!ncrisc_noc_nonposted_atomics_flushed(noc_index));
+#endif
             }
         }
 
@@ -590,31 +594,6 @@ int main() {
                 // message is no longer owned by us.
                 CLEAR_PREVIOUS_LAUNCH_MESSAGE_ENTRY_FOR_WATCHER();
 
-#if defined(ARCH_BLACKHOLE)
-                if (noc_mode == DM_DYNAMIC_NOC) {
-                    noc_fast_atomic_increment<DM_DYNAMIC_NOC>(
-                        noc_index,
-                        NCRISC_AT_CMD_BUF,
-                        dispatch_addr,
-                        NOC_UNICAST_WRITE_VC,
-                        1,
-                        31 /*wrap*/,
-                        false /*linked*/,
-                        post_atomic_increments /*posted*/);
-                    // barrier till the atomic response is back
-                    while (!ncrisc_dynamic_noc_nonposted_atomics_flushed(noc_index));
-                } else {
-                    noc_fast_atomic_increment(
-                        noc_index,
-                        NCRISC_AT_CMD_BUF,
-                        dispatch_addr,
-                        NOC_UNICAST_WRITE_VC,
-                        1,
-                        31 /*wrap*/,
-                        false /*linked*/,
-                        post_atomic_increments /*posted*/);
-                }
-#else
                 noc_fast_atomic_increment(
                     noc_index,
                     NCRISC_AT_CMD_BUF,
@@ -624,6 +603,20 @@ int main() {
                     31 /*wrap*/,
                     false /*linked*/,
                     post_atomic_increments /*posted*/);
+#if defined(ARCH_BLACKHOLE)
+                if (noc_mode == DM_DYNAMIC_NOC) {
+                    // inc dm noc counter for BH as this is non-posted
+                    inc_noc_counter_val<
+                        static_cast<std::underlying_type_t<TensixProcessorTypes>>(TensixProcessorTypes::DM0),
+                        NocBarrierType::NONPOSTED_ATOMICS_ACKED>(noc_index, 1);
+                    // barrier till the atomic response is back
+                    while (!ncrisc_dynamic_noc_nonposted_atomics_flushed(noc_index));
+                    // reset local counters
+                    noc_local_state_init(noc_index);
+                } else {
+                    // flush for BH since this is non-posted, which could cause counter mismatch in the next iter
+                    while (!ncrisc_noc_nonposted_atomics_flushed(noc_index));
+                }
 #endif
                 mailboxes->launch_msg_rd_ptr = (launch_msg_rd_ptr + 1) & (launch_msg_buffer_num_entries - 1);
             }