From 6ce5d7457c57f423eab695bdbea5e4dd2fa90684 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 21 Jun 2023 15:58:39 +0530 Subject: [PATCH] SWDEV-387925 - Remove comparisons that are based on available memory returned from hipMemGetInfo (#262) Change-Id: I3e4b539425482bbda578fd3127f118d47c611db3 --- catch/unit/memory/hipMalloc3DArray.cc | 12 --- catch/unit/memory/hipMallocArray.cc | 13 --- catch/unit/memory/hipMallocConcurrency.cc | 111 +--------------------- catch/unit/memory/hipMallocPitch.cc | 15 --- 4 files changed, 5 insertions(+), 146 deletions(-) diff --git a/catch/unit/memory/hipMalloc3DArray.cc b/catch/unit/memory/hipMalloc3DArray.cc index 73d50ad943..d5b840d7dd 100644 --- a/catch/unit/memory/hipMalloc3DArray.cc +++ b/catch/unit/memory/hipMalloc3DArray.cc @@ -55,8 +55,6 @@ static void Malloc3DArray_DiffSizes(int gpu) { size_t width{size}, height{size}, depth{size}; hipChannelFormatDesc channelDesc = hipCreateChannelDesc(); std::array arr; - size_t pavail, avail; - HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); for (int i = 0; i < ARRAY_LOOP; i++) { HIP_CHECK_THREAD(hipMalloc3DArray(&arr[i], &channelDesc, make_hipExtent(width, height, depth), @@ -65,9 +63,6 @@ static void Malloc3DArray_DiffSizes(int gpu) { for (int i = 0; i < ARRAY_LOOP; i++) { HIP_CHECK_THREAD(hipFreeArray(arr[i])); } - - HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); - REQUIRE_THREAD(pavail == avail); } } @@ -85,7 +80,6 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { std::vector threadlist; int devCnt = 0; devCnt = HipTest::getDeviceCount(); - const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { threadlist.push_back(std::thread(Malloc3DArray_DiffSizes, i)); } @@ -94,12 +88,6 @@ TEST_CASE("Unit_hipMalloc3DArray_MultiThread") { t.join(); } HIP_CHECK_THREAD_FINALIZE(); - const auto avail = getFreeMem(); - - if (pavail != avail) { - WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); - REQUIRE(false); - } } namespace { diff --git a/catch/unit/memory/hipMallocArray.cc b/catch/unit/memory/hipMallocArray.cc index 530eb11077..1ea94d2146 100644 --- a/catch/unit/memory/hipMallocArray.cc +++ b/catch/unit/memory/hipMallocArray.cc @@ -57,9 +57,6 @@ static void MallocArray_DiffSizes(int gpu) { for (const auto& size : runs) { hipChannelFormatDesc desc = hipCreateChannelDesc(); std::array A_d; - size_t pavail, avail; - HIP_CHECK_THREAD(hipMemGetInfo(&pavail, nullptr)); - for (int i = 0; i < ARRAY_LOOP; i++) { HIP_CHECK_THREAD( hipMallocArray(&A_d[i], &desc, std::get<0>(size), std::get<1>(size), hipArrayDefault)); @@ -67,9 +64,6 @@ static void MallocArray_DiffSizes(int gpu) { for (int i = 0; i < ARRAY_LOOP; i++) { HIP_CHECK_THREAD(hipFreeArray(A_d[i])); } - - HIP_CHECK_THREAD(hipMemGetInfo(&avail, nullptr)); - REQUIRE_THREAD(pavail == avail); } } @@ -87,7 +81,6 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { std::vector threadlist; int devCnt = 0; devCnt = HipTest::getDeviceCount(); - const auto pavail = getFreeMem(); for (int i = 0; i < devCnt; i++) { threadlist.push_back(std::thread(MallocArray_DiffSizes, i)); } @@ -96,12 +89,6 @@ TEST_CASE("Unit_hipMallocArray_MultiThread") { t.join(); } HIP_CHECK_THREAD_FINALIZE(); - - const auto avail = getFreeMem(); - if (pavail != avail) { - WARN("Memory leak of hipMalloc3D API in multithreaded scenario"); - REQUIRE(false); - } } // Kernels /////////////////////////////////////// diff --git a/catch/unit/memory/hipMallocConcurrency.cc b/catch/unit/memory/hipMallocConcurrency.cc index 885a13ee7e..17028d232d 100644 --- a/catch/unit/memory/hipMallocConcurrency.cc +++ b/catch/unit/memory/hipMallocConcurrency.cc @@ -85,7 +85,6 @@ static std::atomic g_thTestPassed{true}; static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; - size_t prevAvl, prevTot, curAvl, curTot; bool TestPassed = true; constexpr auto N = 4 * 1024 * 1024; constexpr auto blocksPerCU = 6; // to hide latency @@ -93,16 +92,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { size_t Nbytes = N * sizeof(int); HIP_CHECK(hipSetDevice(gpu)); - HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); - - if (!concurOnOneGPU && (prevAvl < curAvl || prevTot != curTot)) { - //In concurrent calls on one GPU, we cannot verify leaking in this way - printf("%s : Memory allocation mismatch observed." - "Possible memory leak.\n", __func__); - TestPassed &= false; - } unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -121,16 +111,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { TestPassed = false; } - HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HIP_CHECK(hipMemGetInfo(&curAvl, &curTot)); - - if (!concurOnOneGPU && (curAvl < prevAvl || prevTot != curTot)) { - // In concurrent calls on one GPU, we cannot verify leaking in this way - UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed." - << "Possible memory leak."); - TestPassed = false; - } return TestPassed; } @@ -140,8 +121,7 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { * Regress memory allocation and free in loop */ static bool regressAllocInLoop(int gpu) { - bool TestPassed = true; - size_t tot, avail, ptot, pavail, numBytes; + size_t numBytes; int i = 0; int* ptr; @@ -150,41 +130,18 @@ static bool regressAllocInLoop(int gpu) { // Exercise allocation in loop with bigger chunks for (i = 0; i < MaxAllocFree_BigChunks; i++) { - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); HIP_CHECK(hipMalloc(&ptr, numBytes)); - HIP_CHECK(hipMemGetInfo(&avail, &tot)); HIP_CHECK(hipFree(ptr)); - - if (pavail - avail < numBytes) { // We expect pavail-avail >= numBytes - UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << numBytes - << " not matching with hipMemGetInfo - FAIL." - << "pavail=" << pavail << ", ptot=" << ptot - << ", avail=" << avail << ", tot=" << tot - << ", pavail-avail=" << pavail - avail); - TestPassed = false; - break; - } } // Exercise allocation in loop with smaller chunks and maximum iters - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); numBytes = BuffSizeSC; for (i = 0; i < MaxAllocFree_SmallChunks; i++) { HIP_CHECK(hipMalloc(&ptr, numBytes)); - HIP_CHECK(hipFree(ptr)); } - - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - - if ((pavail != avail) || (ptot != tot)) { - UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." - << "Possible memory leak."); - TestPassed &= false; - } - - return TestPassed; + return true; } /** @@ -194,23 +151,13 @@ static bool regressAllocInLoop(int gpu) { static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) { int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; - size_t prevAvl, prevTot, curAvl, curTot; bool TestPassed = true; constexpr auto N = 4 * 1024 * 1024; constexpr auto blocksPerCU = 6; // to hide latency constexpr auto threadsPerBlock = 256; size_t Nbytes = N * sizeof(int); HIPCHECK(hipSetDevice(gpu)); - HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - HIPCHECK(hipMemGetInfo(&curAvl, &curTot)); - - if (!concurOnOneGPU && (prevAvl < curAvl || prevTot != curTot)) { - //In concurrent calls on one GPU, we cannot verify leaking in this way - printf("%s : Memory allocation mismatch observed." - "Possible memory leak.\n", __func__); - TestPassed &= false; - } unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); @@ -229,24 +176,7 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) { TestPassed = false; } - HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot)); HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HIPCHECK(hipMemGetInfo(&curAvl, &curTot)); - - if (!concurOnOneGPU && (curAvl < prevAvl || prevTot != curTot)) { - // In concurrent calls on one GPU, we cannot verify leaking in this way - UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed." - << "Possible memory leak."); - TestPassed = false; - } - - if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) { - // In concurrent calls on one GPU, we cannot verify leaking in this way - UNSCOPED_INFO( - "validateMemoryOnGpuMThread : Memory allocation mismatch observed." - "Possible memory leak."); - TestPassed = false; - } return TestPassed; } @@ -256,8 +186,7 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) { * In Multithreaded Environment */ static bool regressAllocInLoopMthread(int gpu) { - bool TestPassed = true; - size_t tot, avail, ptot, pavail, numBytes; + size_t numBytes; int i = 0; int* ptr; @@ -266,41 +195,18 @@ static bool regressAllocInLoopMthread(int gpu) { // Exercise allocation in loop with bigger chunks for (i = 0; i < MaxAllocFree_BigChunks; i++) { - HIPCHECK(hipMemGetInfo(&pavail, &ptot)); HIPCHECK(hipMalloc(&ptr, numBytes)); - HIPCHECK(hipMemGetInfo(&avail, &tot)); HIPCHECK(hipFree(ptr)); - - if (pavail - avail < numBytes) { // We expect pavail-avail >= numBytes - UNSCOPED_INFO("LoopAllocation " << i << " : Memory allocation of " << numBytes - << " not matching with hipMemGetInfo - FAIL." - << "pavail=" << pavail << ", ptot=" << ptot - << ", avail=" << avail << ", tot=" << tot - << ", pavail-avail=" << pavail - avail); - TestPassed = false; - break; - } } // Exercise allocation in loop with smaller chunks and maximum iters - HIPCHECK(hipMemGetInfo(&pavail, &ptot)); numBytes = BuffSizeSC; - for (i = 0; i < MaxAllocFree_SmallChunks; i++) { HIPCHECK(hipMalloc(&ptr, numBytes)); - HIPCHECK(hipFree(ptr)); } - HIPCHECK(hipMemGetInfo(&avail, &tot)); - - if ((pavail != avail) || (ptot != tot)) { - UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." - << "Possible memory leak."); - TestPassed &= false; - } - - return TestPassed; + return true; } /* @@ -359,7 +265,7 @@ TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") { * of time. */ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") { - size_t avail{0}, tot{0}, pavail{0}, ptot{0}; + size_t avail{0}, tot{0}; bool ret{false}; hipError_t err{}; std::vector ptrlist{}; @@ -370,8 +276,6 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") { HIP_CHECK(hipGetDeviceCount(&devCnt)); REQUIRE(devCnt > 0); - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); - // Allocate small chunks of memory million times for (int i = 0; i < MaxAllocPoolIter; i++) { if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) { @@ -392,13 +296,8 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") { for (auto& t : ptrlist) { HIP_CHECK(hipFree(t)); } - - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - ret = validateMemoryOnGPU(0); REQUIRE(ret == true); - REQUIRE(pavail == avail); - REQUIRE(ptot == tot); } diff --git a/catch/unit/memory/hipMallocPitch.cc b/catch/unit/memory/hipMallocPitch.cc index 5a20671e14..19c803c7c2 100644 --- a/catch/unit/memory/hipMallocPitch.cc +++ b/catch/unit/memory/hipMallocPitch.cc @@ -346,8 +346,6 @@ static void MemoryAllocDiffSizes(int gpu) { } else { width = LARGECHUNK_NUMW * sizeof(T); } - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); for (int i = 0; i < CHUNK_LOOP; i++) { HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d[i]), &pitch_A, width, sizes)); @@ -355,10 +353,6 @@ static void MemoryAllocDiffSizes(int gpu) { for (int i = 0; i < CHUNK_LOOP; i++) { HIP_CHECK(hipFree(A_d[i])); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - if (pavail != avail) { - HIPASSERT(false); - } } } @@ -487,9 +481,6 @@ TEST_CASE("Unit_hipMallocPitch_MultiThread", "") { int devCnt = 0; devCnt = HipTest::getDeviceCount(); - - size_t tot, avail, ptot, pavail; - HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); for (int i = 0; i < devCnt; i++) { threadlist.push_back(std::thread(threadFunc, i)); } @@ -497,12 +488,6 @@ TEST_CASE("Unit_hipMallocPitch_MultiThread", "") { for (auto &t : threadlist) { t.join(); } - HIP_CHECK(hipMemGetInfo(&avail, &tot)); - - if (pavail != avail) { - WARN("Memory leak of hipMallocPitch API in multithreaded scenario"); - REQUIRE(false); - } } /*