diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index cf0b2f3920..22716158d9 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/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; \ diff --git a/projects/hip-tests/catch/multiproc/deviceAllocationMproc.cc b/projects/hip-tests/catch/multiproc/deviceAllocationMproc.cc index 78cb82c47b..93d40ac56a 100644 --- a/projects/hip-tests/catch/multiproc/deviceAllocationMproc.cc +++ b/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(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; } } diff --git a/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc b/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc index e03abd2458..e66cea1a44 100644 --- a/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/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 diff --git a/projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc b/projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc index e53c20ee34..35ac895518 100644 --- a/projects/hip-tests/catch/unit/kernel/hipPrintfKernel.cc +++ b/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()]; diff --git a/projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc b/projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc index 4350618c9f..d88add564a 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc +++ b/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;