From 2bee071530f9ef0677e9a04de4f46b9b07030d5e Mon Sep 17 00:00:00 2001 From: sumanthtg Date: Wed, 2 Sep 2020 21:00:40 +0530 Subject: [PATCH] [dtest] Additional tests for hipOccupancyMax APIs Additional tests for following APIs - hipOccupancyMaxActiveBlocksPerMultiprocessor - hipOccupancyMaxPotentialBlockSize SWDEV-238517 for enhancing hip unit tests Change-Id: Ib9441c1366f46a082e10eb1a572bc7d8ebe1ee37 --- ...upancyMaxActiveBlocksPerMultiprocessor.cpp | 189 +++++++++++++++--- .../hipOccupancyMaxPotentialBlockSize.cpp | 165 +++++++++++++-- 2 files changed, 301 insertions(+), 53 deletions(-) diff --git a/tests/src/runtimeApi/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp b/tests/src/runtimeApi/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp index ad37dac8bc..7fb7c94a3d 100644 --- a/tests/src/runtimeApi/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp +++ b/tests/src/runtimeApi/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp @@ -1,33 +1,51 @@ /* -Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. - +Copyright (c) 2020-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 +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY 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 +FITNNESS 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 +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -// Test the Grid_Launch syntax. + +/** +Testcase Scenarios : + + (TestCase 1):: + 1) Check api behavior by passing numBlocks ptr as NULL. + 2) Pass invalid kernel function/NULL and check the api behavior. + 3) Pass blockSize as zero and check appropriate error-code is returned. + 4) Pass shm as max size_t and validate error-code returned. + + (TestCase 2):: + 5) Validate range by making sure (numBlock * blockSize) doesn't exceed + devProp.maxThreadsPerMultiProcessor. + 6) Check range of out param after passing valid dynSharedMemPerBlk. + + (TestCase 3):: + 7) Test case for using kernel function pointer with template. + +*/ /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM rocclr - * TEST: %t + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST: %t --tests 1 + * TEST: %t --tests 2 + * TEST: %t --tests 3 * HIT_END */ -#include "hip/hip_runtime.h" +#include +#include #include "test_common.h" __global__ void f1(float *a) { *a = 1.0; } @@ -36,29 +54,136 @@ template __global__ void f2(T *a) { *a = 1; } +/** + * Performs argument validation + */ +bool argValidation() { + bool TestPassed = true; + hipError_t ret; + int numBlock = 0, blockSize = 0; + int gridSize = 0; + + // Get potential blocksize + HIPCHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0)); + + // Validate each argument + if ((ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(NULL, f1, + blockSize, 0)) != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for" + " numBlock(NULL). Error: '%s'(%d)\n", hipGetErrorString(ret), ret); + TestPassed &= false; + } + + ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, NULL, + blockSize, 0); + if (ret != hipErrorInvalidValue && ret != hipErrorInvalidDeviceFunction) { + printf("ArgValidation : Inappropritate error value returned for" + " kernelfunc(NULL). numBlk %d, Error: '%s'(%d)\n", numBlock, + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, + f1, 0, 0)) != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for" + " blksize(0), shm(0). numBlk %d, Error: '%s'(%d)\n", numBlock, + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, + f1, 0, std::numeric_limits::max())) + != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for" + " blksize(0), shm(max). numBlk %d, Error: '%s'(%d)\n", numBlock, + hipGetErrorString(ret), ret); + TestPassed &= false; + } + + return TestPassed; +} + + +/** + * Performs range validation on api output + */ +bool rangeValidation() { + hipDeviceProp_t devProp; + bool TestPassed = true; + int numBlock = 0, blockSize = 0; + int gridSize = 0; + + // Get potential blocksize + HIPCHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0)); + + HIPCHECK(hipGetDeviceProperties(&devProp, 0)); + + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, + blockSize, 0)); + + // Check if numBlocks and blockSize are within limits + if ((numBlock <= 0) || + ((numBlock * blockSize) > devProp.maxThreadsPerMultiProcessor)) { + printf("RangeValidation : numBlock %d returned not in range." + "numblk(%d),blocksize(%d) and maxThrdsMP %d", numBlock, numBlock, + blockSize, devProp.maxThreadsPerMultiProcessor); + TestPassed &= false; + } + + // Validate numBlock after passing dynSharedMemPerBlk + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, + blockSize, devProp.sharedMemPerBlock)); + + // Check if numBlocks and blockSize are within limits + if ((numBlock <= 0) || + ((numBlock * blockSize) > devProp.maxThreadsPerMultiProcessor)) { + printf("RangeValidation : numBlock %d returned not in range." + "numblk(%d),blocksize(%d),shm and maxThrdsMP %d", numBlock, numBlock, + blockSize, devProp.maxThreadsPerMultiProcessor); + TestPassed &= false; + } + + return TestPassed; +} + +/** + * Test case for using kernel function pointer with template + */ +bool templateInvocation() { + bool TestPassed = true; + int blockSize = 32; + int numBlock = 0; + + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor + (&numBlock, f2, blockSize, 0)); + if (!numBlock) { + printf("TemplateInvocation : numBlock received as zero"); + TestPassed &= false; + } + + return TestPassed; +} + int main(int argc, char* argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + bool TestPassed = true; - // test case for using kernel function pointer - int gridSize = 0; - int blockSize = 0; - hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0); - assert(gridSize != 0 && blockSize != 0); - - int numBlock = 0; - hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0); - assert(numBlock != 0); - - - // test case for using kernel function pointer with template - gridSize = 0; - blockSize = 0; - hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f2, 0, 0); - assert(gridSize != 0 && blockSize != 0); - - numBlock = 0; - hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f2, blockSize, 0); - assert(numBlock != 0); + if (p_tests == 1) { + TestPassed = argValidation(); + } else if (p_tests == 2) { + TestPassed = rangeValidation(); + } else if (p_tests == 3) { + TestPassed = templateInvocation(); + } else { + printf("Didnt receive any valid option. Try options 1 to 3\n"); + TestPassed = false; + } + if (TestPassed) { passed(); + } else { + failed("hipOccupancyMaxActiveBlocksPerMultiprocessor validation Failed!"); + } } + diff --git a/tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp b/tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp index fc8538df26..39323fa930 100644 --- a/tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp +++ b/tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp @@ -1,33 +1,50 @@ /* -Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. - +Copyright (c) 2020-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 +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY 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 +FITNNESS 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 +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -// Test the Grid_Launch syntax. + +/** +Testcase Scenarios : + + (TestCase 1):: + 1) Pass gridSize as NULL and check appropriate error-code is returned. + 2) Pass blockSize as NULL and check appropriate error-code is returned. + 3) Pass invalid kernel function/NULL and check the api behavior. + + (TestCase 2):: + 4) Validate range by making sure blockSize returned by api doesn't exceed + devProp.maxThreadsPerBlock. + 5) Pass dynSharedMemPerBlk, blockSizeLimit and check out parameter range. + + (TestCase 3):: + 6) Test case for using kernel function pointer with template. + +*/ /* HIT_START * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * TEST: %t --tests 1 + * TEST: %t --tests 2 + * TEST: %t --tests 3 * HIT_END */ -#include "hip/hip_runtime.h" +#include +#include #include "test_common.h" __global__ void f1(float *a) { *a = 1.0; } @@ -36,19 +53,125 @@ template __global__ void f2(T *a) { *a = 1; } +/** + * Performs argument validation + */ +bool argValidation() { + bool TestPassed = true; + hipError_t ret; + int blockSize = 0; + int gridSize = 0; + + // Validate each argument + if ((ret = hipOccupancyMaxPotentialBlockSize(NULL, &blockSize, + f1, 0, 0)) != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for" + " gridSize(NULL). blksize rcvd %d, Error: '%s'(%d)\n", + blockSize, hipGetErrorString(ret), ret); + TestPassed &= false; + } + + if ((ret = hipOccupancyMaxPotentialBlockSize(&gridSize, NULL, + f1, 0, 0)) != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for" + " blockSize(NULL). gridSize rcvd %d, Error: '%s'(%d)\n", + gridSize, hipGetErrorString(ret), ret); + TestPassed &= false; + } + +#ifndef __HIP_PLATFORM_NVCC__ + // nvcc doesnt support kernelfunc(NULL) for api + ret = hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, NULL, 0, 0); + if (ret != hipErrorInvalidValue) { + printf("ArgValidation : Inappropritate error value returned for" + " kernelfunc(NULL). gridSize %d, blkSize %d, Error: '%s'(%d)\n", + gridSize, blockSize, hipGetErrorString(ret), ret); + TestPassed &= false; + } +#endif + + return TestPassed; +} + + +/** + * Performs range validation on api output + */ +bool rangeValidation() { + hipDeviceProp_t devProp; + bool TestPassed = true; + int blockSize = 0; + int gridSize = 0; + + // Get potential blocksize + HIPCHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0)); + + HIPCHECK(hipGetDeviceProperties(&devProp, 0)); + + // Check if blockSize doen't exceed maxThreadsPerBlock + if ((gridSize <= 0) || (blockSize <= 0) || + (blockSize > devProp.maxThreadsPerBlock)) { + printf("RangeValidation : grdSize %d/blkSize %d returned not in range(%d)", + gridSize, blockSize, devProp.maxThreadsPerBlock); + TestPassed &= false; + } + + // Pass dynSharedMemPerBlk, blockSizeLimit and check out param + blockSize = 0; + gridSize = 0; + + HIPCHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, + devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock)); + + if ((gridSize <= 0) || (blockSize <= 0) || + (blockSize > devProp.maxThreadsPerBlock)) { + printf("RangeValidation(Shm,TPB) : grdSize %d/blkSize %d returned" + "not in range(%d)", gridSize, blockSize, devProp.maxThreadsPerBlock); + TestPassed &= false; + } + + + return TestPassed; +} + +/** + * Test case for using kernel function pointer with template + */ +bool templateInvocation() { + bool TestPassed = true; + int gridSize = 0, blockSize = 0; + int numBlock = 0; + + HIPCHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, + &blockSize, f2, 0, 0)); + if (!gridSize || !blockSize) { + printf("TemplateInvocation : gridSize/blockSize received as zero"); + TestPassed &= false; + } + + return TestPassed; +} + int main(int argc, char* argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + bool TestPassed = true; - // test case for using kernel function pointer - int gridSize = 0; - int blockSize = 0; - hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0); - assert(gridSize != 0 && blockSize != 0); + if (p_tests == 1) { + TestPassed = argValidation(); + } else if (p_tests == 2) { + TestPassed = rangeValidation(); + } else if (p_tests == 3) { + TestPassed = templateInvocation(); + } else { + printf("Didnt receive any valid option. Try options 1 to 3\n"); + TestPassed = false; + } - // test case for using kernel function pointer with template - gridSize = 0; - blockSize = 0; - hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f2, 0, 0); - assert(gridSize != 0 && blockSize != 0); + if (TestPassed) { passed(); + } else { + failed("hipOccupancyMaxPotentialBlockSize validation Failed!"); + } } +