From 91de7891a5997dfcc602125c7963e765b617c7ac Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Thu, 23 Oct 2025 21:50:57 +0800 Subject: [PATCH 1/9] Support PDL in sm90_gemm_array_tma_warpspecialized_pingpong.hpp --- ...m90_gemm_array_tma_warpspecialized_pingpong.hpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) 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 d380b2cd..9bd4e300 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 @@ -665,6 +665,9 @@ public: } // Mainloop Producer Warp else if (producer_warp_role == ProducerWarpRole::Mainloop) { + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif int32_t curr_batch = idx2crd(work_tile_info.L_idx, shape<4>(gB_nkl)); // Usually just returns work_tile_info.L_idx; int32_t const mock_l_coord = 0; int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); @@ -766,6 +769,9 @@ public: collective_mainloop.load_tail(mainloop_pipeline, mainloop_pipe_producer_state); } // Mainloop Producer Warp End else if (producer_warp_role == ProducerWarpRole::MainloopAux) { + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif if constexpr (IsMainloopAuxiliaryLoadNeeded) { int32_t curr_batch = idx2crd(work_tile_info.L_idx, shape<4>(gB_nkl)); // Usually just returns work_tile_info.L_idx; int32_t const mock_l_coord = 0; @@ -829,6 +835,9 @@ public: } // Mainloop Auxiliary Load Producer Warp End // Epilogue Producer Warp else if (producer_warp_role == ProducerWarpRole::Epilogue && collective_epilogue.is_producer_load_needed()) { + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); int32_t const sm_count = params.hw_info.sm_count; @@ -1038,6 +1047,11 @@ public: // Get next work tile auto [next_work_tile_info, increment_pipe] = scheduler.fetch_next_work(work_tile_info, tile_scheduler_pipeline, tile_scheduler_pipe_consumer_state); + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + if (!next_work_tile_info.is_valid()) { + cutlass::arch::launch_dependent_grids(); + } + #endif work_tile_info = next_work_tile_info; if (increment_pipe) { ++tile_scheduler_pipe_consumer_state; From acf5990cc236036137087b84cbe4ccd904b7ad83 Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Fri, 24 Oct 2025 13:13:14 +0800 Subject: [PATCH 2/9] Refine position for wait_on_dependent_grids. --- .../kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) 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 9bd4e300..559b9afe 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 @@ -769,10 +769,10 @@ public: collective_mainloop.load_tail(mainloop_pipeline, mainloop_pipe_producer_state); } // Mainloop Producer Warp End else if (producer_warp_role == ProducerWarpRole::MainloopAux) { - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 - cutlass::arch::wait_on_dependent_grids(); - #endif if constexpr (IsMainloopAuxiliaryLoadNeeded) { + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif int32_t curr_batch = idx2crd(work_tile_info.L_idx, shape<4>(gB_nkl)); // Usually just returns work_tile_info.L_idx; int32_t const mock_l_coord = 0; From 1e5f95cbbeb786ff968204493c6f293c4703499f Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Fri, 24 Oct 2025 16:12:14 +0800 Subject: [PATCH 3/9] Support PDL in sm90_gemm_array_tma_warpspecialized_cooperative --- ..._gemm_array_tma_warpspecialized_cooperative.hpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) 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 c69defbf..6e3ce1e0 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 @@ -614,6 +614,9 @@ public: if (producer_warp_role == ProducerWarpRole::Scheduler) { // GroupScheduler requires a producer warp to iterate over the group infos and push // the work tile infos to the downstream pipelines. + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif if constexpr (cute::is_same_v) { do { auto [next_work_tile_info, increment_pipe] = scheduler.advance_to_next_work(tile_scheduler_pipeline, tile_scheduler_pipe_producer_state); @@ -627,6 +630,9 @@ public: } // Mainloop Producer Warp else if (producer_warp_role == ProducerWarpRole::Mainloop) { + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif int32_t curr_batch = idx2crd(work_tile_info.L_idx, shape<4>(gB_nkl)); // Usually just returns work_tile_info.L_idx; int32_t const mock_l_coord = 0; int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); @@ -791,6 +797,9 @@ public: } // Mainloop Auxiliary Load Producer Warp End // Epilogue Producer Warp else if (producer_warp_role == ProducerWarpRole::Epilogue && collective_epilogue.is_producer_load_needed()) { + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); int32_t const sm_count = params.hw_info.sm_count; @@ -994,6 +1003,11 @@ public: // Get next work tile auto [next_work_tile_info, increment_pipe] = scheduler.fetch_next_work(work_tile_info, tile_scheduler_pipeline, tile_scheduler_pipe_consumer_state); + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + if (!next_work_tile_info.is_valid()) { + cutlass::arch::launch_dependent_grids(); + } + #endif work_tile_info = next_work_tile_info; if (increment_pipe) { ++tile_scheduler_pipe_consumer_state; From e1b2ec57e3f1b5c34367c1de27faa0217dc39fcd Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Mon, 17 Nov 2025 09:27:58 +0800 Subject: [PATCH 4/9] Hoist waits above the warp specialized region. --- ...90_gemm_array_tma_warpspecialized_cooperative.hpp | 12 +++--------- .../sm90_gemm_array_tma_warpspecialized_pingpong.hpp | 12 +++--------- 2 files changed, 6 insertions(+), 18 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 6e3ce1e0..9e326d2b 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,15 +608,15 @@ public: // Get pipeline stage increments from tensor shapes auto k_tile_count = size<3>(gA_mkl); + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif if (warp_group_role == WarpGroupRole::Producer) { cutlass::arch::warpgroup_reg_dealloc(); if (producer_warp_role == ProducerWarpRole::Scheduler) { // GroupScheduler requires a producer warp to iterate over the group infos and push // the work tile infos to the downstream pipelines. - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 - cutlass::arch::wait_on_dependent_grids(); - #endif if constexpr (cute::is_same_v) { do { auto [next_work_tile_info, increment_pipe] = scheduler.advance_to_next_work(tile_scheduler_pipeline, tile_scheduler_pipe_producer_state); @@ -630,9 +630,6 @@ public: } // Mainloop Producer Warp else if (producer_warp_role == ProducerWarpRole::Mainloop) { - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 - cutlass::arch::wait_on_dependent_grids(); - #endif int32_t curr_batch = idx2crd(work_tile_info.L_idx, shape<4>(gB_nkl)); // Usually just returns work_tile_info.L_idx; int32_t const mock_l_coord = 0; int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); @@ -797,9 +794,6 @@ public: } // Mainloop Auxiliary Load Producer Warp End // Epilogue Producer Warp else if (producer_warp_role == ProducerWarpRole::Epilogue && collective_epilogue.is_producer_load_needed()) { - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 - cutlass::arch::wait_on_dependent_grids(); - #endif int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); int32_t const sm_count = params.hw_info.sm_count; 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 559b9afe..be70bdff 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,6 +646,9 @@ public: // Get pipeline stage increments from tensor shapes auto k_tile_count = size<3>(gA_mkl); + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + cutlass::arch::wait_on_dependent_grids(); + #endif if (warp_group_role == WarpGroupRole::Producer) { cutlass::arch::warpgroup_reg_dealloc(); @@ -665,9 +668,6 @@ public: } // Mainloop Producer Warp else if (producer_warp_role == ProducerWarpRole::Mainloop) { - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 - cutlass::arch::wait_on_dependent_grids(); - #endif int32_t curr_batch = idx2crd(work_tile_info.L_idx, shape<4>(gB_nkl)); // Usually just returns work_tile_info.L_idx; int32_t const mock_l_coord = 0; int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); @@ -770,9 +770,6 @@ public: } // Mainloop Producer Warp End else if (producer_warp_role == ProducerWarpRole::MainloopAux) { if constexpr (IsMainloopAuxiliaryLoadNeeded) { - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 - cutlass::arch::wait_on_dependent_grids(); - #endif int32_t curr_batch = idx2crd(work_tile_info.L_idx, shape<4>(gB_nkl)); // Usually just returns work_tile_info.L_idx; int32_t const mock_l_coord = 0; @@ -835,9 +832,6 @@ public: } // Mainloop Auxiliary Load Producer Warp End // Epilogue Producer Warp else if (producer_warp_role == ProducerWarpRole::Epilogue && collective_epilogue.is_producer_load_needed()) { - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 - cutlass::arch::wait_on_dependent_grids(); - #endif int32_t const sm_idx = blockIdx.x + (blockIdx.y * gridDim.x); int32_t const sm_count = params.hw_info.sm_count; From b6ad6db219055f33bd20ccb22994ab7e8ec0b9ec Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Mon, 17 Nov 2025 13:21:17 +0800 Subject: [PATCH 5/9] Delete unnecessary #ifdef / #endif. --- .../sm90_gemm_array_tma_warpspecialized_cooperative.hpp | 5 +++-- .../kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp | 5 +++-- 2 files changed, 6 insertions(+), 4 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 9e326d2b..987dd622 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,9 +608,10 @@ public: // Get pipeline stage increments from tensor shapes auto k_tile_count = size<3>(gA_mkl); - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + // Ensure that the kernel does not touch + // unflushed global memory prior to this instruction cutlass::arch::wait_on_dependent_grids(); - #endif + 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 be70bdff..bd979c7b 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,9 +646,10 @@ public: // Get pipeline stage increments from tensor shapes auto k_tile_count = size<3>(gA_mkl); - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + // Ensure that the kernel does not touch + // unflushed global memory prior to this instruction cutlass::arch::wait_on_dependent_grids(); - #endif + if (warp_group_role == WarpGroupRole::Producer) { cutlass::arch::warpgroup_reg_dealloc(); From 929e1e0259d2dcbc0d44f174949f4f8c4182bdc9 Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Tue, 18 Nov 2025 08:45:40 +0800 Subject: [PATCH 6/9] Remove unnecessary #ifdef / #endif for launch_dependent_grids. --- .../sm90_gemm_array_tma_warpspecialized_cooperative.hpp | 4 ++-- .../kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp | 4 ++-- 2 files changed, 4 insertions(+), 4 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 987dd622..25a7e2f7 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 @@ -998,11 +998,11 @@ public: // Get next work tile auto [next_work_tile_info, increment_pipe] = scheduler.fetch_next_work(work_tile_info, tile_scheduler_pipeline, tile_scheduler_pipe_consumer_state); - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + if (!next_work_tile_info.is_valid()) { cutlass::arch::launch_dependent_grids(); } - #endif + work_tile_info = next_work_tile_info; if (increment_pipe) { ++tile_scheduler_pipe_consumer_state; 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 bd979c7b..2c55987a 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 @@ -1042,11 +1042,11 @@ public: // Get next work tile auto [next_work_tile_info, increment_pipe] = scheduler.fetch_next_work(work_tile_info, tile_scheduler_pipeline, tile_scheduler_pipe_consumer_state); - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + if (!next_work_tile_info.is_valid()) { cutlass::arch::launch_dependent_grids(); } - #endif + work_tile_info = next_work_tile_info; if (increment_pipe) { ++tile_scheduler_pipe_consumer_state; From ff7f2dcdfb0415c791d44bbfeab94d2e9d30e023 Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Fri, 21 Nov 2025 11:10:09 +0800 Subject: [PATCH 7/9] 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(); From 17ffd56dfee6132a4eb4aec663c9170cf0d2a73d Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Fri, 21 Nov 2025 11:17:01 +0800 Subject: [PATCH 8/9] Remove unnecessary #ifdef #endif for general gemm. --- .../gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp | 3 +-- .../gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp | 4 ---- 2 files changed, 1 insertion(+), 6 deletions(-) diff --git a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp index 5c04259b..e9908fad 100644 --- a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp +++ b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp @@ -804,7 +804,7 @@ public: // Update starting mainloop pipeline state for the next tile mainloop_pipe_consumer_state.advance(work_k_tile_count); } - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 + if (scheduler.is_last_tile(work_tile_info)) { // Hint on an early release of global memory resources. // The timing of calling this function only influences performance, @@ -812,7 +812,6 @@ public: cutlass::arch::launch_dependent_grids(); } - #endif // Index of warp group within consumer warp groups int consumer_warp_group_idx = canonical_warp_group_idx() - NumLoadWarpGroups; diff --git a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp index 073f3a50..f3d99d4e 100644 --- a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp +++ b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp @@ -800,7 +800,6 @@ public: else if (warp_group_role == WarpGroupRole::Consumer0 || warp_group_role == WarpGroupRole::Consumer1) { cutlass::arch::warpgroup_reg_alloc(); - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 // It is possible to have work tiles start off invalid, // so we have to check that first. if (not work_tile_info.is_valid()) { @@ -811,7 +810,6 @@ public: return; } - #endif if constexpr (IsSchedDynamicPersistent) { // Consumer0's initial tile is static. It starts consuming the 2nd tile. @@ -868,7 +866,6 @@ public: // Update starting mainloop pipeline state for the next tile mainloop_pipe_consumer_state.advance(k_tile_count * NumMmaWarpGroups); - #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 if (scheduler.is_last_tile(work_tile_info, NumMmaWarpGroups)) { // Hint on an early release of global memory resources. // The timing of calling this function only influences performance, @@ -876,7 +873,6 @@ public: cutlass::arch::launch_dependent_grids(); } - #endif // Order two Math WG's Epilogue one after the other math_wg_order_barrier.wait(); From 95f8beb44c3ac0bcfc285c3bb8746ff7854839bb Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Tue, 9 Dec 2025 11:52:23 +0800 Subject: [PATCH 9/9] Revert "Remove unnecessary #ifdef #endif for general gemm." This reverts commit 17ffd56dfee6132a4eb4aec663c9170cf0d2a73d. --- .../gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp | 3 ++- .../gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp | 4 ++++ 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp index e9908fad..5c04259b 100644 --- a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp +++ b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp @@ -804,7 +804,7 @@ public: // Update starting mainloop pipeline state for the next tile mainloop_pipe_consumer_state.advance(work_k_tile_count); } - + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 if (scheduler.is_last_tile(work_tile_info)) { // Hint on an early release of global memory resources. // The timing of calling this function only influences performance, @@ -812,6 +812,7 @@ public: cutlass::arch::launch_dependent_grids(); } + #endif // Index of warp group within consumer warp groups int consumer_warp_group_idx = canonical_warp_group_idx() - NumLoadWarpGroups; diff --git a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp index f3d99d4e..073f3a50 100644 --- a/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp +++ b/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp @@ -800,6 +800,7 @@ public: else if (warp_group_role == WarpGroupRole::Consumer0 || warp_group_role == WarpGroupRole::Consumer1) { cutlass::arch::warpgroup_reg_alloc(); + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 // It is possible to have work tiles start off invalid, // so we have to check that first. if (not work_tile_info.is_valid()) { @@ -810,6 +811,7 @@ public: return; } + #endif if constexpr (IsSchedDynamicPersistent) { // Consumer0's initial tile is static. It starts consuming the 2nd tile. @@ -866,6 +868,7 @@ public: // Update starting mainloop pipeline state for the next tile mainloop_pipe_consumer_state.advance(k_tile_count * NumMmaWarpGroups); + #ifdef CUTLASS_ENABLE_GDC_FOR_SM90 if (scheduler.is_last_tile(work_tile_info, NumMmaWarpGroups)) { // Hint on an early release of global memory resources. // The timing of calling this function only influences performance, @@ -873,6 +876,7 @@ public: cutlass::arch::launch_dependent_grids(); } + #endif // Order two Math WG's Epilogue one after the other math_wg_order_barrier.wait();