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

disable gfx1100 temporarily (#322)

Browse Source 
[ROCm/rocshmem commit: ef3ba6cd45]
This commit is contained in:
Edgar Gabriel
2025-11-14 11:46:19 -06:00
committed by GitHub
parent a500bc8029
commit 2b0bca5c87
2 changed files with 17 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocshmem/CMakeLists.txt
-1
View File
@@ -108,7 +108,6 @@ include(cmake/rocm_local_targets.cmake)
set(DEFAULT_GPUS set(DEFAULT_GPUS
gfx90a:xnack-; gfx90a:xnack-;
gfx90a:xnack+; gfx90a:xnack+;
gfx1100;
gfx1201; gfx1201;
gfx942) gfx942)
projects/rocshmem/src/assembly.hpp
+17 -17
View File
@@ -44,9 +44,9 @@ __device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) {
int ret; int ret;
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile( asm volatile(
"global_load_ubyte %0 %1 off glc slc \n" "global_load_ubyte %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)" "s_waitcnt vmcnt(0)"
@@ -74,9 +74,9 @@ __device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_va
volatile char *read_value) { volatile char *read_value) {
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile( asm volatile(
"global_load_sbyte %0 %1 off glc slc\n " "global_load_sbyte %0 %1 off glc slc\n "
"s_waitcnt vmcnt(0)" "s_waitcnt vmcnt(0)"
@@ -103,9 +103,9 @@ __device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assi
volatile uint64_t *read_value) { volatile uint64_t *read_value) {
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile( asm volatile(
"global_load_dwordx2 %0 %1 off glc slc\n " "global_load_dwordx2 %0 %1 off glc slc\n "
"s_waitcnt vmcnt(0)" "s_waitcnt vmcnt(0)"
@@ -141,9 +141,9 @@ NOWARN(-Wdeprecated-volatile,
case 4: case 4:
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile( asm volatile(
"global_load_dword %0 %1 off glc slc \n" "global_load_dword %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)" "s_waitcnt vmcnt(0)"
@@ -168,9 +168,9 @@ NOWARN(-Wdeprecated-volatile,
case 8: case 8:
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile( asm volatile(
"global_load_dwordx2 %0 %1 off glc slc \n" "global_load_dwordx2 %0 %1 off glc slc \n"
"s_waitcnt vmcnt(0)" "s_waitcnt vmcnt(0)"
@@ -204,7 +204,7 @@ __device__ __forceinline__ void __roc_flush() {
#if not defined USE_HDP_FLUSH #if not defined USE_HDP_FLUSH
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) #if defined(__gfx90a__)
// asm volatile("s_dcache_wb;"); // asm volatile("s_dcache_wb;");
@@ -224,9 +224,9 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
int16_t val16{*(reinterpret_cast<int16_t*>(val))}; int16_t val16{*(reinterpret_cast<int16_t*>(val))};
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16)); asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16));
#endif #endif
#if defined(__gfx942__) || defined(__gfx950__) #if defined(__gfx942__) || defined(__gfx950__)
@@ -241,9 +241,9 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
int32_t val32{*(reinterpret_cast<int32_t*>(val))}; int32_t val32{*(reinterpret_cast<int32_t*>(val))};
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32)); asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32));
#endif #endif
#if defined(__gfx942__) || defined(__gfx950__) #if defined(__gfx942__) || defined(__gfx950__)
@@ -258,9 +258,9 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
int64_t val64{*(reinterpret_cast<int64_t*>(val))}; int64_t val64{*(reinterpret_cast<int64_t*>(val))};
#if defined(__gfx906__) #if defined(__gfx906__)
#endif #endif
#if defined(__gfx908__) #if defined(__gfx908__) || defined(__gfx1100__)
#endif #endif
#if defined(__gfx90a__) || defined (__gfx1100__) #if defined(__gfx90a__)
asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64)); asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64));
#endif #endif
#if defined(__gfx942__) || defined(__gfx950__) #if defined(__gfx942__) || defined(__gfx950__)