initial commit for gfx12 support (#305)

[ROCm/rocshmem commit: d185fe3555]

projects/rocshmem/CMakeLists.txt

@@ -109,6 +109,7 @@ set(DEFAULT_GPUS
     gfx90a:xnack-;
     gfx90a:xnack+;
     gfx1100;
+    gfx1201;
     gfx942)
 
 if(${ROCM_MAJOR_VERSION} GREATER 6)

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;
     }