ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-30 03:37:38 +00:00
Code Issues Packages Projects Releases Wiki Activity

fix async inline assembly

Browse Source
This commit is contained in:
joye
2025-05-19 09:33:17 +08:00
parent 2a9f0fff5d
commit 0f2dfa8d38
1 changed files with 3 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
include/ck_tile/core/arch/amd_buffer_addressing.hpp
View File

@@ -1666,7 +1666,7 @@ struct async_buffer_load<16, pre_nop>
asm volatile("s_nop 4\n"
"buffer_load_dwordx4 %1, %2, 0 offen offset:%3 lds"
: "=r"(smem) /*dummy dependency for smem*/
: "v"(voffset), "s"(rsrc), "n"(ioffset)
: "v"(v_offset), "s"(rsrc), "n"(i_offset)
: "memory");
// __builtin_amdgcn_struct_buffer_load_lds(rsrc, smem, 16, i_offset, v_offset, 0, 0, 0);
#else
@@ -1694,7 +1694,7 @@ struct async_buffer_load<12, pre_nop>
asm volatile("s_nop 4\n"
"buffer_load_dwordx3 %1, %2, 0 offen offset:%3 lds"
: "=r"(smem) /*dummy dependency for smem*/
: "v"(voffset), "s"(rsrc), "n"(ioffset)
: "v"(v_offset), "s"(rsrc), "n"(i_offset)
: "memory");
// __builtin_amdgcn_struct_buffer_load_lds(rsrc, smem, 12, i_offset, v_offset, 0, 0, 0);
#else
@@ -1722,7 +1722,7 @@ struct async_buffer_load<4, pre_nop>
asm volatile("s_nop 4\n"
"buffer_load_dword %1, %2, 0 offen offset:%3 lds"
: "=r"(smem) /*dummy dependency for smem*/
: "v"(voffset), "s"(rsrc), "n"(ioffset)
: "v"(v_offset), "s"(rsrc), "n"(i_offset)
: "memory");
// __builtin_amdgcn_struct_buffer_load_lds(rsrc, smem, 4, i_offset, v_offset, 0, 0, 0);
#else