AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

[SWDEV-454315] - Fix for Atomic bitwise ops testcases

Przeglądaj źródła 
Change-Id: Ib402c6499ad9218fba89e78c8f91a87ca3537b35
This commit is contained in:
Rahul Manocha
2024-03-29 13:01:41 -07:00
zatwierdzone przez Rakesh Roy
rodzic 787f7f8df1
commit 90bafb8ae8
4 zmienionych plików z 46 dodań i 362 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
catch/hipTestMain/config/config_amd_linux
-180
Wyświetl plik
@@ -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",
catch/hipTestMain/config/config_amd_windows
-144
Wyświetl plik
@@ -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",
catch/unit/atomics/atomicXor.cc
+2 -3
Wyświetl plik
@@ -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<TestType, Bitwise::AtomicOperation::kXor>(
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<TestType, Bitwise::AtomicOperation::kXor>(
2, warp_size - 1, cache_line_size);
2, warp_size, cache_line_size);
}
}
}
catch/unit/atomics/bitwise_common.hh
+44 -35
Wyświetl plik
@@ -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 <typename TestType, AtomicOperation operation>
std::tuple<std::vector<TestType>, std::vector<TestType>> 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<TestType, operation>();
const auto mask = kMask;
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 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 <typename TestType, AtomicOperation operation, bool use_shared_mem,
int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
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<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));
@@ -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<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 );
// Test Values on Device
TestType test_value = GetTestValue<TestType, operation>();
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<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);
}
@@ -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 <typename TestType, AtomicOperation operation, int memory_scope = __HIP_MEMORY_SCOPE_AGENT>
@@ -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<TestType, operation, false>(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<TestType, operation, false>(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<TestType, operation, false, __HIP_MEMORY_SCOPE_SYSTEM>(params);