From 7def265d94fe82d289dd6522aaa05e686aa3f3f5 Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Tue, 12 Mar 2024 12:49:42 +0000 Subject: [PATCH] SWDEV-449111 - Lesser CU size tends to have block & thread dim = 0. OOM as PAL uses half of CPU memory once GPU memory exhausts. Change-Id: I2399f7247f5172fdc7980a4652ff68df4b5598c5 --- catch/unit/cooperativeGrps/thread_block.cc | 12 +++++++++++ catch/unit/memory/hipFreeMipmappedArray.cc | 24 ++++++++++++++-------- 2 files changed, 28 insertions(+), 8 deletions(-) diff --git a/catch/unit/cooperativeGrps/thread_block.cc b/catch/unit/cooperativeGrps/thread_block.cc index 0b7a5eff7e..a5018a9cec 100644 --- a/catch/unit/cooperativeGrps/thread_block.cc +++ b/catch/unit/cooperativeGrps/thread_block.cc @@ -80,6 +80,10 @@ static __global__ void thread_block_non_member_thread_rank_getter(unsigned int* TEST_CASE("Unit_Thread_Block_Getters_Positive_Basic") { const auto blocks = GenerateBlockDimensions(); const auto threads = GenerateThreadDimensions(); + if (blocks.x <= 0 || blocks.y <= 0 || blocks.z <= 0 || + threads.x <= 0 || threads.y <= 0 || threads.z <= 0) { + return; + } INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z); INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z); const CPUGrid grid(blocks, threads); @@ -154,6 +158,10 @@ TEST_CASE("Unit_Thread_Block_Getters_Positive_Basic") { TEST_CASE("Unit_Thread_Block_Getters_Via_Base_Type_Positive_Basic") { const auto blocks = GenerateBlockDimensions(); const auto threads = GenerateThreadDimensions(); + if (blocks.x <= 0 || blocks.y <= 0 || blocks.z <= 0 || + threads.x <= 0 || threads.y <= 0 || threads.z <= 0) { + return; + } INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z); INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z); @@ -201,6 +209,10 @@ TEST_CASE("Unit_Thread_Block_Getters_Via_Base_Type_Positive_Basic") { TEST_CASE("Unit_Thread_Block_Getters_Via_Non_Member_Functions_Positive_Basic") { const auto blocks = GenerateBlockDimensions(); const auto threads = GenerateThreadDimensions(); + if (blocks.x <= 0 || blocks.y <= 0 || blocks.z <= 0 || + threads.x <= 0 || threads.y <= 0 || threads.z <= 0) { + return; + } INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z); INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z); diff --git a/catch/unit/memory/hipFreeMipmappedArray.cc b/catch/unit/memory/hipFreeMipmappedArray.cc index df6c4faee0..4fe2fb7436 100644 --- a/catch/unit/memory/hipFreeMipmappedArray.cc +++ b/catch/unit/memory/hipFreeMipmappedArray.cc @@ -107,20 +107,28 @@ TEMPLATE_TEST_CASE("Unit_hipFreeMipmappedArrayMultiTArray", "", char, int) { extent.height = GENERATE(64, 256, 1024); extent.depth = GENERATE(0, 64, 256, 1024); - for (auto& ptr : ptrs) { - HIP_CHECK(hipMallocMipmappedArray(&ptr, &desc, extent, numLevels, flags)); + int i = 0; + for (; i < ptrs.size(); i++) { + if (hipErrorOutOfMemory == hipMallocMipmappedArray(&ptrs[i], &desc, extent, + numLevels, flags)) { + break; + } } - for (auto ptr : ptrs) { - threads.emplace_back(([ptr] { - HIP_CHECK_THREAD(hipFreeMipmappedArray(ptr)); - HIP_CHECK_THREAD(hipStreamQuery(nullptr)); - })); + for (int j = 0; j < i; j++) { + threads.emplace_back([ptrs,j] { + if (hipSuccess != hipFreeMipmappedArray(ptrs[j])) { + return; + } + if (hipSuccess != hipStreamQuery(nullptr)) { + return; + } + }); } for (auto& t : threads) { t.join(); } - + HIP_CHECK_THREAD_FINALIZE(); }