add the sync barrier for persistent kernel (#2977)

Thomas Ning · web-flow · commit ae9f29b7d514 · 2025-10-07T11:54:04.000-07:00
diff --git a/include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp
@@ -1134,6 +1134,7 @@ struct UniversalGemmKernel
 
         while(block_id < num_work)
         {
+            s_waitcnt_barrier();
             // Get the tile index for this block
             const auto tile_idx = amd_wave_read_first_lane(block_id % num_tiles);
             const auto [iM, iN] = TilePartitioner{kargs.M, kargs.N}.GetOutputTileIndex(tile_idx);

Original file line number	Diff line number	Diff line change
`@@ -1134,6 +1134,7 @@ struct UniversalGemmKernel`
`1134`	`1134`
`1135`	`1135`	`while(block_id < num_work)`
`1136`	`1136`	`{`
	`1137`	`+ s_waitcnt_barrier();`
`1137`	`1138`	`// Get the tile index for this block`
`1138`	`1139`	`const auto tile_idx = amd_wave_read_first_lane(block_id % num_tiles);`
`1139`	`1140`	`const auto [iM, iN] = TilePartitioner{kargs.M, kargs.N}.GetOutputTileIndex(tile_idx);`