From 15517aaf18a03cc80ae96245b5eacdb145732a41 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 1/9] 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 --- .../multiproc/hipMemCoherencyTstMProc.cc | 232 +++++++----------- tests/catch/unit/memory/hipMemAdvise.cc | 10 +- tests/catch/unit/memory/hipMemAdviseMmap.cc | 2 +- tests/catch/unit/memory/hipMemCoherencyTst.cc | 48 ++-- 4 files changed, 117 insertions(+), 175 deletions(-) diff --git a/tests/catch/multiproc/hipMemCoherencyTstMProc.cc b/tests/catch/multiproc/hipMemCoherencyTstMProc.cc index 10d6e27f6d..863dfc7178 100644 --- a/tests/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/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/tests/catch/unit/memory/hipMemAdvise.cc b/tests/catch/unit/memory/hipMemAdvise.cc index 96a0318bb8..c0673753a9 100644 --- a/tests/catch/unit/memory/hipMemAdvise.cc +++ b/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/tests/catch/unit/memory/hipMemAdviseMmap.cc b/tests/catch/unit/memory/hipMemAdviseMmap.cc index 7a9bd989d3..f7f4db3fcc 100644 --- a/tests/catch/unit/memory/hipMemAdviseMmap.cc +++ b/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/tests/catch/unit/memory/hipMemCoherencyTst.cc b/tests/catch/unit/memory/hipMemCoherencyTst.cc index 67f455de29..ce1fbdaed8 100644 --- a/tests/catch/unit/memory/hipMemCoherencyTst.cc +++ b/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"); } } From e28f31e1f56d799cf6c977b86829c2592c295bf1 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:49:04 +0530 Subject: [PATCH 2/9] SWDEV-337766 - [catch2][dtest]Fix for MI200 failures for tests: Unit_hipMemcpyPeer_Negative and Unit_hipMemcpyPeerAsync_Negative (#2986) Change-Id: I88aa2580a589f6fcef3a098ee68922e3431aa75a --- tests/catch/unit/memory/hipMemcpyPeer.cc | 6 ++++-- tests/catch/unit/memory/hipMemcpyPeerAsync.cc | 8 ++++---- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/tests/catch/unit/memory/hipMemcpyPeer.cc b/tests/catch/unit/memory/hipMemcpyPeer.cc index 48e881895d..33ee2e9f70 100644 --- a/tests/catch/unit/memory/hipMemcpyPeer.cc +++ b/tests/catch/unit/memory/hipMemcpyPeer.cc @@ -73,11 +73,13 @@ TEST_CASE("Unit_hipMemcpyPeer_Negative") { } SECTION("Passing invalid Destination device ID") { - REQUIRE(hipMemcpyPeer(B_d, 10, A_d, 0, copy_bytes) != hipSuccess); + REQUIRE(hipMemcpyPeer(B_d, numDevices, A_d, 0, copy_bytes) != + hipSuccess); } SECTION("Passing invalid Source device ID") { - REQUIRE(hipMemcpyPeer(B_d, 1, A_d, 10, copy_bytes) != hipSuccess); + REQUIRE(hipMemcpyPeer(B_d, 1, A_d, numDevices, copy_bytes) != + hipSuccess); } HipTest::freeArrays(A_d, B_d, nullptr, A_h, B_h, nullptr, false); } else { diff --git a/tests/catch/unit/memory/hipMemcpyPeerAsync.cc b/tests/catch/unit/memory/hipMemcpyPeerAsync.cc index 56a76cbad2..159c8309e3 100644 --- a/tests/catch/unit/memory/hipMemcpyPeerAsync.cc +++ b/tests/catch/unit/memory/hipMemcpyPeerAsync.cc @@ -82,12 +82,12 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Negative") { } SECTION("Passing invalid Destination device ID") { - REQUIRE(hipMemcpyPeerAsync(B_d, 10, A_d, 0, copy_bytes, + REQUIRE(hipMemcpyPeerAsync(B_d, numDevices, A_d, 0, copy_bytes, stream) != hipSuccess); } SECTION("Passing invalid Source device ID") { - REQUIRE(hipMemcpyPeerAsync(B_d, 10, A_d, 0, copy_bytes, + REQUIRE(hipMemcpyPeerAsync(B_d, 0, A_d, numDevices, copy_bytes, stream) != hipSuccess); } @@ -155,14 +155,14 @@ TEST_CASE("Unit_hipMemcpyPeerAsync_Basic") { stream)); HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamSynchronize(stream)); } SECTION("Calling hipMemcpyPerAsync() using hipStreamPerThread") { HIP_CHECK(hipMemcpyPeerAsync(X_d, 1, A_d, 0, copy_bytes, hipStreamPerThread)); HIP_CHECK(hipMemcpyPeerAsync(Y_d, 1, B_d, 0, copy_bytes, hipStreamPerThread)); - HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); } hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0, static_cast(X_d), From 752acff9b20679faea2ffe2269343a9b3a96e3b1 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:49:18 +0530 Subject: [PATCH 3/9] SWDEV-361129 - [catch2][dtest] Enabling tests hipGraphInstantiateWithFlags for AMD (#2987) Change-Id: I39227ab51417e55f4ab05c48b61e9735856e58ac --- tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc | 6 ------ 1 file changed, 6 deletions(-) diff --git a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc index 5e53d0357c..107bffd296 100644 --- a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc +++ b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc @@ -43,11 +43,9 @@ Mapping is missing for NVIDIA platform hence skipping the testcases #include constexpr size_t N = 1000000; -#if HT_AMD /* This test covers the negative scenarios of hipGraphInstantiateWithFlags API */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") { -#if HT_NVIDIA SECTION("Passing nullptr pGraphExec") { hipGraph_t graph; HIP_CHECK(hipGraphCreate(&graph, 0)); @@ -67,7 +65,6 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_Negative") { hipGraphExec_t graphExec; REQUIRE(hipGraphInstantiateWithFlags(&graphExec, graph, 10) != hipSuccess); } -#endif } /* This function verifies the following scenarios @@ -267,7 +264,6 @@ This testcase verifies hipGraphInstantiateWithFlags API by creating dependency graph on GPU-0 and instantiate, launching and verifying the result on GPU-1 */ -#if HT_NVIDIA TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") { int numDevices = 0; int canAccessPeer = 0; @@ -283,7 +279,6 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") { SUCCEED("skipped the testcase as no of devices is less than 2"); } } -#endif /* This testcase verifies hipGraphInstantiateWithFlags API by creating capture graph and instantiate, launching and verifying @@ -325,4 +320,3 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_StreamCaptureDeviceContextChg") { SUCCEED("skipped the testcase as no of devices is less than 2"); } } -#endif From acf96881652ae83dae6d436f7c53c30dd046715b 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:50:51 +0530 Subject: [PATCH 4/9] SWDEV-356559 - update the note on hipIpcOpenMemHandle showing the different from CUDA (#2988) Change-Id: I4116eaa3a8f31283dda7d7c0c680acdb954036b9 --- include/hip/hip_runtime_api.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index aa36787636..0685876aef 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -1707,8 +1707,9 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr); * hipErrorInvalidHandle, * hipErrorTooManyPeers * - * @note No guarantees are made about the address returned in @p *devPtr. - * In particular, multiple processes may not receive the same address for the same @p handle. + * @note During multiple processes, using the same memory handle opened by the current context, + * there is no guarantee that the same device poiter will be returned in @p *devPtr. + * This is diffrent from CUDA. * */ hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags); From 53599bce3e32e710764f11d893b4d8b73cb4d2e5 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:52:06 +0530 Subject: [PATCH 5/9] SWDEV-318349 - Enable test of graph launched by peer (#2997) - Enable Unit_hipGraphInstantiateWithFlags_DependencyGraph for AMD. - Add stream sync for null stream. - Fix Unit_hipGraphInstantiateWithFlags_DependencyGraph. change-Id: I5774563fdfed1afd29a43927f3630adbfb75c48a --- .../unit/graph/hipGraphInstantiateWithFlags.cc | 17 ++++------------- 1 file changed, 4 insertions(+), 13 deletions(-) diff --git a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc index 107bffd296..14302b84f9 100644 --- a/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc +++ b/tests/catch/unit/graph/hipGraphInstantiateWithFlags.cc @@ -157,6 +157,7 @@ void GraphInstantiateWithFlags_DependencyGraph(bool ctxt_change = false) { // Instantiate and launch the cloned graph HIP_CHECK(hipGraphInstantiateWithFlags(&graphExec, graph, 0)); HIP_CHECK(hipGraphLaunch(graphExec, 0)); + HIP_CHECK(hipStreamSynchronize(0)); // Verify graph execution result HipTest::checkVectorADD(A_h, B_h, C_h, N); @@ -245,20 +246,9 @@ by creating dependency graph and instantiate, launching and verifying the result */ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraph") { - int numDevices = 0; - int canAccessPeer = 0; - HIP_CHECK(hipGetDeviceCount(&numDevices)); - if (numDevices > 1) { - HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, 1)); - if (canAccessPeer) { - GraphInstantiateWithFlags_DependencyGraph(); - } else { - SUCCEED("Machine does not seem to have P2P"); - } - } else { - SUCCEED("skipped the testcase as no of devices is less than 2"); - } + GraphInstantiateWithFlags_DependencyGraph(); } + /* This testcase verifies hipGraphInstantiateWithFlags API by creating dependency graph on GPU-0 and instantiate, launching and verifying @@ -279,6 +269,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithFlags_DependencyGraphDeviceCtxtChg") { SUCCEED("skipped the testcase as no of devices is less than 2"); } } + /* This testcase verifies hipGraphInstantiateWithFlags API by creating capture graph and instantiate, launching and verifying From 8baf649ff597a5914a5d82042758fd61acf9df5f 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:52:24 +0530 Subject: [PATCH 6/9] SWDEV-361383 - Fixing compilation issue for stress test files under catch2 framework (#2998) Change-Id: I4253a6cdb2c10bcb0021f869b7b3c4f1acd28787 --- .../Stress_deviceAllocationStress.cc | 2 +- .../stress/memory/hipMallocManagedStress.cc | 2 +- .../memory/hipMemPrftchAsyncStressTst.cc | 24 +++---- .../stress/memory/hipMemcpyMThreadMSize.cc | 8 +-- tests/catch/stress/memory/memcpy.cc | 10 +-- .../printf/Stress_printf_ComplexKernels.cc | 2 +- tests/catch/stress/stream/streamEnqueue.cc | 63 ++++++++++--------- 7 files changed, 56 insertions(+), 55 deletions(-) diff --git a/tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc b/tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc index e60bf66bb4..5e6c51216e 100644 --- a/tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc +++ b/tests/catch/stress/deviceallocation/Stress_deviceAllocationStress.cc @@ -349,7 +349,7 @@ static bool TestMemoryAllocationInLoop(int test_type, } if (!bPassed) break; } - hipFree(outputVec_d); + HIP_CHECK(hipFree(outputVec_d)); free(outputVec_h); return bPassed; } diff --git a/tests/catch/stress/memory/hipMallocManagedStress.cc b/tests/catch/stress/memory/hipMallocManagedStress.cc index 8a11ab35ef..87950a4f66 100644 --- a/tests/catch/stress/memory/hipMallocManagedStress.cc +++ b/tests/catch/stress/memory/hipMallocManagedStress.cc @@ -158,7 +158,7 @@ TEST_CASE("Stress_hipMallocManaged_MultiSize") { hipStream_t strm; HIP_CHECK(hipStreamCreate(&strm)); dim3 dimBlock(blockSize, 1, 1); - for (int i = 1; i < (1024*1024); ++i) { + for (int i = 1; i < (1024*100); ++i) { HIP_CHECK(hipMallocManaged(&Hmm1, i)); HIP_CHECK(hipMallocManaged(&Hmm2, i)); for (int j = 0; j < i; ++j) { diff --git a/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc b/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc index 0a159b8dbb..a551721b40 100644 --- a/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc +++ b/tests/catch/stress/memory/hipMemPrftchAsyncStressTst.cc @@ -57,9 +57,8 @@ static int HmmAttrPrint() { return managed; } -static void ReleaseResource(int *Hmm, int *Hmm1, hipStream_t *strm) { +static void ReleaseResource(int *Hmm, hipStream_t *strm) { HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipFree(Hmm1)); HIP_CHECK(hipStreamDestroy(*strm)); } @@ -70,11 +69,10 @@ static void ReleaseResource(int *Hmm, int *Hmm1, hipStream_t *strm) { TEST_CASE("Unit_hipMemPrefetchAsyncOneToAll") { int MangdMem = HmmAttrPrint(); if (MangdMem == 1) { - int *Hmm = nullptr, *Hmm1 = nullptr, NumDevs, MemSz = (4096 * 4); + int *Hmm1 = nullptr, NumDevs, MemSz = (4096 * 4); int InitVal = 123, NumElms = MemSz/4; bool IfTestPassed = true; HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MemSz)); HIP_CHECK(hipMallocManaged(&Hmm1, MemSz)); for (int i = 0; i < NumElms; ++i) { Hmm1[i] = InitVal; @@ -93,44 +91,40 @@ TEST_CASE("Unit_hipMemPrefetchAsyncOneToAll") { // Prefetching memory from i to j HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, j, strm)); HIP_CHECK(hipStreamSynchronize(strm)); - MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms); + MemPrftchAsyncKernel1<<<(NumElms/32), 32, 0, strm>>>(Hmm1, NumElms); HIP_CHECK(hipStreamSynchronize(strm)); // Verifying the result for (int m = 0; m < NumElms; ++m) { - if (Hmm[m] != (InitVal * InitVal)) { + if (Hmm1[m] != (InitVal * InitVal)) { IfTestPassed = false; } } if (!IfTestPassed) { - ReleaseResource(Hmm, Hmm1, &strm); + ReleaseResource(Hmm1, &strm); INFO("Did not find expected value!"); REQUIRE(false); } - // Resetting the values in Hmm - HIP_CHECK(hipMemset(Hmm, 0, MemSz)); // Prefetching memory from j to i HIP_CHECK(hipMemPrefetchAsync(Hmm1, MemSz, i, strm)); HIP_CHECK(hipStreamSynchronize(strm)); - MemPrftchAsyncKernel<<<(NumElms/32), 32, 0, strm>>>(Hmm, Hmm1, NumElms); + MemPrftchAsyncKernel1<<<(NumElms/32), 32, 0, strm>>>(Hmm1, NumElms); HIP_CHECK(hipStreamSynchronize(strm)); // Verifying the result for (int m = 0; m < NumElms; ++m) { - if (Hmm[m] != (InitVal * InitVal)) { + if (Hmm1[m] != (InitVal * InitVal)) { IfTestPassed = false; } } if (!IfTestPassed) { - ReleaseResource(Hmm, Hmm1, &strm); + ReleaseResource(Hmm1, &strm); INFO("Did not find expected value!"); REQUIRE(false); } - // Resetting the values in Hmm - HIP_CHECK(hipMemset(Hmm, 0, MemSz)); + HIP_CHECK(hipStreamDestroy(strm)); } } // Releasing the resources in case all the scenarios passed - HIP_CHECK(hipFree(Hmm)); HIP_CHECK(hipFree(Hmm1)); } else { SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " diff --git a/tests/catch/stress/memory/hipMemcpyMThreadMSize.cc b/tests/catch/stress/memory/hipMemcpyMThreadMSize.cc index 3b0e8b7ee1..301f3225ec 100644 --- a/tests/catch/stress/memory/hipMemcpyMThreadMSize.cc +++ b/tests/catch/stress/memory/hipMemcpyMThreadMSize.cc @@ -72,7 +72,7 @@ void Memcpy_And_verify(int NUM_ELM) { for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIP_CHECK(hipMemcpy(A_d[j], A_d[i], NUM_ELM * sizeof(TestType), hipMemcpyDefault)); @@ -122,7 +122,7 @@ void Memcpy_And_verify(int NUM_ELM) { int canAccessPeer = 0; for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIP_CHECK(hipMemcpyHtoD(hipDeviceptr_t(A_d[i]), A_h, NUM_ELM * sizeof(TestType))); @@ -165,7 +165,7 @@ void Memcpy_And_verify(int NUM_ELM) { for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIP_CHECK(hipMemcpyAsync(A_d[j], A_d[i], NUM_ELM * sizeof(TestType), @@ -219,7 +219,7 @@ void Memcpy_And_verify(int NUM_ELM) { for (int i = 0; i < Available_Gpus; ++i) { for (int j = i+1; j < Available_Gpus; ++j) { canAccessPeer = 0; - hipDeviceCanAccessPeer(&canAccessPeer, i, j); + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, i, j)); if (canAccessPeer) { HIP_CHECK(hipSetDevice(j)); HIP_CHECK(hipMemcpyDtoDAsync(hipDeviceptr_t(A_d[j]), diff --git a/tests/catch/stress/memory/memcpy.cc b/tests/catch/stress/memory/memcpy.cc index 21957e0fdb..7ba95b4e77 100644 --- a/tests/catch/stress/memory/memcpy.cc +++ b/tests/catch/stress/memory/memcpy.cc @@ -5,25 +5,25 @@ TEST_CASE("Stress_hipMalloc", "DifferentSizes") { SECTION("Size 10") { auto res = hipMalloc(&d_a, sizeof(10)); REQUIRE(res == hipSuccess); - hipFree(d_a); + HIP_CHECK(hipFree(d_a)); d_a = nullptr; } SECTION("Size 100") { auto res = hipMalloc(&d_a, sizeof(100)); REQUIRE(res == hipSuccess); - hipFree(d_a); + HIP_CHECK(hipFree(d_a)); d_a = nullptr; } SECTION("Size 1000") { auto res = hipMalloc(&d_a, sizeof(1000)); REQUIRE(res == hipSuccess); - hipFree(d_a); + HIP_CHECK(hipFree(d_a)); d_a = nullptr; } SECTION("Size 10000") { auto res = hipMalloc(&d_a, sizeof(10000)); REQUIRE(res == hipSuccess); - hipFree(d_a); + HIP_CHECK(hipFree(d_a)); d_a = nullptr; } SECTION("Size MAX") { @@ -31,4 +31,4 @@ TEST_CASE("Stress_hipMalloc", "DifferentSizes") { REQUIRE(res == hipErrorOutOfMemory); d_a = nullptr; } -} \ No newline at end of file +} diff --git a/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc b/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc index 46bb6f74b1..50ee409475 100644 --- a/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc +++ b/tests/catch/stress/printf/Stress_printf_ComplexKernels.cc @@ -494,7 +494,7 @@ TEST_CASE("Stress_printf_ComplexKernelMultStreamMultGpu") { unsigned int print_limit = 4; // = 4 GB uint32_t iterCount = 1; int numOfGPUs = 0; - hipGetDeviceCount(&numOfGPUs); + HIP_CHECK(hipGetDeviceCount(&numOfGPUs)); if (numOfGPUs < 2) { printf("Skipping test because numOfGPUs < 2\n"); return; diff --git a/tests/catch/stress/stream/streamEnqueue.cc b/tests/catch/stress/stream/streamEnqueue.cc index c44ff96afc..9bbb641930 100644 --- a/tests/catch/stress/stream/streamEnqueue.cc +++ b/tests/catch/stress/stream/streamEnqueue.cc @@ -26,11 +26,13 @@ THE SOFTWARE. #include #include -__global__ void addVal(unsigned long long* ptr, size_t index, unsigned long long val) { +__global__ void addVal(unsigned long long* ptr, size_t index, + unsigned long long val) { atomicAdd(ptr + index, val); } -// Create a copy constructible AtomicWrap around std::atomic so that we can put it in a vector +// Create a copy constructible AtomicWrap around std::atomic so that +// we can put it in a vector template struct AtomicWrap { std::atomic data; @@ -68,18 +70,19 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads") { constexpr size_t maxWork = 10000; constexpr size_t maxVal = 10; - std::uniform_int_distribution genIndex(0, hwThreads - 1); + std::uniform_int_distribution genIndex(0, + hwThreads - 1); std::uniform_int_distribution genWork(0, maxWork); std::uniform_int_distribution genVal(0, maxVal); auto enqueueKernelThread = [&](hipStream_t stream) { auto iter = genWork(engine); // Generate work to be done via thread - for (auto i = 0; i < iter; i++) { + for (unsigned long i = 0; i < iter; i++) { auto index = genIndex(engine); // Generate Index to add to - auto val = genVal(engine); // Generate value to add to the destination + auto val = genVal(engine); // Generate value to add to the destination hostData[index].data += val; // Replicate it on host addVal<<<1, 1, 0, stream>>>(dPtr, static_cast(index), - static_cast(val)); // And on device + static_cast(val)); // And on device } }; @@ -101,8 +104,8 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads") { HIP_CHECK(hipStreamDestroy(stream)); auto hPtr = std::make_unique(hwThreads); - HIP_CHECK( - hipMemcpy(hPtr.get(), dPtr, sizeof(unsigned long long) * hwThreads, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(hPtr.get(), dPtr, sizeof(unsigned long long) * hwThreads, + hipMemcpyDeviceToHost)); HIP_CHECK(hipFree(dPtr)); @@ -113,7 +116,7 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads") { } } -__global__ void doOperation(int* dPtr, size_t size, int val) { +__global__ void doOperation(int* dPtr, int val) { auto i = threadIdx.x; atomicAdd(dPtr + i, val); } @@ -135,14 +138,15 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") { std::vector streamPool{}; streamPool.reserve(deviceCount * streamPerGPU); - - std::map streamToDeviceMemory; // Map of stream and device memory - std::map> streamToHostMemory; // Map of stream and host result - std::map streamToDeviceIndex; // Map of stream and device it was created on - + // Map of stream and device memory + std::map streamToDeviceMemory; + // Map of stream and host result + std::map> streamToHostMemory; + // Map of stream and device it was created on + std::map streamToDeviceIndex; constexpr size_t size = 1024; - for (size_t i = 0; i < deviceCount; i++) { + for (int i = 0; i < deviceCount; i++) { HIP_CHECK(hipSetDevice(i)); for (size_t j = 0; j < streamPerGPU; j++) { @@ -155,8 +159,8 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") { HIP_CHECK(hipMalloc(&dPtr, sizeof(int) * size)); REQUIRE(dPtr != nullptr); HIP_CHECK(hipMemset(dPtr, 0, sizeof(int) * size)); - - streamToDeviceMemory[stream] = dPtr; // All streams work on exclusive memory + // All streams work on exclusive memory + streamToDeviceMemory[stream] = dPtr; streamToHostMemory[stream] = AtomicWrap(0); // CPU result @@ -171,8 +175,10 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") { std::random_device device; std::mt19937 engine(device()); - std::uniform_int_distribution genVal(-maxVal, maxVal); - std::uniform_int_distribution genStream(0, streamPool.size() - 1); + std::uniform_int_distribution genVal(-maxVal, + maxVal); + std::uniform_int_distribution genStream(0, + streamPool.size() - 1); #if HT_NVIDIA std::mutex ness; // On nvidia, current device needs to match stream's device @@ -183,7 +189,8 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") { #if HT_NVIDIA std::unique_lock lock(ness); // Lock on creation #endif - hipStream_t stream = streamPool[genStream(engine)]; // Get a random stream + // Get a random stream + hipStream_t stream = streamPool[genStream(engine)]; // TODO use HIP_CHECK_THREAD when PR#2664 is merged if (hipSuccess != hipSetDevice(streamToDeviceIndex[stream])) { @@ -191,11 +198,10 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") { } int val = genVal(engine); // Generate Value to add/sub to - - streamToHostMemory[stream].data.fetch_add(val); // Replicate result on CPU + // Replicate result on CPU + streamToHostMemory[stream].data.fetch_add(val); auto dPtr = streamToDeviceMemory[stream]; - doOperation<<<1, 1024, 0, stream>>>(dPtr, size, - val); // On GPU + doOperation<<<1, 1024, 0, stream>>>(dPtr, val); // On GPU } }; @@ -219,13 +225,14 @@ TEST_CASE("Stress_StreamEnqueue_DifferentThreads_MultiGPU") { for (auto& i : streamPool) { HIP_CHECK(hipStreamSynchronize(i)); auto dResult = std::make_unique(size); - HIP_CHECK(hipMemcpy(dResult.get(), streamToDeviceMemory[i], sizeof(int) * size, - hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(dResult.get(), streamToDeviceMemory[i], + sizeof(int) * size, hipMemcpyDeviceToHost)); HIP_CHECK(hipFree(streamToDeviceMemory[i])); HIP_CHECK(hipStreamDestroy(i)); auto res = streamToHostMemory[i].data.load(); INFO("Matching CPU: " << res << " GPU: " << dResult[0] << " Dev Ptr: " - << streamToDeviceMemory[i] << " on Device: " << streamToDeviceIndex[i]); - REQUIRE(std::all_of(dResult.get(), dResult.get() + size, [=](int r) { return r == res; })); + << streamToDeviceMemory[i] << " on Device: " << streamToDeviceIndex[i]); + REQUIRE(std::all_of(dResult.get(), dResult.get() + size, + [=](int r) { return r == res; })); } } From c4efb12f1bf3c6adfcbdaea9405b5de246f2c040 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 19 Oct 2022 17:38:03 +0530 Subject: [PATCH 7/9] SWDEV-350781 - Address Unit_hipMemAdvise_TstAlignedAllocMem catch test failure (#2999) Change-Id: I79002ac644667f280ee14a07c28547fb878b320d --- tests/catch/unit/memory/hipMemAdvise.cc | 133 +++++++++--------------- 1 file changed, 49 insertions(+), 84 deletions(-) diff --git a/tests/catch/unit/memory/hipMemAdvise.cc b/tests/catch/unit/memory/hipMemAdvise.cc index c0673753a9..cc2a52a807 100644 --- a/tests/catch/unit/memory/hipMemAdvise.cc +++ b/tests/catch/unit/memory/hipMemAdvise.cc @@ -664,98 +664,63 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { 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); - } + // The following code block checks for gfx90a so as to skip if the device is not MI200 - /* 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}; - p = fork(); + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); - 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] == 1) || (GpuId[0] == 1)) { - WARN("This test is not applicable on MI60 & MI100." - "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, "gfx906"); - if (p) { - WARN("gfx906 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]); - exit(0); - } - int stat = 0; - if (fork() == 0) { - // The below part should be inside fork - int managed = HmmAttrPrint(); - if (managed == 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 ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + int stat = 0; + if (fork() == 0) { + // The below part should be inside fork + int managed = HmmAttrPrint(); + if (managed == 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 } - } - if (DataMismatch != 0) { - WARN("DataMismatch observed!!"); - exit(9); // 9 for failure } else { - exit(10); // 10 for Pass result + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); } } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } - } else { - wait(&stat); - int Result = WEXITSTATUS(stat); - if (Result != 10) { - REQUIRE(false); + wait(&stat); + int Result = WEXITSTATUS(stat); + if (Result != 10) { + REQUIRE(false); + } } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); } + } #endif From fb938a607f66a1aef2289b7e42e4e7020fe03fa0 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 19 Oct 2022 19:12:01 +0530 Subject: [PATCH 8/9] SWDEV-286739 - Support hipDeviceAttributeWallClockRate (#2984) Part 4: Add unit test of wall_clock64() Change-Id: I52cbba6d67d21cde6da19c5ab533159f426a9bf7 --- tests/catch/unit/CMakeLists.txt | 3 + tests/catch/unit/clock/CMakeLists.txt | 29 +++++ tests/catch/unit/clock/hipClockCheck.cc | 116 ++++++++++++++++++ .../module/hipExtLaunchKernelGGL.cpp | 10 +- 4 files changed, 153 insertions(+), 5 deletions(-) create mode 100644 tests/catch/unit/clock/CMakeLists.txt create mode 100644 tests/catch/unit/clock/hipClockCheck.cc diff --git a/tests/catch/unit/CMakeLists.txt b/tests/catch/unit/CMakeLists.txt index 9874ce29d2..0ce9be5ed9 100644 --- a/tests/catch/unit/CMakeLists.txt +++ b/tests/catch/unit/CMakeLists.txt @@ -31,3 +31,6 @@ add_subdirectory(texture) add_subdirectory(streamperthread) add_subdirectory(kernel) add_subdirectory(multiThread) +if(HIP_PLATFORM STREQUAL "amd") +add_subdirectory(clock) +endif() diff --git a/tests/catch/unit/clock/CMakeLists.txt b/tests/catch/unit/clock/CMakeLists.txt new file mode 100644 index 0000000000..c181f1a76b --- /dev/null +++ b/tests/catch/unit/clock/CMakeLists.txt @@ -0,0 +1,29 @@ +# Copyright (c) 2022 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +# THE SOFTWARE. + +# Common Tests - Test independent of all platforms + +set(TEST_SRC + hipClockCheck.cc +) + +hip_add_exe_to_target(NAME ClockCheckTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) diff --git a/tests/catch/unit/clock/hipClockCheck.cc b/tests/catch/unit/clock/hipClockCheck.cc new file mode 100644 index 0000000000..4d817a3764 --- /dev/null +++ b/tests/catch/unit/clock/hipClockCheck.cc @@ -0,0 +1,116 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +#define ONESECOND 1000 // in ms +#define HALFSECOND 500 // in ms + +enum CLOCK_MODE { + CLOCK_MODE_CLOCK64, + CLOCK_MODE_WALL_CLOCK64 +}; + +__global__ void kernel_c(int clockRate, uint64_t wait_t) { + uint64_t start = clock64() / clockRate, cur = 0; // in ms + do { cur = clock64() / clockRate-start;} while (cur < wait_t); +} + +__global__ void kernel_w(int clockRate, uint64_t wait_t) { + uint64_t start = wall_clock64() / clockRate, cur = 0; // in ms + do { cur = wall_clock64() / clockRate-start;} while (cur < wait_t); +} + +bool verifyTimeExecution(CLOCK_MODE m, float time1, float time2, + float expectedTime1, float expectedTime2) { + bool testStatus = false; + float ratio = m == CLOCK_MODE_CLOCK64 ? 0.5 : 0.01; + + if (fabs(time1 - expectedTime1) < ratio * expectedTime1 + && fabs(time2 - expectedTime2) < ratio * expectedTime2) { + WARN("Succeeded: Expected Vs Actual: Kernel1 - " << expectedTime1 << " Vs " << time1 + << ", Kernel2 - " << expectedTime2 << " Vs " << time2); + testStatus = true; + } else { + FAIL_CHECK("Failed: Expected Vs Actual: Kernel1 -" << expectedTime1 << " Vs " << time1 + << ", Kernel2 - " << expectedTime2 << " Vs " << time2); + testStatus = false; + } + return testStatus; +} + +/* + * Launching kernel1 and kernel2 and then we try to + * get the event elapsed time of each kernel using the start and + * end events.The event elapsed time should return us the kernel + * execution time for that particular kernel +*/ +bool kernelTimeExecution(CLOCK_MODE m, int clockRate, + uint64_t expectedTime1, uint64_t expectedTime2) { + hipStream_t stream; + hipEvent_t start_event1, end_event1, start_event2, end_event2; + float time1 = 0, time2 = 0; + HIPCHECK(hipEventCreate(&start_event1)); + HIPCHECK(hipEventCreate(&end_event1)); + HIPCHECK(hipEventCreate(&start_event2)); + HIPCHECK(hipEventCreate(&end_event2)); + HIPCHECK(hipStreamCreate(&stream)); + hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w, + dim3(1), dim3(1), 0, stream, start_event1, end_event1, 0, clockRate, expectedTime1); + hipExtLaunchKernelGGL( m == CLOCK_MODE_CLOCK64 ? kernel_c : kernel_w, + dim3(1), dim3(1), 0, stream, start_event2, end_event2, 0, clockRate, expectedTime2); + HIPCHECK(hipStreamSynchronize(stream)); + HIPCHECK(hipEventElapsedTime(&time1, start_event1, end_event1)); + HIPCHECK(hipEventElapsedTime(&time2, start_event2, end_event2)); + + HIPCHECK(hipStreamDestroy(stream)); + HIPCHECK(hipEventDestroy(start_event1)); + HIPCHECK(hipEventDestroy(end_event1)); + HIPCHECK(hipEventDestroy(start_event2)); + HIPCHECK(hipEventDestroy(end_event2)); + + return verifyTimeExecution(m, time1, time2, expectedTime1, expectedTime2); +} + +TEST_CASE("Unit_hipClock64_Check") { + HIPCHECK(hipSetDevice(0)); + int clockRate = 0; // in KHz + HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeClockRate, 0)); + + SECTION("Verify kernel execution time via clock64()") { + CHECK(kernelTimeExecution(CLOCK_MODE_CLOCK64, clockRate, ONESECOND, HALFSECOND)); + } +} + +TEST_CASE("Unit_hipWallClock64_Check") { + HIPCHECK(hipSetDevice(0)); + int clockRate = 0; // in KHz + HIPCHECK(hipDeviceGetAttribute(&clockRate, hipDeviceAttributeWallClockRate, 0)); + + if(!clockRate) { + INFO("hipDeviceAttributeWallClockRate has not been supported. Skipped"); + return; + } + + SECTION("Verify kernel execution time via wall_clock64()") { + CHECK(kernelTimeExecution(CLOCK_MODE_WALL_CLOCK64, clockRate, ONESECOND, HALFSECOND)); + } +} diff --git a/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp b/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp index e6f5c49a9b..4df6bd8600 100644 --- a/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp +++ b/tests/src/runtimeApi/module/hipExtLaunchKernelGGL.cpp @@ -34,8 +34,8 @@ #include "test_common.h" #include "hip/hip_ext.h" -#define FOURSEC_KERNEL 4999 -#define TWOSEC_KERNEL 2999 +#define FIVESEC_KERNEL 4999 +#define THREESEC_KERNEL 2999 __device__ int globalvar = 1; __global__ void TwoSecKernel(int clockrate) { @@ -163,12 +163,12 @@ bool KernelTimeExecution() { e = hipEventElapsedTime(&time_4sec, start_event1, end_event1); e = hipEventElapsedTime(&time_2sec, start_event2, end_event2); - if ( (time_4sec < static_cast(FOURSEC_KERNEL)) && - (time_2sec < static_cast(TWOSEC_KERNEL))) { + if ( (time_4sec < static_cast(FIVESEC_KERNEL)) && + (time_2sec < static_cast(THREESEC_KERNEL))) { testStatus = true; } else { printf("Expected Vs Actual: Kernel1-<%d Vs %f Kernel2-<%d Vs %f\n", - FOURSEC_KERNEL, time_4sec, TWOSEC_KERNEL, time_2sec); + FIVESEC_KERNEL, time_4sec, THREESEC_KERNEL, time_2sec); testStatus = false; } From 685390469af762268ab766245933033455f8c29d Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi <53337087+satyanveshd@users.noreply.github.com> Date: Thu, 20 Oct 2022 07:26:03 +0530 Subject: [PATCH 9/9] SWDEV-350781 - Remove Pagetable validation tests from HIP (#2970) --- tests/catch/multiproc/CMakeLists.txt | 1 - .../multiproc/hipHostMallocTestsMproc.cc | 318 -------------- .../runtimeApi/memory/hipHostMallocTests.cpp | 401 ++---------------- 3 files changed, 35 insertions(+), 685 deletions(-) delete mode 100644 tests/catch/multiproc/hipHostMallocTestsMproc.cc diff --git a/tests/catch/multiproc/CMakeLists.txt b/tests/catch/multiproc/CMakeLists.txt index 5541067207..85a57a9a71 100644 --- a/tests/catch/multiproc/CMakeLists.txt +++ b/tests/catch/multiproc/CMakeLists.txt @@ -9,7 +9,6 @@ set(LINUX_TEST_SRC hipGetDevicePropertiesMproc.cc hipSetGetDeviceMproc.cc hipIpcMemAccessTest.cc - hipHostMallocTestsMproc.cc hipMallocConcurrencyMproc.cc hipMemCoherencyTstMProc.cc hipIpcEventHandle.cc diff --git a/tests/catch/multiproc/hipHostMallocTestsMproc.cc b/tests/catch/multiproc/hipHostMallocTestsMproc.cc deleted file mode 100644 index b07922ec4b..0000000000 --- a/tests/catch/multiproc/hipHostMallocTestsMproc.cc +++ /dev/null @@ -1,318 +0,0 @@ -/* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/** -Testcase Scenarios : - 1) Validate page table allocation of hipHostMalloc() api when - HIP_VISIBLE_DEVICES set to single device. - 2) Validate page table allocation of hipHostMalloc() api when - HIP_VISIBLE_DEVICES set to list of multiple devices. -*/ - -#include -#include - -#if defined(__linux__) -#include -#include -#include -#include - -#if HT_AMD -/* AMD specific as below tests use rocm-smi */ - -/** - * Defines - */ -#define LIB_ROCMSMI "librocm_smi64.so" -#define ALLOC_SIZE (30*1024*1024) - -/** - * Global variables - */ -static rsmi_status_t (*rsmi_dev_memory_usage_get_fp)(uint32_t, - rsmi_memory_type_t, - uint64_t *); -static rsmi_status_t (*rsmi_init_fp)(uint64_t); -static rsmi_status_t (*rsmi_shut_down_fp)(); -static void *rocm_smi_h; - - -/** - * Fetches Gpu device count - */ -static void getDeviceCount(int *pdevCnt) { - int fd[2], val = 0; - pid_t childpid; - - // create pipe descriptors - pipe(fd); - - // disable visible_devices env from shell - unsetenv("ROCR_VISIBLE_DEVICES"); - unsetenv("HIP_VISIBLE_DEVICES"); - - childpid = fork(); - - if (childpid > 0) { // Parent - close(fd[1]); - // parent will wait to read the device cnt - read(fd[0], &val, sizeof(val)); - - // close the read-descriptor - close(fd[0]); - - // wait for child exit - wait(nullptr); - - *pdevCnt = val; - } else if (!childpid) { // Child - int devCnt = 1; - // writing only, no need for read-descriptor - close(fd[0]); - - HIP_CHECK(hipGetDeviceCount(&devCnt)); - // send the value on the write-descriptor: - write(fd[1], &devCnt, sizeof(devCnt)); - - // close the write descriptor: - close(fd[1]); - exit(0); - } else { // failure - *pdevCnt = 0; - return; - } -} - -/** - * Initializes rocm smi library handles - */ -static bool rocm_smi_init() { - // Open ROCm SMI Library - if (!(rocm_smi_h = dlopen(LIB_ROCMSMI, RTLD_LAZY))) { - printf("Error opening rocm smi library!\n"); - return false; - } - - void* fnsym = dlsym(rocm_smi_h, "rsmi_dev_memory_usage_get"); - if (!fnsym) { - printf("Error getting rsmi_dev_memory_usage_get() function\n"); - dlclose(rocm_smi_h); - return false; - } - rsmi_dev_memory_usage_get_fp = reinterpret_cast(fnsym); - - fnsym = dlsym(rocm_smi_h, "rsmi_init"); - if (!fnsym) { - printf("Error getting rsmi_init() function\n"); - dlclose(rocm_smi_h); - return false; - } - rsmi_init_fp = reinterpret_cast(fnsym); - - fnsym = dlsym(rocm_smi_h, "rsmi_shut_down"); - if (!fnsym) { - printf("Error getting rsmi_shut_down() function\n"); - dlclose(rocm_smi_h); - return false; - } - rsmi_shut_down_fp = reinterpret_cast(fnsym); - - uint64_t init_flags = 0; - rsmi_status_t retsmi_init; - retsmi_init = rsmi_init_fp(init_flags); - if (RSMI_STATUS_SUCCESS != retsmi_init) { - printf("Error when initializing rocm_smi\n"); - dlclose(rocm_smi_h); - return false; - } - - return true; -} - -/** - * Exits rocm smi library - */ -static void rocm_smi_exit() { - rsmi_shut_down_fp(); - dlclose(rocm_smi_h); -} - -/** - * Validates page table memory allocations - * by setting visible devices selected. - */ -static bool validatePageTableAllocations(const char *devList, int visDevCnt) { - int fd[2]; - bool testResult = false; - pid_t pid; - int numdev = 0; - - getDeviceCount(&numdev); - REQUIRE(numdev > 0); - - if (pipe(fd) < 0) { - printf("Pipe system call failed\n"); - return false; - } - - pid = fork(); - - if (!pid) { // Child process - rsmi_status_t ret; - std::vector prev, current; - uint64_t used = 0; - int tmpdev = 0, changeCnt = 0, indx = 0; - char *ptr = nullptr; - bool testPassed = true; - - // Disable visible_devices env from shell - unsetenv("ROCR_VISIBLE_DEVICES"); - unsetenv("HIP_VISIBLE_DEVICES"); - - setenv("HIP_VISIBLE_DEVICES", devList, 1); - - // First Call to initialize hip api - HIP_CHECK(hipGetDeviceCount(&tmpdev)); - - - // Get memory snapshot before hostmalloc - for (indx = 0; indx < numdev; indx++) { - ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used); - if (RSMI_STATUS_SUCCESS != ret) { - printf("Error while running rsmi_dev_memory_usage_get func\n"); - dlclose(rocm_smi_h); - rsmi_shut_down_fp(); - return false; - } - prev.push_back(used); - } - - HIP_CHECK(hipHostMalloc(&ptr, ALLOC_SIZE)); - - // Get memory snapshot after hostmalloc - for (indx = 0; indx < numdev; indx++) { - ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used); - if (RSMI_STATUS_SUCCESS != ret) { - printf("Error while running rsmi_dev_memory_usage_get func\n"); - dlclose(rocm_smi_h); - rsmi_shut_down_fp(); - HIP_CHECK(hipHostFree(ptr)); - return false; - } - current.push_back(used); - } - - for (indx = 0; indx < numdev; indx++) { - // For visible hip devices, there should be increase in VRAM usage - // due to page table allocations - // For NON visible hip devices, there can be reduction in VRAM usage - // due to removal of page tables from them - if (current[indx] > prev[indx]) - changeCnt++; - } - - // Check if memory allocation happened only for visible devices - if (changeCnt == visDevCnt) { - testPassed = true; - } else { - testPassed = false; - } - - HIP_CHECK(hipHostFree(ptr)); - - // writing only, no need for read-descriptor - close(fd[0]); - - // send the value on the write-descriptor: - write(fd[1], &testPassed, sizeof(testPassed)); - - // close the write descriptor: - close(fd[1]); - exit(0); - } else if (pid > 0) { // parent - close(fd[1]); - read(fd[0], &testResult, sizeof(testResult)); - close(fd[0]); - wait(nullptr); - } else { - printf("fork() failed\n"); - HIP_ASSERT(false); - } - - return testResult; -} - -/** - * Test page table allocation when HIP_VISIBLE_DEVICES set to - * single device - */ -TEST_CASE("Unit_hipHostMalloc_SingleVisibleDevicePageAlloc") { - int devCnt; - std::string str; - - if (!rocm_smi_init()) { - WARN("Testcase skipped as rocm smi not initialized/present"); - return; - } - - getDeviceCount(&devCnt); - REQUIRE(devCnt > 0); - - // Select single visible device and validate memory usage - for (int i = 0; i < devCnt; i++) { - str = std::to_string(i); - REQUIRE(validatePageTableAllocations(str.c_str(), 1) == true); - } - - rocm_smi_exit(); -} - -/** - * Test page table allocation when HIP_VISIBLE_DEVICES set to - * multiple devices - */ -TEST_CASE("Unit_hipHostMalloc_MultipleVisibleDevicesPageAlloc") { - int devCnt = 0, vdCnt = 0; - std::string str; - - if (!rocm_smi_init()) { - WARN("Testcase skipped as rocm smi not initialized/present"); - return; - } - - getDeviceCount(&devCnt); - REQUIRE(devCnt > 0); - - // Select multiple visible devices and validate memory usage - for (int i = 0; i < devCnt; i++) { - if (i == 0) - str += std::to_string(i); - else - str += "," + std::to_string(i); - - vdCnt++; - REQUIRE(validatePageTableAllocations(str.c_str(), vdCnt) == true); - } - - rocm_smi_exit(); -} -#endif // HT_AMD -#endif // __linux__ diff --git a/tests/src/runtimeApi/memory/hipHostMallocTests.cpp b/tests/src/runtimeApi/memory/hipHostMallocTests.cpp index 013b533529..8388869a75 100644 --- a/tests/src/runtimeApi/memory/hipHostMallocTests.cpp +++ b/tests/src/runtimeApi/memory/hipHostMallocTests.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2020 - 2022 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -28,323 +28,16 @@ Testcase Scenarios : 4) Pass size as zero for hipHostMalloc() api and check ptr is reset with with return value success. - (TestCase 2):: - 5) Validate memory usage of hipHostMalloc() api when HIP_VISIBLE_DEVICES set - to single device. - 6) Validate memory usage of hipHostMalloc() api when HIP_VISIBLE_DEVICES set - to list of multiple devices. */ -/* Tests 2 and 3 are rocclr specific tests and not supported on nvidia */ /* HIT_START - * BUILD_CMD: %t %hc %S/%s %S/../../test_common.cpp -I%S/../../ -o %T/%t -ldl -std=c++11 - * TEST: %t --tests 1 - * TEST: %t --tests 2 - * TEST: %t --tests 3 + * BUILD_CMD: %t %hc %S/%s %S/../../test_common.cpp -I%S/../../ -o %T/%t -std=c++11 + * TEST: %t * HIT_END */ -#ifdef __linux__ -#include -#include -#include -#include -#endif - -#include -#include -#include -#include "hipHostMallocTests.h" - -/** - * Defines - */ -#define LIB_ROCMSMI "librocm_smi64.so" +#include "test_common.h" #define NUM_BYTES 1000 -#define ALLOC_SIZE (30*1024*1024) - - -/** - * Global variables - */ -rsmi_status_t (*rsmi_dev_memory_usage_get_fp)(uint32_t, rsmi_memory_type_t, - uint64_t *); -rsmi_status_t (*rsmi_init_fp)(uint64_t); -rsmi_status_t (*rsmi_shut_down_fp)(); -void *rocm_smi_h; - - -/** - * Fetches Gpu device count - */ -void getDeviceCount(int *pdevCnt) { -#ifdef __linux__ - int fd[2], val = 0; - pid_t childpid; - - // create pipe descriptors - pipe(fd); - - // disable visible_devices env from shell - unsetenv("ROCR_VISIBLE_DEVICES"); - unsetenv("HIP_VISIBLE_DEVICES"); - - childpid = fork(); - - if (childpid > 0) { // Parent - close(fd[1]); - // parent will wait to read the device cnt - read(fd[0], &val, sizeof(val)); - - // close the read-descriptor - close(fd[0]); - - // wait for child exit - wait(NULL); - - *pdevCnt = val; - } else if (!childpid) { // Child - int devCnt = 1; - // writing only, no need for read-descriptor - close(fd[0]); - - HIPCHECK(hipGetDeviceCount(&devCnt)); - // send the value on the write-descriptor: - write(fd[1], &devCnt, sizeof(devCnt)); - - // close the write descriptor: - close(fd[1]); - exit(0); - } else { // failure - *pdevCnt = 1; - return; - } - -#else - HIPCHECK(hipGetDeviceCount(pdevCnt)); -#endif -} - -#if defined(__linux__) - -/** - * Initializes rocm smi library handles - */ -bool rocm_smi_init() { - // Open ROCm SMI Library - if (!(rocm_smi_h = dlopen(LIB_ROCMSMI, RTLD_LAZY))) { - printf("Error opening rocm smi library!\n"); - return false; - } - - void* fnsym = dlsym(rocm_smi_h, "rsmi_dev_memory_usage_get"); - if (!fnsym) { - printf("Error getting rsmi_dev_memory_usage_get() function\n"); - dlclose(rocm_smi_h); - return false; - } - rsmi_dev_memory_usage_get_fp = reinterpret_cast(fnsym); - - fnsym = dlsym(rocm_smi_h, "rsmi_init"); - if (!fnsym) { - printf("Error getting rsmi_init() function\n"); - dlclose(rocm_smi_h); - return false; - } - rsmi_init_fp = reinterpret_cast(fnsym); - - fnsym = dlsym(rocm_smi_h, "rsmi_shut_down"); - if (!fnsym) { - printf("Error getting rsmi_shut_down() function\n"); - dlclose(rocm_smi_h); - return false; - } - rsmi_shut_down_fp = reinterpret_cast(fnsym); - - uint64_t init_flags = 0; - rsmi_status_t retsmi_init; - retsmi_init = rsmi_init_fp(init_flags); - if (RSMI_STATUS_SUCCESS != retsmi_init) { - printf("Error when initializing rocm_smi\n"); - dlclose(rocm_smi_h); - return false; - } - - return true; -} - -/** - * Exits rocm smi library - */ -void rocm_smi_exit() { - rsmi_shut_down_fp(); - dlclose(rocm_smi_h); -} - -/** - * Validates page table memory allocations - * by setting visible devices selected. - */ -bool validatePageTableAllocations(const char *devList, int devCnt) { - int fd[2]; - bool testResult = false; - pid_t pid; - int numdev = 0; - - getDeviceCount(&numdev); - if (pipe(fd) < 0) { - printf("Pipe system call failed\n"); - return false; - } - - pid = fork(); - - if (!pid) { // Child process - rsmi_status_t ret; - std::vector prev, current; - uint64_t used = 0; - int tmpdev = 0, changeCnt = 0, indx = 0; - char *ptr = NULL; - bool testPassed = true; - - // Disable visible_devices env from shell - unsetenv("ROCR_VISIBLE_DEVICES"); - unsetenv("HIP_VISIBLE_DEVICES"); - - setenv("HIP_VISIBLE_DEVICES", devList, 1); - - // First Call to initialize hip api - hipGetDeviceCount(&tmpdev); - - - // Get memory snapshot before hostmalloc - for (indx = 0; indx < numdev; indx++) { - ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used); - if (RSMI_STATUS_SUCCESS != ret) { - printf("Error while running rsmi_dev_memory_usage_get func\n"); - dlclose(rocm_smi_h); - rsmi_shut_down_fp(); - return false; - } - prev.push_back(used); - } - - HIPCHECK(hipHostMalloc(&ptr, ALLOC_SIZE)); - - // Get memory snapshot after hostmalloc - for (indx = 0; indx < numdev; indx++) { - ret = rsmi_dev_memory_usage_get_fp(indx, RSMI_MEM_TYPE_VRAM, &used); - if (RSMI_STATUS_SUCCESS != ret) { - printf("Error while running rsmi_dev_memory_usage_get func\n"); - dlclose(rocm_smi_h); - rsmi_shut_down_fp(); - hipHostFree(ptr); - return false; - } - current.push_back(used); - } - - for (indx = 0; indx < numdev; indx++) { - // For visible hip devices, there should be increase in VRAM usage - // due to page table allocations - // For NON visible hip devices, there can be reduction in VRAM usage - // due to removal of page tables from them - if (current[indx] > prev[indx]) - changeCnt++; - } - - // Check if memory allocation happened only for visible devices - if (changeCnt == devCnt) { - testPassed = true; - } else { - testPassed = false; - } - - hipHostFree(ptr); - - // writing only, no need for read-descriptor - close(fd[0]); - - // send the value on the write-descriptor: - write(fd[1], &testPassed, sizeof(testPassed)); - - // close the write descriptor: - close(fd[1]); - exit(0); - } else if (pid > 0) { // parent - close(fd[1]); - read(fd[0], &testResult, sizeof(testResult)); - close(fd[0]); - wait(NULL); - } else { - printf("fork() failed\n"); - testResult = false; - } - - return testResult; -} - -/** - * Validate memory usage selecting single visible device - */ -bool validateHostMallocSingleVisibleDevice() { - int devCnt; - std::string str; - bool TestPassed = true; - - if (!rocm_smi_init()) { - printf("%s Testcase skipped as rocm smi not initialized/present\n", - __func__); - return true; - } - getDeviceCount(&devCnt); - - // Select single visible device and validate memory usage - for (int i = 0; i < devCnt; i++) { - str = std::to_string(i); - TestPassed = validatePageTableAllocations(str.c_str(), 1); - if (!TestPassed) - break; - } - - rocm_smi_exit(); - return TestPassed; -} - -/** - * Validate memory usage selecting multiple visible devices - */ -bool validateHostMallocMultipleVisibleDevices() { - int devCnt = 0, vdCnt = 0; - std::string str; - bool TestPassed = true; - - if (!rocm_smi_init()) { - printf("%s Testcase skipped as rocm smi not initialized/present\n", - __func__); - return true; - } - getDeviceCount(&devCnt); - - // Select multiple visible devices and validate memory usage - for (int i = 0; i < devCnt; i++) { - if (i == 0) - str += std::to_string(i); - else - str += "," + std::to_string(i); - - vdCnt++; - TestPassed = validatePageTableAllocations(str.c_str(), vdCnt); - if (!TestPassed) - break; - } - - rocm_smi_exit(); - return TestPassed; -} -#endif - int main(int argc, char *argv[]) { HipTest::parseStandardArguments(argc, argv, true); @@ -353,63 +46,40 @@ int main(int argc, char *argv[]) { size_t allocSize = NUM_BYTES; char *ptr; - if (p_tests == 1) { - // Pass ptr as nullptr. - if ((ret = hipHostMalloc(static_cast(nullptr), allocSize)) - != hipErrorInvalidValue) { - printf("ArgValidation : Inappropritate error value returned for " - "ptr as nullptr. Error: '%s'(%d)\n", - hipGetErrorString(ret), ret); - TestPassed &= false; - } + // Pass ptr as nullptr. + if ((ret = hipHostMalloc(static_cast(nullptr), allocSize)) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for " + "ptr as nullptr. Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } - // Size as max(size_t). - if ((ret = hipHostMalloc(&ptr, - std::numeric_limits::max())) - != hipErrorOutOfMemory) { - printf("ArgValidation : Inappropritate error value returned for " - "max(size_t). Error: '%s'(%d)\n", - hipGetErrorString(ret), ret); - TestPassed &= false; - } + // Size as max(size_t). + if ((ret = hipHostMalloc(&ptr, + std::numeric_limits::max())) + != hipErrorOutOfMemory) { + printf("ArgValidation : Inappropritate error value returned for " + "max(size_t). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } - // Flags as max(uint). - if ((ret = hipHostMalloc(&ptr, allocSize, - std::numeric_limits::max())) - != hipErrorInvalidValue) { - printf("ArgValidation : Inappropritate error value returned for " - "max(uint). Error: '%s'(%d)\n", - hipGetErrorString(ret), ret); - TestPassed &= false; - } + // Flags as max(uint). + if ((ret = hipHostMalloc(&ptr, allocSize, + std::numeric_limits::max())) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for " + "max(uint). Error: '%s'(%d)\n", + hipGetErrorString(ret), ret); + TestPassed &= false; + } - // Pass size as zero and check ptr reset. - HIPCHECK(hipHostMalloc(&ptr, 0)); - if (ptr) { - TestPassed &= false; - printf("ArgValidation : ptr is not reset when size(0)\n"); - } - } else if (p_tests == 2) { - // Test page table allocation when HIP_VISIBLE_DEVICES set to - // single device -#if defined(__linux__) - TestPassed = validateHostMallocSingleVisibleDevice(); -#else - printf("Test validateHostMallocSingleVisibleDevice skipped on" - "non-linux\n"); -#endif - } else if (p_tests == 3) { - // Test page table allocation when HIP_VISIBLE_DEVICES set to - // multiple devices -#if defined(__linux__) - TestPassed = validateHostMallocMultipleVisibleDevices(); -#else - printf("Test validateHostMallocMultipleVisibleDevices skipped on" - "non-linux\n"); -#endif - } else { - printf("Didnt receive any valid option. Try options 1 to 3\n"); - TestPassed = false; + // Pass size as zero and check ptr reset. + HIPCHECK(hipHostMalloc(&ptr, 0)); + if (ptr) { + TestPassed &= false; + printf("ArgValidation : ptr is not reset when size(0)\n"); } if (TestPassed) { @@ -418,4 +88,3 @@ int main(int argc, char *argv[]) { failed("hipHostMallocTests validation Failed!"); } } -