From 929e1e0259d2dcbc0d44f174949f4f8c4182bdc9 Mon Sep 17 00:00:00 2001 From: HydraQYH Date: Tue, 18 Nov 2025 08:45:40 +0800 Subject: [PATCH] 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;