From 95f8beb44c3ac0bcfc285c3bb8746ff7854839bb Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Tue, 9 Dec 2025 11:52:23 +0800 Subject: [PATCH] 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();