From 16a737d8293d011ef4faf9a98e0de487047e2e1c Mon Sep 17 00:00:00 2001 From: vinay birur Date: Mon, 5 Jun 2023 15:09:47 +0530 Subject: [PATCH] SWDEV-389805/SWDEV-403960 - [catch2][dtest] Enable xnack+ check condition Change-Id: I9627d75d0d3258cf261c8e4bfe6c7c3c35c8f9c1 --- catch/multiproc/hipMemCoherencyTstMProc.cc | 306 +++--------------- catch/unit/memory/hipHmmOvrSubscriptionTst.cc | 147 ++------- catch/unit/memory/hipMemAdvise_old.cc | 93 ++---- 3 files changed, 92 insertions(+), 454 deletions(-) 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(&Dptr), Ptr, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&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(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 -#ifdef __linux__ -#include -#include -#include -#include -#include -#include -#include -#include -#include -#endif -#include #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(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 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(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(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