ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-04-20 06:49:15 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add a workaround for a compiler issue for bwd on gfx90a and ROCm 7.1.1 (#3369)

Browse Source 
Sometimes there are not enough wait-states between v_mfma_f32... and v_accvgpr_read_b32 instructions if they are separated by s_cbranch.
The workaround is to read accvgprs to vgpr before branching.
This commit is contained in:
Anton Gorenko
2025-12-08 20:44:17 +05:00
committed by GitHub
parent 878b4e7f46
commit ca6143f0b2
1 changed files with 9 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
include/ck_tile/ops/fmha/pipeline/block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr.hpp
View File

@@ -552,6 +552,15 @@ struct BlockFmhaBwdDQDKDVPipelineKRKTRVR
});
});
}
#if defined(__gfx9__)
else
{
// Workaround for a compiler issue: sometimes there are not enough wait-states
// between v_mfma_f32... and v_accvgpr_read_b32 instructions if they are separated
// by s_cbranch.
tile_elementwise_inout([](auto& x) { asm("; force move to %0" : "+v"(x)); }, s_acc);
}
#endif
{
bool need_perpixel_check = mask.IsEdgeTile(