diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index 680dfa8a0..debfbbde7 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -116,6 +116,13 @@ THE SOFTWARE. * @} */ +/** + * @defgroup OccupancyTest Occupancy + * @{ + * This section describes tests for the occupancy functions of HIP runtime API. + * @} + */ + /** * @defgroup PeerToPeerTest PeerToPeer Device Memory Access * @{ diff --git a/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc b/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc index 65e1b42ab..9f58ed7e6 100644 --- a/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc +++ b/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 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,15 +16,33 @@ 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. */ -/* -Testcase Scenarios : -Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation - Test correct -execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessor for diffrent parameter values -Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters - Test unsuccessful -execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessor api when parameters are invalid -*/ + #include "occupancy_common.hh" +/** + * @addtogroup hipModuleOccupancyMaxActiveBlocksPerMultiprocessor hipModuleOccupancyMaxActiveBlocksPerMultiprocessor + * @{ + * @ingroup OccupancyTest + * `hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, hipFunction_t f, + * int blockSize, size_t dynSharedMemPerBlk)` - + * Returns occupancy for a device function. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the grid size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When block size is 0 + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters") { hipModule_t module; hipFunction_t function; @@ -50,6 +68,21 @@ TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Para HIP_CHECK(hipModuleUnload(module)); } +/** + * Test Description + * ------------------------ + * - Check if grid size and block size are within valid range using basic kernel functions: + * -# When `dynSharedMemPerBlk = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation") { hipDeviceProp_t devProp; hipModule_t module; diff --git a/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc b/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc index 8df78e448..7b60be2a4 100644 --- a/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc +++ b/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 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,17 +16,36 @@ 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. */ -/* -Testcase Scenarios : -Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Positive_RangeValidation - Test -correct execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags for diffrent -parameter values -Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters - Test -unsuccessful execution of hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags api when -parameters are invalid -*/ + #include "occupancy_common.hh" +/** + * @addtogroup hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags + * @{ + * @ingroup OccupancyTest + * `hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, hipFunction_t f, + * int blockSize, size_t dynSharedMemPerBlk, unsigned int flags)` - + * Returns occupancy for a device function. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the grid size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When block size is 0 + * - Expected output: return `hipErrorInvalidValue` + * -# When flag is invalid, because only default flag is supported + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters") { hipModule_t module; hipFunction_t function; @@ -60,6 +79,21 @@ TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Nega HIP_CHECK(hipModuleUnload(module)); } +/** + * Test Description + * ------------------------ + * - Check if grid size and block size are within valid range using basic kernel functions: + * -# When `dynSharedMemPerBlk = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE( "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Positive_RangeValidation") { hipDeviceProp_t devProp; diff --git a/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc b/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc index 6f49b9efe..a9056d7a0 100644 --- a/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc +++ b/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 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,15 +16,33 @@ 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. */ -/* -Testcase Scenarios : -Unit_hipModuleOccupancyMaxPotentialBlockSize_Positive_RangeValidation - Test correct execution of -hipModuleOccupancyMaxPotentialBlockSize for diffrent parameter values -Unit_hipModuleOccupancyMaxPotentialBlockSize_Negative_Parameters - Test unsuccessful execution of -hipModuleOccupancyMaxPotentialBlockSize api when parameters are invalid -*/ + #include "occupancy_common.hh" +/** + * @addtogroup hipModuleOccupancyMaxPotentialBlockSize hipModuleOccupancyMaxPotentialBlockSize + * @{ + * @ingroup OccupancyTest + * `hipModuleOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, + * hipFunction_t f, size_t dynSharedMemPerBlk, int blockSizeLimit)` - + * Determine the grid and block sizes to achieve maximum occupancy for a kernel. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the grid size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the block size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSize_Negative_Parameters") { hipModule_t module; hipFunction_t function; @@ -42,6 +60,21 @@ TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSize_Negative_Parameters") { HIP_CHECK(hipModuleUnload(module)); } +/** + * Test Description + * ------------------------ + * - Check if grid size and block size are within valid range using basic kernel functions: + * -# When `dynSharedMemPerBlk = 0, blockSizeLimit = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSize_Positive_RangeValidation") { hipDeviceProp_t devProp; hipModule_t module; diff --git a/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc b/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc index 50107b465..96b6d8e8d 100644 --- a/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc +++ b/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 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,15 +16,36 @@ 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. */ -/* -Testcase Scenarios : -Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Positive_RangeValidation - Test correct -execution of hipModuleOccupancyMaxPotentialBlockSizeWithFlags for diffrent parameter values -Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters - Test unsuccessful -execution of hipModuleOccupancyMaxPotentialBlockSizeWithFlags api when parameters are invalid -*/ + #include "occupancy_common.hh" +/** + * @addtogroup hipModuleOccupancyMaxPotentialBlockSizeWithFlags hipModuleOccupancyMaxPotentialBlockSizeWithFlags + * @{ + * @ingroup OccupancyTest + * `hipModuleOccupancyMaxPotentialBlockSizeWithFlags(int* gridSize, int* blockSize, + * hipFunction_t f, size_t dynSharedMemPerBlk, int blockSizeLimit, unsigned int flags)` - + * Determine the grid and block sizes to achieves maximum occupancy for a kernel. + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the grid size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the block size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When flag is invalid, because only default flag is supported + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters") { hipModule_t module; hipFunction_t function; @@ -52,6 +73,21 @@ TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parame HIP_CHECK(hipModuleUnload(module)); } +/** + * Test Description + * ------------------------ + * - Check if grid size and block size are within valid range using basic kernel functions: + * -# When `dynSharedMemPerBlk = 0, blockSizeLimit = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Positive_RangeValidation") { hipDeviceProp_t devProp; hipModule_t module; diff --git a/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc b/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc index 6199f45c3..d35d259a1 100644 --- a/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc +++ b/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc @@ -16,21 +16,41 @@ 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. */ -/* -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" +/** + * @addtogroup hipOccupancyMaxActiveBlocksPerMultiprocessor + * hipOccupancyMaxActiveBlocksPerMultiprocessor + * @{ + * @ingroup OccupancyTest + * `hipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, const void* f, + * int blockSize, size_t dynSharedMemPerBlk)` - + * Returns occupancy for a device function. + */ + static __global__ void f1(float* a) { *a = 1.0; } template static __global__ void f2(T* a) { *a = 1; } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the grid size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When block size is 0 + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the function is `nullptr` + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidDeviceFunction` + * Test source + * ------------------------ + * - unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters") { int numBlocks = 0; int blockSize = 0; @@ -53,6 +73,21 @@ TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters } } +/** + * Test Description + * ------------------------ + * - Check if grid size and block size are within valid range using basic kernel functions: + * -# When `dynSharedMemPerBlk = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation") { hipDeviceProp_t devProp; int blockSize = 0; @@ -84,6 +119,21 @@ TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValid } } +/** + * Test Description + * ------------------------ + * - Check is number of blocks is greater than 0 when API is invoked with a template: + * -# When `dynSharedMemPerBlk = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_TemplateInvocation") { hipDeviceProp_t devProp; int blockSize = 0; @@ -116,3 +166,19 @@ TEST_CASE("Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Positive_TemplateIn blockSize, devProp.maxThreadsPerMultiProcessor); } } + +/** + * End doxygen group hipOccupancyMaxActiveBlocksPerMultiprocessor. + * @} + */ + +/** + * @addtogroup hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags + * hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags + * @{ + * @ingroup OccupancyTest + * `hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const void* f, + * int blockSize, size_t dynSharedMemPerBlk, unsigned int flags __dparm(hipOccupancyDefault))` - + * Returns occupancy for a device function. + * @warning Flags ignored currently, skipped tests implementation. + */ diff --git a/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc b/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc index 0fc60cedf..38abafe5d 100644 --- a/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc +++ b/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc @@ -16,21 +16,40 @@ 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. */ -/* -Testcase Scenarios : -Unit_hipOccupancyMaxPotentialBlockSize_Positive_RangeValidation - Test correct execution of -hipOccupancyMaxPotentialBlockSize for diffrent parameter values -Unit_hipOccupancyMaxPotentialBlockSize_Positive_TemplateInvocation - Test correct execution of -hipOccupancyMaxPotentialBlockSize template for diffrent parameter values -Unit_hipOccupancyMaxPotentialBlockSize_Negative_Parameters - Test unsuccessful execution of -hipOccupancyMaxPotentialBlockSize api when parameters are invalid -*/ + #include "occupancy_common.hh" +/** + * @addtogroup hipOccupancyMaxPotentialBlockSize hipOccupancyMaxPotentialBlockSize + * @{ + * @ingroup OccupancyTest + * `hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, + * const void* f, size_t dynSharedMemPerBlk, int blockSizeLimit)` - + * Determine the grid and block sizes to achieves maximum occupancy for a kernel. + */ + static __global__ void f1(float* a) { *a = 1.0; } template static __global__ void f2(T* a) { *a = 1; } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the grid size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When output pointer to the block size is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the function is `nullptr` + * - Platform specific (NVIDIA) + * - Expected output: return `hipErrorInvalidDeviceFunction` + * Test source + * ------------------------ + * - unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Negative_Parameters") { // Common negative tests MaxPotentialBlockSizeNegative([](int* gridSize, int* blockSize) { @@ -50,6 +69,21 @@ TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Negative_Parameters") { #endif } +/** + * Test Description + * ------------------------ + * - Check if grid size and block size are within valid range using basic kernel functions: + * -# When `dynSharedMemPerBlk = 0, blockSizeLimit = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Positive_RangeValidation") { hipDeviceProp_t devProp; @@ -73,6 +107,21 @@ TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Positive_RangeValidation") { } } +/** + * Test Description + * ------------------------ + * - Check is number of blocks is greater than 0 when API is invoked with a template: + * -# When `dynSharedMemPerBlk = 0, blockSizeLimit = 0` + * - Expected output: return `hipSuccess` + * -# When `dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock` + * - Expected output: return `hipSuccess` + * Test source + * ------------------------ + * - unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Positive_TemplateInvocation") { hipDeviceProp_t devProp;