From ff7f2dcdfb0415c791d44bbfeab94d2e9d30e023 Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Fri, 21 Nov 2025 11:10:09 +0800 Subject: [PATCH] Remove duplicated cutlass::arch::wait_on_dependent_grids(); --- .../sm90_gemm_array_tma_warpspecialized_cooperative.hpp | 4 ---- .../kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp | 4 ---- 2 files changed, 8 deletions(-) diff --git a/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp b/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp index 25a7e2f7..e758059e 100644 --- a/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp +++ b/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp @@ -608,10 +608,6 @@ public: // Get pipeline stage increments from tensor shapes auto k_tile_count = size<3>(gA_mkl); - // Ensure that the kernel does not touch - // unflushed global memory prior to this instruction - cutlass::arch::wait_on_dependent_grids(); - if (warp_group_role == WarpGroupRole::Producer) { cutlass::arch::warpgroup_reg_dealloc(); diff --git a/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp b/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp index 2c55987a..7120db78 100644 --- a/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp +++ b/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp @@ -646,10 +646,6 @@ public: // Get pipeline stage increments from tensor shapes auto k_tile_count = size<3>(gA_mkl); - // Ensure that the kernel does not touch - // unflushed global memory prior to this instruction - cutlass::arch::wait_on_dependent_grids(); - if (warp_group_role == WarpGroupRole::Producer) { cutlass::arch::warpgroup_reg_dealloc();