diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index 9dc8bbb073..23d61f3ca6 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -13,6 +13,7 @@ "Unit_hipInit_Negative", "Unit_hipDeviceReset_Positive_Basic", "Unit_hipDeviceReset_Positive_Threaded", + "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Positive_Basic", "Unit_hipKernelNameRef_Negative_Parameters", diff --git a/catch/hipTestMain/config/config_amd_windows_common.json b/catch/hipTestMain/config/config_amd_windows_common.json index c6dcda0c92..c88dfdfe7c 100644 --- a/catch/hipTestMain/config/config_amd_windows_common.json +++ b/catch/hipTestMain/config/config_amd_windows_common.json @@ -95,6 +95,7 @@ "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread", "Note: intermittent Seg fault failure ", "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", + "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Positive_Basic", diff --git a/catch/unit/occupancy/CMakeLists.txt b/catch/unit/occupancy/CMakeLists.txt index eceb8626ca..2ad7eb5a6d 100644 --- a/catch/unit/occupancy/CMakeLists.txt +++ b/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 hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.cc ) diff --git a/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc b/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc index e6fdb3ba04..6199f45c38 100644 --- a/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cc +++ b/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 +/* +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" -static __global__ void f1(float *a) { *a = 1.0; } +static __global__ void f1(float* a) { *a = 1.0; } -template -static __global__ void f2(T *a) { *a = 1; } +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; +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)); + // 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_Positive_RangeValidation") { + hipDeviceProp_t devProp; + int blockSize = 0; + int gridSize = 0; + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); - HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0)); + SECTION("dynSharedMemPerBlk = 0") { + // Get potential blocksize + HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0)); - // Check if numBlocks and blockSize are within limits - REQUIRE(numBlock > 0); - REQUIRE((numBlock * blockSize) <= devProp.maxThreadsPerMultiProcessor); + 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)); - // 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); + 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 - (&numBlock, f2, blockSize, 0)); - REQUIRE(numBlock > 0); + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + + SECTION("dynSharedMemPerBlk = 0") { + // Get potential blocksize + HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f2, 0, 0)); + + MaxActiveBlocksPerMultiprocessor( + [blockSize](int* numBlocks) { + return hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f2, + blockSize, 0); + }, + blockSize, devProp.maxThreadsPerMultiProcessor); + } + + SECTION("dynSharedMemPerBlk = sharedMemPerBlock") { + // Get potential blocksize + HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f2, + devProp.sharedMemPerBlock, 0)); + + MaxActiveBlocksPerMultiprocessor( + [blockSize, devProp](int* numBlocks) { + return hipOccupancyMaxActiveBlocksPerMultiprocessor( + numBlocks, f2, blockSize, devProp.sharedMemPerBlock); + }, + blockSize, devProp.maxThreadsPerMultiProcessor); + } } - diff --git a/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor_old.cc b/catch/unit/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor_old.cc new file mode 100644 index 0000000000..e6fdb3ba04 --- /dev/null +++ b/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 + +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/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc b/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc index e31988a55b..0fc60cedfa 100644 --- a/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.cc +++ b/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize.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,65 +16,82 @@ 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 +/* +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" -static __global__ void f1(float *a) { *a = 1.0; } +static __global__ void f1(float* a) { *a = 1.0; } -template -static __global__ void f2(T *a) { *a = 1; } +template static __global__ void f2(T* a) { *a = 1; } -TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Negative") { - hipError_t ret; - int blockSize = 0; - int gridSize = 0; +TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Negative_Parameters") { + // Common negative tests + MaxPotentialBlockSizeNegative([](int* gridSize, int* blockSize) { + return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f1, 0, 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); +#if HT_AMD +#if 0 // EXSWHTEC-219 + SECTION("Kernel function is NULL") { + int blockSize = 0; + int gridSize = 0; + // nvcc doesnt support kernelfunc(NULL) for api + HIP_CHECK_ERROR(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, NULL, 0, 0), + hipErrorInvalidDeviceFunction); + } +#endif #endif } -TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_rangeValidation") { +TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Positive_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); + SECTION("dynSharedMemPerBlk = 0, blockSizeLimit = 0") { + MaxPotentialBlockSize( + [](int* gridSize, int* blockSize) { + return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f1, 0, 0); + }, + devProp.maxThreadsPerBlock); + } + SECTION("dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock") { + MaxPotentialBlockSize( + [devProp](int* gridSize, int* blockSize) { + return hipOccupancyMaxPotentialBlockSize( + gridSize, blockSize, f1, devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock); + }, + devProp.maxThreadsPerBlock); + } } -TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_templateInvocation") { - int gridSize = 0, blockSize = 0; +TEST_CASE("Unit_hipOccupancyMaxPotentialBlockSize_Positive_TemplateInvocation") { + hipDeviceProp_t devProp; - HIP_CHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, - &blockSize, f2, 0, 0)); - REQUIRE(gridSize > 0); - REQUIRE(blockSize > 0); + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + + SECTION("dynSharedMemPerBlk = 0, blockSizeLimit = 0") { + MaxPotentialBlockSize( + [](int* gridSize, int* blockSize) { + return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f2, 0, 0); + }, + devProp.maxThreadsPerBlock); + } + + SECTION("dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock") { + MaxPotentialBlockSize( + [devProp](int* gridSize, int* blockSize) { + return hipOccupancyMaxPotentialBlockSize( + gridSize, blockSize, f2, devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock); + }, + devProp.maxThreadsPerBlock); + } } - diff --git a/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize_old.cc b/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize_old.cc new file mode 100644 index 0000000000..e31988a55b --- /dev/null +++ b/catch/unit/occupancy/hipOccupancyMaxPotentialBlockSize_old.cc @@ -0,0 +1,80 @@ +/* +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 + +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); +} + diff --git a/catch/unit/occupancy/occupancy_common.hh b/catch/unit/occupancy/occupancy_common.hh new file mode 100644 index 0000000000..d03caad35b --- /dev/null +++ b/catch/unit/occupancy/occupancy_common.hh @@ -0,0 +1,72 @@ +/* +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 +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. +*/ +#pragma once + +#include + +template void MaxPotentialBlockSize(F func, int maxThreadsPerBlock) { + int gridSize = 0; + int blockSize = 0; + + // Get potential blocksize + HIP_CHECK(func(&gridSize, &blockSize)); + + // Check if blockSize doesn't exceed maxThreadsPerBlock + REQUIRE(gridSize > 0); + REQUIRE(blockSize > 0); + REQUIRE(blockSize <= maxThreadsPerBlock); + REQUIRE(gridSize * blockSize < static_cast(std::pow(2, 32))); +} + +template void MaxPotentialBlockSizeNegative(F func) { + int blockSize = 0; + int gridSize = 0; + + // Validate common arguments + SECTION("gridSize is nullptr") { + HIP_CHECK_ERROR(func(nullptr, &blockSize), hipErrorInvalidValue); + } + SECTION("blockSize is nullptr") { + HIP_CHECK_ERROR(func(&gridSize, nullptr), hipErrorInvalidValue); + } +} + +template +void MaxActiveBlocksPerMultiprocessor(F func, int blockSize, int maxThreadsPerMultiProcessor) { + int numBlocks = 0; + + // Validate maximum active block pre multiprocessor + HIP_CHECK(func(&numBlocks)); + + // Check if numBlocks and blockSize are within limits + REQUIRE(numBlocks > 0); + REQUIRE((numBlocks * blockSize) <= maxThreadsPerMultiProcessor); +} + +template void MaxActiveBlocksPerMultiprocessorNegative(F func, int blockSize) { + int numBlocks = 0; + + // Validate common arguments + SECTION("numBlocks is nullptr") { + HIP_CHECK_ERROR(func(nullptr, blockSize, 0), hipErrorInvalidValue); + } + SECTION("Block size is 0") { + HIP_CHECK_ERROR(func(&numBlocks, 0, 0), hipErrorInvalidValue); + } +}