diff --git a/projects/hip/tests/catch/stress/memory/hipMalloc.cc b/projects/hip/tests/catch/stress/memory/hipMalloc.cc new file mode 100644 index 0000000000..a3057a6e4c --- /dev/null +++ b/projects/hip/tests/catch/stress/memory/hipMalloc.cc @@ -0,0 +1,50 @@ +/* + 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 WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING 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 ANNY 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. + */ + +#include + +#include + +// Stress allocation tests +// Try to allocate as much memory as possible +// But since max allocation can fail, we need to be happy with atleast 1/4th of memory +TEST_CASE("Stress_hipMalloc_HighSizeAlloc") { + size_t devMemTotal{0}, devMemFree{0}; + HIP_CHECK(hipMemGetInfo(&devMemFree, &devMemTotal)); + REQUIRE(devMemFree > 0); + REQUIRE(devMemTotal > 0); + + char* d_ptr{nullptr}; + size_t counter{0}; + + INFO("Free Mem Available: " << devMemFree << " bytes out of " << devMemTotal << " bytes!"); + while (hipMalloc(&d_ptr, devMemFree) != hipSuccess && devMemFree > 1) { + counter++; + devMemFree >>= 1; // reduce the memory to be allocated by half + INFO("Attempt to allocate " << devMemFree << " bytes out of " << devMemTotal + << " bytes failed!"); + REQUIRE(counter <= 2); // Make sure that we are atleast able to allocate 1/4th of max memory + } + + HIP_CHECK(hipMemset(d_ptr, 1, devMemFree)); + auto ptr = std::unique_ptr{new unsigned char[devMemFree]}; + HIP_CHECK(hipMemcpy(ptr.get(), d_ptr, devMemFree, hipMemcpyDeviceToHost)); + HIP_CHECK(hipFree(d_ptr)); + REQUIRE(std::all_of(ptr.get(), ptr.get() + devMemFree, [](unsigned char n) { return n == 1; })); +} diff --git a/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc b/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc index 7e992cd152..e582056685 100644 --- a/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc +++ b/projects/hip/tests/catch/stress/memory/hipMallocManagedStress.cc @@ -150,7 +150,7 @@ static int HmmAttrPrint() { // The following test case allocation, host access, device access of HMM // memory from size 1 to 10KB -TEST_CASE("Unit_hipMallocManaged_MultiSize") { +TEST_CASE("Stress_hipMallocManaged_MultiSize") { IfTestPassed = true; int managed = HmmAttrPrint(); if (managed == 1) { @@ -196,7 +196,7 @@ TEST_CASE("Unit_hipMallocManaged_MultiSize") { // The following test case tests the behavior of kernel with a HMM memory and // hipMalloc memory -TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") { +TEST_CASE("Stress_hipMallocManaged_KrnlWth2MemTypes") { IfTestPassed = true; int *Hmm = NULL, *Dptr = NULL, InitVal = 123; size_t NumElms = (1024 * 1024); @@ -241,7 +241,7 @@ TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") { // The following test case tests when the same Hmm memory is used for // launching multiple different kernels will results in any issue -TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") { +TEST_CASE("Stress_hipMallocManaged_MultiKrnlHmmAccess") { int managed = HmmAttrPrint(); if (managed) { int InitVal = 123, NumElms = (1024 * 1024); @@ -253,7 +253,7 @@ TEST_CASE("Unit_hipMallocManaged_MultiKrnlHmmAccess") { } // Testing the allocation of/scenarios around max possible memory -TEST_CASE("Unit_hipMallocManaged_ExtremeSizes") { +TEST_CASE("Stress_hipMallocManaged_ExtremeSizes") { int managed = HmmAttrPrint(); if (managed == 1) { bool IfTestPassed = true; diff --git a/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc b/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc index 98bc2ab014..592ce86897 100644 --- a/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc +++ b/projects/hip/tests/catch/unit/memory/hipMallocConcurrency.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -51,14 +51,13 @@ Testcase Scenarios : #include #include - -#include -#include #include +#include +#include /* Buffer size for bigger chunks in alloc/free cycles */ -static constexpr auto BuffSizeBC = 5*1024*1024; +static constexpr auto BuffSizeBC = 5 * 1024 * 1024; /* Buffer size for smaller chunks in alloc/free cycles */ static constexpr auto BuffSizeSC = 16; @@ -68,19 +67,18 @@ static constexpr auto BuffSizeSC = 16; static constexpr auto NumDiv = 100; /* Max alloc/free iterations for smaller chunks */ -static constexpr auto MaxAllocFree_SmallChunks = (5000000/NumDiv); +static constexpr auto MaxAllocFree_SmallChunks = (5000000 / NumDiv); /* Max alloc/free iterations for bigger chunks */ static constexpr auto MaxAllocFree_BigChunks = 10000; /* Max alloc and pool iterations */ -static constexpr auto MaxAllocPoolIter = (2000000/NumDiv); +static constexpr auto MaxAllocPoolIter = (2000000 / NumDiv); /* Test status shared across threads */ static std::atomic g_thTestPassed{true}; - /** * Validates data consistency on supplied gpu */ @@ -103,9 +101,8 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), - 0, 0, static_cast(A_d), - static_cast(B_d), C_d, N); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, + static_cast(A_d), static_cast(B_d), C_d, N); HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); @@ -121,9 +118,8 @@ static bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) { if (!concurOnOneGPU && (prevAvl != curAvl || 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."); + UNSCOPED_INFO("validateMemoryOnGPU : Memory allocation mismatch observed." + << "Possible memory leak."); TestPassed = false; } @@ -138,7 +134,7 @@ static bool regressAllocInLoop(int gpu) { bool TestPassed = true; size_t tot, avail, ptot, pavail, numBytes; int i = 0; - int *ptr; + int* ptr; HIP_CHECK(hipSetDevice(gpu)); numBytes = BuffSizeBC; @@ -150,11 +146,12 @@ static bool regressAllocInLoop(int gpu) { 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); + 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; } @@ -173,8 +170,8 @@ static bool regressAllocInLoop(int gpu) { HIP_CHECK(hipMemGetInfo(&avail, &tot)); if ((pavail != avail) || (ptot != tot)) { - UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." << - "Possible memory leak."); + UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." + << "Possible memory leak."); TestPassed &= false; } @@ -203,9 +200,8 @@ static bool validateMemoryOnGpuMThread(int gpu, bool concurOnOneGPU = false) { HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); - hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), - 0, 0, static_cast(A_d), - static_cast(B_d), C_d, N); + hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0, + static_cast(A_d), static_cast(B_d), C_d, N); HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); @@ -238,7 +234,7 @@ static bool regressAllocInLoopMthread(int gpu) { bool TestPassed = true; size_t tot, avail, ptot, pavail, numBytes; int i = 0; - int *ptr; + int* ptr; HIPCHECK(hipSetDevice(gpu)); numBytes = BuffSizeBC; @@ -250,11 +246,12 @@ static bool regressAllocInLoopMthread(int gpu) { 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); + 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; } @@ -273,8 +270,8 @@ static bool regressAllocInLoopMthread(int gpu) { HIPCHECK(hipMemGetInfo(&avail, &tot)); if ((pavail != avail) || (ptot != tot)) { - UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." << - "Possible memory leak."); + UNSCOPED_INFO("LoopAllocation : Memory allocation mismatch observed." + << "Possible memory leak."); TestPassed &= false; } @@ -285,18 +282,15 @@ static bool regressAllocInLoopMthread(int gpu) { * Thread func to regress alloc and check data consistency */ static void threadFunc(int gpu) { - g_thTestPassed = regressAllocInLoopMthread(gpu) - && validateMemoryOnGpuMThread(gpu); + g_thTestPassed = regressAllocInLoopMthread(gpu) && validateMemoryOnGpuMThread(gpu); - UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << - g_thTestPassed.load()); + UNSCOPED_INFO("thread execution status on gpu" << gpu << ":" << g_thTestPassed.load()); } /* Performs Argument Validation of api */ TEST_CASE("Unit_hipMalloc_ArgumentValidation") { - int *ptr; - hipError_t ret; + int* ptr{nullptr}; SECTION("hipMalloc() when size(0)") { HIP_CHECK(hipMalloc(&ptr, 0)); @@ -304,21 +298,17 @@ TEST_CASE("Unit_hipMalloc_ArgumentValidation") { REQUIRE(ptr == nullptr); } - SECTION("hipFree() when freeing nullptr ") { - ptr = nullptr; - // api should return success and shudnt crash + SECTION("hipFree() when freeing nullptr") { HIP_CHECK(hipFree(ptr)); } SECTION("hipMalloc() with invalid argument") { - constexpr auto sizeBytes = 100; - ret = hipMalloc(nullptr, sizeBytes); - REQUIRE(ret != hipSuccess); + HIP_CHECK_ERROR(hipMalloc(nullptr, 100), hipErrorInvalidValue); } SECTION("hipMalloc() with max size_t") { - ret = hipMalloc(&ptr, std::numeric_limits::max()); - REQUIRE(ret != hipSuccess); + HIP_CHECK_ERROR(hipMalloc(&ptr, std::numeric_limits::max()), + hipErrorMemoryAllocation); } } @@ -344,12 +334,12 @@ TEST_CASE("Unit_hipMalloc_LoopRegressionAllocFreeCycles") { * of time. */ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") { - size_t avail, tot, pavail, ptot; - bool ret; - hipError_t err; - std::vector ptrlist; + size_t avail{0}, tot{0}, pavail{0}, ptot{0}; + bool ret{false}; + hipError_t err{}; + std::vector ptrlist{}; constexpr auto BuffSize = 10; - int devCnt, *ptr; + int devCnt{0}, *ptr{nullptr}; // Get GPU count HIP_CHECK(hipGetDeviceCount(&devCnt)); @@ -358,14 +348,13 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") { HIP_CHECK(hipMemGetInfo(&pavail, &ptot)); // Allocate small chunks of memory million times - for (int i = 0; i < MaxAllocPoolIter ; i++) { + for (int i = 0; i < MaxAllocPoolIter; i++) { if ((err = hipMalloc(&ptr, BuffSize)) != hipSuccess) { HIP_CHECK(hipMemGetInfo(&avail, &tot)); - INFO("Loop regression pool allocation failure. " << - "Total gpu memory " << tot/(1024.0*1024.0) <<", Free memory " << - avail/(1024.0*1024.0) << " iter " << i << " error " - << hipGetErrorString(err)); + INFO("Loop regression pool allocation failure. " + << "Total gpu memory " << tot / (1024.0 * 1024.0) << ", Free memory " + << avail / (1024.0 * 1024.0) << " iter " << i << " error " << hipGetErrorString(err)); REQUIRE(false); } @@ -375,7 +364,7 @@ TEST_CASE("Unit_hipMalloc_AllocateAndPoolBuffers") { } // Free ptrs at later point of time - for ( auto &t : ptrlist ) { + for (auto& t : ptrlist) { HIP_CHECK(hipFree(t)); } @@ -404,7 +393,7 @@ TEST_CASE("Unit_hipMalloc_Multithreaded_MultiGPU") { threadlist.push_back(std::thread(threadFunc, i)); } - for (auto &t : threadlist) { + for (auto& t : threadlist) { t.join(); }