Skip to content

Commit a3a8e0a

Browse files
committed
Remove duplicated cutlass::arch::wait_on_dependent_grids();
1 parent 79d03b1 commit a3a8e0a

File tree

2 files changed

+0
-8
lines changed

2 files changed

+0
-8
lines changed

include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -608,10 +608,6 @@ class GemmUniversal<
608608
// Get pipeline stage increments from tensor shapes
609609
auto k_tile_count = size<3>(gA_mkl);
610610
611-
// Ensure that the kernel does not touch
612-
// unflushed global memory prior to this instruction
613-
cutlass::arch::wait_on_dependent_grids();
614-
615611
if (warp_group_role == WarpGroupRole::Producer) {
616612
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();
617613

include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -646,10 +646,6 @@ class GemmUniversal<
646646
// Get pipeline stage increments from tensor shapes
647647
auto k_tile_count = size<3>(gA_mkl);
648648
649-
// Ensure that the kernel does not touch
650-
// unflushed global memory prior to this instruction
651-
cutlass::arch::wait_on_dependent_grids();
652-
653649
if (warp_group_role == WarpGroupRole::Producer) {
654650
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();
655651

0 commit comments

Comments
 (0)