Skip to content

GitLab

Commit 69ad91b2 authored by illsilin's avatar illsilin
Browse files

fix the gfx12 assembly syntax

parent bfdda4fe
Hide whitespace changes
Inline Side-by-side
Showing with 8 additions and 9 deletions
include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp
View file @ 69ad91b2
...@@ -481,8 +481,8 @@ struct BlockwiseGemmXdlopsInterwave_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1 ...@@ -481,8 +481,8 @@ struct BlockwiseGemmXdlopsInterwave_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1
{ {
#ifdef __gfx12__ #ifdef __gfx12__
asm volatile("\ asm volatile("\
s_barrier_signal \n \ s_barrier_signal -1 \n \
s_barrier_wait \ s_barrier_wait -1 \
" ::); " ::);
#else #else
asm volatile("s_barrier" ::); asm volatile("s_barrier" ::);
......
include/ck/utility/synchronization.hpp
View file @ 69ad91b2
...@@ -12,9 +12,9 @@ __device__ void block_sync_lds() ...@@ -12,9 +12,9 @@ __device__ void block_sync_lds()
#if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM #if CK_EXPERIMENTAL_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM
#ifdef __gfx12__ #ifdef __gfx12__
asm volatile("\ asm volatile("\
s_wait_idle lgkmcnt(0) \n \ s_wait_idle \n \
s_barrier_signal \n \ s_barrier_signal -1 \n \
s_barrier_wait \ s_barrier_wait -1 \
" ::); " ::);
#else #else
asm volatile("\ asm volatile("\
...@@ -31,10 +31,9 @@ __device__ void block_sync_lds_direct_load() ...@@ -31,10 +31,9 @@ __device__ void block_sync_lds_direct_load()
{ {
#ifdef __gfx12__ #ifdef __gfx12__
asm volatile("\ asm volatile("\
s_wait_idle vmcnt(0) \n \ s_wait_idle \n \
s_wait_idle lgkmcnt(0) \n \ s_barrier_signal -1 \n \
s_barrier_signal \n \ s_barrier_wait -1 \
s_barrier_wait \
" ::); " ::);
#else #else
asm volatile("\ asm volatile("\
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment