From 4a31affb76b9c90b2a3c31e3933088e80cb84ecc Mon Sep 17 00:00:00 2001 From: MachineTom Date: Thu, 25 Sep 2025 10:58:59 -0400 Subject: [PATCH] Users/taosang/SWDEV-510994 - Refractor atomics header and tests (#902) * 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 --- .../include/hip/amd_detail/amd_hip_atomic.h | 228 +++++++++++++----- .../catch/hipTestMain/config/config_amd_linux | 52 +--- .../hipTestMain/config/config_amd_windows | 52 +--- .../catch/include/hip_test_common.hh | 34 +++ .../catch/unit/atomics/arithmetic_common.hh | 45 +--- .../catch/unit/atomics/atomicAdd_system.cc | 6 +- .../hip-tests/catch/unit/atomics/atomicAnd.cc | 2 +- .../catch/unit/atomics/atomicAnd_system.cc | 6 +- .../catch/unit/atomics/atomicCAS_system.cc | 6 +- .../catch/unit/atomics/atomicExch_common.hh | 32 +-- .../catch/unit/atomics/atomicExch_system.cc | 6 +- .../hip-tests/catch/unit/atomics/atomicMax.cc | 2 +- .../catch/unit/atomics/atomicMax_system.cc | 6 +- .../hip-tests/catch/unit/atomics/atomicMin.cc | 2 +- .../catch/unit/atomics/atomicMin_system.cc | 6 +- .../hip-tests/catch/unit/atomics/atomicOr.cc | 2 +- .../catch/unit/atomics/atomicOr_system.cc | 6 +- .../catch/unit/atomics/atomicSub_system.cc | 6 +- .../hip-tests/catch/unit/atomics/atomicXor.cc | 2 +- .../catch/unit/atomics/atomicXor_system.cc | 6 +- .../catch/unit/atomics/bitwise_common.hh | 35 +-- .../catch/unit/atomics/memory_order_common.hh | 16 +- .../catch/unit/atomics/min_max_common.hh | 49 +--- 23 files changed, 286 insertions(+), 321 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h index 299a79765c..d9b79faeb9 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_atomic.h @@ -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 struct Cond_t; template struct Cond_t { @@ -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); + } } diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 4c48ed1076..bc1602a1a7 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -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 ===", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 847d48ac7f..04a8bc3fe5 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -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", diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index 78fb6ec732..7072943081 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -40,6 +40,16 @@ THE SOFTWARE. #include #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 diff --git a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh index 2b4db75dd8..da5f43b3c7 100644 --- a/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh +++ b/projects/hip-tests/catch/unit/atomics/arithmetic_common.hh @@ -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 && 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(p, stream, mem_devs[i].ptr(), - old_vals); + LaunchKernel(p, stream, + mem_devs[i].ptr(), old_vals); } } @@ -479,12 +473,10 @@ void TestCore(const TestParams& p) { Verify(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(params); diff --git a/projects/hip-tests/catch/unit/atomics/atomicAdd_system.cc b/projects/hip-tests/catch/unit/atomics/atomicAdd_system.cc index 3520ff1472..e42a0edbaf 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicAdd_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicAdd_system.cc @@ -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( 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( 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( 2, 2, 1, sizeof(TestType), 4); diff --git a/projects/hip-tests/catch/unit/atomics/atomicAnd.cc b/projects/hip-tests/catch/unit/atomics/atomicAnd.cc index 94213e7548..cec5135086 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicAnd.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicAnd.cc @@ -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); diff --git a/projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc b/projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc index 2c89f504ce..a42eb974a8 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc @@ -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( 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( 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( 2, 2, warp_size, cache_line_size); diff --git a/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc b/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc index 872bfa9dd7..326ba41557 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc @@ -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( 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( 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( 2, 2, 1, sizeof(TestType), 4); diff --git a/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh b/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh index ba3105de63..bfe3c9d11a 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh +++ b/projects/hip-tests/catch/unit/atomics/atomicExch_common.hh @@ -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 @@ -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; diff --git a/projects/hip-tests/catch/unit/atomics/atomicExch_system.cc b/projects/hip-tests/catch/unit/atomics/atomicExch_system.cc index a96fe061ec..eeb19c3661 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicExch_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicExch_system.cc @@ -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(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(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(2, 2, 1, sizeof(TestType), 4); } diff --git a/projects/hip-tests/catch/unit/atomics/atomicMax.cc b/projects/hip-tests/catch/unit/atomics/atomicMax.cc index d390fec698..dbc0bacd63 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicMax.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicMax.cc @@ -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); diff --git a/projects/hip-tests/catch/unit/atomics/atomicMax_system.cc b/projects/hip-tests/catch/unit/atomics/atomicMax_system.cc index 5af7665f9c..8484a3ebd0 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicMax_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicMax_system.cc @@ -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( 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( 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( 2, 2, warp_size, cache_line_size); diff --git a/projects/hip-tests/catch/unit/atomics/atomicMin.cc b/projects/hip-tests/catch/unit/atomics/atomicMin.cc index d3478b398a..3165ac2b67 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicMin.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicMin.cc @@ -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); diff --git a/projects/hip-tests/catch/unit/atomics/atomicMin_system.cc b/projects/hip-tests/catch/unit/atomics/atomicMin_system.cc index a0f4592923..9d7f646b36 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicMin_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicMin_system.cc @@ -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( 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( 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( 2, 2, warp_size, cache_line_size); diff --git a/projects/hip-tests/catch/unit/atomics/atomicOr.cc b/projects/hip-tests/catch/unit/atomics/atomicOr.cc index e01efa9b53..98fda5007d 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicOr.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicOr.cc @@ -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); diff --git a/projects/hip-tests/catch/unit/atomics/atomicOr_system.cc b/projects/hip-tests/catch/unit/atomics/atomicOr_system.cc index 6a9a4d1e35..ba0906560c 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicOr_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicOr_system.cc @@ -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( 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( 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( 2, 2, warp_size, cache_line_size); diff --git a/projects/hip-tests/catch/unit/atomics/atomicSub_system.cc b/projects/hip-tests/catch/unit/atomics/atomicSub_system.cc index f02bf6ea80..a5a1111b52 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicSub_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicSub_system.cc @@ -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( 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( 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( 2, 2, 1, sizeof(TestType), 4); diff --git a/projects/hip-tests/catch/unit/atomics/atomicXor.cc b/projects/hip-tests/catch/unit/atomics/atomicXor.cc index 9bce78886a..f26c6dd5b0 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicXor.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicXor.cc @@ -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); diff --git a/projects/hip-tests/catch/unit/atomics/atomicXor_system.cc b/projects/hip-tests/catch/unit/atomics/atomicXor_system.cc index 4c70d92652..0d3523b382 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicXor_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicXor_system.cc @@ -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( 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( 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( 2, 2, warp_size, cache_line_size); diff --git a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh index 9c7bf0f5da..23e6d9cf78 100644 --- a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh +++ b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh @@ -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(p, stream, mem_devs[i].ptr(), - old_vals); + LaunchKernel(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(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 @@ -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(params); diff --git a/projects/hip-tests/catch/unit/atomics/memory_order_common.hh b/projects/hip-tests/catch/unit/atomics/memory_order_common.hh index 112158782a..8251f8e790 100644 --- a/projects/hip-tests/catch/unit/atomics/memory_order_common.hh +++ b/projects/hip-tests/catch/unit/atomics/memory_order_common.hh @@ -234,8 +234,12 @@ template } template 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 flag(LinearAllocs::hipMallocManaged, sizeof(int)); @@ -426,8 +430,12 @@ template void Test() { } template 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 counter1(LinearAllocs::hipMallocManaged, sizeof(int)); diff --git a/projects/hip-tests/catch/unit/atomics/min_max_common.hh b/projects/hip-tests/catch/unit/atomics/min_max_common.hh index 356030a7e8..3df6658824 100644 --- a/projects/hip-tests/catch/unit/atomics/min_max_common.hh +++ b/projects/hip-tests/catch/unit/atomics/min_max_common.hh @@ -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 && 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 && 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(p, stream, mem_devs[i].ptr(), - old_vals); + LaunchKernel(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(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 @@ -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(params);