diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 51b81bbe1d..3a3268da65 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -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 ===", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 26d0e8dee8..984665dcb6 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -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", 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 13234564d9..bf4893b08b 100644 --- a/projects/hip-tests/catch/unit/atomics/min_max_common.hh +++ b/projects/hip-tests/catch/unit/atomics/min_max_common.hh @@ -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 std::tuple, std::vector> TestKernelHostRef(const TestParams& p) { const auto val = GetTestValue(); - + 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 ? kFloatingPointTestValue : kIntegerTestValue; - std::vector res_vals(p.width, test_value); + std::vector res_vals(p.num_devices * p.width, test_value); std::vector 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 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> old_vals_devs; + std::vector> mem_devs; std::vector 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 mem_dev(p.alloc_type, mem_alloc_size); - + // Host Memory std::vector old_vals(p.num_devices * p.kernel_count * p.ThreadCount()); - std::vector res_vals(p.width); - - TestType* const mem_ptr = - p.alloc_type == LinearAllocs::hipMalloc ? mem_dev.ptr() : mem_dev.host_ptr(); + std::vector res_vals(p.num_devices * p.width); + // Initialize Device Memory TestType test_value = std::is_floating_point_v ? 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(p, stream, mem_dev.ptr(), + LaunchKernel(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(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 @@ -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(params); } } } - + #ifdef __linux__ SECTION("Shared memory") { params.blocks = dim3(1); params.alloc_type = LinearAllocs::hipMalloc; TestCore(params); } + #endif } template @@ -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(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(params); } } } - } // namespace MinMax