ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 09:08:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

change block_sync_lds function to be consistent with gfx12 path in Old CK

Browse Source
This commit is contained in:
Philip Maybank
2025-10-03 10:50:14 -04:00
parent ef43078788
commit eed91c3984
1 changed files with 6 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
include/ck_tile/core/arch/arch.hpp
View File

@@ -191,7 +191,12 @@ CK_TILE_DEVICE void s_waitcnt_barrier()
template <index_t lgkmcnt = 0>
CK_TILE_DEVICE void block_sync_lds()
{
s_waitcnt_barrier<waitcnt_arg::kMaxVmCnt, waitcnt_arg::kMaxExpCnt, lgkmcnt>();
// s_waitcnt_barrier<waitcnt_arg::kMaxVmCnt, waitcnt_arg::kMaxExpCnt, lgkmcnt>();
asm volatile("\
s_wait_dscnt 0x0 \n \
s_barrier_signal -1 \n \
s_barrier_wait -1 \
" ::);
}
template <index_t vmcnt = 0>