From f7249e092bb71beffdb7c93b1f866b9f78add1fc Mon Sep 17 00:00:00 2001 From: venkatesh-amd Date: Fri, 14 Nov 2025 15:41:45 +0530 Subject: [PATCH] =?UTF-8?q?SWDEV-533237=20:=20Added=20test=20cases=20for?= =?UTF-8?q?=20hipOccupancyAvailableDynamicSMemPer=E2=80=A6=20(#716)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * SWDEV-533237 Added test cases for hipOccupancyAvailableDynamicSMemPerBlock API * SWDEV-533237 : Added test cases for hipOccupancyAvailableDynamicSMemPerBlock * SWDEV-533237 : Addressed review comments for hipOccupancyAvailableDynamicSMemPerBlock aip test cases --------- Co-authored-by: jainprad <92369414+jainprad@users.noreply.github.com> --- .../catch/unit/occupancy/CMakeLists.txt | 1 + ...ipOccupancyAvailableDynamicSMemPerBlock.cc | 154 ++++++++++++++++++ 2 files changed, 155 insertions(+) create mode 100644 projects/hip-tests/catch/unit/occupancy/hipOccupancyAvailableDynamicSMemPerBlock.cc diff --git a/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt b/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt index 3f5a3b9827..ba92c278ae 100644 --- a/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt @@ -10,6 +10,7 @@ set(TEST_SRC hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.cc hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc + hipOccupancyAvailableDynamicSMemPerBlock.cc ) add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/simple_kernel.code diff --git a/projects/hip-tests/catch/unit/occupancy/hipOccupancyAvailableDynamicSMemPerBlock.cc b/projects/hip-tests/catch/unit/occupancy/hipOccupancyAvailableDynamicSMemPerBlock.cc new file mode 100644 index 0000000000..35b34a6169 --- /dev/null +++ b/projects/hip-tests/catch/unit/occupancy/hipOccupancyAvailableDynamicSMemPerBlock.cc @@ -0,0 +1,154 @@ + +/* +Copyright (c) 2025 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. +*/ + +/** + * @brief Returns dynamic shared memory available per block when launching + * numBlocks blocks on SM. + * + * @ingroup Occupancy + * Returns in dynamicSmemSize the maximum size of dynamic shared memory + * to allow numBlocks blocks per SM. + * + * @param [out] dynamicSmemSize Returned maximum dynamic shared memory. + * @param [in] f Kernel function for which occupancy is + * calculated. + * @param [in] numBlocks Number of blocks to fit on SM + * @param [in] blockSize Size of the block + * + * @return #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidDeviceFunction, + * #hipErrorInvalidValue, #hipErrorUnknown + */ + +#include +#include +#include + +constexpr size_t SIZE = 1024; + +static __global__ void dynamicReverse(int *d, int n) { + __shared__ int s[SIZE]; + + int t = threadIdx.x; // unique index it starts from 0 + int tr = n - t - 1; // it is reverse index of t, if t = 0, then tr = n-1 + + s[t] = d[t]; + __syncthreads(); + d[t] = s[tr]; +} + +static __global__ void f1(float *a) { *a = 1.0; } + +/** + * Test Description + * ------------------------ + * - This test will verify hipOccupancyAvailableDynamicSMemPerBlock api with + * invalid parameters + * Test source + * ------------------------ + * - occupancy/hipOccupancyAvailableDynamicSMemPerBlock.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipOccupancyAvailableDynamicSMemPerBlock_Negative") { + size_t dynamicSmemSize; + int numBlocks = 1; + SECTION("Number of blocks are zero") { + HIP_CHECK_ERROR( + hipOccupancyAvailableDynamicSMemPerBlock(&dynamicSmemSize, f1, 0, SIZE), + hipErrorInvalidValue); + } + + SECTION("Block size is zero") { + HIP_CHECK_ERROR(hipOccupancyAvailableDynamicSMemPerBlock(&dynamicSmemSize, + f1, numBlocks, 0), + hipErrorInvalidValue); + } + +#if HT_AMD + SECTION("Invalid driver funtion") { + HIP_CHECK_ERROR(hipOccupancyAvailableDynamicSMemPerBlock( + &dynamicSmemSize, NULL, numBlocks, SIZE), + hipErrorInvalidDeviceFunction); + } + + SECTION("dynamicSmemSize is null") { + HIP_CHECK_ERROR( + hipOccupancyAvailableDynamicSMemPerBlock(nullptr, f1, numBlocks, SIZE), + hipErrorInvalidValue); + } +#endif +} + +/** + * Test Description + * ------------------------ + * - This test will verify funtionality of + * hipOccupancyAvailableDynamicSMemPerBlock api Test source + * ------------------------ + * - occupancy/hipOccupancyAvailableDynamicSMemPerBlock.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipOccupancyAvailableDynamicSMemPerBlock_Positive") { + size_t dynamicSmemSize = 0; + int numBlocks = 1; + int inputArray[SIZE], expectedOutput[SIZE], actualOutput[SIZE]; + + for (int i = 0; i < SIZE; i++) { + inputArray[i] = i; + expectedOutput[i] = SIZE - i - 1; // Expected reversed array + actualOutput[i] = 0; + } + + int *deviceArray{}; // Device pointer + HIP_CHECK( + hipMalloc(&deviceArray, SIZE * sizeof(int))); // Allocate device memory + + HIP_CHECK(hipMemcpy(deviceArray, inputArray, SIZE * sizeof(int), + hipMemcpyHostToDevice)); + + HIP_CHECK(hipOccupancyAvailableDynamicSMemPerBlock( + &dynamicSmemSize, dynamicReverse, numBlocks, SIZE)); + hipDeviceProp_t devProp; + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + INFO("Available Dynamic shared memory size : " + << dynamicSmemSize + << ", Dynamic shared memory calculated from device properties : " + << devProp.sharedMemPerBlock - SIZE * sizeof(int)); + REQUIRE(dynamicSmemSize == devProp.sharedMemPerBlock - SIZE * sizeof(int)); + dynamicReverse<<>>(deviceArray, SIZE); + + HIP_CHECK(hipMemcpy(actualOutput, deviceArray, SIZE * sizeof(int), + hipMemcpyDeviceToHost)); + // Verify results + for (int i = 0; i < SIZE; i++) { + INFO("Results mismatched : actualOutput[" << i << "]!=expectedOutput[" << i + << "],(" << actualOutput[i] << "," + << expectedOutput[i] << ")"); + REQUIRE(actualOutput[i] == expectedOutput[i]); + } +} + +/** + * End doxygen group occupancyTest. + * @} + */