diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh
index 47b7bb6183a3..ca4f10198e55 100644
--- a/projects/hip-tests/catch/include/hip_test_common.hh
+++ b/projects/hip-tests/catch/include/hip_test_common.hh
@@ -524,13 +524,18 @@ inline bool checkConcurrentKernels(int num_devices) {
   return true;
 }
 
+// The runtime reports pageable memory access iff XNACK is on; stall-only XNACK GPUs never
+// report xnack+ in gcnArchName.
 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<int *>(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<int *>(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..bd2e5ca430b6 100644
--- a/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc
+++ b/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc
@@ -242,19 +242,9 @@ 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) {
-    int pageMemAccess = 0;
-    HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, hipDeviceAttributePageableMemoryAccess, 0));
-    WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess);
-    if (HipTest::isManagedMemorySupportedOnDevice(0) && (pageMemAccess == 1)) {
+  // Skip if XNACK is off
+  if (HipTest::isXnackOn()) {
+    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<int*>(aligned_alloc(4096, 4096 * 4));
@@ -290,13 +280,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;