ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-14 02:57:45 +00:00
Code Issues Packages Projects Releases Wiki Activity

forward compatible

Browse Source
This commit is contained in:
zanzhang
2025-12-24 14:51:24 +08:00
parent 1c3151963b
commit 925b245697
2 changed files with 9 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
include/ck_tile/core/arch/amd_buffer_coherence.hpp
View File

@@ -83,6 +83,9 @@ enum struct amd_buffer_coherence_enum
slc = SYSTEM_NT,
glc_slc = DEVICE_NT | SYSTEM_NT,
// To forward compatible
WAVE_NT0 = INT_MAX,
WAVE_NT1 = INT_MAX - 1,
// gfx94: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
// SC[1:0] System Cache level: 0=wave, 1=group, 2=device, 3=system
// NT Non-Temporal: 0=expect temporal reuse; 1=do not expect temporal reuse
@@ -118,6 +121,12 @@ enum struct amd_buffer_coherence_enum
SYSTEM_NT0 = 0,
DEVICE_NT1 = glc,
SYSTEM_NT1 = slc,
// To forward compatible
GROUP_NT0 = INT_MAX - 2,
GROUP_NT1 = INT_MAX - 3,
WAVE_NT0 = INT_MAX,
WAVE_NT1 = INT_MAX - 1,
#endif
};

2
include/ck_tile/ops/flatmm/pipeline/flatmm_pipeline_agmem_bgmem_creg_v1.hpp
View File

@@ -31,9 +31,7 @@ struct BaseFlatmmPipelineAGmemBGmemCRegV1
ck_tile::ignore = K;
if(M <= 416)
{
#if defined(__gfx942__) || defined(__gfx950__)
return ck_tile::amd_buffer_coherence_enum::WAVE_NT1;
#endif
}
return ck_tile::amd_buffer_coherence_enum::coherence_default;
}