SWDEV-532473 - Add Pcie atomic support check to failing tests (#575)

Co-authored-by: Rahul Manocha <rmanocha@amd.com>

projects/hip-tests/catch/include/hip_test_common.hh

@@ -334,13 +334,12 @@ inline bool isImageSupported() {
   return imageSupport != 0;
 }
 
-inline bool isPcieAtomicsSupported() {
-  int pcieAtomics = 1;
+inline bool isPcieAtomicSupported() {
+  int pcieAtomic = 1;
   int device;
   HIP_CHECK(hipGetDevice(&device));
-  HIPCHECK(hipDeviceGetAttribute(&pcieAtomics, hipDeviceAttributeHostNativeAtomicSupported,
-           device));
-  return pcieAtomics != 0;
+  HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, device));
+  return pcieAtomic;
 }
 
 inline bool isP2PSupported(int& d1, int& d2) {
@@ -549,12 +548,11 @@ class BlockingContext {
     return;                                                                                        \
   }
 
-// This must be called in host-device memory conherency tests
-#define CHECK_PCIE_ATOMICS_SUPPORT                                                                 \
-  if (!HipTest::isPcieAtomicsSupported()) {                                                        \
-    INFO("Pcie atomics is not support on the device. Skipped.");                                   \
+#define CHECK_PCIE_ATOMIC_SUPPORT                                                                 \
+  if (!HipTest::isPcieAtomicSupported()) {                                                        \
+    HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped");                         \
     return;                                                                                        \
-  }
+  }   
 
 #define CHECK_P2P_SUPPORT                                                                          \
   int d1, d2;                                                                                      \

projects/hip-tests/catch/multiproc/deviceAllocationMproc.cc

@@ -43,7 +43,6 @@ static __global__ void kerTestDeviceMalloc(size_t size) {
   if (myId == 0) {
     dev_common_ptr = reinterpret_cast<char*>(malloc(size));
     if (dev_common_ptr == nullptr) {
-      printf("Device Allocation Failed! \n");
       return;
     }
   }
@@ -57,7 +56,6 @@ static __global__ void kerTestDeviceWrite() {
   int myId = threadIdx.x + blockDim.x * blockIdx.x;
   // Allocate
   if (dev_common_ptr == nullptr) {
-    printf("Device Allocation Failed! \n");
     return;
   }
   *(dev_common_ptr + myId) = SCHAR_MAX;
@@ -95,7 +93,6 @@ static __global__ void kerTestDeviceNew(size_t size) {
   if (myId == 0) {
     dev_common_ptr = new char[size];
     if (dev_common_ptr == nullptr) {
-      printf("Device Allocation Failed! \n");
       return;
     }
   }

projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc

@@ -102,7 +102,7 @@ bool static TstCoherency(int* Ptr, bool HmmMem) {
 // The following test is failing on Nvidia platform hence disabled it for now
 #if HT_AMD
 TEST_CASE("Unit_malloc_CoherentTst") {
-  CHECK_PCIE_ATOMICS_SUPPORT
+  CHECK_PCIE_ATOMIC_SUPPORT
   hipDeviceProp_t prop;
   HIPCHECK(hipGetDeviceProperties(&prop, 0));
   char* p = NULL;
@@ -164,7 +164,7 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
 // The following test is failing on Nvidia platform hence disabling it for now
 #if HT_AMD
 TEST_CASE("Unit_mmap_CoherentTst") {
-  CHECK_PCIE_ATOMICS_SUPPORT
+  CHECK_PCIE_ATOMIC_SUPPORT
   hipDeviceProp_t prop;
   HIPCHECK(hipGetDeviceProperties(&prop, 0));
   char* p = NULL;
@@ -408,7 +408,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
   }
   int stat = 0;
   if (fork() == 0) {  // child process
-    CHECK_PCIE_ATOMICS_SUPPORT
+    CHECK_PCIE_ATOMIC_SUPPORT;
     int *Ptr = nullptr, SIZE = sizeof(int);
     bool HmmMem = false;
     // Allocating hipHostMalloc() memory
@@ -437,7 +437,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
   }
   int stat = 0;
   if (fork() == 0) {  // child process
-    CHECK_PCIE_ATOMICS_SUPPORT
+    CHECK_PCIE_ATOMIC_SUPPORT
     int *Ptr = nullptr, SIZE = sizeof(int);
     bool HmmMem = false;
     // Allocating hipHostMalloc() memory
@@ -465,7 +465,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
   }
   int stat = 0;
   if (fork() == 0) {  // child process
-    CHECK_PCIE_ATOMICS_SUPPORT
+    CHECK_PCIE_ATOMIC_SUPPORT
     int *Ptr = nullptr, SIZE = sizeof(int);
     bool HmmMem = false;
     // Allocating hipHostMalloc() memory
@@ -493,7 +493,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") {
   }
   int stat = 0;
   if (fork() == 0) {  // child process
-    CHECK_PCIE_ATOMICS_SUPPORT
+    CHECK_PCIE_ATOMIC_SUPPORT
     int *Ptr = nullptr, SIZE = sizeof(int);
     bool HmmMem = false;
     // Allocating hipHostMalloc() memory

projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc

@@ -53,6 +53,7 @@ TEST_CASE("Unit_kernel_ChkPrintf") {
   const char* check = st.c_str();
   for (int i = 0; i < device_count; ++i) {
     HIP_CHECK(hipSetDevice(i));
+    if (!HipTest::isPcieAtomicSupported()) continue;
     hipLaunchKernelGGL(run_printf, dim3(1), dim3(1), 0, 0);
     HIP_CHECK(hipDeviceSynchronize());
     char* data = new char[st.size()];

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

@@ -106,13 +106,7 @@ static void TstCoherency(int* ptr, MemoryType type) {
    behavior is observed or not with memory allocated using hipHostMalloc()*/
 TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
   HIP_CHECK(hipSetDevice(0));
-  int pcieAtomic = 0;
-  HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0));
-  if (!pcieAtomic) {
-    fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n");
-    REQUIRE(true);
-    return;
-  }
+  CHECK_PCIE_ATOMIC_SUPPORT;
 
   int *Ptr = nullptr, SIZE = sizeof(int);
   YES_COHERENT = false;
@@ -138,6 +132,8 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
 #if HT_AMD
 TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
   HIP_CHECK(hipSetDevice(0));
+  CHECK_PCIE_ATOMIC_SUPPORT;
+
   int *Ptr = nullptr, SIZE = sizeof(int), managed = 0;
   YES_COHERENT = false;