From 885c9dcf896d6c4f59c986480b358edbe15c447e Mon Sep 17 00:00:00 2001 From: Ioannis Assiouras Date: Wed, 14 Oct 2026 01:30:34 +0000 Subject: [PATCH 1/2] test(hip-tests): detect XNACK via the pageable memory access attribute MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude --- .../catch/include/hip_test_common.hh | 12 ++++++---- .../multiproc/hipMemCoherencyTstMProc.cc | 24 ++++--------------- .../stress/memory/hipHmmOvrSubscriptionTst.cc | 9 ++----- .../stress/memory/hipHostRegisterStress.cc | 3 +-- .../catch/unit/memory/hipMemAdvise_old.cc | 19 +++------------ .../hipGetProcAddressVmmApis.cc | 11 +-------- 6 files changed, 19 insertions(+), 59 deletions(-) diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index 47b7bb6183a3..dcafce8fe612 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -524,13 +524,17 @@ inline bool checkConcurrentKernels(int num_devices) { return true; } +// Stall-only XNACK GPUs never report xnack+ in gcnArchName; ask for pageable access instead. inline bool isXnackOn() { - hipDeviceProp_t prop; +#if HT_AMD int device = 0; + int pageable = 0; HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - std::string gfxName(prop.gcnArchName); - return gfxName.find("xnack+") != std::string::npos; + HIP_CHECK(hipDeviceGetAttribute(&pageable, hipDeviceAttributePageableMemoryAccess, device)); + return pageable != 0; +#else + return false; +#endif } inline bool areWarpMatchFunctionsSupported() { diff --git a/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc b/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc index 10dc66f44e6f..b34e496ec192 100644 --- a/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/projects/hip-tests/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -91,11 +91,7 @@ bool static TstCoherency(int* Ptr, bool HmmMem) { HIP_TEST_CASE(Unit_malloc_CoherentTst) { CHECK_PCIE_ATOMIC_SUPPORT CHECK_MANAGED_MEMORY_SUPPORT - hipDeviceProp_t prop; - HIPCHECK(hipGetDeviceProperties(&prop, 0)); - char* p = NULL; - p = strstr(prop.gcnArchName, "xnack+"); - if (p) { + if (HipTest::isXnackOn()) { int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = true; @@ -116,11 +112,7 @@ HIP_TEST_CASE(Unit_malloc_CoherentTst) { #if HT_AMD HIP_TEST_CASE(Unit_malloc_CoherentTstWthAdvise) { CHECK_MANAGED_MEMORY_SUPPORT - hipDeviceProp_t prop; - HIPCHECK(hipGetDeviceProperties(&prop, 0)); - char* p = NULL; - p = strstr(prop.gcnArchName, "xnack+"); - if (p) { + if (HipTest::isXnackOn()) { int *Ptr = nullptr, SIZE = sizeof(int); // Allocating hipMallocManaged() memory @@ -146,11 +138,7 @@ HIP_TEST_CASE(Unit_malloc_CoherentTstWthAdvise) { HIP_TEST_CASE(Unit_mmap_CoherentTst) { CHECK_PCIE_ATOMIC_SUPPORT CHECK_MANAGED_MEMORY_SUPPORT - hipDeviceProp_t prop; - HIPCHECK(hipGetDeviceProperties(&prop, 0)); - char *p = NULL; - p = strstr(prop.gcnArchName, "xnack+"); - if (p) { + if (HipTest::isXnackOn()) { bool HmmMem = true; int *Ptr = reinterpret_cast(mmap(NULL, sizeof(int), PROT_READ | PROT_WRITE, @@ -170,11 +158,7 @@ HIP_TEST_CASE(Unit_mmap_CoherentTst) { #if HT_AMD HIP_TEST_CASE(Unit_mmap_CoherentTstWthAdvise) { CHECK_MANAGED_MEMORY_SUPPORT - hipDeviceProp_t prop; - HIPCHECK(hipGetDeviceProperties(&prop, 0)); - char *p = NULL; - p = strstr(prop.gcnArchName, "xnack+"); - if (p) { + if (HipTest::isXnackOn()) { int SIZE = sizeof(int); int *Ptr = reinterpret_cast(mmap(NULL, SIZE, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, 0, 0)); diff --git a/projects/hip-tests/catch/stress/memory/hipHmmOvrSubscriptionTst.cc b/projects/hip-tests/catch/stress/memory/hipHmmOvrSubscriptionTst.cc index 91dcc74612dd..5cc33fb939eb 100644 --- a/projects/hip-tests/catch/stress/memory/hipHmmOvrSubscriptionTst.cc +++ b/projects/hip-tests/catch/stress/memory/hipHmmOvrSubscriptionTst.cc @@ -22,13 +22,8 @@ HIP_TEST_CASE(Stress_HMM_OverSubscriptionTst) { const bool hmm = HipTest::isManagedMemorySupportedOnDevice(0); bool shouldRun = []() -> bool { -#if HT_AMD // For AMD this gcn arch needs to have xnack+ - int device = 0; - hipDeviceProp_t props{}; - HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&props, device)); - std::string arch(props.gcnArchName); - return arch.find("xnack+") != std::string::npos; +#if HT_AMD // For AMD this needs XNACK on + return HipTest::isXnackOn(); #else // For CUDA this depends on SM and attribute check should be fine return true; #endif diff --git a/projects/hip-tests/catch/stress/memory/hipHostRegisterStress.cc b/projects/hip-tests/catch/stress/memory/hipHostRegisterStress.cc index 456923c23828..67e5a3cc4617 100644 --- a/projects/hip-tests/catch/stress/memory/hipHostRegisterStress.cc +++ b/projects/hip-tests/catch/stress/memory/hipHostRegisterStress.cc @@ -42,9 +42,8 @@ static __global__ void Inc(uint8_t* Ad) { HIP_TEST_CASE(Stress_hipHostRegister_Oversubscription) { hipDeviceProp_t prop; HIP_CHECK(hipGetDeviceProperties(&prop, 0)); - std::string arch = prop.gcnArchName; #if HT_AMD - if (std::string::npos == arch.find("xnack+")) { + if (!HipTest::isXnackOn()) { HIP_SKIP_TEST(HipTest::SkipReason::kGpuXnackNotEnabled); } #endif diff --git a/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc b/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc index f74aa2068085..73448391214d 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc @@ -242,15 +242,8 @@ HIP_TEST_CASE(Unit_hipMemAdvise_TstAccessedByFlg4) { successfully without hang or segfault*/ #if __linux__ && HT_AMD HIP_TEST_CASE(Unit_hipMemAdvise_TstAlignedAllocMem) { - // The following code block checks for xnack+ - // so as to skip if the device is not xnack+ - hipDeviceProp_t prop; - int device; - HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - std::string gfxName(prop.gcnArchName); - - if (gfxName.find("xnack+") != std::string::npos) { + // Skip if XNACK is off + if (HipTest::isXnackOn()) { int pageMemAccess = 0; HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, hipDeviceAttributePageableMemoryAccess, 0)); WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess); @@ -290,13 +283,7 @@ HIP_TEST_CASE(Unit_hipMemAdvise_TstAlignedAllocMem_XNACK) { HIP_SKIP_TEST("cannot set XNACK via environment."); } - hipDeviceProp_t prop; - int device; - HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - std::string gfxName(prop.gcnArchName); - - if (gfxName.find("xnack+") != std::string::npos) { + if (HipTest::isXnackOn()) { hip::SpawnProc proc("hipMemAdviseTstAlignedAllocMem", true); REQUIRE(proc.run() == 0); } else { diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc index ac05a86e2e50..b51d86a552d3 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipGetProcAddressVmmApis.cc @@ -31,16 +31,7 @@ HIP_TEST_CASE(Unit_hipGetProcAddress_VMM) { HIP_SKIP_TEST(HipTest::SkipReason::kVmmUnsupported); } - hipDeviceProp_t devProp; - int device; - bool xnackEnabled = false; - HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&devProp, device)); - std::string gfxName(devProp.gcnArchName); - - if (gfxName.find("xnack+") != std::string::npos) { - xnackEnabled = true; - } + const bool xnackEnabled = HipTest::isXnackOn(); void* hipMemGetAllocationGranularity_ptr = nullptr; void* hipMemCreate_ptr = nullptr; From 00da8b48ecd285cfc4a3e57a376011a0e5b805e6 Mon Sep 17 00:00:00 2001 From: Ioannis Assiouras Date: Wed, 14 Oct 2026 10:25:40 +0000 Subject: [PATCH 2/2] test(hip-tests): drop redundant pageable query and document isXnackOn MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude --- projects/hip-tests/catch/include/hip_test_common.hh | 3 ++- projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc | 5 +---- 2 files changed, 3 insertions(+), 5 deletions(-) diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index dcafce8fe612..ca4f10198e55 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -524,7 +524,8 @@ inline bool checkConcurrentKernels(int num_devices) { return true; } -// Stall-only XNACK GPUs never report xnack+ in gcnArchName; ask for pageable access instead. +// The runtime reports pageable memory access iff XNACK is on; stall-only XNACK GPUs never +// report xnack+ in gcnArchName. inline bool isXnackOn() { #if HT_AMD int device = 0; diff --git a/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc b/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc index 73448391214d..bd2e5ca430b6 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc @@ -244,10 +244,7 @@ HIP_TEST_CASE(Unit_hipMemAdvise_TstAccessedByFlg4) { HIP_TEST_CASE(Unit_hipMemAdvise_TstAlignedAllocMem) { // Skip if XNACK is off if (HipTest::isXnackOn()) { - int pageMemAccess = 0; - HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, hipDeviceAttributePageableMemoryAccess, 0)); - WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess); - if (HipTest::isManagedMemorySupportedOnDevice(0) && (pageMemAccess == 1)) { + if (HipTest::isManagedMemorySupportedOnDevice(0)) { int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123; // Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz)); Mllc = reinterpret_cast(aligned_alloc(4096, 4096 * 4));