From cf5ebc5990a298092ccac7d29bd753e5941d9e41 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 28 Dec 2023 18:33:28 +0100 Subject: [PATCH] EXSWHTEC-280 - Implement Unit Tests for launch bounds #209 Change-Id: I43af8c66dfcb4926bc29fe936a37d91fef6a0650 --- catch/include/hip_test_defgroups.hh | 58 ++++++ catch/unit/CMakeLists.txt | 1 + catch/unit/launchBounds/CMakeLists.txt | 47 +++++ catch/unit/launchBounds/launch_bounds.cc | 173 ++++++++++++++++++ .../launch_bounds_compiler_error_kernels.cc | 35 ++++ .../launch_bounds_negative_kernels_rtc.hh | 64 +++++++ .../launch_bounds_parse_error_kernels.cc | 30 +++ 7 files changed, 408 insertions(+) create mode 100644 catch/unit/launchBounds/CMakeLists.txt create mode 100644 catch/unit/launchBounds/launch_bounds.cc create mode 100644 catch/unit/launchBounds/launch_bounds_compiler_error_kernels.cc create mode 100644 catch/unit/launchBounds/launch_bounds_negative_kernels_rtc.hh create mode 100644 catch/unit/launchBounds/launch_bounds_parse_error_kernels.cc diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index e108f296fd..c1f72bdf6c 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -307,3 +307,61 @@ TEST_CASE("Unit_atomicDec_Negative_Parameters") {} * This section describes tests for the Complex type functions. * @} */ + +/** + * @defgroup DeviceLanguageTest Device Language + * @{ + * This section describes tests for the Device Language API. + */ + +/** + * @addtogroup launch_bounds launch_bounds + * @{ + * @ingroup DeviceLanguageTest + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# Compiles kernels that are not created appropriately: + * - Maximum number of threads is 0 + * - Maximum number of threads is not integer value + * - Mimimum number of warps is not integer value + * -# Expected output: compiler error + * Test source + * ------------------------ + * - unit/launch_bounds/CMakeLists.txt + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_CompilerError") {} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# Compiles kernels that are not created appropriately: + * - Maximum number of threads is negative + * - Mimimum number of warps is negative + * - Validates handling of invalid arguments: + * -# Expected output: parse error + * Test source + * ------------------------ + * - unit/launch_bounds/CMakeLists.txt + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_ParseError") {} + +/** + * End doxygen group launch_bounds. + * @} + */ + +/** + * End doxygen group DeviceLanguageTest. + * @} + */ diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 304016410e..9d158dad4d 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -58,3 +58,4 @@ add_subdirectory(vulkan_interop) add_subdirectory(gl_interop) # Disabled on NVIDIA due to defect - EXSWHTEC-246 endif() add_subdirectory(synchronization) +add_subdirectory(launchBounds) diff --git a/catch/unit/launchBounds/CMakeLists.txt b/catch/unit/launchBounds/CMakeLists.txt new file mode 100644 index 0000000000..589fe4d8e8 --- /dev/null +++ b/catch/unit/launchBounds/CMakeLists.txt @@ -0,0 +1,47 @@ +# 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 +# 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. + +set(TEST_SRC + launch_bounds.cc +) + +if(HIP_PLATFORM MATCHES "nvidia") + hip_add_exe_to_target(NAME LaunchBoundsTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS nvrtc) +elseif(HIP_PLATFORM MATCHES "amd") + hip_add_exe_to_target(NAME LaunchBoundsTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS hiprtc) +endif() + +add_test(NAME Unit_Kernel_Launch_bounds_Negative_Parameters_CompilerError + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + launch_bounds_compiler_error_kernels.cc -1) + +if(HIP_PLATFORM MATCHES "amd") + add_test(NAME Unit_Kernel_Launch_bounds_Negative_Parameters_ParseError + COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py + ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH} + launch_bounds_parse_error_kernels.cc -1) +endif() diff --git a/catch/unit/launchBounds/launch_bounds.cc b/catch/unit/launchBounds/launch_bounds.cc new file mode 100644 index 0000000000..72b087331b --- /dev/null +++ b/catch/unit/launchBounds/launch_bounds.cc @@ -0,0 +1,173 @@ +/* +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 +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 +#include "launch_bounds_negative_kernels_rtc.hh" + +/** + * @addtogroup launch_bounds launch_bounds + * @{ + * @ingroup DeviceLanguageTest + * `__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_WARPS_PER_EXECUTION_UNIT)` - + * allows the application to provide usage hints that influence the resources (primarily registers) + * used by the generated code. It is a function attribute that must be attached to a global + * function. + */ + +constexpr int kMaxThreadsPerBlock = 128; +constexpr int kMinWarpsPerMultiprocessor = 2; + +__launch_bounds__(kMaxThreadsPerBlock, kMinWarpsPerMultiprocessor) __global__ + void SumKernel(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); +} + +template void LaunchBoundsWrapper(const int threads_per_block) { + auto block_size = GENERATE(1, 32, 128); + int* A_d; + int* A_h; + int sum{0}; + + A_h = static_cast(malloc(sizeof(int))); + memset(A_h, 0, sizeof(int)); + HIP_CHECK(hipMalloc(&A_d, sizeof(int))); + HIP_CHECK(hipMemcpy(A_d, A_h, sizeof(int), hipMemcpyHostToDevice)); + SumKernel<<>>(A_d); + + if constexpr (out_of_bounds) { + if (threads_per_block < 0) { + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidConfiguration); + } else { +#if HT_AMD + HIP_CHECK_ERROR(hipGetLastError(), hipErrorLaunchFailure); +#else + HIP_CHECK_ERROR(hipGetLastError(), hipErrorInvalidValue); +#endif + } + } else { + HIP_CHECK(hipGetLastError()); + } + + HIP_CHECK(hipMemcpy(A_h, A_d, sizeof(int), hipMemcpyDeviceToHost)); + + if constexpr (!out_of_bounds) { + for (int i = 0; i < threads_per_block * block_size; ++i) { + sum += i; + } + REQUIRE(*A_h == sum); + } + + free(A_h); + HIP_CHECK(hipFree(A_d)); +} + +/** + * Test Description + * ------------------------ + * - Executes simple addition kernel and validates results. + * - The number of threads per block used to launch the kernel + * are complied with the `__launch_bounds__`: + * -# Number of threads per block are less than or equal to the configured maximum value. + * -# Different values are assigned and kernel functionality is validated. + * Test source + * ------------------------ + * - unit/launch_bounds/launch_bounds.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Kernel_Launch_bounds_Positive_Basic") { + auto threads_per_block = GENERATE(1, kMaxThreadsPerBlock / 2, kMaxThreadsPerBlock); + LaunchBoundsWrapper(threads_per_block); +} + +/** + * Test Description + * ------------------------ + * - Validates that the kernels will not be launched if the number of threads + * per block is larger than configured with `__launch_bounds__`: + * -# Expected output: + * - return `hipErrorLaunchFailure` on AMD. + * - return `hipErrorInvalidValue` on NVIDIA. + * Test source + * ------------------------ + * - unit/launch_bounds/launch_bounds.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Kernel_Launch_bounds_Negative_OutOfBounds") { + auto threads_per_block = + GENERATE(-1 * kMaxThreadsPerBlock, -1, kMaxThreadsPerBlock + 1, 2 * kMaxThreadsPerBlock); + LaunchBoundsWrapper(threads_per_block); +} + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# Compiles kernels that are not created appropriately: + * - Maximum number of threads is 0 + * - Maximum number of threads is negative + * - Minimum number of warps is negative + * - Maximum number of threads is not integer value + * - Mimimum number of warps is not integer value + * -# Expected output: compiler error + * - Uses RTC for compilation. + * Test source + * ------------------------ + * - unit/launch_bounds/launch_bounds.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_RTC") { + hiprtcProgram program{}; + +#if HT_AMD + const auto program_source = GENERATE(kMaxThreadsZero, kMaxThreadsNegative, kMinWarpsNegative, + kMaxThreadsNotInt, kMinWarpsNotInt); +#else + // Aligned with CUDA behavior and expected behavior on NVIDIA + const auto program_source = GENERATE(kMaxThreadsNotInt, kMinWarpsNotInt); +#endif + + HIPRTC_CHECK(hiprtcCreateProgram(&program, program_source, "launch_bounds_negative.cc", 0, + nullptr, nullptr)); + hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)}; + + // Get the compile log. + size_t log_size{}; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size)); + std::string log(log_size, ' '); + HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data())); + int error_count{0}; + std::string error_message{"error:"}; + + size_t n_pos = log.find(error_message, 0); + while (n_pos != std::string::npos) { + ++error_count; + n_pos = log.find(error_message, n_pos + 1); + } + + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + REQUIRE(error_count > 0); + HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION); +} diff --git a/catch/unit/launchBounds/launch_bounds_compiler_error_kernels.cc b/catch/unit/launchBounds/launch_bounds_compiler_error_kernels.cc new file mode 100644 index 0000000000..4a8ec9885f --- /dev/null +++ b/catch/unit/launchBounds/launch_bounds_compiler_error_kernels.cc @@ -0,0 +1,35 @@ +/* +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 +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 + +__launch_bounds__(0) __global__ void MaxThreadsZero(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); +} + +__launch_bounds__(1.5) __global__ void MaxThreadsNotInt(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); +} + +__launch_bounds__(128, 1.5) __global__ void MinWarpsNotInt(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); +} diff --git a/catch/unit/launchBounds/launch_bounds_negative_kernels_rtc.hh b/catch/unit/launchBounds/launch_bounds_negative_kernels_rtc.hh new file mode 100644 index 0000000000..a341ce67ff --- /dev/null +++ b/catch/unit/launchBounds/launch_bounds_negative_kernels_rtc.hh @@ -0,0 +1,64 @@ +/* +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 +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 + +/* +Negative kernels used for the launch bounds negative Test Cases that are using RTC. +*/ + +static constexpr auto kMaxThreadsZero{ + R"( + __launch_bounds__(0) __global__ void SumKernel(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); + } + )"}; + +static constexpr auto kMaxThreadsNegative{ + R"( + __launch_bounds__(-1) __global__ void SumKernel(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); + } + )"}; + +static constexpr auto kMinWarpsNegative{ + R"( + __launch_bounds__(128, -1) __global__ void SumKernel(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); + } + )"}; + +static constexpr auto kMaxThreadsNotInt{ + R"( + __launch_bounds__(1.5) __global__ void SumKernel(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); + } + )"}; + +static constexpr auto kMinWarpsNotInt{ + R"( + __launch_bounds__(128, 1.5) __global__ void SumKernel(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); + } + )"}; diff --git a/catch/unit/launchBounds/launch_bounds_parse_error_kernels.cc b/catch/unit/launchBounds/launch_bounds_parse_error_kernels.cc new file mode 100644 index 0000000000..e0ed6093f5 --- /dev/null +++ b/catch/unit/launchBounds/launch_bounds_parse_error_kernels.cc @@ -0,0 +1,30 @@ +/* +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 +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 + +__launch_bounds__(-1) __global__ void MaxThreadsNegative(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); +} + +__launch_bounds__(128, -1) __global__ void MinWarpsNegative(int* sum) { + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + atomicAdd(sum, tid); +}