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

Fix

Browse Source
This commit is contained in:
Enrico Degregori
2026-05-12 07:39:55 +00:00
parent a170e2bd9d
commit 066d21ec32
1 changed files with 12 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp
View File

@@ -1281,14 +1281,20 @@ CK_TILE_DEVICE void async_buffer_load_dwordxn_v(void* smem,
index_t /*flag*/ = 0,
bool_constant<pre_nop> = {})
{
// #define CK_TILE_ASYNC_LOAD_WITH_INSTR(instr) \
// if constexpr(pre_nop) \
// asm volatile("s_nop 4\n" instr " %1, %2, 0 offen offset:%3 lds" \
// : "=r"(smem) /*dummy dependency for smem*/ \
// : "v"(voffset), "s"(rsrc), "n"(ioffset) \
// : "memory"); \
// else \
// asm volatile(instr " %1, %2, 0 offen offset:%3 lds" \
// : "=r"(smem) /*dummy dependency for smem*/ \
// : "v"(voffset), "s"(rsrc), "n"(ioffset) \
// : "memory");
#define CK_TILE_ASYNC_LOAD_WITH_INSTR(instr) \
if constexpr(pre_nop) \
asm volatile("s_nop 4\n" instr " %1, %2, 0 offen offset:%3 lds" \
: "=r"(smem) /*dummy dependency for smem*/ \
: "v"(voffset), "s"(rsrc), "n"(ioffset) \
: "memory"); \
else \
asm volatile(instr " %1, %2, 0 offen offset:%3 lds" \
: "=r"(smem) /*dummy dependency for smem*/ \
: "v"(voffset), "s"(rsrc), "n"(ioffset) \
: "memory");