SWDEV-289406 - Occuancy test migration

Change-Id: Ic6bce9bbad908bd210727f9981f7a3fc750a91ed

tests/catch/hipTestMain/CMakeLists.txt

@@ -12,6 +12,7 @@ endif()
 target_link_libraries(UnitTests PRIVATE UnitDeviceTests
                                         MemoryTest
                                         StreamTest
+                                        OccupancyTest
                                         stdc++fs)
 
 # Add AMD Only Tests

tests/catch/unit/CMakeLists.txt

@@ -1,5 +1,6 @@
 add_subdirectory(memory)
 add_subdirectory(deviceLib)
 add_subdirectory(stream)
+add_subdirectory(occupancy)
 # Disable Saxpy test temporarily to see if CI Passes
 # add_subdirectory(rtc)

tests/catch/unit/occupancy/CMakeLists.txt

@@ -0,0 +1,11 @@
+# Common Tests - Test independent of all platforms
+set(TEST_SRC
+  hipOccupancyMaxActiveBlocksPerMultiprocessor.cc
+  hipOccupancyMaxPotentialBlockSize.cc
+)
+
+# Create shared lib of all tests
+add_library(OccupancyTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+
+# Add dependency on build_tests to build it on this custom target
+add_dependencies(build_tests OccupancyTest)

tests/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc

@@ -0,0 +1,91 @@
+/*
+Copyright (c) 2021-Present 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);
+}
+

tests/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc

@@ -0,0 +1,80 @@
+/*
+Copyright (c) 2021-Present 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; }
+
+TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Negative") {
+  hipError_t ret;
+  int blockSize = 0;
+  int gridSize = 0;
+
+  // Validate each argument
+  ret = hipOccupancyMaxPotentialBlockSize(NULL, &blockSize, f1, 0, 0);
+  REQUIRE(ret != hipSuccess);
+
+  ret = hipOccupancyMaxPotentialBlockSize(&gridSize, NULL, f1, 0, 0);
+  REQUIRE(ret != hipSuccess);
+
+#ifndef __HIP_PLATFORM_NVIDIA__
+  // nvcc doesnt support kernelfunc(NULL) for api
+  ret = hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, NULL, 0, 0);
+  REQUIRE(ret != hipSuccess);
+#endif
+}
+
+TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_rangeValidation") {
+  hipDeviceProp_t devProp;
+  int blockSize = 0;
+  int gridSize = 0;
+
+  // Get potential blocksize
+  HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0));
+
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  // Check if blockSize doen't exceed maxThreadsPerBlock
+  REQUIRE(gridSize > 0); REQUIRE(blockSize > 0);
+  REQUIRE(blockSize <= devProp.maxThreadsPerBlock);
+
+  // Pass dynSharedMemPerBlk, blockSizeLimit and check out param
+  blockSize = 0;
+  gridSize = 0;
+
+  HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1,
+           devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock));
+
+  // Check if blockSize doen't exceed maxThreadsPerBlock
+  REQUIRE(gridSize > 0); REQUIRE(blockSize > 0);
+  REQUIRE(blockSize <= devProp.maxThreadsPerBlock);
+
+}
+
+TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_templateInvocation") {
+  int gridSize = 0, blockSize = 0;
+
+  HIP_CHECK(hipOccupancyMaxPotentialBlockSize<void(*)(int *)>(&gridSize,
+                       &blockSize, f2, 0, 0));
+  REQUIRE(gridSize > 0);
+  REQUIRE(blockSize > 0);
+}
+