SWDEV-425830 - Add pcie atomic query in atomic tests

Add code to query pcie-atomics in atomic tests.
If pcie-atomics is not supported, skip the tests.
Enable the tests for Nvidia GPUs as well.

Change-Id: I7bfc8600511d78d8c73ab526f9846cd268651278


[ROCm/hip-tests commit: fafcad9162]

projects/hip-tests/catch/unit/memory/CMakeLists.txt

@@ -72,14 +72,14 @@ set(TEST_SRC
     hipMemAdviseMmap.cc
     hipMallocManaged.cc
     hipMemRangeGetAttribute.cc
-    hipMemRangeGetAttribute_old.cc)
+    hipMemRangeGetAttribute_old.cc
+    hipMemCoherencyTst.cc)
 
 if(HIP_PLATFORM MATCHES "amd")
   set(TEST_SRC
       ${TEST_SRC}
       hipMemPtrGetInfo.cc
       hipPointerGetAttributes.cc
-      hipMemCoherencyTst.cc
       hipExtMallocWithFlags.cc
       hipMallocMngdMultiThread.cc
       hipMemVmm.cc
@@ -101,6 +101,7 @@ hip_add_exe_to_target(NAME MemoryTest1
 if(HIP_PLATFORM MATCHES "amd")
   set_source_files_properties(hipHostRegister.cc PROPERTIES COMPILE_FLAGS -std=c++17)
   add_executable(hipHostRegisterPerf EXCLUDE_FROM_ALL hipHostRegister_exe.cc)
+  add_dependencies(build_tests hipHostRegisterPerf)
 if(UNIX)
   add_executable(hipMemAdviseTstAlignedAllocMem EXCLUDE_FROM_ALL hipMemAdvise_AlignedAllocMem_Exe.cc)
   add_dependencies(MemoryTest1 hipMemAdviseTstAlignedAllocMem)
@@ -207,17 +208,18 @@ hip_add_exe_to_target(NAME MemoryTest2
   TEST_SRC ${TEST_SRC}
   TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
 
-if(HIP_PLATFORM MATCHES "amd")
-  set(TEST_SRC
-    hipSVMTestByteGranularity.cpp
-    hipSVMTestFineGrainMemoryConsistency.cpp
-    hipSVMTestFineGrainSyncBuffers.cpp
-    hipSVMTestSharedAddressSpaceFineGrain.cpp
-  )
+set(TEST_SRC
+  hipSVMTestByteGranularity.cpp
+  hipSVMTestFineGrainMemoryConsistency.cpp
+  hipSVMTestFineGrainSyncBuffers.cpp
+  hipSVMTestSharedAddressSpaceFineGrain.cpp
+)
 
-  hip_add_exe_to_target(NAME SVMAtomicTest
-    TEST_SRC ${TEST_SRC}
-    TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
+hip_add_exe_to_target(NAME SVMAtomicTest
+  TEST_SRC ${TEST_SRC}
+  TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
 
-  add_dependencies(build_tests hipHostRegisterPerf)
-endif()
+if(HIP_PLATFORM MATCHES "nvidia")
+  set_target_properties(SVMAtomicTest PROPERTIES COMPILE_FLAGS -arch=sm_70)
+  set_target_properties(MemoryTest1 PROPERTIES COMPILE_FLAGS -arch=sm_70)
+endif()

projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc

@@ -94,10 +94,15 @@ static void TstCoherency(int* ptr, bool hmmMem) {
 
 /* Test case description: The following test validates if fine grain
    behavior is observed or not with memory allocated using hipHostMalloc()*/
-// The following tests are disabled for Nvidia as they are not consistently
-// passing
-#if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
+  int pcieAtomic = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0));
+  if (!pcieAtomic) {
+    fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n");
+    REQUIRE(true);
+    return;
+  }
+
   int *Ptr = nullptr, SIZE = sizeof(int);
   bool HmmMem = false;
   YES_COHERENT = false;
@@ -117,8 +122,6 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
   HIP_CHECK(hipHostFree(Ptr));
   REQUIRE(YES_COHERENT);
 }
-#endif
-
 
 /* Test case description: The following test validates if fine grain
    behavior is observed or not with memory allocated using hipMallocManaged()*/

projects/hip-tests/catch/unit/memory/hipSVMTestByteGranularity.cpp

@@ -74,6 +74,13 @@ __global__ void sum_neighbor_locations(char* a, unsigned int num_devices,
 *  - HIP_VERSION >= 5.7
 */
 TEST_CASE("test_svm_byte_granularity") {
+  int pcieAtomic = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0));
+  if (!pcieAtomic) {
+    fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n");
+    REQUIRE(true);
+    return;
+  }
   const int num_elements = 2048;
   int num_devices = 0;
   HIP_CHECK(hipGetDeviceCount(&num_devices));

projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainMemoryConsistency.cpp

@@ -234,6 +234,13 @@ void launch_kernels_and_verify(std::vector<hipStream_t> &streams, unsigned int n
 *  - HIP_VERSION >= 5.7
 */
 TEST_CASE("test_svm_fine_grain_memory_consistency") {
+  int pcieAtomic = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0));
+  if (!pcieAtomic) {
+    fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n");
+    REQUIRE(true);
+    return;
+  }
   const int num_elements = 2167;
   int num_devices = 0;
   HIP_CHECK(hipGetDeviceCount(&num_devices));

projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainSyncBuffers.cpp

@@ -76,6 +76,13 @@ void spawnAnalysisTask(int location)
 *  - HIP_VERSION >= 5.7
 */
 TEST_CASE("test_svm_fine_grain_sync_buffers") {
+  int pcieAtomic = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0));
+  if (!pcieAtomic) {
+    fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n");
+    REQUIRE(true);
+    return;
+  }
   size_t num_pixels = 1024 * 1024 * 2;
   hipStream_t stream;
   HIP_CHECK(hipSetDevice(0));

projects/hip-tests/catch/unit/memory/hipSVMTestSharedAddressSpaceFineGrain.cpp

@@ -129,7 +129,8 @@ void verify_linked_lists_on_device(hipStream_t stream, Node* pNodes,
   int correct_count = *pNumCorrect;
   if(correct_count != ListLength * numLists)
   {
-    fprintf(stderr,"Failed\n");
+    fprintf(stderr, "Failed: correct_count = %d, ListLength=%u, numLists = %u\n", correct_count,
+            ListLength, numLists);
     REQUIRE(false);
   }
 }
@@ -239,6 +240,7 @@ TEST_CASE("test_svm_shared_address_space_fine_grain_buffers") {
 * Test requirements
 * ------------------------
 *  - Host specific (WINDOWS and LINUX)
+*  - Unified address supported on devices
 *  - System fine grain access supported on devices
 *  - HIP_VERSION >= 5.7
 */

projects/hip-tests/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc

@@ -36,4 +36,4 @@ TEST_CASE("Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters") {
     HIP_CHECK_ERROR(hipDestroyExternalSemaphore(ext_semaphore), hipErrorInvalidValue);
   }
 #endif
-}
+}