AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Remove support code for GFX940 and GFX941 targets (#55)

Browse Source 
[ROCm/rocshmem commit: 1210b6419f]
This commit is contained in:
Avinash Kethineedi
2025-03-21 14:31:49 -05:00
committed by GitHub
parent 033253fbdf
commit 6f78b4300f
1 changed files with 11 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocshmem/src/assembly.hpp
+11 -11
View File
@@ -51,7 +51,7 @@ __device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) {
: "=v"(ret)
: "v"(src));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile(
"global_load_ubyte %0 %1 off sc0 sc1 \n"
"s_waitcnt vmcnt(0)"
@@ -74,7 +74,7 @@ __device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_va
: "=v"(*assigned_value)
: "v"(read_value));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile(
"global_load_sbyte %0 %1 off sc0 sc1\n "
"s_waitcnt vmcnt(0)"
@@ -96,7 +96,7 @@ __device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assi
: "=v"(*assigned_value)
: "v"(read_value));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile(
"global_load_dwordx2 %0 %1 off sc0 sc1\n "
"s_waitcnt vmcnt(0)"
@@ -127,7 +127,7 @@ NOWARN(-Wdeprecated-volatile,
: "=v"(ret)
: "v"(src));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile(
"global_load_dword %0 %1 off sc0 sc1 \n"
"s_waitcnt vmcnt(0)"
@@ -147,7 +147,7 @@ NOWARN(-Wdeprecated-volatile,
: "=v"(ret)
: "v"(src));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile(
"global_load_dwordx2 %0 %1 off sc0 sc1 \n"
"s_waitcnt vmcnt(0)"
@@ -172,7 +172,7 @@ __device__ __forceinline__ void __roc_inv() {
#if defined(__gfx90a__)
// asm volatile("buffer_wbinvl1;");
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
// asm volatile("buffer_inv sc0 sc1;");
#endif
#endif
@@ -188,7 +188,7 @@ __device__ __forceinline__ void __roc_flush() {
// asm volatile("s_dcache_wb;");
// asm volatile("buffer_wbl2;");
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
// asm volatile("s_dcache_wb;");
// asm volatile("buffer_wbl2;");
#endif
@@ -207,7 +207,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
#if defined(__gfx90a__)
asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile("flat_store_short %0 %1 sc0 sc1" : : "v"(dst), "v"(val16));
#endif
break;
@@ -221,7 +221,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
#if defined(__gfx90a__)
asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile("flat_store_dword %0 %1 sc0 sc1" : : "v"(dst), "v"(val32));
#endif
break;
@@ -235,7 +235,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
#if defined(__gfx90a__)
asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile("flat_store_dwordx2 %0 %1 sc0 sc1" : : "v"(dst), "v"(val64));
#endif
break;
@@ -257,7 +257,7 @@ __device__ __forceinline__ uint64_t __read_clock() {
"s_waitcnt lgkmcnt(0)\n"
: "=s"(clock));
#endif
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
#if defined(__gfx942__)
asm volatile(
"s_memrealtime %0\n"
"s_waitcnt lgkmcnt(0)\n"