ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-11 17:00:18 +00:00
Code Issues Packages Projects Releases Wiki Activity

tempsave, update the blocksync functions

Browse Source
This commit is contained in:
aska-0096
2025-08-11 14:21:09 +00:00
parent 7b8052d7ca
commit b86f7786e2
1 changed files with 12 additions and 22 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -89,21 +89,6 @@ CK_TILE_DEVICE index_t get_thread_id() { return threadIdx.x; }
CK_TILE_DEVICE index_t get_block_id() { return blockIdx.x; }
template <index_t lgkmcnt = 0>
CK_TILE_DEVICE void block_sync_lds()
{
if constexpr(lgkmcnt > 15)
{
__builtin_amdgcn_s_waitcnt(CK_TILE_S_CNT_MAX & CK_TILE_LGKMCNT(15));
}
else
{
__builtin_amdgcn_s_waitcnt(CK_TILE_S_CNT_MAX & CK_TILE_LGKMCNT(lgkmcnt));
}
__builtin_amdgcn_s_barrier();
}
CK_TILE_DEVICE void block_sync_load_raw(index_t cnt = 0)
{
#ifdef __gfx12__
@@ -122,13 +107,6 @@ CK_TILE_DEVICE void block_sync_load_raw(index_t cnt = 0)
#endif
}
template <index_t vmcnt>
CK_TILE_DEVICE void block_sync_lds_direct_load()
{
// We don't sync the lds insts here.
__builtin_amdgcn_s_waitcnt(CK_TILE_S_CNT_MAX & CK_TILE_VMCNT(vmcnt));
__builtin_amdgcn_s_barrier();
}
// https://llvm.org/docs/AMDGPU/gfx9_waitcnt.html
struct waitcnt_arg
{
@@ -181,6 +159,18 @@ CK_TILE_DEVICE void s_waitcnt_barrier()
__builtin_amdgcn_s_barrier();
}
template <index_t lgkmcnt = 0>
CK_TILE_DEVICE void block_sync_lds()
{
s_waitcnt_barrier<waitcnt_arg::kMaxVmCnt, waitcnt_arg::kMaxExpCnt, lgkmcnt>();
}
template <index_t vmcnt = 0>
CK_TILE_DEVICE void block_sync_lds_direct_load()
{
s_waitcnt_barrier<vmcnt, waitcnt_arg::kMaxExpCnt, waitcnt_arg::kMaxLgkmCnt>();
}
CK_TILE_DEVICE void s_nop(index_t cnt = 0)
{
#if 1