EXSWHTEC-106 - Reimplement tests for hipOccupancyMaxActiveBlocksPerMultiprocessor and hipOccupancyMaxPotentialBlockSize APIs

catch/hipTestMain/config/config_amd_linux_common.json

@@ -14,6 +14,7 @@
 	   "Unit_hipInit_Negative",
 	   "Unit_hipMemset_Negative_OutOfBoundsPtr",
 	   "Unit_hipDeviceReset_Positive_Basic",
-	   "Unit_hipDeviceReset_Positive_Threaded"
+	   "Unit_hipDeviceReset_Positive_Threaded",
+	   "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters"
 	]
 }

catch/hipTestMain/config/config_amd_windows_common.json

@@ -93,6 +93,7 @@
 	   "Note: needs to be enabled when streamPerThread issues are fixed",
 	   "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread",
 	   "Note: intermittent Seg fault failure ",
-	   "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags"
+	   "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags",
+	   "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters"
 	]
 }

catch/unit/occupancy/CMakeLists.txt

@@ -1,7 +1,9 @@
 # Common Tests - Test independent of all platforms
 set(TEST_SRC
   hipOccupancyMaxActiveBlocksPerMultiprocessor.cc
+  hipOccupancyMaxActiveBlocksPerMultiprocessor_old.cc
   hipOccupancyMaxPotentialBlockSize.cc
+  hipOccupancyMaxPotentialBlockSize_old.cc
 )
 
 hip_add_exe_to_target(NAME OccupancyTest

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>
-
-static __global__ void f1(float *a) { *a = 1.0; }
+/*
+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"
 
-template <typename T>
-static __global__ void f2(T *a) { *a = 1; }
+static __global__ void f1(float* a) { *a = 1.0; }
 
-/**
- * Defines
- */
-#define OccupancyDisableCachingOverride 0x01
+template <typename T> static __global__ void f2(T* a) { *a = 1; }
 
-TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative") {
-  hipError_t ret;
-  int numBlock = 0, blockSize = 0;
-  int gridSize = 0, defBlkSize = 32;
+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));
 
-  // 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);
+  // 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_rangeValidation") {
+TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation") {
   hipDeviceProp_t devProp;
-  int numBlock = 0, blockSize = 0;
+  int blockSize = 0;
   int gridSize = 0;
 
-  // Get potential blocksize
-  HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0));
-
   HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
 
-  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0));
-
-  // Check if numBlocks and blockSize are within limits
-  REQUIRE(numBlock > 0);
-  REQUIRE((numBlock * blockSize) <= devProp.maxThreadsPerMultiProcessor);
-
-  // 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);
+  SECTION("dynSharedMemPerBlk = 0") {
+    // Get potential blocksize
+    HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0));
+
+    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));
+
+    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);
+  }
+}

catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor_old.cc

@@ -0,0 +1,91 @@
+/*
+Copyright (c) 2021 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 WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+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>
+
+static __global__ void f1(float *a) { *a = 1.0; }
+
+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;
+  int gridSize = 0;
+
+  // Get potential blocksize
+  HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0));
+
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0));
+
+  // Check if numBlocks and blockSize are within limits
+  REQUIRE(numBlock > 0);
+  REQUIRE((numBlock * blockSize) <= devProp.maxThreadsPerMultiProcessor);
+
+  // 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);
+}
+
+TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_templateInvocation") {
+  int blockSize = 32;
+  int numBlock = 0;
+
+  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>
+                                                  (&numBlock, f2, blockSize, 0));
+  REQUIRE(numBlock > 0);
+}
+