diff --git a/tests/catch/multiproc/hipMemCoherencyTstMProc.cc b/tests/catch/multiproc/hipMemCoherencyTstMProc.cc
index 8e36c4331d..10d6e27f6d 100644
--- a/tests/catch/multiproc/hipMemCoherencyTstMProc.cc
+++ b/tests/catch/multiproc/hipMemCoherencyTstMProc.cc
@@ -113,9 +113,8 @@ TEST_CASE("Unit_malloc_CoherentTst") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no*/
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -128,8 +127,8 @@ TEST_CASE("Unit_malloc_CoherentTst") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -138,16 +137,11 @@ TEST_CASE("Unit_malloc_CoherentTst") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -208,9 +202,8 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -223,8 +216,8 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -233,16 +226,11 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -305,9 +293,8 @@ TEST_CASE("Unit_mmap_CoherentTst") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -320,8 +307,8 @@ TEST_CASE("Unit_mmap_CoherentTst") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is not applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -330,16 +317,11 @@ TEST_CASE("Unit_mmap_CoherentTst") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -403,9 +385,8 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
     REQUIRE(false);
   }
 
-  /* GpuId[0] for gfx906 exists--> 1 for yes and 0 for no
-     GpuId[0] for gfx908 exists--> 1 for yes and 0 for no*/
-  int GpuId[2] = {0, 0};
+  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
+  int GpuId[1] = {0};
   p = fork();
 
   if (p < 0) {
@@ -418,8 +399,8 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
     //  Read string from child and close reading end.
     read(fd1[0], GpuId, 2 * sizeof(int));
     close(fd1[0]);
-    if ((GpuId[0] == 1) || (GpuId[1] == 1)) {
-      WARN("This test is not applicable on MI60 & MI100."
+    if (GpuId[0] == 0) {
+      WARN("This test is applicable for MI200."
               "Skipping the test!!");
       exit(0);
     }
@@ -428,16 +409,11 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
     hipDeviceProp_t prop;
     HIPCHECK(hipGetDeviceProperties(&prop, 0));
     char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx906");
+    p = strstr(prop.gcnArchName, "gfx90a");
     if (p) {
-      WARN("gfx906 gpu found on this system!!");
+      WARN("gfx90a gpu found on this system!!");
       GpuId[0] = 1;
     }
-    p = strstr(prop.gcnArchName, "gfx908");
-    if (p) {
-      WARN("gfx908 gpu found on this system!!");
-      GpuId[1] = 1;
-    }
     // Write concatenated string and close writing end
     write(fd1[1], GpuId, 2 * sizeof(int));
     close(fd1[1]);
@@ -497,7 +473,7 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -538,7 +514,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -579,7 +555,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -587,7 +563,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
     int *Ptr = nullptr, *PtrD = nullptr, SIZE = sizeof(int);
     YES_COHERENT = false;
     // Allocating hipHostMalloc() memory
-     HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
+    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
     *Ptr = 4;
     hipStream_t strm;
     HIP_CHECK(hipStreamCreate(&strm));
@@ -620,7 +596,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") {
   if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
   int stat = 0;
@@ -662,33 +638,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
-    *Ptr = 4;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
+      *Ptr = 4;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif
@@ -700,33 +686,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
-    *Ptr = 1;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocPortable));
+      *Ptr = 1;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif
@@ -737,33 +733,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
-    *Ptr = 4;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocWriteCombined));
+      *Ptr = 4;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif
@@ -775,33 +781,43 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") {
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") {
   if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
+      WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!");
       REQUIRE(false);
   }
-  int stat = 0;
-  if (fork() == 0) {  // child process
-    int *Ptr = nullptr, SIZE = sizeof(int);
-    bool HmmMem = false;
-    YES_COHERENT = false;
-    // Allocating hipHostMalloc() memory
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
-    *Ptr = 1;
-    TstCoherency(Ptr, HmmMem);
-    if (YES_COHERENT) {
-      // exit() with code 10 which indicates pass
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(10);
-    } else {
-      // exit() with code 9 which indicates fail
-      HIP_CHECK(hipHostFree(Ptr));
-      exit(9);
-    }
-  } else {  // parent process
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
+  int stat = 0, Pageable = 0;
+
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable) {
+    if (fork() == 0) {  // child process
+      int *Ptr = nullptr, SIZE = sizeof(int);
+      bool HmmMem = false;
+      YES_COHERENT = false;
+      // Allocating hipHostMalloc() memory
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocNumaUser));
+      *Ptr = 1;
+      TstCoherency(Ptr, HmmMem);
+      if (YES_COHERENT) {
+        // exit() with code 10 which indicates pass
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(10);
+      } else {
+        // exit() with code 9 which indicates fail
+        HIP_CHECK(hipHostFree(Ptr));
+        exit(9);
+      }
+    } else {  // parent process
+      wait(&stat);
+      int Result = WEXITSTATUS(stat);
+      if (Result != 10) {
+        REQUIRE(false);
+      }
     }
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+               "attribute. Hence skipping the test with Pass result.\n"); 
   }
 }
 #endif
diff --git a/tests/catch/unit/memory/hipMemCoherencyTst.cc b/tests/catch/unit/memory/hipMemCoherencyTst.cc
index f04f9835d0..73bf4aeb09 100644
--- a/tests/catch/unit/memory/hipMemCoherencyTst.cc
+++ b/tests/catch/unit/memory/hipMemCoherencyTst.cc
@@ -91,54 +91,39 @@ static void TstCoherency(int *Ptr, bool HmmMem) {
   }
 }
 
-static int HmmAttrPrint() {
-  int managed = 0;
-  INFO("The following are the attribute values related to HMM for"
-         " device 0:\n");
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-              hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
-  INFO("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributeConcurrentManagedAccess, 0));
-  INFO("hipDeviceAttributeConcurrentManagedAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributePageableMemoryAccess, 0));
-  INFO("hipDeviceAttributePageableMemoryAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
-  INFO("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
-         << managed);
-
-  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
-                                  0));
-  INFO("hipDeviceAttributeManagedMemory: " << managed);
-  return managed;
-}
-
-
 /* Test case description: The following test validates if fine grain
    behavior is observed or not with memory allocated using hipHostMalloc()*/
 // The following tests are disabled for Nvidia as they are not consistently
 // passing
 #if HT_AMD
 TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
-  int *Ptr = nullptr,  SIZE = sizeof(int);
+  int *Ptr = nullptr,  SIZE = sizeof(int), Pageable = 0;
   bool HmmMem = false;
   YES_COHERENT = false;
-  // Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
-  SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
-  }
-  SECTION("hipHostMalloc with Default flag") {
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
-  }
-  SECTION("hipHostMalloc with hipHostMallocMapped flag") {
-    HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
-  }
 
-  TstCoherency(Ptr, HmmMem);
-  HIP_CHECK(hipHostFree(Ptr));
-  REQUIRE(YES_COHERENT);
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  if (Pageable == 1) {
+    // Allocating hipHostMalloc() memory with hipHostMallocCoherent flag
+    SECTION("hipHostMalloc with hipHostMallocCoherent flag") {
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocCoherent));
+    }
+    SECTION("hipHostMalloc with Default flag") {
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE));
+    }
+    SECTION("hipHostMalloc with hipHostMallocMapped flag") {
+      HIP_CHECK(hipHostMalloc(&Ptr, SIZE, hipHostMallocMapped));
+    }
+
+    TstCoherency(Ptr, HmmMem);
+    HIP_CHECK(hipHostFree(Ptr));
+    REQUIRE(YES_COHERENT);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess "
+             "attribute. Hence skipping the test with Pass result.\n");
+  }
 }
 #endif
 
@@ -149,12 +134,19 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") {
 // passing
 #if HT_AMD
 TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
-  int *Ptr = nullptr, SIZE = sizeof(int);
+  int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0, managed = 0;
   bool HmmMem = true;
   YES_COHERENT = false;
 
-  int managed =  HmmAttrPrint();
-  if (managed == 1) {
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+
+  if (managed == 1 && Pageable == 1) {
     // Allocating hipMallocManaged() memory
     SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
       HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
@@ -166,8 +158,8 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
     HIP_CHECK(hipFree(Ptr));
     REQUIRE(YES_COHERENT);
   } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
+    SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
+           "device attribute. Hence skipping the test with Pass result.\n");
   }
 }
 #endif
@@ -175,30 +167,40 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") {
 /* Test case description: The following test validates if memory access is fine
    with memory allocated using hipMallocManaged() and CoarseGrain Advise*/
 TEST_CASE("Unit_hipMallocManaged_CoherentTstWthAdvise") {
-  int *Ptr = nullptr, SIZE = sizeof(int);
+  int *Ptr = nullptr, SIZE = sizeof(int), managed = 0;
   YES_COHERENT = false;
-  // Allocating hipMallocManaged() memory
-  SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
-    HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
-  }
-  SECTION("hipMallocManaged with hipMemAttachHost flag") {
-    HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost));
-  }
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+
+  if (managed == 1) {
+    // Allocating hipMallocManaged() memory
+    SECTION("hipMallocManaged with hipMemAttachGlobal flag") {
+      HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal));
+    }
+    SECTION("hipMallocManaged with hipMemAttachHost flag") {
+      HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachHost));
+    }
 #if HT_AMD
-  HIP_CHECK(hipMemAdvise(Ptr, SIZE, hipMemAdviseSetCoarseGrain, 0));
+    HIP_CHECK(hipMemAdvise(Ptr, SIZE, hipMemAdviseSetCoarseGrain, 0));
 #endif
-  // Initializing Ptr memory with 9
-  *Ptr = 9;
-  hipStream_t strm;
-  HIP_CHECK(hipStreamCreate(&strm));
-  SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
-  HIP_CHECK(hipStreamSynchronize(strm));
-  if (*Ptr == 81) {
-    YES_COHERENT = true;
+    // Initializing Ptr memory with 9
+    *Ptr = 9;
+    hipStream_t strm;
+    HIP_CHECK(hipStreamCreate(&strm));
+    SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
+    HIP_CHECK(hipStreamSynchronize(strm));
+    if (*Ptr == 81) {
+      YES_COHERENT = true;
+    }
+    HIP_CHECK(hipFree(Ptr));
+    HIP_CHECK(hipStreamDestroy(strm));
+    REQUIRE(YES_COHERENT);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+             "attribute. Hence skipping the test with Pass result.\n");
   }
-  HIP_CHECK(hipFree(Ptr));
-  HIP_CHECK(hipStreamDestroy(strm));
-  REQUIRE(YES_COHERENT);
 }
 
 
@@ -226,12 +228,18 @@ TEST_CASE("Unit_hipMalloc_CoherentTst") {
    hipExtMallocWithFlags()*/
 #if HT_AMD
 TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") {
-  int *Ptr = nullptr, SIZE = sizeof(int), InitVal = 9;
+  int *Ptr = nullptr, SIZE = sizeof(int), InitVal = 9, Pageable = 0, managed = 0;
   bool FineGrain = true;
   YES_COHERENT = false;
 
-  int managed =  HmmAttrPrint();
-  if (managed == 1) {
+  HIP_CHECK(hipDeviceGetAttribute(&Pageable,
+                                 hipDeviceAttributePageableMemoryAccess, 0));
+  INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable);
+
+  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
+                                  0));
+  INFO("hipDeviceAttributeManagedMemory: " << managed);
+  if (managed == 1 && Pageable == 1) {
     // Allocating hipExtMallocWithFlags() memory with flags
     SECTION("hipExtMallocWithFlags with hipDeviceMallocFinegrained flag") {
       HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast<void**>(&Ptr), SIZE*2,
@@ -264,8 +272,8 @@ TEST_CASE("Unit_hipExtMallocWithFlags_CoherentTst") {
     HIP_CHECK(hipFree(Ptr));
     REQUIRE(YES_COHERENT);
   } else {
-    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-           "attribute. Hence skipping the testing with Pass result.\n");
+    SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess"
+           "device attribute. Hence skipping the test with Pass result.\n");
   }
 }
 #endif