From bb08c70c4eba1c7773e4c889a4f8ad9f7ad3cc09 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 14:47:48 +0100 Subject: [PATCH] EXSWHTEC-350 - Implement tests for hipModuleLaunchCooperativeKernel #429 Change-Id: I42342c7d4cceed67990f603cd83473fe76e7f856 [ROCm/hip-tests commit: 303836a7fda08afe8e57f01212dd8f5afc12002b] --- .../catch/include/hip_test_defgroups.hh | 7 + .../catch/unit/module/CMakeLists.txt | 2 + .../hipModuleLaunchCooperativeKernel.cc | 211 ++++++++++++++++ ...oduleLaunchCooperativeKernelMultiDevice.cc | 227 ++++++++++++++++++ .../catch/unit/module/launch_kernel_module.cc | 6 + 5 files changed, 453 insertions(+) create mode 100644 projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc create mode 100644 projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 0a56d94239..8191c8b96f 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -179,6 +179,13 @@ THE SOFTWARE. * @} */ +/** + * @defgroup ModuleTest Module Management + * @{ + * This section describes the module management types & functions of HIP runtime API. + * @} + */ + /** * @defgroup TextureTest Texture Management * @{ diff --git a/projects/hip-tests/catch/unit/module/CMakeLists.txt b/projects/hip-tests/catch/unit/module/CMakeLists.txt index 5d951cf27c..76ca9e9ec6 100644 --- a/projects/hip-tests/catch/unit/module/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/module/CMakeLists.txt @@ -29,6 +29,8 @@ set(TEST_SRC hipModuleLaunchKernel.cc hipModuleGetGlobal.cc hipModuleGetTexRef.cc + hipModuleLaunchCooperativeKernel.cc + hipModuleLaunchCooperativeKernelMultiDevice.cc ) add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code diff --git a/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc b/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc new file mode 100644 index 0000000000..0ca6a31293 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernel.cc @@ -0,0 +1,211 @@ +/* +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. +*/ + +/** + * @addtogroup hipModuleLaunchCooperativeKernel hipModuleLaunchCooperativeKernel + * @{ + * @ingroup ModuleTest + * `hipModuleLaunchCooperativeKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, + * unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, + * unsigned int sharedMemBytes, hipStream_t stream, void ** kernelParams)` - + * Launches kernel f with launch parameters and shared memory on stream with arguments passed to + * kernelParams, where thread blocks can cooperate and synchronize as they execute. + */ + +#include +#include +#include + +#include "hip_module_launch_kernel_common.hh" + +/** + * Test Description + * ------------------------ + * - Tests `hipModuleLaunchCooperativeKernel` for a cooperative kernel with no parameters, and for + * a normal kernel with parameters. + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchCooperativeKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ +TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Positive_Basic") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + SECTION("Cooperative kernel with no arguments") { + hipFunction_t f = GetKernel(mg.module(), "CoopKernel"); + HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 2, 2, 1, 1, 1, 1, 0, nullptr, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("Kernel with arguments using kernelParams") { + hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + + LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); + + int* result_ptr = result_dev.ptr(); + void* kernel_args[1] = {&result_ptr}; + HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_args)); + + int result = 0; + HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 42); + } +} + +/** + * Test Description + * ------------------------ + * - Positive parameters test for `hipModuleLaunchCooperativeKernel`. + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchCooperativeKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ +TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Positive_Parameters") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + + SECTION("blockDim.x == maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX); + HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, x, 1, 1, 0, nullptr, nullptr)); + } + + SECTION("blockDim.y == maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY); + HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, y, 1, 1, 0, nullptr, nullptr)); + } + + SECTION("blockDim.z == maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ); + HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, z, 1, 1, 0, nullptr, nullptr)); + } +} + +/** + * Test Description + * ------------------------ + * - Negative parameters test for `hipModuleLaunchCooperativeKernel`. + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchCooperativeKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ +TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Negative_Parameters") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + + SECTION("f == nullptr") { + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernel(nullptr, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr), + hipErrorInvalidResourceHandle); + } + + SECTION("gridDim.x == 0") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 0, 1, 1, 1, 1, 1, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("gridDim.y == 0") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 0, 1, 1, 1, 1, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("gridDim.z == 0") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 0, 1, 1, 1, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.x == 0") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 0, 1, 1, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.y == 0") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 0, 1, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.z == 0") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, 0, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.x > maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u; + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, x, 1, 1, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.y > maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u; + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, y, 1, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.z > maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u; + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, z, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { + const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock); + const unsigned int dim = std::ceil(std::cbrt(max)); + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernel(f, 1, 1, 1, dim, dim, dim, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + +#if HT_AMD // Disabled due to defect EXSWHTEC-351 + SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { + const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u; + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, 1, max, nullptr, nullptr), + hipErrorInvalidValue); + } + + SECTION("Invalid stream") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, 1, 0, stream, nullptr), + hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc b/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc new file mode 100644 index 0000000000..1deaae02c0 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc @@ -0,0 +1,227 @@ +/* +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. +*/ + +/** + * @addtogroup hipModuleLaunchCooperativeKernelMultiDevice + * hipModuleLaunchCooperativeKernelMultiDevice + * @{ + * @ingroup ModuleTest + * `hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams* launchParamsList, unsigned + * int numDevices, unsigned int flags)` - + * Launches kernels on multiple devices where thread blocks can cooperate and synchronize as they + * execute. + */ + +#include +#include +#include + +#include "hip_module_launch_kernel_common.hh" + +/** + * Test Description + * ------------------------ + * - Tests `hipModuleLaunchCooperativeKernel` for a cooperative kernel with no parameters. + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ +TEST_CASE("Unit_hipModuleLaunchCooperativeKernelMultiDevice_Positive_Basic") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipFunction_t f = GetKernel(mg.module(), "CoopKernel"); + + const auto device_count = HipTest::getDeviceCount(); + + std::vector params_list(device_count); + + int device = 0; + for (auto& params : params_list) { + params.function = f; + params.gridDimX = 1; + params.gridDimY = 1; + params.gridDimZ = 1; + params.blockDimX = 1; + params.blockDimY = 1; + params.blockDimZ = 1; + params.kernelParams = nullptr; + params.sharedMemBytes = 0; + HIP_CHECK(hipSetDevice(device++)); + HIP_CHECK(hipStreamCreate(¶ms.hStream)); + } + + HIP_CHECK(hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u)); + + for (const auto params : params_list) { + HIP_CHECK(hipStreamSynchronize(params.hStream)); + } + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.hStream)); + } +} + +/** + * Test Description + * ------------------------ + * - Negative parameters test for `hipModuleLaunchCooperativeKernelMultiDevice`. + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ +TEST_CASE("Unit_hipModuleLaunchCooperativeKernelMultiDevice_Negative_Parameters") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipFunction_t f = GetKernel(mg.module(), "CoopKernel"); + + const auto device_count = HipTest::getDeviceCount(); + + std::vector params_list(device_count); + + int device = 0; + for (auto& params : params_list) { + params.function = f; + params.gridDimX = 1; + params.gridDimY = 1; + params.gridDimZ = 1; + params.blockDimX = 1; + params.blockDimY = 1; + params.blockDimZ = 1; + params.kernelParams = nullptr; + params.sharedMemBytes = 0; + HIP_CHECK(hipSetDevice(device++)); + HIP_CHECK(hipStreamCreate(¶ms.hStream)); + } + + SECTION("launchParamsList == nullptr") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernelMultiDevice(nullptr, device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("numDevices == 0") { + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), 0, 0u), + hipErrorInvalidValue); + } + + SECTION("numDevices > device count") { + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count + 1, 0u), + hipErrorInvalidValue); + } + + SECTION("invalid flags") { + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 999), + hipErrorInvalidValue); + } + + if (device_count > 1) { + SECTION("launchParamsList.func doesn't match across all devices") { + params_list[1].function = GetKernel(mg.module(), "NOPKernel"); + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("launchParamsList.gridDim doesn't match across all kernels") { + params_list[1].gridDimX = 2; + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("launchParamsList.blockDim doesn't match across all kernels") { + params_list[1].blockDimX = 2; + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("launchParamsList.sharedMem doesn't match across all kernels") { + params_list[1].sharedMemBytes = 1024; + HIP_CHECK_ERROR( + hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + } + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.hStream)); + } +} + +/** + * Test Description + * ------------------------ + * - Tries running `hipModuleLaunchCooperativeKernelMultiDevice` with multiple kernels on the same + * device. + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.5 + */ +TEST_CASE("Unit_hipModuleLaunchCooperativeKernelMultiDevice_Negative_MultiKernelSameDevice") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipFunction_t f = GetKernel(mg.module(), "CoopKernel"); + + HIP_CHECK(hipSetDevice(0)); + + std::vector params_list(2); + + for (auto& params : params_list) { + params.function = f; + params.gridDimX = 1; + params.gridDimY = 1; + params.gridDimZ = 1; + params.blockDimX = 1; + params.blockDimY = 1; + params.blockDimZ = 1; + params.kernelParams = nullptr; + params.sharedMemBytes = 0; + HIP_CHECK(hipStreamCreate(¶ms.hStream)); + } + + HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), 2, 0u), + hipErrorInvalidValue); + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.hStream)); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/module/launch_kernel_module.cc b/projects/hip-tests/catch/unit/module/launch_kernel_module.cc index 01c04b45d6..12821da450 100644 --- a/projects/hip-tests/catch/unit/module/launch_kernel_module.cc +++ b/projects/hip-tests/catch/unit/module/launch_kernel_module.cc @@ -20,6 +20,7 @@ THE SOFTWARE. */ #include +#include extern "C" { __global__ void NOPKernel() {} @@ -34,4 +35,9 @@ __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) { } } } + +__global__ void CoopKernel() { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + grid.sync(); +} } \ No newline at end of file