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

Merge pull request #2719 from HydraQYH/dev_support_pdl_for_sm90_gemm_array_tma_ws

Browse Source 
Support PDL for SM90 Array TMA GEMM
This commit is contained in:
Junkai-Wu
2025-12-09 18:32:15 +08:00
committed by GitHub
parent 1e7d3e030b 95f8beb44c
commit c6dfdf8375
2 changed files with 10 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp
View File

@@ -994,6 +994,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);
if (!next_work_tile_info.is_valid()) {
cutlass::arch::launch_dependent_grids();
}
work_tile_info = next_work_tile_info;
if (increment_pipe) {
++tile_scheduler_pipe_consumer_state;

5
include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp
View File

@@ -1038,6 +1038,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);
if (!next_work_tile_info.is_valid()) {
cutlass::arch::launch_dependent_grids();
}
work_tile_info = next_work_tile_info;
if (increment_pipe) {
++tile_scheduler_pipe_consumer_state;