diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index b5a83301df..63df9a75bd 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -289,117 +289,9 @@ "Unit_AtomicsWithRandomActiveLanesInWavefront_DivergentInteger", "Unit_hipGraphAddMemcpyNodeToSymbol_Positive_Basic", "Unit_hipStreamBeginCapture_Positive_Functional", - "Unit_atomicAnd_Positive_SameAddress - int", - "Unit_atomicAnd_Positive_SameAddress - unsigned int", - "Unit_atomicAnd_Positive_SameAddress - unsigned long", - "Unit_atomicAnd_Positive_SameAddress - unsigned long long", - "Unit_atomicAnd_Positive_Adjacent_Addresses - int", - "Unit_atomicAnd_Positive_Adjacent_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Adjacent_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Adjacent_Addresses - unsigned long long", - "Unit_atomicAnd_Positive_Scattered_Addresses - int", - "Unit_atomicAnd_Positive_Scattered_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Scattered_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Scattered_Addresses - unsigned long long", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - int", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - unsigned int", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - unsigned long", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - unsigned long long", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - int", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - int", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long", "Unit_atomicAnd_Negative_Parameters_RTC", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Same_Address - int", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Same_Address - unsigned int", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Same_Address - unsigned long", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Same_Address - unsigned long long", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Adjacent_Addresses - int", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned int", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long long", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Scattered_Addresses - int", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned int", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long", - "Unit_atomicAnd_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long long", - "Unit_atomicOr_Positive_SameAddress - int", - "Unit_atomicOr_Positive_SameAddress - unsigned int", - "Unit_atomicOr_Positive_SameAddress - unsigned long", - "Unit_atomicOr_Positive_SameAddress - unsigned long long", - "Unit_atomicOr_Positive_Adjacent_Addresses - int", - "Unit_atomicOr_Positive_Adjacent_Addresses - unsigned int", - "Unit_atomicOr_Positive_Adjacent_Addresses - unsigned long", - "Unit_atomicOr_Positive_Adjacent_Addresses - unsigned long long", - "Unit_atomicOr_Positive_Scattered_Addresses - int", - "Unit_atomicOr_Positive_Scattered_Addresses - unsigned int", - "Unit_atomicOr_Positive_Scattered_Addresses - unsigned long", - "Unit_atomicOr_Positive_Scattered_Addresses - unsigned long long", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - int", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - unsigned int", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - unsigned long", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - unsigned long long", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - int", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - int", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - unsigned int", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - unsigned long", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long", "Unit_atomicOr_Negative_Parameters_RTC", - "Unit_atomicOr_system_Positive_Peer_GPUs_Same_Address - int", - "Unit_atomicOr_system_Positive_Peer_GPUs_Same_Address - unsigned int", - "Unit_atomicOr_system_Positive_Peer_GPUs_Same_Address - unsigned long", - "Unit_atomicOr_system_Positive_Peer_GPUs_Same_Address - unsigned long long", - "Unit_atomicOr_system_Positive_Peer_GPUs_Adjacent_Addresses - int", - "Unit_atomicOr_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned int", - "Unit_atomicOr_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long", - "Unit_atomicOr_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long long", - "Unit_atomicOr_system_Positive_Peer_GPUs_Scattered_Addresses - int", - "Unit_atomicOr_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned int", - "Unit_atomicOr_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long", - "Unit_atomicOr_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long long", - "Unit_atomicXor_Positive_SameAddress - int", - "Unit_atomicXor_Positive_SameAddress - unsigned int", - "Unit_atomicXor_Positive_SameAddress - unsigned long", - "Unit_atomicXor_Positive_SameAddress - unsigned long long", - "Unit_atomicXor_Positive_Adjacent_Addresses - int", - "Unit_atomicXor_Positive_Adjacent_Addresses - unsigned int", - "Unit_atomicXor_Positive_Adjacent_Addresses - unsigned long", - "Unit_atomicXor_Positive_Adjacent_Addresses - unsigned long long", - "Unit_atomicXor_Positive_Scattered_Addresses - int", - "Unit_atomicXor_Positive_Scattered_Addresses - unsigned int", - "Unit_atomicXor_Positive_Scattered_Addresses - unsigned long", - "Unit_atomicXor_Positive_Scattered_Addresses - unsigned long long", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - int", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - unsigned int", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - unsigned long", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - unsigned long long", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - int", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - int", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - unsigned int", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - unsigned long", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long", "Unit_atomicXor_Negative_Parameters_RTC", - "Unit_atomicXor_system_Positive_Peer_GPUs_Same_Address - int", - "Unit_atomicXor_system_Positive_Peer_GPUs_Same_Address - unsigned int", - "Unit_atomicXor_system_Positive_Peer_GPUs_Same_Address - unsigned long", - "Unit_atomicXor_system_Positive_Peer_GPUs_Same_Address - unsigned long long", - "Unit_atomicXor_system_Positive_Peer_GPUs_Adjacent_Addresses - int", - "Unit_atomicXor_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned int", - "Unit_atomicXor_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long", - "Unit_atomicXor_system_Positive_Peer_GPUs_Adjacent_Addresses - unsigned long long", - "Unit_atomicXor_system_Positive_Peer_GPUs_Scattered_Addresses - int", - "Unit_atomicXor_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned int", - "Unit_atomicXor_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long", - "Unit_atomicXor_system_Positive_Peer_GPUs_Scattered_Addresses - unsigned long long", "Unit_atomicMin_Positive_SameAddress - int", "Unit_atomicMin_Positive_SameAddress - unsigned int", "Unit_atomicMin_Positive_SameAddress - unsigned long", @@ -600,78 +492,6 @@ "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___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - unsigned long long", "Unit_Kernel_Launch_bounds_Negative_OutOfBounds", "Unit_Kernel_Launch_bounds_Negative_Parameters_RTC", "Unit___threadfence_block_Positive_Basic_Peer", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index d29ae0b663..9507716096 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -777,80 +777,8 @@ "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_Positive_SameAddress - int", - "Unit_atomicAnd_Positive_SameAddress - unsigned int", - "Unit_atomicAnd_Positive_SameAddress - unsigned long", - "Unit_atomicAnd_Positive_SameAddress - unsigned long long", - "Unit_atomicAnd_Positive_Adjacent_Addresses - int", - "Unit_atomicAnd_Positive_Adjacent_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Adjacent_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Adjacent_Addresses - unsigned long long", - "Unit_atomicAnd_Positive_Scattered_Addresses - int", - "Unit_atomicAnd_Positive_Scattered_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Scattered_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Scattered_Addresses - unsigned long long", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - int", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - unsigned int", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - unsigned long", - "Unit_atomicAnd_Positive_Multi_Kernel_Same_Address - unsigned long long", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - int", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - int", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - unsigned int", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - unsigned long", - "Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long", "Unit_atomicAnd_Negative_Parameters_RTC", - "Unit_atomicOr_Positive_SameAddress - int", - "Unit_atomicOr_Positive_SameAddress - unsigned int", - "Unit_atomicOr_Positive_SameAddress - unsigned long", - "Unit_atomicOr_Positive_SameAddress - unsigned long long", - "Unit_atomicOr_Positive_Adjacent_Addresses - int", - "Unit_atomicOr_Positive_Adjacent_Addresses - unsigned int", - "Unit_atomicOr_Positive_Adjacent_Addresses - unsigned long", - "Unit_atomicOr_Positive_Adjacent_Addresses - unsigned long long", - "Unit_atomicOr_Positive_Scattered_Addresses - int", - "Unit_atomicOr_Positive_Scattered_Addresses - unsigned int", - "Unit_atomicOr_Positive_Scattered_Addresses - unsigned long", - "Unit_atomicOr_Positive_Scattered_Addresses - unsigned long long", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - int", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - unsigned int", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - unsigned long", - "Unit_atomicOr_Positive_Multi_Kernel_Same_Address - unsigned long long", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - int", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long", - "Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - int", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - unsigned int", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - unsigned long", - "Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long", "Unit_atomicOr_Negative_Parameters_RTC", - "Unit_atomicXor_Positive_SameAddress - int", - "Unit_atomicXor_Positive_SameAddress - unsigned int", - "Unit_atomicXor_Positive_SameAddress - unsigned long", - "Unit_atomicXor_Positive_SameAddress - unsigned long long", - "Unit_atomicXor_Positive_Adjacent_Addresses - int", - "Unit_atomicXor_Positive_Adjacent_Addresses - unsigned int", - "Unit_atomicXor_Positive_Adjacent_Addresses - unsigned long", - "Unit_atomicXor_Positive_Adjacent_Addresses - unsigned long long", - "Unit_atomicXor_Positive_Scattered_Addresses - int", - "Unit_atomicXor_Positive_Scattered_Addresses - unsigned int", - "Unit_atomicXor_Positive_Scattered_Addresses - unsigned long", - "Unit_atomicXor_Positive_Scattered_Addresses - unsigned long long", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - int", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - unsigned int", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - unsigned long", - "Unit_atomicXor_Positive_Multi_Kernel_Same_Address - unsigned long long", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - int", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - unsigned int", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long", - "Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses - unsigned long long", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - int", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - unsigned int", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - unsigned long", - "Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses - unsigned long long", "Unit_atomicXor_Negative_Parameters_RTC", "Unit_atomicMin_Positive_SameAddress - int", "Unit_atomicMin_Positive_SameAddress - unsigned int", @@ -1017,78 +945,6 @@ "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___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Wavefront_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_and_Positive_Workgroup_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Wavefront_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_or_Positive_Workgroup_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Wavefront_Scattered_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_SameAddress - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Adjacent_Addresses - unsigned long long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - unsigned int", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - unsigned long", - "Unit___hip_atomic_fetch_xor_Positive_Workgroup_Scattered_Addresses - unsigned long long", "=== Below tests cause timeout in stress test of 09/02/24 ===", "Unit_Device___half2half2_Accuracy_Positive", "Unit_Device_make_half2_Accuracy_Positive", diff --git a/catch/unit/atomics/atomicXor.cc b/catch/unit/atomics/atomicXor.cc index be7de6a44b..994b89020a 100644 --- a/catch/unit/atomics/atomicXor.cc +++ b/catch/unit/atomics/atomicXor.cc @@ -144,11 +144,10 @@ TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses", "" unsigned long, unsigned long long) { int warp_size = 0; HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0)); - for (auto current = 0; current < cmd_options.iterations; ++current) { DYNAMIC_SECTION("Adjacent address " << current) { Bitwise::SingleDeviceMultipleKernelTest( - 2, warp_size - 1, sizeof(TestType)); + 2, warp_size, sizeof(TestType)); } } } @@ -174,7 +173,7 @@ TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses", " for (auto current = 0; current < cmd_options.iterations; ++current) { DYNAMIC_SECTION("Scattered address " << current) { Bitwise::SingleDeviceMultipleKernelTest( - 2, warp_size - 1, cache_line_size); + 2, warp_size, cache_line_size); } } } diff --git a/catch/unit/atomics/bitwise_common.hh b/catch/unit/atomics/bitwise_common.hh index 887d25d4f9..dfd9d96148 100644 --- a/catch/unit/atomics/bitwise_common.hh +++ b/catch/unit/atomics/bitwise_common.hh @@ -135,7 +135,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); @@ -174,29 +174,34 @@ struct TestParams { template std::tuple, std::vector> TestKernelHostRef(const TestParams& p) { + const auto thread_count_per_kernel = p.ThreadCount(); const auto thread_count = p.num_devices * p.kernel_count * p.ThreadCount(); TestType test_value = GetTestValue(); const auto mask = kMask; - 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 tid = 0u; tid < thread_count; ++tid) { - auto& res = res_vals[tid % p.width]; - old_vals.push_back(res); + 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); - if constexpr (operation == AtomicOperation::kAnd || operation == AtomicOperation::kAndSystem || - operation == AtomicOperation::kBuiltinAnd) { - res = res & mask; - } else if constexpr (operation == AtomicOperation::kOr || - operation == AtomicOperation::kOrSystem || - operation == AtomicOperation::kBuiltinOr) { - res = res | mask; - } else if constexpr (operation == AtomicOperation::kXor || - operation == AtomicOperation::kXorSystem || - operation == AtomicOperation::kBuiltinXor) { - res = res ^ mask; + if constexpr (operation == AtomicOperation::kAnd || operation == AtomicOperation::kAndSystem || + operation == AtomicOperation::kBuiltinAnd) { + res = res & mask; + } else if constexpr (operation == AtomicOperation::kOr || + operation == AtomicOperation::kOrSystem || + operation == AtomicOperation::kBuiltinOr) { + res = res | mask; + } else if constexpr (operation == AtomicOperation::kXor || + operation == AtomicOperation::kXorSystem || + operation == AtomicOperation::kBuiltinXor) { + res = res ^ mask; + } + } } } @@ -237,7 +242,10 @@ template void TestCore(const TestParams& p) { const auto old_vals_alloc_size = p.kernel_count * p.ThreadCount() * sizeof(TestType); + const auto mem_alloc_size = p.width * p.pitch; + // Device Memory 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)); @@ -245,39 +253,40 @@ 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 ); + // Test Values on Device TestType test_value = GetTestValue(); - 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 and get back old vals 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); } @@ -287,7 +296,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 GENERATE_COPY(dim3(sm_count)); } template @@ -326,7 +335,7 @@ 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); @@ -361,7 +370,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); @@ -401,7 +410,7 @@ 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::hipMalloc}) { params.alloc_type = alloc_type; DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) { TestCore(params);