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

@@ -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);
 

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();
 }