AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

SWDEV-454316 - Fix for Atomic Min Max Testcases

Bläddra i källkod 
Change-Id: Ieca965e6e17f4b17769938228340791b9bbe45ab


[ROCm/hip-tests commit: 0c230d3cf0]
This commit is contained in:
Rahul Manocha
2024-03-29 14:03:49 -07:00
committad av Rahul Manocha
förälder 66a1b8d8e4
incheckning 572eb13c5c
3 ändrade filer med 48 tillägg och 397 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/hipTestMain/config/config_amd_linux
+4 -204
Visa fil
@@ -222,12 +222,6 @@
"Unit_hipExtModuleLaunchKernel_Negative_Parameters",
"Unit_hipLaunchKernel_Negative_Parameters",
"Unit_Device_modf_modff_Negative_RTC",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - double",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Same_Address - float",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Same_Address - double",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Same_Address - float",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - double",
"SWDEV-446588 - Disable graph multi gpu testcases until graph has support for it",
"Unit_hipGraphExecUpdate_Negative_MultiDevice_Context_Changed",
"Unit_hipGraphMem_Alloc_Free_NodeGetParams_Functional_MultiDevice",
@@ -247,206 +241,8 @@
"Unit_atomicAnd_Negative_Parameters_RTC",
"Unit_atomicOr_Negative_Parameters_RTC",
"Unit_atomicXor_Negative_Parameters_RTC",
"Unit_atomicMin_Positive_SameAddress - int",
"Unit_atomicMin_Positive_SameAddress - unsigned int",
"Unit_atomicMin_Positive_SameAddress - unsigned long",
"Unit_atomicMin_Positive_SameAddress - unsigned long long",
"Unit_atomicMin_Positive_SameAddress - float",
"Unit_atomicMin_Positive_SameAddress - double",
"Unit_atomicMin_Positive_Adjacent_Addresses - int",
"Unit_atomicMin_Positive_Adjacent_Addresses - unsigned int",
"Unit_atomicMin_Positive_Adjacent_Addresses - unsigned long",
"Unit_atomicMin_Positive_Adjacent_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Adjacent_Addresses - float",
"Unit_atomicMin_Positive_Adjacent_Addresses - double",
"Unit_atomicMin_Positive_Scattered_Addresses - int",
"Unit_atomicMin_Positive_Scattered_Addresses - unsigned int",
"Unit_atomicMin_Positive_Scattered_Addresses - unsigned long",
"Unit_atomicMin_Positive_Scattered_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Scattered_Addresses - float",
"Unit_atomicMin_Positive_Scattered_Addresses - double",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - int",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - unsigned int",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - unsigned long",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - unsigned long long",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - float",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - double",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - int",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - int",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - unsigned int",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - unsigned long",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_atomicMin_Negative_Parameters_RTC",
"Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address - float",
"Unit_atomicMin_system_Positive_Peer_GPUs_Same_Address - double",
"Unit_atomicMin_system_Positive_Peer_GPUs_Adjacent_Addresses - int",
"Unit_atomicMin_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned int",
"Unit_atomicMin_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long",
"Unit_atomicMin_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long long",
"Unit_atomicMin_system_Positive_Peer_GPUs_Adjacent_Addresses - float",
"Unit_atomicMin_system_Positive_Peer_GPUs_Adjacent_Addresses - double",
"Unit_atomicMin_system_Positive_Peer_GPUs_Scattered_Addresses - int",
"Unit_atomicMin_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned int",
"Unit_atomicMin_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long",
"Unit_atomicMin_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long long",
"Unit_atomicMin_system_Positive_Peer_GPUs_Scattered_Addresses - float",
"Unit_atomicMin_system_Positive_Peer_GPUs_Scattered_Addresses - double",
"Unit_atomicMax_Positive_SameAddress - int",
"Unit_atomicMax_Positive_SameAddress - unsigned int",
"Unit_atomicMax_Positive_SameAddress - unsigned long",
"Unit_atomicMax_Positive_SameAddress - unsigned long long",
"Unit_atomicMax_Positive_SameAddress - double",
"Unit_atomicMax_Positive_Adjacent_Addresses - int",
"Unit_atomicMax_Positive_Adjacent_Addresses - unsigned int",
"Unit_atomicMax_Positive_Adjacent_Addresses - unsigned long",
"Unit_atomicMax_Positive_Adjacent_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Adjacent_Addresses - float",
"Unit_atomicMax_Positive_Adjacent_Addresses - double",
"Unit_atomicMax_Positive_Scattered_Addresses - int",
"Unit_atomicMax_Positive_Scattered_Addresses - unsigned int",
"Unit_atomicMax_Positive_Scattered_Addresses - unsigned long",
"Unit_atomicMax_Positive_Scattered_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Scattered_Addresses - float",
"Unit_atomicMax_Positive_Scattered_Addresses - double",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - int",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - unsigned int",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - unsigned long",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - unsigned long long",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - float",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - int",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - int",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - unsigned int",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - unsigned long",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_atomicMax_Negative_Parameters_RTC",
"Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address - float",
"Unit_atomicMax_system_Positive_Peer_GPUs_Same_Address - double",
"Unit_atomicMax_system_Positive_Peer_GPUs_Adjacent_Addresses - int",
"Unit_atomicMax_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned int",
"Unit_atomicMax_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long",
"Unit_atomicMax_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long long",
"Unit_atomicMax_system_Positive_Peer_GPUs_Adjacent_Addresses - float",
"Unit_atomicMax_system_Positive_Peer_GPUs_Adjacent_Addresses - double",
"Unit_atomicMax_system_Positive_Peer_GPUs_Scattered_Addresses - int",
"Unit_atomicMax_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned int",
"Unit_atomicMax_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long",
"Unit_atomicMax_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long long",
"Unit_atomicMax_system_Positive_Peer_GPUs_Scattered_Addresses - float",
"Unit_atomicMax_system_Positive_Peer_GPUs_Scattered_Addresses - double",
"Unit_safeAtomicMin_Positive_Adjacent_Addresses - float",
"Unit_safeAtomicMin_Positive_Adjacent_Addresses - double",
"Unit_safeAtomicMin_Positive_Scattered_Addresses - float",
"Unit_safeAtomicMin_Positive_Scattered_Addresses - double",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Same_Address - float",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Same_Address - double",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_unsafeAtomicMin_Positive_SameAddress - double",
"Unit_unsafeAtomicMin_Positive_Adjacent_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Adjacent_Addresses - double",
"Unit_unsafeAtomicMin_Positive_Scattered_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Scattered_Addresses - double",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Same_Address - double",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_safeAtomicMax_Positive_Adjacent_Addresses - float",
"Unit_safeAtomicMax_Positive_Adjacent_Addresses - double",
"Unit_safeAtomicMax_Positive_Scattered_Addresses - float",
"Unit_safeAtomicMax_Positive_Scattered_Addresses - double",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Same_Address - float",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_unsafeAtomicMax_Positive_SameAddress - double",
"Unit_unsafeAtomicMax_Positive_Adjacent_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Adjacent_Addresses - double",
"Unit_unsafeAtomicMax_Positive_Scattered_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Scattered_Addresses - double",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Same_Address - double",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - double",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - double",
"Unit_Kernel_Launch_bounds_Negative_OutOfBounds",
"Unit_Kernel_Launch_bounds_Negative_Parameters_RTC",
"Unit___threadfence_block_Positive_Basic_Peer",
@@ -1059,6 +855,10 @@
"Unit_safeAtomicMin_Positive_SameAddress - float",
"=== SWDEV-454220 : Below test hanged in stress test on 22/03/24 ===",
"Unit_hipExtLaunchKernel_Positive_Basic",
"=== 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",
#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
+4 -163
Visa fil
@@ -736,171 +736,8 @@
"Unit_atomicAnd_Negative_Parameters_RTC",
"Unit_atomicOr_Negative_Parameters_RTC",
"Unit_atomicXor_Negative_Parameters_RTC",
"Unit_atomicMin_Positive_SameAddress - int",
"Unit_atomicMin_Positive_SameAddress - unsigned int",
"Unit_atomicMin_Positive_SameAddress - unsigned long",
"Unit_atomicMin_Positive_SameAddress - unsigned long long",
"Unit_atomicMin_Positive_SameAddress - float",
"Unit_atomicMin_Positive_SameAddress - double",
"Unit_atomicMin_Positive_Adjacent_Addresses - int",
"Unit_atomicMin_Positive_Adjacent_Addresses - unsigned int",
"Unit_atomicMin_Positive_Adjacent_Addresses - unsigned long",
"Unit_atomicMin_Positive_Adjacent_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Adjacent_Addresses - float",
"Unit_atomicMin_Positive_Adjacent_Addresses - double",
"Unit_atomicMin_Positive_Scattered_Addresses - int",
"Unit_atomicMin_Positive_Scattered_Addresses - unsigned int",
"Unit_atomicMin_Positive_Scattered_Addresses - unsigned long",
"Unit_atomicMin_Positive_Scattered_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Scattered_Addresses - float",
"Unit_atomicMin_Positive_Scattered_Addresses - double",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - int",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - unsigned int",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - unsigned long",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - unsigned long long",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - float",
"Unit_atomicMin_Positive_Multi_Kernel_Same_Address - double",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - int",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_atomicMin_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - int",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - unsigned int",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - unsigned long",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_atomicMin_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_atomicMin_Negative_Parameters_RTC",
"Unit_atomicMax_Positive_SameAddress - int",
"Unit_atomicMax_Positive_SameAddress - unsigned int",
"Unit_atomicMax_Positive_SameAddress - unsigned long",
"Unit_atomicMax_Positive_SameAddress - unsigned long long",
"Unit_atomicMax_Positive_Adjacent_Addresses - int",
"Unit_atomicMax_Positive_Adjacent_Addresses - unsigned int",
"Unit_atomicMax_Positive_Adjacent_Addresses - unsigned long",
"Unit_atomicMax_Positive_Adjacent_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Adjacent_Addresses - float",
"Unit_atomicMax_Positive_Adjacent_Addresses - double",
"Unit_atomicMax_Positive_Scattered_Addresses - int",
"Unit_atomicMax_Positive_Scattered_Addresses - unsigned int",
"Unit_atomicMax_Positive_Scattered_Addresses - unsigned long",
"Unit_atomicMax_Positive_Scattered_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Scattered_Addresses - float",
"Unit_atomicMax_Positive_Scattered_Addresses - double",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - int",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - unsigned int",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - unsigned long",
"Unit_atomicMax_Positive_Multi_Kernel_Same_Address - unsigned long long",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - int",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_atomicMax_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - int",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - unsigned int",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - unsigned long",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_atomicMax_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_atomicMax_Negative_Parameters_RTC",
"Unit_safeAtomicMin_Positive_Adjacent_Addresses - float",
"Unit_safeAtomicMin_Positive_Adjacent_Addresses - double",
"Unit_safeAtomicMin_Positive_Scattered_Addresses - float",
"Unit_safeAtomicMin_Positive_Scattered_Addresses - double",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_safeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_unsafeAtomicMin_Positive_Adjacent_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Adjacent_Addresses - double",
"Unit_unsafeAtomicMin_Positive_Scattered_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Scattered_Addresses - double",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_unsafeAtomicMin_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_safeAtomicMax_Positive_Adjacent_Addresses - float",
"Unit_safeAtomicMax_Positive_Adjacent_Addresses - double",
"Unit_safeAtomicMax_Positive_Scattered_Addresses - float",
"Unit_safeAtomicMax_Positive_Scattered_Addresses - double",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_safeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit_unsafeAtomicMax_Positive_Adjacent_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Adjacent_Addresses - double",
"Unit_unsafeAtomicMax_Positive_Scattered_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Scattered_Addresses - double",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Adjacent_Addresses - double",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - float",
"Unit_unsafeAtomicMax_Positive_Multi_Kernel_Scattered_Addresses - double",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_min_Positive_Wavefront_Scattered_Addresses - double",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_min_Positive_Workgroup_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Wavefront_Scattered_Addresses - double",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_SameAddress - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Adjacent_Addresses - double",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - unsigned int",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - unsigned long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - unsigned long long",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - float",
"Unit___hip_atomic_fetch_max_Positive_Workgroup_Scattered_Addresses - double",
"=== Below tests cause timeout in stress test of 09/02/24 ===",
"Unit_Device___half2half2_Accuracy_Positive",
"Unit_Device_make_half2_Accuracy_Positive",
@@ -1127,6 +964,10 @@
"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",
"=== SWDEV-475482 - Disable tests to merge clr change",
"Unit_hipCreateTextureObject_LinearResource",
"Unit_hipCreateTextureObject_Pitch2DResource",
projects/hip-tests/catch/unit/atomics/min_max_common.hh
+40 -30
Visa fil
@@ -142,7 +142,7 @@ __global__ void TestKernel(TestType* const global_mem, TestType* const old_vals,
__syncthreads();
}
const auto n = cooperative_groups::this_grid().size() - width;
const auto n = cooperative_groups::this_grid().size();
TestType* atomic_addr = PitchedOffset(mem, pitch, tid % width);
@@ -182,31 +182,34 @@ struct TestParams {
template <typename TestType, AtomicOperation operation>
std::tuple<std::vector<TestType>, std::vector<TestType>> TestKernelHostRef(const TestParams& p) {
const auto val = GetTestValue<TestType, operation>();
const auto thread_count_per_kernel = p.ThreadCount();
const auto thread_count = p.num_devices * p.kernel_count * p.ThreadCount();
TestType test_value =
std::is_floating_point_v<TestType> ? kFloatingPointTestValue : kIntegerTestValue;
std::vector<TestType> res_vals(p.width, test_value);
std::vector<TestType> res_vals(p.num_devices * p.width, test_value);
std::vector<TestType> old_vals;
old_vals.reserve(thread_count);
for (auto i = 0u; i < p.num_devices; ++i) {
for (auto j = 0u; j < p.kernel_count; ++j) {
for (auto tid = 0u; tid < thread_count_per_kernel; ++tid) {
auto& res = res_vals[tid % p.width + (i * p.width)];
old_vals.push_back(res);
for (auto tid = 0u; tid < thread_count; ++tid) {
auto& res = res_vals[tid % p.width];
old_vals.push_back(res);
if constexpr (operation == AtomicOperation::kMin || operation == AtomicOperation::kMinSystem ||
if constexpr (operation == AtomicOperation::kMin || operation == AtomicOperation::kMinSystem ||
operation == AtomicOperation::kUnsafeMin ||
operation == AtomicOperation::kSafeMin ||
operation == AtomicOperation::kBuiltinMin) {
res = std::min(res, val);
} else if constexpr (operation == AtomicOperation::kMax ||
res = std::min(res, val);
} else if constexpr (operation == AtomicOperation::kMax ||
operation == AtomicOperation::kMaxSystem ||
operation == AtomicOperation::kUnsafeMax ||
operation == AtomicOperation::kSafeMax ||
operation == AtomicOperation::kBuiltinMax) {
res = std::max(res, val);
res = std::max(res, val);
}
}
}
}
@@ -246,8 +249,12 @@ void LaunchKernel(const TestParams& p, hipStream_t stream, TestType* const mem_p
template <typename TestType, AtomicOperation operation, bool use_shared_mem,
int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
void TestCore(const TestParams& p) {
// Device Memory Allocation
const auto old_vals_alloc_size = p.kernel_count * p.ThreadCount() * sizeof(TestType);
const auto mem_alloc_size = p.width * p.pitch;
std::vector<LinearAllocGuard<TestType>> old_vals_devs;
std::vector<LinearAllocGuard<TestType>> mem_devs;
std::vector<StreamGuard> streams;
for (auto i = 0; i < p.num_devices; ++i) {
HIP_CHECK(hipSetDevice(i));
@@ -255,40 +262,43 @@ void TestCore(const TestParams& p) {
for (auto j = 0; j < p.kernel_count; ++j) {
streams.emplace_back(Streams::created);
}
mem_devs.emplace_back(p.alloc_type, mem_alloc_size);
}
const auto mem_alloc_size = p.width * p.pitch;
LinearAllocGuard<TestType> mem_dev(p.alloc_type, mem_alloc_size);
// Host Memory
std::vector<TestType> old_vals(p.num_devices * p.kernel_count * p.ThreadCount());
std::vector<TestType> res_vals(p.width);
TestType* const mem_ptr =
p.alloc_type == LinearAllocs::hipMalloc ? mem_dev.ptr() : mem_dev.host_ptr();
std::vector<TestType> res_vals(p.num_devices * p.width);
// Initialize Device Memory
TestType test_value =
std::is_floating_point_v<TestType> ? kFloatingPointTestValue : kIntegerTestValue;
HIP_CHECK(hipMemset(mem_ptr, 0, mem_alloc_size));
for (int i = 0; i < p.width * p.pitch / sizeof(TestType); ++i) {
HIP_CHECK(hipMemcpy(&mem_ptr[i], &test_value, sizeof(TestType), hipMemcpyHostToDevice));
for (auto i = 0u; i < p.num_devices; ++i) {
TestType* const mem_ptr =
p.alloc_type == LinearAllocs::hipMalloc ? mem_devs[i].ptr() : mem_devs[i].host_ptr();
HIP_CHECK(hipMemset(mem_ptr, 0, mem_alloc_size));
for (int i = 0; i < p.width * p.pitch / sizeof(TestType); ++i) {
HIP_CHECK(hipMemcpy(&mem_ptr[i], &test_value, sizeof(TestType), hipMemcpyHostToDevice));
}
}
// Launch kernel
for (auto i = 0u; i < p.num_devices; ++i) {
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_dev.ptr(),
LaunchKernel<TestType, operation, use_shared_mem, memory_scope>(p, stream, mem_devs[i].ptr(),
old_vals);
}
}
// Copy Results back to Host
for (auto i = 0u; i < p.num_devices; ++i) {
const auto device_offset = i * p.kernel_count * p.ThreadCount();
HIP_CHECK(hipMemcpy(old_vals.data() + device_offset, old_vals_devs[i].ptr(),
old_vals_alloc_size, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy2D(res_vals.data() + i*p.width, sizeof(TestType), mem_devs[i].ptr(),
p.pitch, sizeof(TestType), p.width, hipMemcpyDeviceToHost));
}
HIP_CHECK(hipMemcpy2D(res_vals.data(), sizeof(TestType), mem_ptr, p.pitch, sizeof(TestType),
p.width, hipMemcpyDeviceToHost));
Verify<TestType, operation>(p, res_vals, old_vals);
}
@@ -298,7 +308,7 @@ inline dim3 GenerateThreadDimensions() { return GENERATE(dim3(16), dim3(1024));
inline dim3 GenerateBlockDimensions() {
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(sm_count);
}
template <typename TestType, AtomicOperation operation, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
@@ -334,19 +344,20 @@ void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int p
}
using LA = LinearAllocs;
for (const auto alloc_type :
{LA::hipMalloc, LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) {
{LA::hipMalloc}) {
params.alloc_type = alloc_type;
DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
TestCore<TestType, operation, false>(params);
}
}
}
#ifdef __linux__
SECTION("Shared memory") {
params.blocks = dim3(1);
params.alloc_type = LinearAllocs::hipMalloc;
TestCore<TestType, operation, true>(params);
}
#endif
}
template <typename TestType, AtomicOperation operation>
@@ -369,7 +380,7 @@ void SingleDeviceMultipleKernelTest(const unsigned int kernel_count, const unsig
using LA = LinearAllocs;
for (const auto alloc_type :
{LA::hipMalloc, LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) {
{LA::hipMalloc}) {
params.alloc_type = alloc_type;
DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
TestCore<TestType, operation, false>(params);
@@ -409,12 +420,11 @@ void MultipleDeviceMultipleKernelTest(const unsigned int num_devices,
params.pitch = pitch;
using LA = LinearAllocs;
for (const auto alloc_type : {LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) {
for (const auto alloc_type : {LA::hipHostMalloc}) {
params.alloc_type = alloc_type;
DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
TestCore<TestType, operation, false, __HIP_MEMORY_SCOPE_SYSTEM>(params);
}
}
}
} // namespace MinMax