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

Users/taosang/SWDEV-510994 - Refractor atomics header and tests (#902)

Browse Source 
* SWDEV-550626 - Refactor atomics header and tests

1. Introduce __HIP_ATOMIC_BACKWARD_COMPAT.
By default we define __HIP_ATOMIC_BACKWARD_COMPAT=1 to
let hip atomic functions maintain old assumptions. if
users want to adopt the new behavior, that is , by default
assume no-fine-grained no-remote-memory, then they can
define __HIP_ATOMIC_BACKWARD_COMPAT=0 and get the new
behaviour.

2. Use  __HIP_ATOMIC_BACKWARD_COMPAT_MEMORY to replace
original __HIP_FINE_GRAINED_MEMORY  in atomic header.
And apply __HIP_FINE_GRAINED_MEMORY onto all 
atomicXXX_system() functions to prevent failure on memory
allocated by hipHostMalloc().

3. Replace HIP_TEST_FINE_GRAINED_MEMORY with
HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY in hip-tests.

4. Fix negative test errors.
    Fix managed memory test error on memory order.
    some other minor changes.
    As a result  all originally disabled tests are enabled.

5. Add more atomics tests in some cases.

6. Reduce test time in each case.
     Reduce iteration number to 1 for tests that cost too much time.

8. Put common codes into hip_test_common.hh
This commit is contained in:
MachineTom
2025-09-25 10:58:59 -04:00
committed by GitHub
parent becb4646bd
commit 4a31affb76
23 changed files with 286 additions and 321 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h
+163 -65
View File
@@ -26,6 +26,16 @@ THE SOFTWARE.
#include "amd_device_functions.h"
#endif
#if !defined(__HIP_ATOMIC_BACKWARD_COMPAT)
#define __HIP_ATOMIC_BACKWARD_COMPAT 1
#endif
#if defined(__has_extension) && __has_extension(clang_atomic_attributes) && __HIP_ATOMIC_BACKWARD_COMPAT
#define __HIP_ATOMIC_BACKWARD_COMPAT_MEMORY [[clang::atomic(fine_grained_memory, remote_memory)]]
#else
#define __HIP_ATOMIC_BACKWARD_COMPAT_MEMORY
#endif
template <bool B, typename T, typename F> struct Cond_t;
template <typename T, typename F> struct Cond_t<true, T, F> {
@@ -106,8 +116,10 @@ __device__ inline unsigned short int atomicCAS(unsigned short int* address,
__device__ inline unsigned short int atomicCAS_system(unsigned short int* address,
unsigned short int compare,
unsigned short int val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
}
return compare;
}
@@ -118,8 +130,10 @@ __device__ inline int atomicCAS(int* address, int compare, int val) {
}
__device__ inline int atomicCAS_system(int* address, int compare, int val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
}
return compare;
}
@@ -132,8 +146,10 @@ __device__ inline unsigned int atomicCAS(unsigned int* address, unsigned int com
__device__ inline unsigned int atomicCAS_system(unsigned int* address, unsigned int compare,
unsigned int val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
}
return compare;
}
@@ -146,8 +162,10 @@ __device__ inline unsigned long atomicCAS(unsigned long* address, unsigned long
__device__ inline unsigned long atomicCAS_system(unsigned long* address, unsigned long compare,
unsigned long val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
}
return compare;
}
@@ -161,8 +179,10 @@ __device__ inline unsigned long long atomicCAS(unsigned long long* address,
__device__ inline unsigned long long atomicCAS_system(unsigned long long* address,
unsigned long long compare,
unsigned long long val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
}
return compare;
}
@@ -173,8 +193,10 @@ __device__ inline float atomicCAS(float* address, float compare, float val) {
}
__device__ inline float atomicCAS_system(float* address, float compare, float val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
}
return compare;
}
@@ -185,8 +207,10 @@ __device__ inline double atomicCAS(double* address, double compare, double val)
}
__device__ inline double atomicCAS_system(double* address, double compare, double val) {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
__hip_atomic_compare_exchange_strong(address, &compare, val, __ATOMIC_RELAXED, __ATOMIC_RELAXED,
__HIP_MEMORY_SCOPE_SYSTEM);
}
return compare;
}
@@ -195,7 +219,9 @@ __device__ inline int atomicAdd(int* address, int val) {
}
__device__ inline int atomicAdd_system(int* address, int val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicAdd(unsigned int* address, unsigned int val) {
@@ -203,7 +229,9 @@ __device__ inline unsigned int atomicAdd(unsigned int* address, unsigned int val
}
__device__ inline unsigned int atomicAdd_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicAdd(unsigned long* address, unsigned long val) {
@@ -211,7 +239,9 @@ __device__ inline unsigned long atomicAdd(unsigned long* address, unsigned long
}
__device__ inline unsigned long atomicAdd_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicAdd(unsigned long long* address,
@@ -221,27 +251,25 @@ __device__ inline unsigned long long atomicAdd(unsigned long long* address,
__device__ inline unsigned long long atomicAdd_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
#if defined(__has_extension) && __has_extension(clang_atomic_attributes)
#define __HIP_FINE_GRAINED_MEMORY [[clang::atomic(fine_grained_memory)]]
#else
#define __HIP_FINE_GRAINED_MEMORY
#endif
__device__ inline float atomicAdd(float* address, float val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicAdd(address, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
}
__device__ inline float atomicAdd_system(float* address, float val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
#if !defined(__HIPCC_RTC__)
@@ -253,14 +281,16 @@ __device__ inline double atomicAdd(double* address, double val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicAdd(address, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
}
__device__ inline double atomicAdd_system(double* address, double val) {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline int atomicSub(int* address, int val) {
@@ -268,7 +298,9 @@ __device__ inline int atomicSub(int* address, int val) {
}
__device__ inline int atomicSub_system(int* address, int val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicSub(unsigned int* address, unsigned int val) {
@@ -276,7 +308,9 @@ __device__ inline unsigned int atomicSub(unsigned int* address, unsigned int val
}
__device__ inline unsigned int atomicSub_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicSub(unsigned long* address, unsigned long val) {
@@ -284,7 +318,9 @@ __device__ inline unsigned long atomicSub(unsigned long* address, unsigned long
}
__device__ inline unsigned long atomicSub_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicSub(unsigned long long* address,
@@ -294,35 +330,41 @@ __device__ inline unsigned long long atomicSub(unsigned long long* address,
__device__ inline unsigned long long atomicSub_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline float atomicSub(float* address, float val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicAdd(address, -val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
}
__device__ inline float atomicSub_system(float* address, float val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline double atomicSub(double* address, double val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicAdd(address, -val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
}
__device__ inline double atomicSub_system(double* address, double val) {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(address, -val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline int atomicExch(int* address, int val) {
@@ -330,7 +372,9 @@ __device__ inline int atomicExch(int* address, int val) {
}
__device__ inline int atomicExch_system(int* address, int val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicExch(unsigned int* address, unsigned int val) {
@@ -338,7 +382,9 @@ __device__ inline unsigned int atomicExch(unsigned int* address, unsigned int va
}
__device__ inline unsigned int atomicExch_system(unsigned int* address, unsigned int val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicExch(unsigned long* address, unsigned long val) {
@@ -346,7 +392,9 @@ __device__ inline unsigned long atomicExch(unsigned long* address, unsigned long
}
__device__ inline unsigned long atomicExch_system(unsigned long* address, unsigned long val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicExch(unsigned long long* address,
@@ -356,7 +404,9 @@ __device__ inline unsigned long long atomicExch(unsigned long long* address,
__device__ inline unsigned long long atomicExch_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline float atomicExch(float* address, float val) {
@@ -364,7 +414,9 @@ __device__ inline float atomicExch(float* address, float val) {
}
__device__ inline float atomicExch_system(float* address, float val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline double atomicExch(double* address, double val) {
@@ -372,7 +424,9 @@ __device__ inline double atomicExch(double* address, double val) {
}
__device__ inline double atomicExch_system(double* address, double val) {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_exchange(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline int atomicMin(int* address, int val) {
@@ -380,7 +434,9 @@ __device__ inline int atomicMin(int* address, int val) {
}
__device__ inline int atomicMin_system(int* address, int val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicMin(unsigned int* address, unsigned int val) {
@@ -388,7 +444,9 @@ __device__ inline unsigned int atomicMin(unsigned int* address, unsigned int val
}
__device__ inline unsigned int atomicMin_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicMin(unsigned long* address, unsigned long val) {
@@ -396,7 +454,9 @@ __device__ inline unsigned long atomicMin(unsigned long* address, unsigned long
}
__device__ inline unsigned long atomicMin_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicMin(unsigned long long* address,
@@ -406,7 +466,9 @@ __device__ inline unsigned long long atomicMin(unsigned long long* address,
__device__ inline unsigned long long atomicMin_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline long long atomicMin(long long* address, long long val) {
@@ -414,14 +476,16 @@ __device__ inline long long atomicMin(long long* address, long long val) {
}
__device__ inline long long atomicMin_system(long long* address, long long val) {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline float atomicMin(float* addr, float val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMin(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
@@ -431,7 +495,7 @@ __device__ inline float atomicMin_system(float* addr, float val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMin(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
#endif
@@ -441,7 +505,7 @@ __device__ inline double atomicMin(double* addr, double val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMin(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
@@ -451,7 +515,7 @@ __device__ inline double atomicMin_system(double* addr, double val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMin(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
#endif
@@ -462,7 +526,9 @@ __device__ inline int atomicMax(int* address, int val) {
}
__device__ inline int atomicMax_system(int* address, int val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicMax(unsigned int* address, unsigned int val) {
@@ -470,7 +536,9 @@ __device__ inline unsigned int atomicMax(unsigned int* address, unsigned int val
}
__device__ inline unsigned int atomicMax_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicMax(unsigned long* address, unsigned long val) {
@@ -478,7 +546,9 @@ __device__ inline unsigned long atomicMax(unsigned long* address, unsigned long
}
__device__ inline unsigned long atomicMax_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicMax(unsigned long long* address,
@@ -488,21 +558,25 @@ __device__ inline unsigned long long atomicMax(unsigned long long* address,
__device__ inline unsigned long long atomicMax_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline long long atomicMax(long long* address, long long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
__device__ inline long long atomicMax_system(long long* address, long long val) {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline float atomicMax(float* addr, float val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMax(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
@@ -512,7 +586,7 @@ __device__ inline float atomicMax_system(float* addr, float val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMax(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
#endif
@@ -522,7 +596,7 @@ __device__ inline double atomicMax(double* addr, double val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMax(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
}
#endif
@@ -532,7 +606,7 @@ __device__ inline double atomicMax_system(double* addr, double val) {
#if defined(__AMDGCN_UNSAFE_FP_ATOMICS__)
return unsafeAtomicMax(addr, val);
#else
__HIP_FINE_GRAINED_MEMORY {
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(addr, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
#endif
@@ -551,7 +625,9 @@ __device__ inline int atomicAnd(int* address, int val) {
}
__device__ inline int atomicAnd_system(int* address, int val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicAnd(unsigned int* address, unsigned int val) {
@@ -559,7 +635,9 @@ __device__ inline unsigned int atomicAnd(unsigned int* address, unsigned int val
}
__device__ inline unsigned int atomicAnd_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicAnd(unsigned long* address, unsigned long val) {
@@ -567,7 +645,9 @@ __device__ inline unsigned long atomicAnd(unsigned long* address, unsigned long
}
__device__ inline unsigned long atomicAnd_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicAnd(unsigned long long* address,
@@ -577,7 +657,9 @@ __device__ inline unsigned long long atomicAnd(unsigned long long* address,
__device__ inline unsigned long long atomicAnd_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_and(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline int atomicOr(int* address, int val) {
@@ -585,7 +667,9 @@ __device__ inline int atomicOr(int* address, int val) {
}
__device__ inline int atomicOr_system(int* address, int val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicOr(unsigned int* address, unsigned int val) {
@@ -593,7 +677,9 @@ __device__ inline unsigned int atomicOr(unsigned int* address, unsigned int val)
}
__device__ inline unsigned int atomicOr_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicOr(unsigned long* address, unsigned long val) {
@@ -601,7 +687,9 @@ __device__ inline unsigned long atomicOr(unsigned long* address, unsigned long v
}
__device__ inline unsigned long atomicOr_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicOr(unsigned long long* address, unsigned long long val) {
@@ -610,7 +698,9 @@ __device__ inline unsigned long long atomicOr(unsigned long long* address, unsig
__device__ inline unsigned long long atomicOr_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_or(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline int atomicXor(int* address, int val) {
@@ -618,7 +708,9 @@ __device__ inline int atomicXor(int* address, int val) {
}
__device__ inline int atomicXor_system(int* address, int val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned int atomicXor(unsigned int* address, unsigned int val) {
@@ -626,7 +718,9 @@ __device__ inline unsigned int atomicXor(unsigned int* address, unsigned int val
}
__device__ inline unsigned int atomicXor_system(unsigned int* address, unsigned int val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long atomicXor(unsigned long* address, unsigned long val) {
@@ -634,7 +728,9 @@ __device__ inline unsigned long atomicXor(unsigned long* address, unsigned long
}
__device__ inline unsigned long atomicXor_system(unsigned long* address, unsigned long val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
__device__ inline unsigned long long atomicXor(unsigned long long* address,
@@ -644,5 +740,7 @@ __device__ inline unsigned long long atomicXor(unsigned long long* address,
__device__ inline unsigned long long atomicXor_system(unsigned long long* address,
unsigned long long val) {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
__HIP_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_xor(address, val, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SYSTEM);
}
}
projects/hip-tests/catch/hipTestMain/config/config_amd_linux
+10 -42
View File
@@ -74,19 +74,6 @@
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters",
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation",
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Positive_RangeValidation",
"=== SWDEV-435667: Below tests failing randomly in stress test on 01/12/23 ===",
"Unit_atomicExch_system_Positive_Peer_GPUs - int",
"Unit_atomicExch_system_Positive_Peer_GPUs - unsigned int",
"Unit_atomicExch_system_Positive_Peer_GPUs - unsigned long",
"Unit_atomicExch_system_Positive_Peer_GPUs - unsigned long long",
"Unit_atomicExch_system_Positive_Peer_GPUs - float",
"Unit_atomicExch_system_Positive_Peer_GPUs - double",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - int",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned int",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned long",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned long long",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - float",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - double",
"=== SWDEV-439004: Below tests failing randomly in CQE staging ===",
"Unit_hipGraphicsMapResources_Negative_Parameters",
"Unit_hipGraphicsSubResourceGetMappedArray_Negative_Parameters",
@@ -111,11 +98,10 @@
"Unit_Coalesced_Group_Tiled_Partition_Sync_Positive_Basic - uint16_t",
"Unit_Coalesced_Group_Tiled_Partition_Sync_Positive_Basic - uint32_t",
"=== SWDEV-444987 - Below tests fail in stress testing on 25/01/2023 ===",
"Unit_atomicAnd_Negative_Parameters_RTC",
"Unit_atomicOr_Negative_Parameters_RTC",
"Unit_atomicXor_Negative_Parameters_RTC",
"Unit_atomicMin_Negative_Parameters_RTC",
"Unit_atomicMax_Negative_Parameters_RTC",
"Unit_floatTM",
"Unit_TestMathFuncComplex",
"Unit_hipGraphAddMemcpyNodeToSymbol_Positive_Basic",
"Unit_hipStreamBeginCapture_Positive_Functional",
"Unit_Kernel_Launch_bounds_Negative_OutOfBounds",
"Unit_Kernel_Launch_bounds_Negative_Parameters_RTC",
"Unit_Device_sin_Accuracy_Positive - float",
@@ -577,11 +563,6 @@
"Unit_Device___longlong_as_double_Negative_RTC",
"Unit_Device___hiloint2double_Positive",
"Unit_Device___hiloint2double_Negative_RTC",
"Unit_atomicAdd_Negative_Parameters_RTC",
"Unit_atomicSub_Negative_Parameters_RTC",
"Unit_atomicInc_Negative_Parameters_RTC",
"Unit_atomicDec_Negative_Parameters_RTC",
"Unit_atomicCAS_Negative_Parameters_RTC",
"SWDEV-447384, SWDEV-447932: These tests fail in gfx1100, gfx1101 & gfx1102",
"Unit_hipFreeAsync_Negative_Parameters",
"SWDEV-445928: These tests fail in PSDB stress test on 09/02/2024",
@@ -589,24 +570,12 @@
"Unit_Device___float2half_rd_Accuracy_Limited_Positive",
"Unit_Device___float2half_ru_Accuracy_Limited_Positive",
"Unit_Device___float2half_rz_Accuracy_Limited_Positive",
"=== These tests fail on linux PSDB 21/11/24 ===",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - double",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - float",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - double",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - float",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Same_Address - double",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Same_Address - float",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Same_Address - double",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Same_Address - float",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Same_Address - double",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Same_Address - float",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Same_Address - double",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Same_Address - float",
"=== SWDEV-454316 : Below tests fail in stress test ===",
"Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address - float",
"Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address - double",
"Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address - float",
"Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address - double",
"Unit_hipGraphInstantiateWithFlags_StreamCaptureDeviceContextChg",
"=== SWDEV-457316 Below tests are disabled temporarily to avoid combined PSDB ===",
"Unit_hipGraphAddMemFreeNode_Negative_NotSupported",
"=== SWDEV-475482 - Disable tests to merge clr change ===",
"Unit_hipCreateTextureObject_LinearResource",
"Unit_hipCreateTextureObject_Pitch2DResource",
"=== SWDEV-511679 : Below tests fail in stress test ===",
"Unit_hipIpcOpenMemHandle_Negative_Open_In_Two_Contexts_Same_Device",
"Unit_hipIpcCloseMemHandle_Positive_Reference_Counting",
@@ -665,7 +634,6 @@
"Unit_Device___float2half_rd_SmallVals_Sanity_Positive",
"Unit_Device___float2half_ru_SmallVals_Sanity_Positive",
"Unit_Device___float2half_rz_SmallVals_Sanity_Positive",
"Unit_safeAtomicMin_Positive_SameAddress - float",
#endif
#if defined gfx1030
"=== SWDEV-445961: These tests hang in PSDB stress test on 09/02/2024 ===",
projects/hip-tests/catch/hipTestMain/config/config_amd_windows
+2 -50
View File
@@ -276,19 +276,6 @@
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters",
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation",
"Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Positive_RangeValidation",
"=== SWDEV-435667: Below tests failing randomly in stress test on 01/12/23 ===",
"Unit_atomicExch_system_Positive_Peer_GPUs - int",
"Unit_atomicExch_system_Positive_Peer_GPUs - unsigned int",
"Unit_atomicExch_system_Positive_Peer_GPUs - unsigned long",
"Unit_atomicExch_system_Positive_Peer_GPUs - unsigned long long",
"Unit_atomicExch_system_Positive_Peer_GPUs - float",
"Unit_atomicExch_system_Positive_Peer_GPUs - double",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - int",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned int",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned long",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned long long",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - float",
"Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - double",
"=== SWDEV-435667: Below tests failing randomly in stress test on 08/12/23 ===",
"Unit_hipMemPoolSetAccess_Negative_Parameters",
"SWDEV-438524: Below tests taking long time to run in stress test on 15/12/23 ===",
@@ -597,7 +584,6 @@
"Unit_hipModuleGetTexRef_Positive_Basic",
"Unit_Kernel_Launch_bounds_Negative_OutOfBounds",
"Unit_Kernel_Launch_bounds_Negative_Parameters_RTC",
"Unit_AtomicBuiltins_Negative_Parameters_RTC",
"Unit_hipMemcpy2D_H2D-D2D-D2H - int",
"Unit_hipMemcpy2D_H2D-D2D-D2H - float",
"Unit_hipMemcpy2D_H2D-D2D-D2H - double",
@@ -618,25 +604,7 @@
"Unit_hipMemset2DAsync_capturehipMemset2DAsync",
"Unit_hipOccupancyMaxPotBlkSizeVariableSMemWithFlags_Functional",
"Unit_hipDynamicShared",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - double",
"Unit_atomicExch_Positive - int",
"Unit_atomicExch_Positive - unsigned int",
"Unit_atomicExch_Positive - unsigned long",
"Unit_atomicExch_Positive - unsigned long long",
"Unit_atomicExch_Positive - float",
"Unit_atomicExch_Positive - double",
"Unit___hip_atomic_exchange_Positive_Wavefront - int",
"Unit___hip_atomic_exchange_Positive_Wavefront - unsigned int",
"Unit___hip_atomic_exchange_Positive_Wavefront - unsigned long",
"Unit___hip_atomic_exchange_Positive_Wavefront - unsigned long long",
"Unit___hip_atomic_exchange_Positive_Wavefront - float",
"Unit___hip_atomic_exchange_Positive_Wavefront - double",
"Unit___hip_atomic_exchange_Positive_Workgroup - int",
"Unit___hip_atomic_exchange_Positive_Workgroup - unsigned int",
"Unit___hip_atomic_exchange_Positive_Workgroup - unsigned long",
"Unit___hip_atomic_exchange_Positive_Workgroup - unsigned long long",
"Unit___hip_atomic_exchange_Positive_Workgroup - float",
"Unit___hip_atomic_exchange_Positive_Workgroup - double",
"Unit___syncthreads_Positive_Basic",
"Unit___syncthreads_count_Positive_Basic",
"Unit___syncthreads_and_Positive_Basic",
@@ -725,12 +693,7 @@
"Unit_Coalesced_Group_Tiled_Partition_Sync_Positive_Basic - uint8_t",
"Unit_Coalesced_Group_Tiled_Partition_Sync_Positive_Basic - uint16_t",
"Unit_Coalesced_Group_Tiled_Partition_Sync_Positive_Basic - uint32_t",
"Below tests failed in stress test of 25/01/24 ===",
"Unit_atomicAnd_Negative_Parameters_RTC",
"Unit_atomicOr_Negative_Parameters_RTC",
"Unit_atomicXor_Negative_Parameters_RTC",
"Unit_atomicMin_Negative_Parameters_RTC",
"Unit_atomicMax_Negative_Parameters_RTC",
"=== Below tests cause timeout in stress test of 09/02/24 ===",
"Unit_Device___half2half2_Accuracy_Positive",
"Unit_Device_make_half2_Accuracy_Positive",
@@ -924,11 +887,6 @@
"Unit_Device___dmul_rn_Accuracy_Positive",
"Unit_Device___ddiv_rn_Accuracy_Positive",
"Unit_Device___fma_rn_Accuracy_Positive",
"Unit_atomicAdd_Negative_Parameters_RTC",
"Unit_atomicSub_Negative_Parameters_RTC",
"Unit_atomicInc_Negative_Parameters_RTC",
"Unit_atomicDec_Negative_Parameters_RTC",
"Unit_atomicCAS_Negative_Parameters_RTC",
"SWDEV-450909: Test failed in stress testing",
"Unit_RTC_LinkDestroy_Default",
"=== SWDEV-457316 Below tests are disabled temporarily to avoid combined PSDB ===",
@@ -954,12 +912,6 @@
"Unit_Coalesced_Group_Tiled_Partition_Shfl_Down_Positive_Basic - unsigned long long",
"Unit_Coalesced_Group_Tiled_Partition_Shfl_Down_Positive_Basic - float",
"Unit_Coalesced_Group_Tiled_Partition_Shfl_Down_Positive_Basic - double",
"=== SWDEV-454316 : Below tests fail in stress test ===",
"Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address - float",
"Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address - double",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Same_Address - float",
"Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address - float",
"Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address - double",
"=== SWDEV-475482 - Disable tests to merge clr change",
"Unit_hipCreateTextureObject_LinearResource",
"Unit_hipCreateTextureObject_Pitch2DResource",
projects/hip-tests/catch/include/hip_test_common.hh
+34
View File
@@ -40,6 +40,16 @@ THE SOFTWARE.
#include <sys/resource.h>
#endif
#if !defined(__HIP_ATOMIC_BACKWARD_COMPAT)
#define __HIP_ATOMIC_BACKWARD_COMPAT 1
#endif
#if defined(__has_extension) && __has_extension(clang_atomic_attributes) && __HIP_ATOMIC_BACKWARD_COMPAT
#define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY [[clang::atomic(fine_grained_memory, remote_memory)]]
#else
#define HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY
#endif
#ifdef TEST_CLOCK_CYCLE
#define clock_function() clock64()
#else
@@ -388,6 +398,30 @@ inline bool isP2PSupported(int& d1, int& d2) {
return supported;
}
inline bool checkConcurrentKernels(int num_devices) {
for (auto i = 0; i < num_devices; ++i) {
HIP_CHECK(hipSetDevice(i));
int concurrent_kernels = 0;
HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i));
if (!concurrent_kernels) {
return false;
}
}
if (num_devices > 1) {
HIP_CHECK(hipSetDevice(0));
}
return true;
}
inline bool isXnackOn() {
hipDeviceProp_t prop;
int device = 0;
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
return gfxName.find("xnack+") != std::string::npos;
}
inline bool areWarpMatchFunctionsSupported() {
int matchFunctionsSupported = 1;
#if HT_NVIDIA
projects/hip-tests/catch/unit/atomics/arithmetic_common.hh
+9 -36
View File
@@ -29,12 +29,6 @@ THE SOFTWARE.
namespace cg = cooperative_groups;
#if defined(__has_extension) && __has_extension(clang_atomic_attributes)
#define HIP_TEST_FINE_GRAINED_MEMORY [[clang::atomic(fine_grained_memory)]]
#else
#define HIP_TEST_FINE_GRAINED_MEMORY
#endif
// Atomic operations for which the tests in this file apply for
enum class AtomicOperation {
kAdd = 0,
@@ -141,7 +135,7 @@ __device__ TestType PerformAtomicOperation(TestType* const mem, const LinearAllo
return CASAtomicAddSystem(mem, val);
} else if constexpr (operation == AtomicOperation::kBuiltinAdd) {
if (std::is_floating_point_v<TestType> && allocType == LinearAllocs::hipHostMalloc) {
HIP_TEST_FINE_GRAINED_MEMORY {
HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_add(mem, val, __ATOMIC_RELAXED, memory_scope);
}
} else {
@@ -450,8 +444,8 @@ void TestCore(const TestParams& p) {
for (auto j = 0u; j < p.kernel_count; ++j) {
const auto& stream = streams[i * p.kernel_count + j].stream();
const auto old_vals = old_vals_devs[i].ptr() + j * p.ThreadCount();
LaunchKernel<TestType, operation, use_shared_mem, memory_scope>(p, stream, mem_devs[i].ptr(),
old_vals);
LaunchKernel<TestType, operation, use_shared_mem, memory_scope>(p, stream,
mem_devs[i].ptr(), old_vals);
}
}
@@ -479,12 +473,10 @@ void TestCore(const TestParams& p) {
Verify<TestType, operation>(p, res_vals, old_vals);
}
inline dim3 GenerateThreadDimensions() { return dim3(16); }
inline dim3 GenerateThreadDimensions() { return dim3(1024); }
inline dim3 GenerateBlockDimensions() {
int sm_count = 0;
HIP_CHECK(hipDeviceGetAttribute(&sm_count, hipDeviceAttributeMultiprocessorCount, 0));
return dim3(sm_count);
return dim3(8);
}
// Configures and creates the TestCore for a single device, and a single kernel launch
@@ -581,28 +573,9 @@ void MultipleDeviceMultipleKernelAndHostTest(const unsigned int num_devices,
}
}
CHECK_P2P_SUPPORT
if (kernel_count > 1) {
for (auto i = 0u; i < num_devices; ++i) {
int canAccess = 0;
for (auto j = 0u; j < num_devices; ++j) {
if (i != j) {
HIP_CHECK(hipDeviceCanAccessPeer(&canAccess, i, j));
if(canAccess == 0) {
std::string msg = "P2P access check failed between dev1:" + std::to_string(i) + ",dev2:" + std::to_string(j);
HipTest::HIP_SKIP_TEST(msg.c_str());
return;
}
}
}
int concurrent_kernels = 0;
HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i));
if (!concurrent_kernels) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
}
if (!HipTest::checkConcurrentKernels(num_devices)) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
TestParams params;
@@ -615,7 +588,7 @@ void MultipleDeviceMultipleKernelAndHostTest(const unsigned int num_devices,
params.host_thread_count = host_thread_count;
using LA = LinearAllocs;
for (const auto alloc_type : {LA::hipMalloc}) {
for (const auto alloc_type : {LA::hipMalloc, LA::hipHostMalloc}) {
params.alloc_type = alloc_type;
DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
TestCore<TestType, operation, false, __HIP_MEMORY_SCOPE_SYSTEM>(params);
projects/hip-tests/catch/unit/atomics/atomicAdd_system.cc
+3 -3
View File
@@ -60,7 +60,7 @@ TEMPLATE_TEST_CASE("Unit_atomicAdd_system_Positive_Peer_GPUs", "", int, unsigned
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kAddSystem>(
2, 2, 1, sizeof(TestType));
@@ -109,7 +109,7 @@ TEMPLATE_TEST_CASE("Unit_atomicAdd_system_Positive_Host_And_GPU", "", int, unsig
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kAddSystem>(
1, 1, 1, sizeof(TestType), 4);
@@ -158,7 +158,7 @@ TEMPLATE_TEST_CASE("Unit_atomicAdd_system_Positive_Host_And_Peer_GPUs", "", int,
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kAddSystem>(
2, 2, 1, sizeof(TestType), 4);
projects/hip-tests/catch/unit/atomics/atomicAnd.cc
+1 -1
View File
@@ -207,7 +207,7 @@ TEST_CASE("Unit_atomicAnd_Negative_Parameters_RTC") {
HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
int error_count{0};
// Please check the content of negative_kernels_rtc.hh
int expected_error_count{9};
int expected_error_count{10};
std::string error_message{"error:"};
size_t n_pos = log.find(error_message, 0);
projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc
+3 -3
View File
@@ -47,7 +47,7 @@ THE SOFTWARE.
*/
TEMPLATE_TEST_CASE("Unit_atomicAnd_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
unsigned long, unsigned long long) {
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAndSystem>(
2, 2, 1, sizeof(TestType));
@@ -73,7 +73,7 @@ TEMPLATE_TEST_CASE("Unit_atomicAnd_system_Positive_Peer_GPUs_Adjacent_Addresses"
int warp_size = 0;
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Adjacent address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAndSystem>(
2, 2, warp_size, sizeof(TestType));
@@ -100,7 +100,7 @@ TEMPLATE_TEST_CASE("Unit_atomicAnd_system_Positive_Peer_GPUs_Scattered_Addresses
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Scattered address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAndSystem>(
2, 2, warp_size, cache_line_size);
projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc
+3 -3
View File
@@ -66,7 +66,7 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Peer_GPUs", "", int, unsigned
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kCASAddSystem>(
2, 2, 1, sizeof(TestType));
@@ -116,7 +116,7 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Host_And_GPU", "", int, unsig
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kCASAddSystem>(
1, 1, 1, sizeof(TestType), 4);
@@ -166,7 +166,7 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Host_And_Peer_GPUs", "", int,
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kCASAddSystem>(
2, 2, 1, sizeof(TestType), 4);
projects/hip-tests/catch/unit/atomics/atomicExch_common.hh
+6 -26
View File
@@ -217,6 +217,7 @@ class AtomicExchCRTP {
const auto shared_mem_size = use_shared_mem ? mem_alloc_size : 0u;
for (auto i = 0u; i < p.num_devices; ++i) {
HIP_CHECK(hipSetDevice(i));
const auto device_offset = i * p.kernel_count * thread_count;
for (auto j = 0u; j < p.kernel_count; ++j) {
const auto& stream = streams[i * p.kernel_count + j].stream();
@@ -301,12 +302,10 @@ class AtomicExch
}
};
inline dim3 GenerateAtomicExchThreadDimensions() { return GENERATE(dim3(16), dim3(1024)); }
inline dim3 GenerateAtomicExchThreadDimensions() { return dim3(1024); }
inline dim3 GenerateAtomicExchBlockDimensions() {
int sm_count = 0;
HIP_CHECK(hipDeviceGetAttribute(&sm_count, hipDeviceAttributeMultiprocessorCount, 0));
return GENERATE_COPY(dim3(sm_count), dim3(sm_count + sm_count / 2));
return dim3(8);
}
template <typename TestType, AtomicScopes scope, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
@@ -395,28 +394,9 @@ void AtomicExchMultipleDeviceMultipleKernelAndHostTest(const unsigned int num_de
}
}
CHECK_P2P_SUPPORT
if (kernel_count > 1) {
for (auto i = 0u; i < num_devices; ++i) {
int canAccess = 0;
for (auto j = 0u; j < num_devices; ++j) {
if (i != j) {
HIP_CHECK(hipDeviceCanAccessPeer(&canAccess, i, j));
if(canAccess == 0) {
std::string msg = "P2P access check failed between dev1:" + std::to_string(i) + ",dev2:" + std::to_string(j);
HipTest::HIP_SKIP_TEST(msg.c_str());
return;
}
}
}
int concurrent_kernels = 0;
HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i));
if (!concurrent_kernels) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
}
if (!HipTest::checkConcurrentKernels(num_devices)) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
AtomicExchParams params;
projects/hip-tests/catch/unit/atomics/atomicExch_system.cc
+3 -3
View File
@@ -65,7 +65,7 @@ TEMPLATE_TEST_CASE("Unit_atomicExch_system_Positive_Peer_GPUs", "", int, unsigne
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(2, 2, 1, sizeof(TestType));
}
@@ -119,7 +119,7 @@ TEMPLATE_TEST_CASE("Unit_atomicExch_system_Positive_Host_And_GPU", "", int, unsi
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(1, 1, 1, sizeof(TestType), 4);
}
@@ -174,7 +174,7 @@ TEMPLATE_TEST_CASE("Unit_atomicExch_system_Positive_Host_And_Peer_GPUs", "", int
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
AtomicExchMultipleDeviceMultipleKernelAndHostTest<TestType>(2, 2, 1, sizeof(TestType), 4);
}
projects/hip-tests/catch/unit/atomics/atomicMax.cc
+1 -1
View File
@@ -207,7 +207,7 @@ TEST_CASE("Unit_atomicMax_Negative_Parameters_RTC") {
HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
int error_count{0};
// Please check the content of negative_kernels_rtc.hh
int expected_error_count{8};
int expected_error_count{7};
std::string error_message{"error:"};
size_t n_pos = log.find(error_message, 0);
projects/hip-tests/catch/unit/atomics/atomicMax_system.cc
+3 -3
View File
@@ -52,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address", "",
TEMPLATE_TEST_CASE("Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
unsigned long, unsigned long long) {
#endif
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MinMax::MultipleDeviceMultipleKernelTest<TestType, MinMax::AtomicOperation::kMaxSystem>(
2, 2, 1, sizeof(TestType));
@@ -83,7 +83,7 @@ TEMPLATE_TEST_CASE("Unit_atomicMax_system_Positive_Peer_GPUs_Adjacent_Addresses"
int warp_size = 0;
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Adjacent address " << current) {
MinMax::MultipleDeviceMultipleKernelTest<TestType, MinMax::AtomicOperation::kMaxSystem>(
2, 2, warp_size, sizeof(TestType));
@@ -115,7 +115,7 @@ TEMPLATE_TEST_CASE("Unit_atomicMax_system_Positive_Peer_GPUs_Scattered_Addresses
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Scattered address " << current) {
MinMax::MultipleDeviceMultipleKernelTest<TestType, MinMax::AtomicOperation::kMaxSystem>(
2, 2, warp_size, cache_line_size);
projects/hip-tests/catch/unit/atomics/atomicMin.cc
+1 -1
View File
@@ -207,7 +207,7 @@ TEST_CASE("Unit_atomicMin_Negative_Parameters_RTC") {
HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
int error_count{0};
// Please check the content of negative_kernels_rtc.hh
int expected_error_count{8};
int expected_error_count{7};
std::string error_message{"error:"};
size_t n_pos = log.find(error_message, 0);
projects/hip-tests/catch/unit/atomics/atomicMin_system.cc
+3 -3
View File
@@ -52,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address", "",
TEMPLATE_TEST_CASE("Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
unsigned long, unsigned long long) {
#endif
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MinMax::MultipleDeviceMultipleKernelTest<TestType, MinMax::AtomicOperation::kMinSystem>(
2, 2, 1, sizeof(TestType));
@@ -83,7 +83,7 @@ TEMPLATE_TEST_CASE("Unit_atomicMin_system_Positive_Peer_GPUs_Adjacent_Addresses"
int warp_size = 0;
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Adjacent address " << current) {
MinMax::MultipleDeviceMultipleKernelTest<TestType, MinMax::AtomicOperation::kMinSystem>(
2, 2, warp_size, sizeof(TestType));
@@ -115,7 +115,7 @@ TEMPLATE_TEST_CASE("Unit_atomicMin_system_Positive_Peer_GPUs_Scattered_Addresses
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Scattered address " << current) {
MinMax::MultipleDeviceMultipleKernelTest<TestType, MinMax::AtomicOperation::kMinSystem>(
2, 2, warp_size, cache_line_size);
projects/hip-tests/catch/unit/atomics/atomicOr.cc
+1 -1
View File
@@ -207,7 +207,7 @@ TEST_CASE("Unit_atomicOr_Negative_Parameters_RTC") {
HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
int error_count{0};
// Please check the content of negative_kernels_rtc.hh
int expected_error_count{9};
int expected_error_count{10};
std::string error_message{"error:"};
size_t n_pos = log.find(error_message, 0);
projects/hip-tests/catch/unit/atomics/atomicOr_system.cc
+3 -3
View File
@@ -47,7 +47,7 @@ THE SOFTWARE.
*/
TEMPLATE_TEST_CASE("Unit_atomicOr_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
unsigned long, unsigned long long) {
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOrSystem>(
2, 2, 1, sizeof(TestType));
@@ -73,7 +73,7 @@ TEMPLATE_TEST_CASE("Unit_atomicOr_system_Positive_Peer_GPUs_Adjacent_Addresses",
int warp_size = 0;
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Adjacent address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOrSystem>(
2, 2, warp_size, sizeof(TestType));
@@ -100,7 +100,7 @@ TEMPLATE_TEST_CASE("Unit_atomicOr_system_Positive_Peer_GPUs_Scattered_Addresses"
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Scattered address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOrSystem>(
2, 2, warp_size, cache_line_size);
projects/hip-tests/catch/unit/atomics/atomicSub_system.cc
+3 -3
View File
@@ -60,7 +60,7 @@ TEMPLATE_TEST_CASE("Unit_atomicSub_system_Positive_Peer_GPUs", "", int, unsigned
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kSubSystem>(
2, 2, 1, sizeof(TestType));
@@ -109,7 +109,7 @@ TEMPLATE_TEST_CASE("Unit_atomicSub_system_Positive_Host_And_GPU", "", int, unsig
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kSubSystem>(
1, 1, 1, sizeof(TestType), 4);
@@ -158,7 +158,7 @@ TEMPLATE_TEST_CASE("Unit_atomicSub_system_Positive_Host_And_Peer_GPUs", "", int,
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
MultipleDeviceMultipleKernelAndHostTest<TestType, AtomicOperation::kSubSystem>(
2, 2, 1, sizeof(TestType), 4);
projects/hip-tests/catch/unit/atomics/atomicXor.cc
+1 -1
View File
@@ -206,7 +206,7 @@ TEST_CASE("Unit_atomicXor_Negative_Parameters_RTC") {
HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
int error_count{0};
// Please check the content of negative_kernels_rtc.hh
int expected_error_count{9};
int expected_error_count{10};
std::string error_message{"error:"};
size_t n_pos = log.find(error_message, 0);
projects/hip-tests/catch/unit/atomics/atomicXor_system.cc
+3 -3
View File
@@ -47,7 +47,7 @@ THE SOFTWARE.
*/
TEMPLATE_TEST_CASE("Unit_atomicXor_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
unsigned long, unsigned long long) {
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Same address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXorSystem>(
2, 2, 1, sizeof(TestType));
@@ -73,7 +73,7 @@ TEMPLATE_TEST_CASE("Unit_atomicXor_system_Positive_Peer_GPUs_Adjacent_Addresses"
int warp_size = 0;
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Adjacent address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXorSystem>(
2, 2, warp_size, sizeof(TestType));
@@ -100,7 +100,7 @@ TEMPLATE_TEST_CASE("Unit_atomicXor_system_Positive_Peer_GPUs_Scattered_Addresses
HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
const auto cache_line_size = 128u;
for (auto current = 0; current < cmd_options.iterations; ++current) {
for (auto current = 0; current < 1; ++current) {
DYNAMIC_SECTION("Scattered address " << current) {
Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXorSystem>(
2, 2, warp_size, cache_line_size);
projects/hip-tests/catch/unit/atomics/bitwise_common.hh
+8 -27
View File
@@ -276,8 +276,8 @@ void TestCore(const TestParams& p) {
for (auto j = 0u; j < p.kernel_count; ++j) {
const auto& stream = streams[i * p.kernel_count + j].stream();
const auto old_vals = old_vals_devs[i].ptr() + j * p.ThreadCount();
LaunchKernel<TestType, operation, use_shared_mem, memory_scope>(p, stream, mem_devs[i].ptr(),
old_vals);
LaunchKernel<TestType, operation, use_shared_mem, memory_scope>(p, stream,
mem_devs[i].ptr(), old_vals);
}
}
for (auto i = 0; i < p.num_devices; ++i) {
@@ -296,12 +296,10 @@ void TestCore(const TestParams& p) {
Verify<TestType, operation>(p, res_vals, old_vals);
}
inline dim3 GenerateThreadDimensions() { return GENERATE(dim3(16), dim3(1024)); }
inline dim3 GenerateThreadDimensions() { return dim3(1024); }
inline dim3 GenerateBlockDimensions() {
int sm_count = 0;
HIP_CHECK(hipDeviceGetAttribute(&sm_count, hipDeviceAttributeMultiprocessorCount, 0));
return GENERATE_COPY(dim3(sm_count));
return dim3(8);
}
template <typename TestType, AtomicOperation operation, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
@@ -396,26 +394,9 @@ void MultipleDeviceMultipleKernelTest(const unsigned int num_devices,
}
}
if (kernel_count > 1) {
for (auto i = 0u; i < num_devices; ++i) {
int canAccess = 0;
for (auto j = 0u; j < num_devices; ++j) {
if (i != j) {
HIP_CHECK(hipDeviceCanAccessPeer(&canAccess, i, j));
if(canAccess == 0) {
std::string msg = "P2P access check failed between dev1:" + std::to_string(i) + ",dev2:" + std::to_string(j);
HipTest::HIP_SKIP_TEST(msg.c_str());
return;
}
}
}
int concurrent_kernels = 0;
HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i));
if (!concurrent_kernels) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
}
if (!HipTest::checkConcurrentKernels(num_devices)) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
TestParams params;
@@ -427,7 +408,7 @@ void MultipleDeviceMultipleKernelTest(const unsigned int num_devices,
params.pitch = pitch;
using LA = LinearAllocs;
for (const auto alloc_type : {LA::hipMalloc}) {
for (const auto alloc_type : {LA::hipMalloc, LA::hipHostMalloc}) {
params.alloc_type = alloc_type;
DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
TestCore<TestType, operation, false, __HIP_MEMORY_SCOPE_SYSTEM>(params);
projects/hip-tests/catch/unit/atomics/memory_order_common.hh
+12 -4
View File
@@ -234,8 +234,12 @@ template <BuiltinAtomicOperation operation, int memory_order, int memory_scope>
}
template <BuiltinAtomicOperation operation, int memory_order> void SystemTest() {
HipTest::HIP_SKIP_TEST("Skip system scope tests due to random failures!!");
return;
// The test need xnack on to make managed memory be host coherent.
// If xnack is off, managed memory has coarse grained access.
if (!HipTest::isXnackOn()) {
std::cerr << "GPU is not xnack enabled hence skipping system test\n";
return;
}
std::thread host_thread;
LinearAllocGuard<int> flag(LinearAllocs::hipMallocManaged, sizeof(int));
@@ -426,8 +430,12 @@ template <BuiltinAtomicOperation operation, int memory_scope> void Test() {
}
template <BuiltinAtomicOperation operation> void SystemTest() {
HipTest::HIP_SKIP_TEST("Skip system scope tests due to random failures!!");
return;
// The test need xnack on to make managed memory be host coherent.
// If xnack is off, managed memory has coarse grained access.
if (!HipTest::isXnackOn()) {
std::cerr << "GPU is not xnack enabled hence skipping system test\n";
return;
}
std::thread host_producer, host_consumer;
LinearAllocGuard<int> counter1(LinearAllocs::hipMallocManaged, sizeof(int));
projects/hip-tests/catch/unit/atomics/min_max_common.hh
+10 -39
View File
@@ -29,12 +29,6 @@ THE SOFTWARE.
namespace cg = cooperative_groups;
#if defined(__has_extension) && __has_extension(clang_atomic_attributes)
#define HIP_TEST_FINE_GRAINED_MEMORY [[clang::atomic(fine_grained_memory)]]
#else
#define HIP_TEST_FINE_GRAINED_MEMORY
#endif
namespace MinMax {
enum class AtomicOperation {
kMin = 0,
@@ -88,14 +82,14 @@ __device__ TestType PerformAtomicOperation(TestType* const mem, const LinearAllo
return safeAtomicMax(mem, val);
} else if constexpr (operation == AtomicOperation::kBuiltinMin) {
if (std::is_floating_point_v<TestType> && allocType == LinearAllocs::hipHostMalloc)
HIP_TEST_FINE_GRAINED_MEMORY {
HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_min(mem, val, __ATOMIC_RELAXED, memory_scope);
} else {
return __hip_atomic_fetch_min(mem, val, __ATOMIC_RELAXED, memory_scope);
}
} else if constexpr (operation == AtomicOperation::kBuiltinMax) {
if (std::is_floating_point_v<TestType> && allocType == LinearAllocs::hipHostMalloc)
HIP_TEST_FINE_GRAINED_MEMORY {
HIP_TEST_ATOMIC_BACKWARD_COMPAT_MEMORY {
return __hip_atomic_fetch_max(mem, val, __ATOMIC_RELAXED, memory_scope);
} else {
return __hip_atomic_fetch_max(mem, val, __ATOMIC_RELAXED, memory_scope);
@@ -306,8 +300,8 @@ void TestCore(const TestParams& p) {
for (auto j = 0u; j < p.kernel_count; ++j) {
const auto& stream = streams[i * p.kernel_count + j].stream();
const auto old_vals = old_vals_devs[i].ptr() + j * p.ThreadCount();
LaunchKernel<TestType, operation, use_shared_mem, memory_scope>(p, stream, mem_devs[i].ptr(),
old_vals);
LaunchKernel<TestType, operation, use_shared_mem, memory_scope>(p, stream,
mem_devs[i].ptr(), old_vals);
}
}
for (auto i = 0; i < p.num_devices; ++i) {
@@ -326,12 +320,10 @@ void TestCore(const TestParams& p) {
Verify<TestType, operation>(p, res_vals, old_vals);
}
inline dim3 GenerateThreadDimensions() { return GENERATE(dim3(16), dim3(1024)); }
inline dim3 GenerateThreadDimensions() { return dim3(1024); }
inline dim3 GenerateBlockDimensions() {
int sm_count = 0;
HIP_CHECK(hipDeviceGetAttribute(&sm_count, hipDeviceAttributeMultiprocessorCount, 0));
return dim3(sm_count);
return dim3(8);
}
template <typename TestType, AtomicOperation operation, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
@@ -422,31 +414,10 @@ void MultipleDeviceMultipleKernelTest(const unsigned int num_devices,
return;
}
}
CHECK_P2P_SUPPORT
if (kernel_count > 1) {
for (auto i = 0u; i < num_devices; ++i) {
int canAccess = 0;
for (auto j = 0u; j < num_devices; ++j) {
if (i != j) {
HIP_CHECK(hipDeviceCanAccessPeer(&canAccess, i, j));
if(canAccess == 0) {
std::string msg = "P2P access check failed between dev1:" + std::to_string(i) + ",dev2:" + std::to_string(j);
HipTest::HIP_SKIP_TEST(msg.c_str());
return;
}
}
}
int concurrent_kernels = 0;
HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i));
if (!concurrent_kernels) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
}
if (!HipTest::checkConcurrentKernels(num_devices)) {
HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
return;
}
TestParams params;
params.num_devices = num_devices;
params.kernel_count = kernel_count;
@@ -456,7 +427,7 @@ void MultipleDeviceMultipleKernelTest(const unsigned int num_devices,
params.pitch = pitch;
using LA = LinearAllocs;
for (const auto alloc_type : {LA::hipMalloc}) {
for (const auto alloc_type : {LA::hipMalloc, LA::hipHostMalloc}) {
params.alloc_type = alloc_type;
DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
TestCore<TestType, operation, false, __HIP_MEMORY_SCOPE_SYSTEM>(params);