ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 01:27:34 +00:00
Code Issues Packages Projects Releases Wiki Activity

Instruction scheduling in block gemm instead of scheduler

Browse Source 
Compiler generates better assembly decreasing spilling for large tiles
AB Quant is unchanged
This commit is contained in:
Enrico Degregori
2026-07-01 13:34:22 +00:00
parent 0523204792
commit d408192bfe
2 changed files with 6 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
include/ck_tile/ops/gemm/block/block_gemm_areg_breg_creg_eight_waves_v1.hpp
View File

@@ -248,6 +248,12 @@ struct BlockGemmARegBRegCRegEightWavesV1
merge_sequences(c_iter_idx{}, c_warp_y_index_zeros),
merge_sequences(sequence<1, 1>{}, c_warp_y_lengths),
c_warp_tensor.get_thread_buffer());
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0);
if constexpr(nIter == 0 && mIter == MIterPerWarp - 1 && kIter == 0)
{
s_waitcnt_lgkm<4>();
__builtin_amdgcn_sched_barrier(0);
}
});
});
}

8
include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async_eight_waves.hpp
View File

@@ -183,14 +183,6 @@ struct GemmPipelineAgBgCrCompAsyncEightWaves : public BaseGemmPipelineAgBgCrComp
// Hot loop scheduler
// ------------------
auto hot_loop_scheduler = [&]() {
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
s_waitcnt_lgkm<4>();
__builtin_amdgcn_sched_group_barrier(0x004, 1, 0); // lgkmcnt / SALU
static_for<0, MFMA_INST - 3, 1>{}([&](auto) {
__builtin_amdgcn_sched_group_barrier(0x008, 1, 0); // MFMA
});
__builtin_amdgcn_sched_barrier(0);
};