diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 6de829e0f4..973b2a9743 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -54,6 +54,8 @@ "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported", "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", + "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters", + "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters", "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphExecMemcpyNodeSetParamsFromSymbol_Positive_Basic", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 0938722c5e..80e0caec46 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -119,6 +119,8 @@ "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported", "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", + "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters", + "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters", "Unit_hipGraphMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphExecMemcpyNodeSetParamsToSymbol_Positive_Basic", "Unit_hipGraphMemcpyNodeSetParamsFromSymbol_Positive_Basic", diff --git a/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt b/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt index 2ad7eb5a6d..ba3d64c981 100644 --- a/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/occupancy/CMakeLists.txt @@ -4,9 +4,21 @@ set(TEST_SRC hipOccupancyMaxActiveBlocksPerMultiprocessor_old.cc hipOccupancyMaxPotentialBlockSize.cc hipOccupancyMaxPotentialBlockSize_old.cc + hipModuleOccupancyMaxPotentialBlockSize.cc + hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc + hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc + hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags.cc ) +add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/simple_kernel.code + COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/simple_kernel.cc -o simple_kernel.code + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/simple_kernel.cc) + +add_custom_target(simple_kernel ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/simple_kernel.code) + hip_add_exe_to_target(NAME OccupancyTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) + +add_dependencies(OccupancyTest simple_kernel) diff --git a/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc new file mode 100644 index 0000000000..65e1b42ab3 --- /dev/null +++ b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessor.cc @@ -0,0 +1,92 @@ +/* +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. +*/ +/* +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" + +TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters") { + hipModule_t module; + hipFunction_t function; + int blockSize = 0; + int gridSize = 0; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + // Get potential blocksize + HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0)); + + // Common negative tests + MaxActiveBlocksPerMultiprocessorNegative( + [&function](int* numBlocks, int blockSize, size_t dynSharedMemPerBlk) { + return hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, function, blockSize, + dynSharedMemPerBlk); + }, + blockSize); + + HIP_CHECK(hipModuleUnload(module)); +} + +TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessor_Positive_RangeValidation") { + hipDeviceProp_t devProp; + hipModule_t module; + hipFunction_t function; + int blockSize = 0; + int gridSize = 0; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + + SECTION("dynSharedMemPerBlk = 0") { + // Get potential blocksize + HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0)); + + MaxActiveBlocksPerMultiprocessor( + [blockSize, &function](int* numBlocks) { + return hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, function, blockSize, + 0); + }, + blockSize, devProp.maxThreadsPerMultiProcessor); + } + SECTION("dynSharedMemPerBlk = sharedMemPerBlock") { + // Get potential blocksize + HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, + devProp.sharedMemPerBlock, 0)); + + MaxActiveBlocksPerMultiprocessor( + [blockSize, devProp, &function](int* numBlocks) { + return hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, function, blockSize, + devProp.sharedMemPerBlock); + }, + blockSize, devProp.maxThreadsPerMultiProcessor); + } + + HIP_CHECK(hipModuleUnload(module)); +} diff --git a/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc new file mode 100644 index 0000000000..8df78e4481 --- /dev/null +++ b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags.cc @@ -0,0 +1,103 @@ +/* +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. +*/ +/* +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" + +TEST_CASE("Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters") { + hipModule_t module; + hipFunction_t function; + int numBlocks = 0; + int blockSize = 0; + int gridSize = 0; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + // Get potential blocksize + HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0)); + + // Common negative tests + MaxActiveBlocksPerMultiprocessorNegative( + [&function](int* numBlocks, int blockSize, size_t dynSharedMemPerBlk) { + return hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + numBlocks, function, blockSize, dynSharedMemPerBlk, hipOccupancyDefault); + }, + blockSize); + + SECTION("Flag is invalid") { + // Only default flag is supported + HIP_CHECK_ERROR(hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + &numBlocks, function, blockSize, 0, 2), + hipErrorInvalidValue); + } + + HIP_CHECK(hipModuleUnload(module)); +} + +TEST_CASE( + "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Positive_RangeValidation") { + hipDeviceProp_t devProp; + hipModule_t module; + hipFunction_t function; + int blockSize = 0; + int gridSize = 0; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + + SECTION("dynSharedMemPerBlk = 0") { + // Get potential blocksize + HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, 0, 0)); + + MaxActiveBlocksPerMultiprocessor( + [blockSize, &function](int* numBlocks) { + return hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + numBlocks, function, blockSize, 0, hipOccupancyDefault); + }, + blockSize, devProp.maxThreadsPerMultiProcessor); + } + SECTION("dynSharedMemPerBlk = sharedMemPerBlock") { + // Get potential blocksize + HIP_CHECK(hipModuleOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, function, + devProp.sharedMemPerBlock, 0)); + + MaxActiveBlocksPerMultiprocessor( + [blockSize, devProp, &function](int* numBlocks) { + return hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + numBlocks, function, blockSize, devProp.sharedMemPerBlock, hipOccupancyDefault); + }, + blockSize, devProp.maxThreadsPerMultiProcessor); + } + + HIP_CHECK(hipModuleUnload(module)); +} diff --git a/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc new file mode 100644 index 0000000000..6f49b9efe8 --- /dev/null +++ b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSize.cc @@ -0,0 +1,75 @@ +/* +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. +*/ +/* +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" + +TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSize_Negative_Parameters") { + hipModule_t module; + hipFunction_t function; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + // Common negative tests + MaxPotentialBlockSizeNegative([&function](int* gridSize, int* blockSize) { + return hipModuleOccupancyMaxPotentialBlockSize(gridSize, blockSize, function, 0, 0); + }); + + HIP_CHECK(hipModuleUnload(module)); +} + +TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSize_Positive_RangeValidation") { + hipDeviceProp_t devProp; + hipModule_t module; + hipFunction_t function; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + + SECTION("dynSharedMemPerBlk = 0, blockSizeLimit = 0") { + MaxPotentialBlockSize( + [&function](int* gridSize, int* blockSize) { + return hipModuleOccupancyMaxPotentialBlockSize(gridSize, blockSize, function, 0, 0); + }, + devProp.maxThreadsPerBlock); + } + + SECTION("dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock") { + MaxPotentialBlockSize( + [&function, devProp](int* gridSize, int* blockSize) { + return hipModuleOccupancyMaxPotentialBlockSize( + gridSize, blockSize, function, devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock); + }, + devProp.maxThreadsPerBlock); + } + + HIP_CHECK(hipModuleUnload(module)); +} diff --git a/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc new file mode 100644 index 0000000000..50107b465f --- /dev/null +++ b/projects/hip-tests/catch/unit/occupancy/hipModuleOccupancyMaxPotentialBlockSizeWithFlags.cc @@ -0,0 +1,87 @@ +/* +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. +*/ +/* +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" + +TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters") { + hipModule_t module; + hipFunction_t function; + int blockSize = 0; + int gridSize = 0; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + // Common negative tests + MaxPotentialBlockSizeNegative([&function](int* gridSize, int* blockSize) { + return hipModuleOccupancyMaxPotentialBlockSizeWithFlags(gridSize, blockSize, function, 0, 0, + hipOccupancyDefault); + }); + + SECTION("Flag is invalid") { + // Only default flag is supported + HIP_CHECK_ERROR( + hipModuleOccupancyMaxPotentialBlockSizeWithFlags(&gridSize, &blockSize, function, 0, 0, 2), + hipErrorInvalidValue); + } + + HIP_CHECK(hipModuleUnload(module)); +} + +TEST_CASE("Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Positive_RangeValidation") { + hipDeviceProp_t devProp; + hipModule_t module; + hipFunction_t function; + + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK(hipModuleLoad(&module, "simple_kernel.code")); + HIPCHECK(hipModuleGetFunction(&function, module, "SimpleKernel")); + + HIP_CHECK(hipGetDeviceProperties(&devProp, 0)); + + SECTION("dynSharedMemPerBlk = 0, blockSizeLimit = 0") { + MaxPotentialBlockSize( + [&function](int* gridSize, int* blockSize) { + return hipModuleOccupancyMaxPotentialBlockSizeWithFlags(gridSize, blockSize, function, 0, + 0, hipOccupancyDefault); + }, + devProp.maxThreadsPerBlock); + } + + SECTION("dynSharedMemPerBlk = sharedMemPerBlock, blockSizeLimit = maxThreadsPerBlock") { + MaxPotentialBlockSize( + [&function, devProp](int* gridSize, int* blockSize) { + return hipModuleOccupancyMaxPotentialBlockSizeWithFlags( + gridSize, blockSize, function, devProp.sharedMemPerBlock, devProp.maxThreadsPerBlock, + hipOccupancyDefault); + }, + devProp.maxThreadsPerBlock); + } + + HIP_CHECK(hipModuleUnload(module)); +} diff --git a/projects/hip-tests/catch/unit/occupancy/occupancy_common.hh b/projects/hip-tests/catch/unit/occupancy/occupancy_common.hh index d03caad35b..5d71c7dac0 100644 --- a/projects/hip-tests/catch/unit/occupancy/occupancy_common.hh +++ b/projects/hip-tests/catch/unit/occupancy/occupancy_common.hh @@ -66,7 +66,5 @@ template void MaxActiveBlocksPerMultiprocessorNegative(F func, int 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); - } + SECTION("Block size is 0") { HIP_CHECK_ERROR(func(&numBlocks, 0, 0), hipErrorInvalidValue); } } diff --git a/projects/hip-tests/catch/unit/occupancy/simple_kernel.cc b/projects/hip-tests/catch/unit/occupancy/simple_kernel.cc new file mode 100644 index 0000000000..e1bf415e95 --- /dev/null +++ b/projects/hip-tests/catch/unit/occupancy/simple_kernel.cc @@ -0,0 +1,25 @@ +/* +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. +*/ + +#include "hip/hip_runtime.h" + +extern "C" __global__ void SimpleKernel(int* a, int* b) { + int tx = threadIdx.x; + b[tx] = a[tx]; +}