EXSWHTEC-106 - Reimplement tests for hipOccupancyMaxActiveBlocksPerMultiprocessor and hipOccupancyMaxPotentialBlockSize APIs (#46)

- Reimplement tests for hipOccupancyMaxActiveBlocksPerMultiprocessor and hipOccupancyMaxPotentialBlockSize APIs
- Add helper file occupancy_common.hh with parameterized templates
- Expand all positive and negative tests to use the templates
- Change section disable macros
- Disable AMD specific test due to defect
- Disable negative test using json file
- Fix disabled files

catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.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
@@ -16,76 +16,103 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
-#include <hip_test_common.hh>
+/*
+Testcase Scenarios :
+Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation - Test correct execution
+of hipOccupancyMaxActiveBlocksPerMultiprocessor for diffrent parameter values
+Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_TemplateInvocation - Test correct
+execution of hipOccupancyMaxActiveBlocksPerMultiprocessor template for diffrent parameter values
+Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters - Test unsuccessful execution
+of hipOccupancyMaxActiveBlocksPerMultiprocessor api when parameters are invalid
+*/
+#include "occupancy_common.hh"
 
-static __global__ void f1(float *a) { *a = 1.0; }
+static __global__ void f1(float* a) { *a = 1.0; }
 
-template <typename T>
-static __global__ void f2(T *a) { *a = 1; }
+template <typename T> static __global__ void f2(T* a) { *a = 1; }
 
-/**
- * Defines
- */
-#define OccupancyDisableCachingOverride 0x01
-
-TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative") {
-  hipError_t ret;
-  int numBlock = 0, blockSize = 0;
-  int gridSize = 0, defBlkSize = 32;
-
-  // Get potential blocksize
-  HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0));
-
-  // Validate each argument
-  ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(NULL, f1, blockSize, 0);
-  REQUIRE(ret != hipSuccess);
-
-  ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, NULL, blockSize, 0);
-  REQUIRE(ret != hipSuccess);
-
-  ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, 0, 0);
-  REQUIRE(ret != hipSuccess);
-
-  ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, 0,
-                                         std::numeric_limits<std::size_t>::max());
-  REQUIRE(ret != hipSuccess);
-
-  ret = hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(&numBlock, f1,
-                                  defBlkSize, 0, OccupancyDisableCachingOverride);
-  REQUIRE(ret == hipSuccess);
-}
-
-TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_rangeValidation") {
-  hipDeviceProp_t devProp;
-  int numBlock = 0, blockSize = 0;
+TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters") {
+  int numBlocks = 0;
+  int blockSize = 0;
   int gridSize = 0;
 
   // Get potential blocksize
   HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0));
 
+  // Common negative tests
+  MaxActiveBlocksPerMultiprocessorNegative(
+      [](int* numBlocks, int blockSize, size_t dynSharedMemPerBlk) {
+        return hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f1, blockSize,
+                                                            dynSharedMemPerBlk);
+      },
+      blockSize);
+
+  SECTION("Kernel function is NULL") {
+    HIP_CHECK_ERROR(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, NULL, blockSize, 0),
+                    hipErrorInvalidDeviceFunction);
+  }
+}
+
+TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation") {
+  hipDeviceProp_t devProp;
+  int blockSize = 0;
+  int gridSize = 0;
+
   HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
 
-  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0));
+  SECTION("dynSharedMemPerBlk = 0") {
+    // Get potential blocksize
+    HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0));
 
-  // Check if numBlocks and blockSize are within limits
-  REQUIRE(numBlock > 0);
-  REQUIRE((numBlock * blockSize) <= devProp.maxThreadsPerMultiProcessor);
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize](int* numBlocks) {
+          return hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f1, blockSize, 0);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
+  SECTION("dynSharedMemPerBlk = sharedMemPerBlock") {
+    // Get potential blocksize
+    HIP_CHECK(
+        hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, devProp.sharedMemPerBlock, 0));
 
-  // Validate numBlock after passing dynSharedMemPerBlk
-  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize,
-                                                           devProp.sharedMemPerBlock));
-
-  // Check if numBlocks and blockSize are within limits
-  REQUIRE(numBlock > 0);
-  REQUIRE((numBlock * blockSize) <= devProp.maxThreadsPerMultiProcessor);
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize, devProp](int* numBlocks) {
+          return hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f1, blockSize,
+                                                              devProp.sharedMemPerBlock);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
 }
 
-TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_templateInvocation") {
-  int blockSize = 32;
-  int numBlock = 0;
+TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_TemplateInvocation") {
+  hipDeviceProp_t devProp;
+  int blockSize = 0;
+  int gridSize = 0;
 
-  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>
-                                                  (&numBlock, f2, blockSize, 0));
-  REQUIRE(numBlock > 0);
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  SECTION("dynSharedMemPerBlk = 0") {
+    // Get potential blocksize
+    HIP_CHECK(hipOccupancyMaxPotentialBlockSize<void (*)(int*)>(&gridSize, &blockSize, f2, 0, 0));
+
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize](int* numBlocks) {
+          return hipOccupancyMaxActiveBlocksPerMultiprocessor<void (*)(int*)>(numBlocks, f2,
+                                                                              blockSize, 0);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
+
+  SECTION("dynSharedMemPerBlk = sharedMemPerBlock") {
+    // Get potential blocksize
+    HIP_CHECK(hipOccupancyMaxPotentialBlockSize<void (*)(int*)>(&gridSize, &blockSize, f2,
+                                                                devProp.sharedMemPerBlock, 0));
+
+    MaxActiveBlocksPerMultiprocessor(
+        [blockSize, devProp](int* numBlocks) {
+          return hipOccupancyMaxActiveBlocksPerMultiprocessor<void (*)(int*)>(
+              numBlocks, f2, blockSize, devProp.sharedMemPerBlock);
+        },
+        blockSize, devProp.maxThreadsPerMultiProcessor);
+  }
 }
-