[rocm-libraries] ROCm/rocm-libraries#4608 (commit 2b85a8e)

Fully disable the atomic 128 functions when targetting
 unknown or spirv (#4608)

## Motivation

The reported issue, ROCm/rocm-libraries#1112.
Functions were already not used for these unknown targets and spirv, but
still showed the error message. These functions should be fully disabled
for these targets.

## Technical Details

Disabled the 128 bit atomic functions.

## Test Plan

Build with unknown and spirv target.

## Test Result
Both were able to be build.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Author: NB4444

rocprim/include/rocprim/intrinsics/atomic.hpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2017-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -162,6 +162,8 @@ namespace detail
         return __hip_atomic_load(address, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
     }
 
+// atomic_load for uint128 is not supported for unknown targets or spirv
+#if !defined(ROCPRIM_TARGET_UNKNOWN) && !defined(ROCPRIM_TARGET_SPIRV)
     ROCPRIM_DEVICE ROCPRIM_INLINE
     __uint128_t atomic_load(const __uint128_t* address)
     {
@@ -197,25 +199,24 @@ namespace detail
     #define ROCPRIM_ATOMIC_LOAD_SHARED(ptr) \
         ROCPRIM_ATOMIC_LOAD("ds_read_b128", "", "s_waitcnt lgkmcnt(0)", ptr)
     // This architecture doesn't support atomics on the global AS.
-    #define ROCPRIM_ATOMIC_LOAD_GLOBAL(ptr) ROCPRIM_ATOMIC_LOAD_FLAT(ptr)
-#elif ROCPRIM_TARGET_RDNA3 || ROCPRIM_TARGET_CDNA2 || ROCPRIM_TARGET_CDNA1 || ROCPRIM_TARGET_GCN5 \
-    || ROCPRIM_TARGET_SPIRV
-    // We don't really know what architecture we are on when targeting
-    // SPIR-V. Lets just assume it's one of these.
-    #define ROCPRIM_ATOMIC_LOAD_FLAT(ptr) \
-        ROCPRIM_ATOMIC_LOAD("flat_load_dwordx4", "glc", "s_waitcnt vmcnt(0)", ptr)
-    #define ROCPRIM_ATOMIC_LOAD_SHARED(ptr) \
-        ROCPRIM_ATOMIC_LOAD("ds_read_b128", "", "s_waitcnt lgkmcnt(0)", ptr)
-    #define ROCPRIM_ATOMIC_LOAD_GLOBAL(ptr) \
-        ROCPRIM_ATOMIC_LOAD("global_load_dwordx4", "off glc", "s_waitcnt vmcnt(0)", ptr)
-#elif defined(__HIP_DEVICE_COMPILE__)
-    // Please submit an issue or pull request!
-    #error support for 128-bit atomics not implemented for current architecture
-#endif
+        #define ROCPRIM_ATOMIC_LOAD_GLOBAL(ptr) ROCPRIM_ATOMIC_LOAD_FLAT(ptr)
+    #elif ROCPRIM_TARGET_RDNA3 || ROCPRIM_TARGET_CDNA2 || ROCPRIM_TARGET_CDNA1 \
+        || ROCPRIM_TARGET_GCN5
+        #define ROCPRIM_ATOMIC_LOAD_FLAT(ptr) \
+            ROCPRIM_ATOMIC_LOAD("flat_load_dwordx4", "glc", "s_waitcnt vmcnt(0)", ptr)
+        #define ROCPRIM_ATOMIC_LOAD_SHARED(ptr) \
+            ROCPRIM_ATOMIC_LOAD("ds_read_b128", "", "s_waitcnt lgkmcnt(0)", ptr)
+        #define ROCPRIM_ATOMIC_LOAD_GLOBAL(ptr) \
+            ROCPRIM_ATOMIC_LOAD("global_load_dwordx4", "off glc", "s_waitcnt vmcnt(0)", ptr)
+    #elif defined(__HIP_DEVICE_COMPILE__)
+        // Please submit an issue or pull request!
+        #error support for 128-bit atomics not implemented for current architecture
+    #endif
 
-#ifdef __HIP_DEVICE_COMPILE__
-    #if !ROCPRIM_TARGET_SPIRV && defined(__has_builtin) \
-        && __has_builtin(__builtin_amdgcn_is_shared) && __has_builtin(__builtin_amdgcn_is_private)
+    #ifdef __HIP_DEVICE_COMPILE__
+        #if !ROCPRIM_TARGET_SPIRV && defined(__has_builtin) \
+            && __has_builtin(__builtin_amdgcn_is_shared)    \
+            && __has_builtin(__builtin_amdgcn_is_private)
 
         auto* ptr = (const __attribute__((address_space(0 /*flat*/))) __uint128_t*)address;
         if(__builtin_amdgcn_is_shared(ptr))
@@ -234,15 +235,15 @@ namespace detail
                 = (const __attribute__((address_space(1 /*global*/))) __uint128_t*)address;
             ROCPRIM_ATOMIC_LOAD_GLOBAL(global_ptr);
         }
-    #else
+        #else
         // SPIR-V does not like the address-space checks. For now
         // lets just do flat loading/storing.
         ROCPRIM_ATOMIC_LOAD_FLAT(address);
-    #endif
-#else
+        #endif
+    #else
         (void)address;
         result = 0;
-#endif
+    #endif
 
         return result;
 
@@ -251,6 +252,7 @@ namespace detail
 #undef ROCPRIM_ATOMIC_LOAD_SHARED
 #undef ROCPRIM_ATOMIC_LOAD_GLOBAL
     }
+#endif
 
     ROCPRIM_DEVICE ROCPRIM_INLINE
     void atomic_store(unsigned char* address, unsigned char value)
@@ -279,6 +281,8 @@ namespace detail
         __hip_atomic_store(address, value, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
     }
 
+// atomic_store for uint128 is not supported for unknown targets or spirv
+#if !defined(ROCPRIM_TARGET_UNKNOWN) && !defined(ROCPRIM_TARGET_SPIRV)
     ROCPRIM_DEVICE ROCPRIM_INLINE
     void atomic_store(__uint128_t* address, const __uint128_t value)
     {
@@ -355,6 +359,7 @@ namespace detail
 #undef ROCPRIM_ATOMIC_STORE_SHARED
 #undef ROCPRIM_ATOMIC_STORE_GLOBAL
     }
+#endif
 
     /// \brief Wait for all vector memory operations to complete
     ///

test/rocprim/test_intrinsics_atomic.cpp

@@ -1,4 +1,4 @@
-// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -29,6 +29,9 @@
 
 #include <cstdint>
 
+// The 128-bit atomic store and atomic load are disabled for these targets.
+#if !defined(ROCPRIM_TARGET_UNKNOWN) && !defined(ROCPRIM_TARGET_SPIRV)
+
 template<typename GetPtr>
 ROCPRIM_INLINE ROCPRIM_DEVICE
 void test_atomics(
@@ -183,3 +186,4 @@ TEST(RocprimAtomicTests, Flat128Bits)
             HIP_CHECK(hipGetLastError());
         });
 }
+#endif