SWDEV-393910 - Port gfx94x changes to mainline.

Change-Id: Ibf727223bbe5230b132b47c39e0fc1d87cbd3b9c

hipamd/include/hip/amd_detail/amd_hip_unsafe_atomics.h

@@ -53,10 +53,7 @@ THE SOFTWARE.
  * @return Original value contained in \p addr.
  */
 __device__ inline float unsafeAtomicAdd(float* addr, float value) {
-#if defined(__gfx940__) &&                                                     \
-    __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f32)
-  return __builtin_amdgcn_flat_atomic_fadd_f32(addr, value);
-#elif defined(__gfx90a__) &&                                                   \
+#if defined(__gfx90a__) &&                                                   \
     __has_builtin(__builtin_amdgcn_is_shared) &&                               \
     __has_builtin(__builtin_amdgcn_is_private) &&                              \
     __has_builtin(__builtin_amdgcn_ds_atomic_fadd_f32) &&                      \
@@ -178,8 +175,7 @@ __device__ inline float unsafeAtomicMin(float* addr, float val) {
  * @return Original value contained in \p addr.
  */
 __device__ inline double unsafeAtomicAdd(double* addr, double value) {
-#if (defined(__gfx90a__) || defined(__gfx940__)) &&                              \
-    __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f64)
+#if defined(__gfx90a__) && __has_builtin(__builtin_amdgcn_flat_atomic_fadd_f64)
   return __builtin_amdgcn_flat_atomic_fadd_f64(addr, value);
 #elif defined (__hip_atomic_fetch_add)
   return __hip_atomic_fetch_add(addr, value, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
@@ -215,7 +211,7 @@ __device__ inline double unsafeAtomicAdd(double* addr, double value) {
  * @return Original value contained at \p addr.
  */
 __device__ inline double unsafeAtomicMax(double* addr, double val) {
-#if (defined(__gfx90a__) || defined(__gfx940__)) && \
+#if (defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) &&  \
     __has_builtin(__builtin_amdgcn_flat_atomic_fmax_f64)
   return __builtin_amdgcn_flat_atomic_fmax_f64(addr, val);
 #else
@@ -268,7 +264,7 @@ __device__ inline double unsafeAtomicMax(double* addr, double val) {
  * @return Original value contained at \p addr.
  */
 __device__ inline double unsafeAtomicMin(double* addr, double val) {
-#if (defined(__gfx90a__) || defined(__gfx940__)) && \
+#if (defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) &&  \
     __has_builtin(__builtin_amdgcn_flat_atomic_fmin_f64)
   return __builtin_amdgcn_flat_atomic_fmin_f64(addr, val);
 #else
@@ -309,12 +305,15 @@ __device__ inline double unsafeAtomicMin(double* addr, double val) {
  * @return Original value contained in \p addr.
  */
 __device__ inline float safeAtomicAdd(float* addr, float value) {
-#if defined(__gfx908__) ||                                                    \
-    (defined(__gfx90a__) && !__has_builtin(__hip_atomic_fetch_add))
+#if defined(__gfx908__) || defined(__gfx941__)                                \
+    || ((defined(__gfx90a__) || defined(__gfx940__) || defined(__gfx942__))   \
+         && !__has_builtin(__hip_atomic_fetch_add))
   // On gfx908, we can generate unsafe FP32 atomic add that does not follow all
   // IEEE rules when -munsafe-fp-atomics is passed. Do a CAS loop emulation instead.
-  // On gfx90a, if we do not have the __hip_atomic_fetch_add builtin, we need to
-  // force a CAS loop here.
+  // On gfx941, we can generate unsafe FP32 atomic add that may not always happen atomically,
+  // so we need to force a CAS loop emulation to ensure safety.
+  // On gfx90a, gfx940 and gfx942 if we do not have the __hip_atomic_fetch_add builtin, we
+  // need to force a CAS loop here.
   float old_val;
 #if __has_builtin(__hip_atomic_load)
   old_val = __hip_atomic_load(addr, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
@@ -434,8 +433,7 @@ __device__ inline float safeAtomicMin(float* addr, float val) {
  * @return Original value contained in \p addr.
  */
 __device__ inline double safeAtomicAdd(double* addr, double value) {
-#if (defined(__gfx90a__) || defined(__gfx940__)) &&                                                    \
-    __has_builtin(__hip_atomic_fetch_add)
+#if defined(__gfx90a__) &&  __has_builtin(__hip_atomic_fetch_add)
   // On gfx90a, with the __hip_atomic_fetch_add builtin, relaxed system-scope
   // atomics will produce safe CAS loops, but are otherwise not different than
   // agent-scope atomics. This logic is only applicable for gfx90a, and should

hipamd/src/amd_hsa_elf.hpp

@@ -95,13 +95,15 @@ enum : unsigned {
   EF_AMDGPU_MACH_AMDGCN_GFX1034       = 0x03e,
   EF_AMDGPU_MACH_AMDGCN_GFX90A        = 0x03f,
   EF_AMDGPU_MACH_AMDGCN_GFX940        = 0x040,
-  EF_AMDGPU_MACH_AMDGCN_GFX1100       = 0x041,
-  EF_AMDGPU_MACH_AMDGCN_GFX1013       = 0x042,
+  EF_AMDGPU_MACH_AMDGCN_GFX941        = 0x041,
+  EF_AMDGPU_MACH_AMDGCN_GFX942        = 0x042,
   EF_AMDGPU_MACH_AMDGCN_RESERVED_0X43 = 0x043,
-  EF_AMDGPU_MACH_AMDGCN_GFX1103       = 0x044,
-  EF_AMDGPU_MACH_AMDGCN_GFX1036       = 0x045,
-  EF_AMDGPU_MACH_AMDGCN_GFX1101       = 0x046,
-  EF_AMDGPU_MACH_AMDGCN_GFX1102       = 0x047,
+  EF_AMDGPU_MACH_AMDGCN_GFX1100       = 0x044,
+  EF_AMDGPU_MACH_AMDGCN_GFX1013       = 0x045,
+  EF_AMDGPU_MACH_AMDGCN_GFX1103       = 0x046,
+  EF_AMDGPU_MACH_AMDGCN_GFX1036       = 0x047,
+  EF_AMDGPU_MACH_AMDGCN_GFX1101       = 0x048,
+  EF_AMDGPU_MACH_AMDGCN_GFX1102       = 0x049,
 
   // First/last AMDGCN-based processors.
   EF_AMDGPU_MACH_AMDGCN_FIRST = EF_AMDGPU_MACH_AMDGCN_GFX600,

hipamd/src/hip_code_object.cpp

@@ -175,6 +175,16 @@ static bool getProcName(uint32_t EFlags, std::string& proc_name, bool& xnackSupp
       sramEccSupported = true;
       proc_name = "gfx940";
       break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX941:
+      xnackSupported = true;
+      sramEccSupported = true;
+      proc_name = "gfx941";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX942:
+      xnackSupported = true;
+      sramEccSupported = true;
+      proc_name = "gfx942";
+      break;
     case EF_AMDGPU_MACH_AMDGCN_GFX1010:
       xnackSupported = true;
       sramEccSupported = false;

hipamd/src/hiprtc/hiprtcComgrHelper.cpp

@@ -162,6 +162,16 @@ static bool getProcName(uint32_t EFlags, std::string& proc_name, bool& xnackSupp
       sramEccSupported = true;
       proc_name = "gfx940";
       break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX941:
+      xnackSupported = true;
+      sramEccSupported = true;
+      proc_name = "gfx941";
+      break;
+    case EF_AMDGPU_MACH_AMDGCN_GFX942:
+      xnackSupported = true;
+      sramEccSupported = true;
+      proc_name = "gfx942";
+      break;
     case EF_AMDGPU_MACH_AMDGCN_GFX1010:
       xnackSupported = true;
       sramEccSupported = false;