diff --git a/catch/multiproc/hipMemCoherencyTstMProc.cc b/catch/multiproc/hipMemCoherencyTstMProc.cc
index 8579aabc2..b9b7da5a1 100644
--- a/catch/multiproc/hipMemCoherencyTstMProc.cc
+++ b/catch/multiproc/hipMemCoherencyTstMProc.cc
@@ -80,13 +80,16 @@ static void TstCoherency(int *Ptr, bool HmmMem) {
 
   // Getting gpu frequency
   if (IsGfx11()) {
-    HIPCHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeWallClockRate, 0));
+    HIPCHECK(hipDeviceGetAttribute(&peak_clk,
+                                   hipDeviceAttributeWallClockRate, 0));
   } else {
-    HIPCHECK(hipDeviceGetAttribute(&peak_clk, hipDeviceAttributeClockRate, 0));
+    HIPCHECK(hipDeviceGetAttribute(&peak_clk,
+                                   hipDeviceAttributeClockRate, 0));
   }
 
   if (!HmmMem) {
-    HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void **>(&Dptr), Ptr, 0));
+    HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void **>(&Dptr),
+                                      Ptr, 0));
     if (IsGfx11()) {
       CoherentTst_gfx11<<<1, 1, 0, strm>>>(Dptr, peak_clk);
     } else {
@@ -121,57 +124,13 @@ static void 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") {
-  if ((setenv("HSA_XNACK", "1", 1)) != 0) {
-    WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
-    REQUIRE(false);
-  }
-  // The following code block is used to check for gfx906/8 so as to skip if
-  // any of the gpus available
-  int fd1[2];  // Used to store two ends of first pipe
-  pid_t p;
-  if (pipe(fd1) == -1) {
-    fprintf(stderr, "Pipe Failed");
-    REQUIRE(false);
-  }
-
-  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no*/
-  int GpuId[1] = {0};
-  p = fork();
-
-  if (p < 0) {
-    fprintf(stderr, "fork Failed");
-    REQUIRE(false);
-  } else if (p > 0) {  // parent process
-    close(fd1[1]);  // Close writing end of first pipe
-    // Wait for child to send a string
-    wait(NULL);
-    //  Read string from child and close reading end.
-    read(fd1[0], GpuId, 2 * sizeof(int));
-    close(fd1[0]);
-    if (GpuId[0] == 0) {
-      WARN("This test is applicable for MI200."
-              "Skipping the test!!");
-      exit(0);
-    }
-  } else {    // child process
-    close(fd1[0]);  // Close read end of first pipe
-    hipDeviceProp_t prop;
-    HIPCHECK(hipGetDeviceProperties(&prop, 0));
-    char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx90a");
-    if (p) {
-      WARN("gfx90a gpu found on this system!!");
-      GpuId[0] = 1;
-    }
-    // Write concatenated string and close writing end
-    write(fd1[1], GpuId, 2 * sizeof(int));
-    close(fd1[1]);
-    exit(0);
-  }
-
-  // Test Case execution begins from here
-  int stat = 0;
-  if (fork() == 0) {
+  hipDeviceProp_t prop;
+  HIPCHECK(hipGetDeviceProperties(&prop, 0));
+  char *p = NULL;
+  p = strstr(prop.gcnArchName, "xnack+");
+  if (p) {
+    // Test Case execution begins from here
+    int stat = 0;
     int managed = 0;
     HIPCHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));
@@ -183,23 +142,10 @@ TEST_CASE("Unit_malloc_CoherentTst") {
       Ptr = reinterpret_cast<int*>(malloc(SIZE));
       TstCoherency(Ptr, HmmMem);
       free(Ptr);
-      if (YES_COHERENT) {
-        // exit() with code 10 which indicates pass
-        exit(10);
-      } else {
-        // exit() with code 9 which indicates fail
-        exit(9);
-      }
-    } else {
-      SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-             "attribute. Hence skipping the testing with Pass result.\n");
-    }
+      REQUIRE(YES_COHERENT);
+    } 
   } else {
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
-    }
+    HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n");
   }
 }
 #endif
@@ -210,55 +156,12 @@ TEST_CASE("Unit_malloc_CoherentTst") {
 // The following test is failing on Nvidia platform hence disabling it for now
 #if HT_AMD
 TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
-  if ((setenv("HSA_XNACK", "1", 1)) != 0) {
-      WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
-      REQUIRE(false);
-  }
-  // The following code block is used to check for gfx906/8 so as to skip if
-  // any of the gpus available
-  int fd1[2];  // Used to store two ends of first pipe
-  pid_t p;
-  if (pipe(fd1) == -1) {
-    fprintf(stderr, "Pipe Failed");
-    REQUIRE(false);
-  }
-
-  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
-  int GpuId[1] = {0};
-  p = fork();
-
-  if (p < 0) {
-    fprintf(stderr, "fork Failed");
-    REQUIRE(false);
-  } else if (p > 0) {  // parent process
-    close(fd1[1]);  // Close writing end of first pipe
-    // Wait for child to send a string
-    wait(NULL);
-    //  Read string from child and close reading end.
-    read(fd1[0], GpuId, 2 * sizeof(int));
-    close(fd1[0]);
-    if (GpuId[0] == 0) {
-      WARN("This test is applicable for MI200."
-              "Skipping the test!!");
-      exit(0);
-    }
-  } else {    // child process
-    close(fd1[0]);  // Close read end of first pipe
-    hipDeviceProp_t prop;
-    HIPCHECK(hipGetDeviceProperties(&prop, 0));
-    char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx90a");
-    if (p) {
-      WARN("gfx90a gpu found on this system!!");
-      GpuId[0] = 1;
-    }
-    // Write concatenated string and close writing end
-    write(fd1[1], GpuId, 2 * sizeof(int));
-    close(fd1[1]);
-    exit(0);
-  }
-  int stat = 0;
-  if (fork() == 0) {
+  hipDeviceProp_t prop;
+  HIPCHECK(hipGetDeviceProperties(&prop, 0));
+  char *p = NULL;
+  p = strstr(prop.gcnArchName, "xnack+");
+  if (p) {
+    int stat = 0;
     int managed = 0;
     HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));
@@ -273,25 +176,10 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") {
       SquareKrnl<<<1, 1, 0, strm>>>(Ptr);
       HIP_CHECK(hipStreamSynchronize(strm));
       HIP_CHECK(hipStreamDestroy(strm));
-      if (*Ptr == 16) {
-        // exit() with code 10 which indicates pass
-        free(Ptr);
-        exit(10);
-      } else {
-        // exit() with code 9 which indicates fail
-        free(Ptr);
-        exit(9);
-      }
-    } else {
-      SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-              "attribute. Hence skipping the testing with Pass result.\n");
+      REQUIRE (*Ptr == 16);
     }
   } else {
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
-    }
+    HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n");
   }
 }
 #endif
@@ -301,55 +189,12 @@ 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") {
-  if ((setenv("HSA_XNACK", "1", 1)) != 0) {
-    WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
-    REQUIRE(false);
-  }
-  // The following code block is used to check for gfx906/8 so as to skip if
-  // any of the gpus available
-  int fd1[2];  // Used to store two ends of first pipe
-  pid_t p;
-  if (pipe(fd1) == -1) {
-    fprintf(stderr, "Pipe Failed");
-    REQUIRE(false);
-  }
-
-  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
-  int GpuId[1] = {0};
-  p = fork();
-
-  if (p < 0) {
-    fprintf(stderr, "fork Failed");
-    REQUIRE(false);
-  } else if (p > 0) {  // parent process
-    close(fd1[1]);  // Close writing end of first pipe
-    // Wait for child to send a string
-    wait(NULL);
-    //  Read string from child and close reading end.
-    read(fd1[0], GpuId, 2 * sizeof(int));
-    close(fd1[0]);
-    if (GpuId[0] == 0) {
-      WARN("This test is not applicable for MI200."
-              "Skipping the test!!");
-      exit(0);
-    }
-  } else {    // child process
-    close(fd1[0]);  // Close read end of first pipe
-    hipDeviceProp_t prop;
-    HIPCHECK(hipGetDeviceProperties(&prop, 0));
-    char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx90a");
-    if (p) {
-      WARN("gfx90a gpu found on this system!!");
-      GpuId[0] = 1;
-    }
-    // Write concatenated string and close writing end
-    write(fd1[1], GpuId, 2 * sizeof(int));
-    close(fd1[1]);
-    exit(0);
-  }
-  int stat = 0;
-  if (fork() == 0) {
+  hipDeviceProp_t prop;
+  HIPCHECK(hipGetDeviceProperties(&prop, 0));
+  char *p = NULL;
+  p = strstr(prop.gcnArchName, "xnack+");
+  if (p) {
+    int stat = 0;
     int managed = 0;
     HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));
@@ -369,21 +214,10 @@ TEST_CASE("Unit_mmap_CoherentTst") {
       if (err != 0) {
         WARN("munmap failed\n");
       }
-      if (YES_COHERENT) {
-        exit(10);
-      } else {
-        exit(9);
-      }
-    } else {
-      SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-             "attribute. Hence skipping the testing with Pass result.\n");
-    }
+      REQUIRE(YES_COHERENT);
+    } 
   } else {
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
-    }
+    HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n");
   }
 }
 #endif
@@ -393,55 +227,12 @@ TEST_CASE("Unit_mmap_CoherentTst") {
 // The following test is failing on Nvidia platform hence disabling it for now
 #if HT_AMD
 TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
-  if ((setenv("HSA_XNACK", "1", 1)) != 0) {
-    WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
-    REQUIRE(false);
-  }
-  // The following code block is used to check for gfx906/8 so as to skip if
-  // any of the gpus available
-  int fd1[2];  // Used to store two ends of first pipe
-  pid_t p;
-  if (pipe(fd1) == -1) {
-    fprintf(stderr, "Pipe Failed");
-    REQUIRE(false);
-  }
-
-  /* GpuId[0] for gfx90a exists--> 1 for yes and 0 for no */
-  int GpuId[1] = {0};
-  p = fork();
-
-  if (p < 0) {
-    fprintf(stderr, "fork Failed");
-    REQUIRE(false);
-  } else if (p > 0) {  // parent process
-    close(fd1[1]);  // Close writing end of first pipe
-    // Wait for child to send a string
-    wait(NULL);
-    //  Read string from child and close reading end.
-    read(fd1[0], GpuId, 2 * sizeof(int));
-    close(fd1[0]);
-    if (GpuId[0] == 0) {
-      WARN("This test is applicable for MI200."
-              "Skipping the test!!");
-      exit(0);
-    }
-  } else {    // child process
-    close(fd1[0]);  // Close read end of first pipe
-    hipDeviceProp_t prop;
-    HIPCHECK(hipGetDeviceProperties(&prop, 0));
-    char *p = NULL;
-    p = strstr(prop.gcnArchName, "gfx90a");
-    if (p) {
-      WARN("gfx90a gpu found on this system!!");
-      GpuId[0] = 1;
-    }
-    // Write concatenated string and close writing end
-    write(fd1[1], GpuId, 2 * sizeof(int));
-    close(fd1[1]);
-    exit(0);
-  }
-  int stat = 0;
-  if (fork() == 0) {
+  hipDeviceProp_t prop;
+  HIPCHECK(hipGetDeviceProperties(&prop, 0));
+  char *p = NULL;
+  p = strstr(prop.gcnArchName, "xnack+");
+  if (p) {
+    int stat = 0;
     int managed = 0;
     HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
                                     0));
@@ -464,26 +255,15 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") {
       bool IfTstPassed = false;
       if (*Ptr == 81) {
         IfTstPassed = true;
-      }
+      } 
       int err = munmap(Ptr, SIZE);
       if (err != 0) {
         WARN("munmap failed\n");
       }
-      if (IfTstPassed) {
-        exit(10);
-      } else {
-        exit(9);
-      }
-    } else {
-      SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-             "attribute. Hence skipping the testing with Pass result.\n");
-    }
+      REQUIRE(IfTstPassed);
+    } 
   } else {
-    wait(&stat);
-    int Result = WEXITSTATUS(stat);
-    if (Result != 10) {
-      REQUIRE(false);
-    }
+    HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n");
   }
 }
 #endif
diff --git a/catch/unit/memory/hipHmmOvrSubscriptionTst.cc b/catch/unit/memory/hipHmmOvrSubscriptionTst.cc
index de5d48ee4..5ff9e469b 100644
--- a/catch/unit/memory/hipHmmOvrSubscriptionTst.cc
+++ b/catch/unit/memory/hipHmmOvrSubscriptionTst.cc
@@ -24,26 +24,12 @@ THE SOFTWARE.
    feature which is part of HMM.*/
 
 #include <hip_test_common.hh>
-#ifdef __linux__
-#include <sys/types.h>
-#include <sys/ipc.h>
-#include <sys/shm.h>
-#include <sys/stat.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-#include <unistd.h>
-#include <sys/wait.h>
-#endif
-#include <list>
 
 #define INIT_VAL 2.5
 #define NUM_ELMS 268435456  // 268435456 * 4 = 1GB
 #define ITERATIONS 10
 #define ONE_GB 1024 * 1024 * 1024
 
-static void GetTotGpuMem(int *TotMem);
-static void DisplayHmmFlgs(int *Signal);
 // Kernel function
 __global__ void Square(int n, float *x) {
   int index = blockIdx.x * blockDim.x + threadIdx.x;
@@ -90,124 +76,29 @@ static void OneGBMemTest(int dev) {
   HIP_CHECK(hipStreamDestroy(strm));
 }
 
-static void GetTotGpuMem(int *TotMem) {
-  size_t FreeMem, TotGpuMem;
-  HIP_CHECK(hipMemGetInfo(&FreeMem, &TotGpuMem));
-  TotMem[0] = (TotGpuMem/(ONE_GB));
-  TotMem[1] = 1;
-}
-
-static void DisplayHmmFlgs(int *Signal) {
-  int managed = 0;
-  WARN("The following are the attribute values related to HMM for"
-         " device 0:\n");
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                    hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
-  WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributeConcurrentManagedAccess, 0));
-  WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-                                 hipDeviceAttributePageableMemoryAccess, 0));
-  WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
-  HIP_CHECK(hipDeviceGetAttribute(&managed,
-              hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
-  WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
-          << managed);
-
-  HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
-                                  0));
-  WARN("hipDeviceAttributeManagedMemory: " << managed);
-
-  // Checking for Vega20 or MI100
+TEST_CASE("Unit_HMM_OverSubscriptionTst") {
+  // Checking if xnack is enabled
   hipDeviceProp_t prop;
   HIP_CHECK(hipGetDeviceProperties(&prop, 0));
   char *p = NULL;
-  p = strstr(prop.gcnArchName, "gfx906");
+  p = strstr(prop.gcnArchName, "xnack+");
   if (p) {
-    WARN("This system has MI60 gpu hence OverSubscription test will be");
-    WARN(" skipped");
-    Signal[2] = 1;
-  }
-  p = strstr(prop.gcnArchName, "gfx908");
-  if (p) {
-    WARN("This system has MI100 gpu hence OverSubscription test will be");
-    WARN(" skipped");
-    Signal[2] = 1;
-  }
-  Signal[1] = managed;
-  Signal[0] = 1;
-}
-
-TEST_CASE("Unit_HMM_OverSubscriptionTst") {
-  int HmmEnabled = 0;
-  // The following Shared Mem is to get Max GPU Mem
-  // The size requested is for three ints
-  // 1) To get Max GPU Mem in GB
-  // 2) To Signal parent that req. info is available to consume
-  // 3) To know if MI60 or MI100 gpu are there in the system
-  key_t key = ftok("shmTotMem", 66);
-  int shmid = shmget(key, (3 * sizeof(int)), 0666|IPC_CREAT);
-  int *TotGpuMem = reinterpret_cast<int*>(shmat(shmid, NULL, 0));
-  TotGpuMem[0] = 0; TotGpuMem[1] = 0;
-  // The following function DisplayHmmFlgs() displays the flag values related
-  // to HMM and also sends us ManagedMemory attribute value
-  if (fork() == 0) {
-    DisplayHmmFlgs(TotGpuMem);
-    exit(1);
-  }
-  while (TotGpuMem[0] == 0) {
-      sleep(2);
-  }
-  // The following if block will skip test if either of MI60 or MI100 is found
-  if (TotGpuMem[2] == 1) {
-    SUCCEED("Test is skipped!!");
-    REQUIRE(true);
+    size_t FreeMem, TotGpuMem;
+    HIP_CHECK(hipMemGetInfo(&FreeMem, &TotGpuMem));
+    int NumGB = (TotGpuMem/(ONE_GB));
+    int TotalThreads = (NumGB + 10);
+    WARN("Launching " << TotalThreads);
+    WARN(" processes to test OverSubscription.");
+
+	std::thread Thrds[NumGB];
+	
+	for (int k = 0; k < TotalThreads; ++k) {
+		Thrds[k] = std::thread(OneGBMemTest, 0);
+	}
+	for (int k = 0; k < TotalThreads; ++k) {
+		Thrds[k].join();
+	}
   } else {
-    HmmEnabled = TotGpuMem[1];
-
-    // Re-setting the shared memory values for further usage
-    TotGpuMem[0] = 0;
-    TotGpuMem[1] = 0;
-
-    std::list<pid_t> PidLst;
-    // The following function gets the MaxGpu memory in GBs and also launches
-    // OverSubscription test
-    if (HmmEnabled) {
-      if ((setenv("HSA_XNACK", "1", 1)) != 0) {
-        WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
-        REQUIRE(false);
-      }
-      if (fork() == 0) {
-        GetTotGpuMem(TotGpuMem);
-      }
-      while (TotGpuMem[1] == 0) {
-        sleep(2);
-      }
-      int NumGB = TotGpuMem[0], TotalThreads = (NumGB + 10);
-      WARN("Launching " << TotalThreads);
-      WARN(" processes to test OverSubscription.");
-      pid_t pid;
-      for (int k = 0; k < TotalThreads; ++k) {
-        pid = fork();
-        PidLst.push_back(pid);
-        if (pid == 0) {
-          OneGBMemTest(0);
-          exit(10);
-        }
-      }
-    } else {
-      SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
-              "attribute. Hence skipping the testing with Pass result.\n");
-    }
-    int status;
-    for (pid_t pd : PidLst) {
-      waitpid(pd, &status, 0);
-      if (!(WIFEXITED(status))) {
-        REQUIRE(false);
-      }
-    }
+	HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n");
   }
-  shmdt(TotGpuMem);
-  shmctl(shmid, IPC_RMID, NULL);
 }
diff --git a/catch/unit/memory/hipMemAdvise_old.cc b/catch/unit/memory/hipMemAdvise_old.cc
index b95199805..30510b5b2 100644
--- a/catch/unit/memory/hipMemAdvise_old.cc
+++ b/catch/unit/memory/hipMemAdvise_old.cc
@@ -651,87 +651,54 @@ TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") {
   }
 }
 
-
 /* Allocate memory using aligned_alloc(), assign PreferredLocation flag to
    the allocated memory and launch a kernel. Kernel should get executed
    successfully without hang or segfault*/
 #if __linux__ && HT_AMD
 TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
-  if ((setenv("HSA_XNACK", "1", 1)) != 0) {
-    WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
-    REQUIRE(false);
-  }
-  // The following code block checks for gfx90a so as to skip if the device is not MI200
-
+  // 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 == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
+  if (gfxName.find("xnack+") != std::string::npos) {
     int stat = 0;
-    if (fork() == 0) {
-      // The below part should be inside fork
-      int managedMem = 0, pageMemAccess = 0;
-      HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess,
-                hipDeviceAttributePageableMemoryAccess, 0));
-      WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess);
-
-      HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0));
-      WARN("hipDeviceAttributeManagedMemory: " << managedMem);
-      if ((managedMem == 1) && (pageMemAccess == 1)) {
-        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));
-        for (int i = 0; i < NumElms; ++i) {
-          Mllc[i] = InitVal;
-        }
-        hipStream_t strm;
-        int DataMismatch = 0;
-        HIP_CHECK(hipStreamCreate(&strm));
-        // The following hipMemAdvise() call is made to know if advise on
-        // aligned_alloc() is causing any issue
-        HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0));
-        HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm));
-        HIP_CHECK(hipStreamSynchronize(strm));
-        MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms);
-        HIP_CHECK(hipStreamSynchronize(strm));
+    int managedMem = 0, pageMemAccess = 0;
+    HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess,
+              hipDeviceAttributePageableMemoryAccess, 0));
+    WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess);
+    HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0));
+    WARN("hipDeviceAttributeManagedMemory: " << managedMem);
+    if ((managedMem == 1) && (pageMemAccess == 1)) {
+      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));
+      for (int i = 0; i < NumElms; ++i) {
+        Mllc[i] = InitVal;
+      }
+      hipStream_t strm;
+      int DataMismatch = 0;
+      HIP_CHECK(hipStreamCreate(&strm));
+      // The following hipMemAdvise() call is made to know if advise on
+      // aligned_alloc() is causing any issue
+      HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0));
+      HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm));
+      HIP_CHECK(hipStreamSynchronize(strm));
+      MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms);
+      HIP_CHECK(hipStreamSynchronize(strm));
         for (int i = 0; i < NumElms; ++i) {
           if (Mllc[i] != (InitVal + 10)) {
             DataMismatch++;
           }
         }
-        if (DataMismatch != 0) {
-          WARN("DataMismatch observed!!");
-          exit(9);  // 9 for failure
-        } else {
-          exit(10);  // 10 for Pass result
-        }
-      } else {
-        SUCCEED("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess "
-                "attribute. Hence skipping the testing with Pass result.\n");
-        exit(Catch::ResultDisposition::ContinueOnFailure);
-      }
-    } else {
-      wait(&stat);
-      int Result = WEXITSTATUS(stat);
-      if (Result == Catch::ResultDisposition::ContinueOnFailure) {
-        WARN("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess "
-             "attribute. Hence skipping the testing with Pass result.\n");
-      } else {
-        if (Result != 10) {
-          REQUIRE(false);
-        }
-      }
-    }
+      REQUIRE(DataMismatch == 0);
+      } 
   } else {
-      SUCCEED("Memory model feature is only supported for gfx90a, Hence"
-              "skipping the testcase for this GPU " << device);
-      WARN("Memory model feature is only supported for gfx90a, Hence"
-              "skipping the testcase for this GPU " << device);
-  }
-
+      HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n");
+    }
 }
 #endif