diff --git a/tests/catch/hipTestMain/CMakeLists.txt b/tests/catch/hipTestMain/CMakeLists.txt index 5102046a1e..c24c63bfa4 100644 --- a/tests/catch/hipTestMain/CMakeLists.txt +++ b/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 diff --git a/tests/catch/unit/CMakeLists.txt b/tests/catch/unit/CMakeLists.txt index 5aca994a4c..3c9b3cb2b3 100644 --- a/tests/catch/unit/CMakeLists.txt +++ b/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) diff --git a/tests/catch/unit/occupancy/CMakeLists.txt b/tests/catch/unit/occupancy/CMakeLists.txt new file mode 100644 index 0000000000..a349698932 --- /dev/null +++ b/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) diff --git a/tests/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc b/tests/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc new file mode 100644 index 0000000000..8b91ca2f7e --- /dev/null +++ b/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 + +static __global__ void f1(float *a) { *a = 1.0; } + +template +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::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 + (&numBlock, f2, blockSize, 0)); + REQUIRE(numBlock > 0); +} + diff --git a/tests/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc b/tests/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc new file mode 100644 index 0000000000..ccc2c93940 --- /dev/null +++ b/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 + +static __global__ void f1(float *a) { *a = 1.0; } + +template +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(&gridSize, + &blockSize, f2, 0, 0)); + REQUIRE(gridSize > 0); + REQUIRE(blockSize > 0); +} +