From 6a8719947427338f0d8ddd2d5560cb35ee2f6d75 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 19 Oct 2022 16:48:44 +0530 Subject: [PATCH] SWDEV-308368 - Fix coherent test bugs (#2985) 1. Remove hipDeviceAttributePageableMemoryAccess checking as it is used to check xnack+ rather than fine-grained synchronization. Note: hipHostMalloc() will alloc memory with fine-grained access and atomic by default when env HIP_HOST_COHERENT=1. For more information, see hip_programming_guide. 2. Fix some other bugs. Change-Id: Ieaaed1b15b4d258193d834104b17d9f03a9e4781 [ROCm/hip commit: 15517aaf18a03cc80ae96245b5eacdb145732a41] --- .../multiproc/hipMemCoherencyTstMProc.cc | 232 +++++++----------- .../tests/catch/unit/memory/hipMemAdvise.cc | 10 +- .../catch/unit/memory/hipMemAdviseMmap.cc | 2 +- .../catch/unit/memory/hipMemCoherencyTst.cc | 48 ++-- 4 files changed, 117 insertions(+), 175 deletions(-) diff --git a/projects/hip/tests/catch/multiproc/hipMemCoherencyTstMProc.cc b/projects/hip/tests/catch/multiproc/hipMemCoherencyTstMProc.cc index 10d6e27f6d..863dfc7178 100644 --- a/projects/hip/tests/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/projects/hip/tests/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -40,15 +40,12 @@ __global__ void CoherentTst(int *ptr, int PeakClk) { // Incrementing the value by 1 - int64_t GpuFrq = (PeakClk * 1000); + int64_t GpuFrq = int64_t(PeakClk) * 1000; int64_t StrtTck = clock64(); atomicAdd(ptr, 1); // The following while loop checks the value in ptr for around 3-4 seconds while ((clock64() - StrtTck) <= (3 * GpuFrq)) { - if (*ptr == 3) { - atomicAdd(ptr, 1); - return; - } + if (atomicCAS(ptr, 3, 4) == 3) break; } } @@ -57,8 +54,6 @@ __global__ void SquareKrnl(int *ptr) { *ptr = (*ptr) * (*ptr); } - - // The variable below will work as signal to decide pass/fail static bool YES_COHERENT = false; @@ -641,40 +636,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); REQUIRE(false); } - int stat = 0, Pageable = 0; + int stat = 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); - } + 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 @@ -689,40 +675,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); REQUIRE(false); } - int stat = 0, Pageable = 0; + int stat = 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); - } + 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 @@ -736,45 +713,35 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); REQUIRE(false); } - int stat = 0, Pageable = 0; + int stat = 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); - } + 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 - /* Test Case Description: The following test checks if the memory exhibits fine grain behavior when HIP_HOST_COHERENT is set to 1*/ // The following test is AMD specific test hence skipping for Nvidia @@ -784,42 +751,31 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); REQUIRE(false); } - int stat = 0, Pageable = 0; + int stat = 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); - } + 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/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc b/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc index 96a0318bb8..c0673753a9 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemAdvise.cc @@ -777,7 +777,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") { REQUIRE(false); } } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " "attribute. Hence skipping the testing with Pass result.\n"); } } @@ -809,7 +809,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") { REQUIRE(false); } } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " "attribute. Hence skipping the testing with Pass result.\n"); } } else { @@ -847,7 +847,7 @@ TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") { } HIP_CHECK(hipFree(Hmm)); } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " "attribute. Hence skipping the testing with Pass result.\n"); } } @@ -929,7 +929,7 @@ TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") { HIP_CHECK(hipFree(Hmm)); HIP_CHECK(hipStreamDestroy(strm)); } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " "attribute. Hence skipping the testing with Pass result.\n"); } } @@ -955,7 +955,7 @@ TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") { REQUIRE(false); } } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributePageableMemoryAccess " + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " "attribute. Hence skipping the testing with Pass result.\n"); } } diff --git a/projects/hip/tests/catch/unit/memory/hipMemAdviseMmap.cc b/projects/hip/tests/catch/unit/memory/hipMemAdviseMmap.cc index 7a9bd989d3..f7f4db3fcc 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemAdviseMmap.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemAdviseMmap.cc @@ -50,7 +50,7 @@ TEST_CASE("Unit_hipMemAdvise_MmapMem") { HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, 0)); INFO("hipDeviceAttributeManagedMemory: " << managed); - if ((managed == 1) && (PageableMem == 1)) { + if (PageableMem == 1) { #ifdef __linux__ // For now this test is enabled only for linux platforms FILE *fptr; diff --git a/projects/hip/tests/catch/unit/memory/hipMemCoherencyTst.cc b/projects/hip/tests/catch/unit/memory/hipMemCoherencyTst.cc index 67f455de29..ce1fbdaed8 100644 --- a/projects/hip/tests/catch/unit/memory/hipMemCoherencyTst.cc +++ b/projects/hip/tests/catch/unit/memory/hipMemCoherencyTst.cc @@ -89,33 +89,24 @@ static void TstCoherency(int* ptr, bool hmmMem) { // passing #if HT_AMD TEST_CASE("Unit_hipHostMalloc_CoherentTst") { - int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0; + int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = false; YES_COHERENT = false; - 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"); + // 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); } #endif @@ -126,19 +117,14 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") { // passing #if HT_AMD TEST_CASE("Unit_hipMallocManaged_CoherentTst") { - int *Ptr = nullptr, SIZE = sizeof(int), Pageable = 0, managed = 0; + int *Ptr = nullptr, SIZE = sizeof(int), managed = 0; bool HmmMem = true; YES_COHERENT = false; - HIP_CHECK(hipDeviceGetAttribute(&Pageable, - hipDeviceAttributePageableMemoryAccess, 0)); - INFO("hipDeviceAttributePageableMemoryAccess: " << Pageable); - HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, 0)); INFO("hipDeviceAttributeManagedMemory: " << managed); - - if (managed == 1 && Pageable == 1) { + if (managed == 1) { // Allocating hipMallocManaged() memory SECTION("hipMallocManaged with hipMemAttachGlobal flag") { HIP_CHECK(hipMallocManaged(&Ptr, SIZE, hipMemAttachGlobal)); @@ -150,7 +136,7 @@ TEST_CASE("Unit_hipMallocManaged_CoherentTst") { HIP_CHECK(hipFree(Ptr)); REQUIRE(YES_COHERENT); } else { - SUCCEED("GPU 0 doesn't support ManagedMemory or PageableMemoryAccess" + SUCCEED("GPU 0 doesn't support ManagedMemory " "device attribute. Hence skipping the test with Pass result.\n"); } }