CK build from windows gfx1100 #2854
Closed
yanite
started this conversation in
Show and tell
Replies: 0 comments
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Uh oh!
There was an error while loading. Please reload this page.
Uh oh!
There was an error while loading. Please reload this page.
-
U:\python_AI\xformers\third_party\composable_kernel_tiled\include\ck_tile/core/arch/amd_buffer_addressing_builtins.hpp:1160:23:error: invalid operand for instruction1160 | "buffer_load_dword %1, %2, 0 offen offset:%3 lds"
| ^
:2:24: note: instantiated into assembly here
2 | buffer_load_dword v24, s[44:47], 0 offen offset:0 lds
| ^
`
**Who knows how this assembly should be written? by gfx1100 **
source code
template <bool pre_nop = false> CK_TILE_DEVICE void async_buffer_load_dword_v(void* smem, int32x4_t rsrc, index_t voffset, index_t /*soffset*/, index_t ioffset /*max 0xFFF*/, index_t /*flag*/ = 0, bool_constant<pre_nop> = {}) { if constexpr(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) : "memory"); else asm volatile("buffer_load_dword %1, %2, 0 offen offset:%3 lds" : "=r"(smem) /*dummy dependency for smem*/ : "v"(voffset), "s"(rsrc), "n"(ioffset) : "memory"); }$ clang --version AMD clang version 22.0.0 //github.com/ROCm/llvm-project.git 8e85e3138dd485c4221cc12aff9eb60ab48ed3b5) Target: x86_64-pc-windows-msvc Thread model: posix InstalledDir: U:\ROCm\TheRock\build\compiler\amd-llvm\build\bin$ rocm-sdk version 7.11.0a20260102Beta Was this translation helpful? Give feedback.
All reactions