fix: add nvlink barrier to ensure the workspace is cleaned before next stage (#238)

Co-authored-by: Chenhao Xu <xch@deepseek.com>

Author: bucket-xv

deep_gemm/include/deep_gemm/impls/sm100_fp8_fp4_mega_moe.cuh

@@ -635,6 +635,14 @@ sm100_fp8_fp4_mega_moe_impl(void* y,
                 }
             }
         }
+
+        // Wait for all ranks to finish cleaning
+        comm::nvlink_barrier<kNumRanks, kNumSMs, kNumDispatchThreads, kDispatchGridSyncIndex>(
+            workspace, sym_buffer, sm_idx, thread_idx,
+            [=]() { cutlass::arch::NamedBarrier::sync(kNumDispatchThreads, kDispatchBarrierIdx); },
+            /* Before the NVLink barrier, there is a grid sync */ true,
+            /* One block is sufficient to block the next stage from launching */ false
+        );
     } else if (warp_idx == kNumDispatchWarps) {
         // Adjust registers
         cutlass::arch::warpgroup_reg_dealloc<kNumNonEpilogueRegisters>();