leavelet/cutlass
1
0
Fork 0
Code Issues Pull Requests Actions 30 Packages Projects Releases Wiki Activity

Remove duplicated cutlass::arch::wait_on_dependent_grids();

Browse Source
This commit is contained in:
HydraQYH
2025-12-06 10:20:35 +08:00
parent 929e1e0259
commit ff7f2dcdfb
2 changed files with 0 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp
-4
View File
@@ -608,10 +608,6 @@ public:
// Get pipeline stage increments from tensor shapes // Get pipeline stage increments from tensor shapes
auto k_tile_count = size<3>(gA_mkl); 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) { if (warp_group_role == WarpGroupRole::Producer) {
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>(); cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();
include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp
-4
View File
@@ -646,10 +646,6 @@ public:
// Get pipeline stage increments from tensor shapes // Get pipeline stage increments from tensor shapes
auto k_tile_count = size<3>(gA_mkl); 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) { if (warp_group_role == WarpGroupRole::Producer) {
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>(); cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();