SWDEV-387925 - Remove comparisons that are based on available memory returned from hipMemGetInfo (#262)
Change-Id: I3e4b539425482bbda578fd3127f118d47c611db3
Tento commit je obsažen v:
odevzdal
GitHub
rodič
e7348b1f04
revize
6ce5d7457c
@@ -55,8 +55,6 @@ static void Malloc3DArray_DiffSizes(int gpu) {
|
||||
size_t width{size}, height{size}, depth{size};
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
|
||||
std::array<hipArray_t, ARRAY_LOOP> 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<std::thread> 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 {
|
||||
|
||||
@@ -57,9 +57,6 @@ static void MallocArray_DiffSizes(int gpu) {
|
||||
for (const auto& size : runs) {
|
||||
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
|
||||
std::array<hipArray_t, ARRAY_LOOP> 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<std::thread> 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 ///////////////////////////////////////
|
||||
|
||||
@@ -85,7 +85,6 @@ static std::atomic<bool> 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<int*> 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);
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -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<void**>(&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);
|
||||
}
|
||||
}
|
||||
|
||||
/*
|
||||
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele