[Bug Fix]Bypass launch grids for SM120 Kernel with SM90 Mainloop & SM100 TileScheduler

include/cutlass/gemm/kernel/sm100_tile_scheduler.hpp

@@ -449,6 +449,11 @@ class PersistentTileSchedulerSm100 {
     return clc_pipe_producer_state;
   }
 
+  CUTLASS_DEVICE
+  bool is_last_tile(WorkTileInfo& work_tile_info, uint32_t advance_count = 1) const {
+    return false;
+  }
+
   // Kernel helper function to get next work tile
   template <class TileSchedulerPipeline, class TileSchedulerPipelineState>
   CUTLASS_HOST_DEVICE

include/cutlass/gemm/kernel/sm100_tile_scheduler_stream_k.hpp

@@ -226,7 +226,12 @@ class PersistentTileSchedulerSm100StreamK {
   PipelineState<Stages> 
   advance_to_next_work(Pipeline& clc_pipeline, PipelineState<Stages> clc_pipe_producer_state) const {
     return sm100_scheduler_.advance_to_next_work(clc_pipeline, clc_pipe_producer_state);
- }
+  }
+
+  CUTLASS_DEVICE
+  bool is_last_tile(WorkTileInfo work_tile_info, uint32_t advance_count = 1) const {
+    return false;
+  }
 
   // Given the inputs, computes the total number of output blocks this problem will compute over
   template<class ProblemShape>

include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_cooperative.hpp

@@ -804,15 +804,14 @@ class GemmUniversal<
           // 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,
           // not functional correctness.
           cutlass::arch::launch_dependent_grids();
 
         }
-        #endif
 
         // Index of warp group within consumer warp groups
         int consumer_warp_group_idx = canonical_warp_group_idx() - NumLoadWarpGroups;

include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp

@@ -868,15 +868,13 @@ class GemmUniversal<
         // 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,
           // not functional correctness.
           cutlass::arch::launch_dependent_grids();
 
         }
-        #endif
 
         // Order two Math WG's Epilogue one after the other
         math_wg_order_barrier.wait();