diff --git a/projects/rocshmem/CMakeLists.txt b/projects/rocshmem/CMakeLists.txt index de49a680b6..88d2341dce 100644 --- a/projects/rocshmem/CMakeLists.txt +++ b/projects/rocshmem/CMakeLists.txt @@ -109,6 +109,7 @@ set(DEFAULT_GPUS gfx90a:xnack-; gfx90a:xnack+; gfx1100; + gfx1201; gfx942) if(${ROCM_MAJOR_VERSION} GREATER 6) diff --git a/projects/rocshmem/src/assembly.hpp b/projects/rocshmem/src/assembly.hpp index 2dd5e8dea0..de9c339c22 100644 --- a/projects/rocshmem/src/assembly.hpp +++ b/projects/rocshmem/src/assembly.hpp @@ -59,6 +59,13 @@ __device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) { "s_waitcnt vmcnt(0)" : "=v"(ret) : "v"(src)); +#endif +#if defined(__gfx1201__) + asm volatile( + "global_load_u8 %0 %1 off scope:SCOPE_SYS \n" + "s_wait_loadcnt 0x0" + : "=v"(ret) + : "v"(src)); #endif return ret; } @@ -83,6 +90,13 @@ __device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_va : "=v"(*assigned_value) : "v"(read_value)); #endif +#if defined(__gfx1201__) + asm volatile( + "global_load_i8 %0 %1 off scope:SCOPE_SYS \n" + "s_wait_loadcnt 0x0" + : "=v"(*assigned_value) + : "v"(read_value)); +#endif } __device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assigned_value, @@ -104,6 +118,13 @@ __device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assi "s_waitcnt vmcnt(0)" : "=v"(*assigned_value) : "v"(read_value)); +#endif + #if defined(__gfx1201__) + asm volatile( + "global_load_b64 %0 %1 off scope:SCOPE_SYS \n" + "s_wait_loadcnt 0x0" + : "=v"(*assigned_value) + : "v"(read_value)); #endif } @@ -135,6 +156,13 @@ NOWARN(-Wdeprecated-volatile, "s_waitcnt vmcnt(0)" : "=v"(ret) : "v"(src)); +#endif +#if defined(__gfx1201__) + asm volatile( + "global_load_b32 %0 %1 off scope:SCOPE_SYS \n" + "s_wait_loadcnt 0x0" + : "=v"(ret) + : "v"(src)); #endif break; case 8: @@ -155,6 +183,13 @@ NOWARN(-Wdeprecated-volatile, "s_waitcnt vmcnt(0)" : "=v"(ret) : "v"(src)); +#endif +#if defined(__gfx1201__) + asm volatile( + "global_load_b64 %0 %1 off scope:SCOPE_SYS \n" + "s_wait_loadcnt 0x0" + : "=v"(ret) + : "v"(src)); #endif break; default: @@ -196,6 +231,9 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, #endif #if defined(__gfx942__) || defined(__gfx950__) asm volatile("flat_store_short %0 %1 sc0 sc1" : : "v"(dst), "v"(val16)); +#endif +#if defined(__gfx1201__) + asm volatile("flat_store_b16 %0 %1 scope:SCOPE_SYS" : : "v"(dst), "v"(val16)); #endif break; } @@ -210,6 +248,9 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, #endif #if defined(__gfx942__) || defined(__gfx950__) asm volatile("flat_store_dword %0 %1 sc0 sc1" : : "v"(dst), "v"(val32)); +#endif +#if defined(__gfx1201__) + asm volatile("flat_store_b32 %0 %1 scope:SCOPE_SYS" : : "v"(dst), "v"(val32)); #endif break; } @@ -224,6 +265,9 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst, #endif #if defined(__gfx942__) || defined(__gfx950__) asm volatile("flat_store_dwordx2 %0 %1 sc0 sc1" : : "v"(dst), "v"(val64)); +#endif +#if defined(__gfx1201__) + asm volatile("flat_store_b64 %0 %1 scope:SCOPE_SYS" : : "v"(dst), "v"(val64)); #endif break; }