From 599e4c9c68fa8f3741ada5bfa8310c852f28c5b0 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 5 Mar 2026 09:07:00 +0200 Subject: [PATCH] Add canMapHostMemory and managedMemory skip guards to catch tests Tests using hipHostMalloc, hipMallocManaged, or hipMemPrefetchAsync now check device capabilities and skip with HIP_SKIP_THIS_TEST when the device doesn't support host memory mapping or managed memory. This avoids false failures on backends like clvk with BufferDevAddr. Co-Authored-By: Claude Opus 4.6 --- tests/catch/unit/compiler/hipClassKernel.cc | 6 ++++++ .../catch/unit/kernel/hipMemFaultStackAllocation.cc | 7 +++++++ tests/catch/unit/memory/hipHostGetDevicePointer.cc | 6 ++++++ tests/catch/unit/memory/hipHostMalloc.cc | 12 ++++++++++++ tests/catch/unit/memory/hipMallocManaged.cc | 5 ++--- .../catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc | 7 +++++++ tests/catch/unit/memory/hipMemcpy.cc | 6 ++++++ tests/catch/unit/memory/hipMemcpyAsync.cc | 6 ++++++ tests/catch/unit/memory/hipMemoryAllocateCoherent.cc | 6 ++++++ tests/catch/unit/memory/hipMemsetFunctional.cc | 6 ++++++ tests/catch/unit/memory/memset.cc | 6 ++++++ 11 files changed, 70 insertions(+), 3 deletions(-) diff --git a/tests/catch/unit/compiler/hipClassKernel.cc b/tests/catch/unit/compiler/hipClassKernel.cc index 71286b706a..68442ea0d4 100644 --- a/tests/catch/unit/compiler/hipClassKernel.cc +++ b/tests/catch/unit/compiler/hipClassKernel.cc @@ -202,6 +202,12 @@ passByValueKernel(testPassByValue obj, bool* result_ecd) { } TEST_CASE("Unit_hipClassKernel_Value") { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } bool *result_ecd,*result_ech; result_ech = AllocateHostMemory(); result_ecd = AllocateDeviceMemory(); diff --git a/tests/catch/unit/kernel/hipMemFaultStackAllocation.cc b/tests/catch/unit/kernel/hipMemFaultStackAllocation.cc index 629afb208c..1cc4766ee5 100644 --- a/tests/catch/unit/kernel/hipMemFaultStackAllocation.cc +++ b/tests/catch/unit/kernel/hipMemFaultStackAllocation.cc @@ -64,6 +64,13 @@ static bool verify(const int* C_d, const int* A_d) { } TEST_CASE("Unit_hipMemFaultStackAllocation_Check") { + int managedSupport = 0; + HIP_CHECK(hipDeviceGetAttribute(&managedSupport, + hipDeviceAttributeManagedMemory, 0)); + if (!managedSupport) { + HipTest::HIP_SKIP_TEST("Device does not support managed memory"); + return; + } hipError_t ret; int *A_d, *C_d; const size_t Nbytes = N * sizeof(int); diff --git a/tests/catch/unit/memory/hipHostGetDevicePointer.cc b/tests/catch/unit/memory/hipHostGetDevicePointer.cc index 7c3e689e05..4e849f2bea 100644 --- a/tests/catch/unit/memory/hipHostGetDevicePointer.cc +++ b/tests/catch/unit/memory/hipHostGetDevicePointer.cc @@ -43,6 +43,12 @@ TEST_CASE("Unit_hipHostGetDevicePointer_Negative") { template __global__ void set(T* ptr, T val) { *ptr = val; } TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } int* hPtr{nullptr}; HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int))); diff --git a/tests/catch/unit/memory/hipHostMalloc.cc b/tests/catch/unit/memory/hipHostMalloc.cc index efe26ed669..33d1aba009 100644 --- a/tests/catch/unit/memory/hipHostMalloc.cc +++ b/tests/catch/unit/memory/hipHostMalloc.cc @@ -172,6 +172,12 @@ This testcase verifies the hipHostMalloc API by 3. validates the result. */ TEST_CASE("Unit_hipHostMalloc_NonCoherent") { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } int* A = nullptr; HIP_CHECK(hipHostMalloc(reinterpret_cast(&A), sizeBytes, hipHostMallocNonCoherent)); @@ -224,6 +230,12 @@ This testcase verifies the hipHostMalloc API by 3. validates the result. */ TEST_CASE("Unit_hipHostMalloc_Default") { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } int* A = nullptr; HIP_CHECK(hipHostMalloc(reinterpret_cast(&A), sizeBytes)); const char* ptrType = "default"; diff --git a/tests/catch/unit/memory/hipMallocManaged.cc b/tests/catch/unit/memory/hipMallocManaged.cc index 1b9766ffa9..8dd7bb0d0d 100644 --- a/tests/catch/unit/memory/hipMallocManaged.cc +++ b/tests/catch/unit/memory/hipMallocManaged.cc @@ -60,9 +60,8 @@ static unsigned threadsPerBlock{256}; TEST_CASE("Unit_hipMallocManaged_Basic") { auto managed = HmmAttrPrint(); if (managed != 1) { - WARN( - "GPU doesn't support hipDeviceAttributeManagedMemory attribute so defaulting to system " - "memory."); + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory"); + return; } float *A, *B, *C; diff --git a/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc b/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc index 482a73cd02..da3ba4f151 100644 --- a/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc +++ b/tests/catch/unit/memory/hipMemPrefetchAsyncExtTsts.cc @@ -322,6 +322,13 @@ TEST_CASE("Unit_hipMemPrefetchAsyncNegativeTst") { which is not multiple of page Size, but still trying to launch kernel and see if we are getting values as expected.*/ TEST_CASE("Unit_hipMemPrefetchAsync_NonPageSz") { + int managedSupport = 0; + HIP_CHECK(hipDeviceGetAttribute(&managedSupport, + hipDeviceAttributeManagedMemory, 0)); + if (!managedSupport) { + HipTest::HIP_SKIP_TEST("Device does not support managed memory"); + return; + } int *Hmm = nullptr, NumElms = 4096*2, InitVal = 123; hipStream_t strm; bool IfTestPassed = true; diff --git a/tests/catch/unit/memory/hipMemcpy.cc b/tests/catch/unit/memory/hipMemcpy.cc index 405bbf0e62..52d586b59c 100644 --- a/tests/catch/unit/memory/hipMemcpy.cc +++ b/tests/catch/unit/memory/hipMemcpy.cc @@ -438,6 +438,12 @@ This testcase verifies the following scenarios */ TEMPLATE_TEST_CASE("Unit_hipMemcpy_H2H-H2D-D2H-H2PinMem", "", int, float, double) { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } TestType *A_d{nullptr}, *B_d{nullptr}; TestType *A_h{nullptr}, *B_h{nullptr}; TestType *A_Ph{nullptr}, *B_Ph{nullptr}; diff --git a/tests/catch/unit/memory/hipMemcpyAsync.cc b/tests/catch/unit/memory/hipMemcpyAsync.cc index 4c4e08ec2d..cc4a05787f 100644 --- a/tests/catch/unit/memory/hipMemcpyAsync.cc +++ b/tests/catch/unit/memory/hipMemcpyAsync.cc @@ -142,6 +142,12 @@ This testcase verifies the following scenarios */ TEMPLATE_TEST_CASE("Unit_hipMemcpyAsync_H2H-H2D-D2H-H2PinMem", "", char, int, float, double) { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } TestType *A_d{nullptr}, *B_d{nullptr}; TestType *A_h{nullptr}, *B_h{nullptr}; TestType *A_Ph{nullptr}, *B_Ph{nullptr}; diff --git a/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc b/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc index 9b356c0210..69bcd96bc1 100644 --- a/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc +++ b/tests/catch/unit/memory/hipMemoryAllocateCoherent.cc @@ -44,6 +44,12 @@ __global__ void Kernel(float* hostRes, int clkRate) { } TEST_CASE("Unit_hipHostMalloc_CoherentAccess") { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } int blocks = 2; float* hostRes; HIP_CHECK(hipHostMalloc(&hostRes, blocks * sizeof(float), diff --git a/tests/catch/unit/memory/hipMemsetFunctional.cc b/tests/catch/unit/memory/hipMemsetFunctional.cc index 2f2b478ae3..0423e327e3 100644 --- a/tests/catch/unit/memory/hipMemsetFunctional.cc +++ b/tests/catch/unit/memory/hipMemsetFunctional.cc @@ -90,6 +90,12 @@ void checkMemset(T value, size_t count, MemsetType memsetType, bool async = fals if (mallocType == hipDeviceMalloc_t) { HIP_CHECK(hipMalloc(&devPtr, count * sizeof(T))); } else if (mallocType == hipHostMalloc_t) { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } HIP_CHECK(hipHostMalloc(&devPtr, count * sizeof(T))); } diff --git a/tests/catch/unit/memory/memset.cc b/tests/catch/unit/memory/memset.cc index 91f2d04601..217b16c31c 100644 --- a/tests/catch/unit/memory/memset.cc +++ b/tests/catch/unit/memory/memset.cc @@ -11,6 +11,12 @@ TEST_CASE("Unit_hipMemset_4bytes") { } TEST_CASE("Unit_hipMemset_4bytes_hostMem") { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + if (!prop.canMapHostMemory) { + HipTest::HIP_SKIP_TEST("Test requires canMapHostMemory support"); + return; + } int* d_a; auto res = hipHostMalloc(&d_a, sizeof(int), 0); REQUIRE(res == hipSuccess);