From f62fdb73804a99bb0df0275732a18922a2c6d687 Mon Sep 17 00:00:00 2001 From: "Gollamandala, Srinivasarao" Date: Tue, 3 Jun 2025 14:37:25 +0530 Subject: [PATCH] SWDEV-526022-[catch2][dtest]-Functional and Negative tests for hipLaunchKernelExC, hipLaunchKernelEx and hipDrvLaunchKernelEx APIs (#156) * Fntl tests of hipLaunchKernelExC, hipLaunchKernelEx and hipDrvLaunchKernelEx APIs * Fixed the test failures on windows [ROCm/hip-tests commit: 60b3aae3c0a69b2f6d79fb95b150279090cce7ec] --- .../catch/unit/kernel/CMakeLists.txt | 1 + .../catch/unit/kernel/hipLaunchKernelEx.cc | 585 ++++++++++++++++++ .../catch/unit/module/CMakeLists.txt | 13 +- .../catch/unit/module/coopKernel.cpp | 96 +++ .../catch/unit/module/hipDrvLaunchKernelEx.cc | 450 ++++++++++++++ 5 files changed, 1144 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/catch/unit/kernel/hipLaunchKernelEx.cc create mode 100644 projects/hip-tests/catch/unit/module/coopKernel.cpp create mode 100644 projects/hip-tests/catch/unit/module/hipDrvLaunchKernelEx.cc diff --git a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt index 6b8b8129f4..ca11dd9358 100644 --- a/projects/hip-tests/catch/unit/kernel/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/kernel/CMakeLists.txt @@ -34,6 +34,7 @@ set(TEST_SRC hipTestMemKernel.cc launch_bounds.cc hipAutoThreadIdx.cc + hipLaunchKernelEx.cc ) if(UNIX) set(TEST_SRC ${TEST_SRC} diff --git a/projects/hip-tests/catch/unit/kernel/hipLaunchKernelEx.cc b/projects/hip-tests/catch/unit/kernel/hipLaunchKernelEx.cc new file mode 100644 index 0000000000..a41fb2d4bd --- /dev/null +++ b/projects/hip-tests/catch/unit/kernel/hipLaunchKernelEx.cc @@ -0,0 +1,585 @@ +/* +Copyright (c) 2025 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 +#include +#include +#include + +static constexpr int N = 10; +static __device__ int devArr[N]; +//------------------------------------------------------------------------------ +// Kernel using hipLaunchKernelEx +//------------------------------------------------------------------------------ +__global__ void cooperativeKernelEx(int *output, int totalThreads) { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + int tid = threadIdx.x + blockDim.x * blockIdx.x; + if (tid < totalThreads) { + output[tid] = tid * 3; + } + grid.sync(); + if (tid == 0) { + output[0] = 2222; + } +} +//------------------------------------------------------------------------------ +// Kernel using hipLaunchKernelExC +//------------------------------------------------------------------------------ +__global__ void cooperativeKernelExC(int *output, int totalThreads) { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + int tid = threadIdx.x + blockDim.x * blockIdx.x; + if (tid < totalThreads) { + output[tid] = tid; + } + grid.sync(); + if (tid == 0) { + output[0] = 1111; + } +} +/* + * Kernel which doesn't use cooperative groups and without any arguments + */ +static __global__ void emptyKernel() {} + +/* + * Kernel which doesn't use cooperative groups and takes an argument + * and updates the value with 100 + */ +static __global__ void argKernel(int *val) { *val = 100; } + +/* + * Kernel which uses cooperative groups and without any arguments + */ +static __global__ void coopEmptykernel() { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + grid.sync(); +} + +/* + * Kernel which uses cooperative groups and takes an array as argument + * and performs below operations + * 1) All the threads in the block fills the element in the devArr array + * based on the blockIdx.x + * 2) Wait for all the blocks completes it's operations + * 3) Fill each element in the output array with sum of elements in devArr + */ +static __global__ void coopFillArrayKernel(int *output) { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + + if (blockIdx.x == 0) + devArr[0] = 10; + else if (blockIdx.x == 1) + devArr[1] = 20; + else if (blockIdx.x == 2) + devArr[2] = 30; + else if (blockIdx.x == 3) + devArr[3] = 40; + else if (blockIdx.x == 4) + devArr[4] = 50; + else if (blockIdx.x == 5) + devArr[5] = 60; + else if (blockIdx.x == 6) + devArr[6] = 70; + else if (blockIdx.x == 7) + devArr[7] = 80; + else if (blockIdx.x == 8) + devArr[8] = 90; + else if (blockIdx.x == 9) + devArr[9] = 100; + + grid.sync(); + + for (int i = 0; i < N; i++) { + output[blockIdx.x] = output[blockIdx.x] + devArr[i]; + } +} + +__global__ void normalKernel(int *output, int totalThreads) { + int tid = threadIdx.x + blockDim.x * blockIdx.x; + if (tid < totalThreads) { + output[tid] = tid * 3; + } +} +/** + * @addtogroup hipLaunchKernelExC hipLaunchKernelExC + * @{ + * @ingroup KernelTest + * `hipError_t hipLaunchKernelExC(const hipLaunchConfig_t* config, const void* + * fPtr, void** args)` - Launches a HIP kernel using a generic function pointer + * and the specified configuration. + */ + +/** + * Test Description + * ------------------------ + * This test verfies the different negative scenarios of hipLaunchKernelExC + + * Test source + * ------------------------ + * - catch/unit/kernel/hipLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipLaunchKernelExC_NegetiveTsts") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + int blockSize = 16; + int totalThreads = 64; + const int numBlocks = (totalThreads + blockSize - 1) / blockSize; + hipLaunchConfig_t config = {}; + config.gridDim = dim3{(uint32_t)numBlocks, 1, 1}; + config.blockDim = dim3{(uint32_t)blockSize, 1, 1}; + config.dynamicSmemBytes = 0; + config.stream = 0; + + hipLaunchAttribute attr; + attr.id = hipLaunchAttributeCooperative; + attr.val.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + int *d_output = nullptr; + HIP_CHECK(hipMalloc(&d_output, totalThreads * sizeof(int))); + HIP_CHECK(hipMemset(d_output, 0, totalThreads * sizeof(int))); + void *kernelArgs[] = {&d_output, (void *)&totalThreads}; + + SECTION("Kernel function as nullptr") { + HIP_CHECK_ERROR(hipLaunchKernelExC(&config, nullptr, kernelArgs), + hipErrorInvalidDeviceFunction); + } + + SECTION("Kernel args as nullptr") { + HIP_CHECK_ERROR( + hipLaunchKernelExC(&config, (void *)cooperativeKernelExC, nullptr), + hipErrorInvalidValue); + } + + SECTION("Non Cooparative Kernel") { + HIP_CHECK_ERROR( + hipLaunchKernelExC(&config, (void *)normalKernel, kernelArgs), + hipSuccess); + } + + hipLaunchConfig_t invalidConfig = {}; + invalidConfig.gridDim = dim3{0, 1, 1}; + invalidConfig.blockDim = dim3{0, 1, 1}; + invalidConfig.dynamicSmemBytes = 0; + invalidConfig.stream = 0; + + hipLaunchAttribute invalidAttr; + invalidAttr.id = hipLaunchAttributeCooperative; + invalidAttr.val.cooperative = 1; + invalidConfig.attrs = &invalidAttr; + invalidConfig.numAttrs = 1; + + SECTION("Invalid Kernel Config") { + HIP_CHECK_ERROR(hipLaunchKernelExC(&invalidConfig, + (void *)cooperativeKernelExC, + kernelArgs), + hipErrorInvalidConfiguration); + } +} +/** + * Test Description + * ------------------------ + * This test verfies the different negative scenarios of hipLaunchKernelEx + + * Test source + * ------------------------ + * - catch/unit/kernel/hipLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ + +TEST_CASE("Unit_hipLaunchKernelEx_NegetiveTsts") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + int blockSize = 16; + int totalThreads = 64; + const int numBlocks = (totalThreads + blockSize - 1) / blockSize; + hipLaunchConfig_t config = {}; + config.gridDim = dim3{(uint32_t)numBlocks, 1, 1}; + config.blockDim = dim3{(uint32_t)blockSize, 1, 1}; + config.dynamicSmemBytes = 0; + config.stream = 0; + + hipLaunchAttribute attr; + attr.id = hipLaunchAttributeCooperative; + attr.val.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + int *d_output = nullptr; + HIP_CHECK(hipMalloc(&d_output, totalThreads * sizeof(int))); + HIP_CHECK(hipMemset(d_output, 0, totalThreads * sizeof(int))); + + SECTION("Empty Kernel Args") { + HIP_CHECK_ERROR(hipLaunchKernelEx(&config, emptyKernel), hipSuccess); + } + + SECTION("Non Cooparative Kernel") { + HIP_CHECK_ERROR(hipLaunchKernelEx(&config, + (void (*)(int *, int))normalKernel, + d_output, totalThreads), + hipSuccess); + } + + hipLaunchConfig_t invalidConfig = {}; + invalidConfig.gridDim = dim3{0, 1, 1}; + invalidConfig.blockDim = dim3{0, 1, 1}; + invalidConfig.dynamicSmemBytes = 0; + invalidConfig.stream = 0; + + hipLaunchAttribute attr_1; + attr_1.id = hipLaunchAttributeCooperative; + attr_1.val.cooperative = 1; + invalidConfig.attrs = &attr_1; + invalidConfig.numAttrs = 1; + + SECTION("Invalid Kernel Config") { + HIP_CHECK_ERROR(hipLaunchKernelEx(&invalidConfig, + (void (*)(int *, int))cooperativeKernelEx, + d_output, totalThreads), + hipErrorInvalidConfiguration); + } +} + +bool runTest(const char *testName, const void *kernelFunc, int totalThreads, + int blockSize, int flagValue, bool useTemplate) { + const int numBlocks = (totalThreads + blockSize - 1) / blockSize; + + int *d_output = nullptr; + HIP_CHECK(hipMalloc(&d_output, totalThreads * sizeof(int))); + HIP_CHECK(hipMemset(d_output, 0, totalThreads * sizeof(int))); + + hipLaunchConfig_t config = {}; + config.gridDim = dim3{(uint32_t)numBlocks, 1, 1}; + config.blockDim = dim3{(uint32_t)blockSize, 1, 1}; + config.dynamicSmemBytes = 0; + config.stream = 0; + + hipLaunchAttribute attr; + attr.id = hipLaunchAttributeCooperative; + attr.val.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + // For a kernel parameter declared as "int* output", pass the address of the + // device pointer. + void *kernelArgs[] = {&d_output, (void *)&totalThreads}; + + if (useTemplate) { + HIP_CHECK(hipLaunchKernelEx(&config, (void (*)(int *, int))kernelFunc, + d_output, totalThreads)); + } else { + HIP_CHECK(hipLaunchKernelExC(&config, kernelFunc, kernelArgs)); + } + + HIP_CHECK(hipDeviceSynchronize()); + + int *h_output = (int *)malloc(totalThreads * sizeof(int)); + HIP_CHECK(hipMemcpy(h_output, d_output, totalThreads * sizeof(int), + hipMemcpyDeviceToHost)); + + // Verify results. + bool success = true; + if (h_output[0] != flagValue) { + printf("%s test failed: Expected flag %d at index 0, got %d\n", testName, + flagValue, h_output[0]); + success = false; + } + for (int i = 1; i < totalThreads; i++) { + int expectedValue = (flagValue == 1111) ? i : (i * 3); + if (h_output[i] != expectedValue) { + printf("%s test failed at index %d: Expected %d, got %d\n", testName, i, + expectedValue, h_output[i]); + success = false; + break; + } + } + + HIP_CHECK(hipFree(d_output)); + free(h_output); + return success; +} +/** + * Test Description + * ------------------------ + * This test verfies basic positive test scenarios of hipLaunchKernelEx, + hipLaunchKernelExC + + * Test source + * ------------------------ + * - catch/unit/kernel/hipLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipLaunchKernelEx_Functional") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + std::string api_type = GENERATE("hipLaunchKernelEx", "hipLaunchKernelExC"); + if (api_type == "hipLaunchKernelEx") { + REQUIRE(runTest(api_type.c_str(), (void *)cooperativeKernelEx, 64, 16, 2222, + true) == true); + } + if (api_type == "hipLaunchKernelExC") { + REQUIRE(runTest(api_type.c_str(), (void *)cooperativeKernelExC, 64, 16, + 1111, false) == true); + } +} +/** + * Test Description + * ------------------------ + * - This testcase checks following scenarios + * - 1) Launch Normal kernel which has no arguments using hipLaunchKernelEx + * - 2) Launch Normal kernel which has no arguments using hipLaunchKernelExC + * - 3) Launch Normal kernel which has arguments by using hipLaunchKernelEx + * - 4) Launch Normal kernel which has arguments by using hipLaunchKernelExC + * - 5) Launch Cooperative kernel which has no arguments by using + * - hipLaunchKernelEx + * - 6) Launch Cooperative kernel which has no arguments by using + * - hipLaunchKernelExC + * Test source + * ------------------------ + * - unit/kernel/hipLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipLaunchKernelEx_With_Different_Kernels") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipLaunchConfig_t config = {}; + config.gridDim = dim3{2, 2, 1}; + config.blockDim = dim3{2, 2, 1}; + config.dynamicSmemBytes = 0; + config.stream = 0; + + hipLaunchAttribute attr; + attr.id = hipLaunchAttributeCooperative; + memset(attr.pad, 0, sizeof(attr.pad)); + attr.val.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + SECTION("Normal kernel with no arguments") { + SECTION("hipLaunchKernelEx") { + HIP_CHECK(hipLaunchKernelEx(&config, emptyKernel)); + } + + SECTION("hipLaunchKernelExC") { + HIP_CHECK(hipLaunchKernelExC( + &config, reinterpret_cast(emptyKernel), nullptr)); + } + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("Kernel with arguments using kernelParams") { + config.gridDim = dim3{1, 1, 1}; + config.blockDim = dim3{1, 1, 1}; + + int *devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, sizeof(int))); + + SECTION("hipLaunchKernelEx") { + HIP_CHECK(hipLaunchKernelEx(&config, argKernel, devMem)); + } + + SECTION("hipLaunchKernelExC") { + void *kernel_args[1] = {&devMem}; + HIP_CHECK(hipLaunchKernelExC(&config, reinterpret_cast(argKernel), + kernel_args)); + } + HIP_CHECK(hipDeviceSynchronize()); + + int result = 0; + HIP_CHECK(hipMemcpy(&result, devMem, sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 100); + } + + SECTION("Cooperative kernel with no arguments") { + SECTION("hipLaunchKernelEx") { + HIP_CHECK(hipLaunchKernelEx(&config, coopEmptykernel)); + } + + SECTION("hipLaunchKernelExC") { + HIP_CHECK(hipLaunchKernelExC( + &config, reinterpret_cast(coopEmptykernel), nullptr)); + } + HIP_CHECK(hipDeviceSynchronize()); + } +} + +/** + * Test Description + * ------------------------ + * - This testcase launches a kernel which uses the cooperative groups and + * - with N number of blocks and one thread in each block by using the + * - hipLaunchKernelEx and hipLaunchKernelExC and validates the kernel output + * Test source + * ------------------------ + * - unit/kernel/hipLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipLaunchKernelEx_With_CooperativeKernelWithArgs") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipLaunchConfig_t config = {}; + config.gridDim = dim3{N, 1, 1}; + config.blockDim = dim3{1, 1, 1}; + config.dynamicSmemBytes = 0; + config.stream = 0; + + hipLaunchAttribute attr; + attr.id = hipLaunchAttributeCooperative; + memset(attr.pad, 0, sizeof(attr.pad)); + attr.val.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + int hostMem[N]; + for (int i = 0; i < N; i++) { + hostMem[i] = 0; + } + + int *devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, N * sizeof(int))); + HIP_CHECK(hipMemcpy(devMem, hostMem, N * sizeof(int), hipMemcpyDefault)); + + SECTION("hipLaunchKernelEx") { + HIP_CHECK(hipLaunchKernelEx(&config, coopFillArrayKernel, devMem)); + } + + SECTION("hipLaunchKernelExC") { + void *kernel_args[1] = {&devMem}; + HIP_CHECK(hipLaunchKernelExC( + &config, reinterpret_cast(coopFillArrayKernel), kernel_args)); + } + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(hostMem, devMem, N * sizeof(int), hipMemcpyDefault)); + for (int i = 0; i < N; i++) { + REQUIRE(hostMem[i] == 550); + } + + HIP_CHECK(hipFree(devMem)); +} + +/** + * Test Description + * ------------------------ + * - This testcase checks following scenarios + * - 1) Launch a kernel with one block and maximum allowed threads in a block + * - in the x direction by using hipLaunchKernelEx and hipLaunchKernelExC + * - 2) Launch a kernel with one block and maximum allowed threads in a block + * - in the y direction by using hipLaunchKernelEx and hipLaunchKernelExC + * - 3) Launch a kernel with one block and maximum allowed threads in a block + * - in the z direction by using hipLaunchKernelEx and hipLaunchKernelExC + * Test source + * ------------------------ + * - unit/kernel/hipLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipLaunchKernelEx_With_MaxBlockDims") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipLaunchConfig_t config = {}; + config.gridDim = dim3{1, 1, 1}; + config.blockDim = dim3{1, 1, 1}; + config.dynamicSmemBytes = 0; + config.stream = 0; + + hipLaunchAttribute attr; + attr.id = hipLaunchAttributeCooperative; + memset(attr.pad, 0, sizeof(attr.pad)); + attr.val.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + SECTION("blockDim.x == maxBlockDimX") { + const unsigned int x = + GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); + config.blockDim = dim3{x, 1, 1}; + + SECTION("hipLaunchKernelEx") { + HIP_CHECK(hipLaunchKernelEx(&config, emptyKernel)); + } + + SECTION("hipLaunchKernelExC") { + HIP_CHECK(hipLaunchKernelExC( + &config, reinterpret_cast(emptyKernel), nullptr)); + } + } + + SECTION("blockDim.y == maxBlockDimY") { + const unsigned int y = + GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0); + config.blockDim = dim3{1, y, 1}; + + SECTION("hipLaunchKernelEx") { + HIP_CHECK(hipLaunchKernelEx(&config, emptyKernel)); + } + + SECTION("hipLaunchKernelExC") { + HIP_CHECK(hipLaunchKernelExC( + &config, reinterpret_cast(emptyKernel), nullptr)); + } + } + + SECTION("blockDim.z == maxBlockDimZ") { + const unsigned int z = + GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0); + config.blockDim = dim3{1, 1, z}; + + SECTION("hipLaunchKernelEx") { + HIP_CHECK(hipLaunchKernelEx(&config, emptyKernel)); + } + + SECTION("hipLaunchKernelExC") { + HIP_CHECK(hipLaunchKernelExC( + &config, reinterpret_cast(emptyKernel), nullptr)); + } + } + HIP_CHECK(hipDeviceSynchronize()); +} +/** + * End doxygen group KernelTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/module/CMakeLists.txt b/projects/hip-tests/catch/unit/module/CMakeLists.txt index 3139f221c1..67884f2b49 100644 --- a/projects/hip-tests/catch/unit/module/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/module/CMakeLists.txt @@ -1,4 +1,4 @@ -# Copyright (c) 2023-2024 Advanced Micro Devices, Inc. All Rights Reserved. +# Copyright (c) 2023-2025 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 @@ -33,6 +33,7 @@ set(TEST_SRC hipModuleLaunchCooperativeKernelMultiDevice.cc hipFuncGetAttribute.cc hipGetFuncBySymbol.cc + hipDrvLaunchKernelEx.cc ) add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code @@ -63,11 +64,19 @@ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_tex_ref_module.cc) add_custom_target(get_tex_ref_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code) +add_custom_target(coopKernel.code + COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} + ${CMAKE_CURRENT_SOURCE_DIR}/coopKernel.cpp + -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/coopKernel.code + -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} + -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code + ${CMAKE_CURRENT_BINARY_DIR}/coopKernel.code ) # Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900 --offload-arch=gfx906" # having space at the start/end of OFFLOAD_ARCH_STR can cause build failures @@ -213,6 +222,7 @@ if(HIP_PLATFORM MATCHES "amd") else() set(RTCLIB "nvrtc") endif() + hip_add_exe_to_target(NAME ModuleTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests @@ -220,6 +230,7 @@ hip_add_exe_to_target(NAME ModuleTest COMMON_SHARED_SRC ${COMMON_SHARED_SRC} COMPILE_OPTIONS -std=c++17) +add_dependencies(ModuleTest coopKernel.code) add_dependencies(ModuleTest get_function_module) add_dependencies(ModuleTest launch_kernel_module) add_dependencies(ModuleTest get_global_test_module) diff --git a/projects/hip-tests/catch/unit/module/coopKernel.cpp b/projects/hip-tests/catch/unit/module/coopKernel.cpp new file mode 100644 index 0000000000..d0095288aa --- /dev/null +++ b/projects/hip-tests/catch/unit/module/coopKernel.cpp @@ -0,0 +1,96 @@ + +/* +Copyright (c) 2025 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +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 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. +*/ + +#include "hip/hip_runtime.h" +#include + +using namespace cooperative_groups; +extern "C" { +__global__ void cooperativeKernelEx(int* output, int totalThreads) { + grid_group grid = this_grid(); + int tid = threadIdx.x + blockDim.x * blockIdx.x; + if (tid < totalThreads) { + output[tid] = tid * 3; + } + grid.sync(); + if (tid == 0) { + output[0] = 2222; + } +} + +/* + * Kernel which doesn't use cooperative groups and without any arguments + */ +__global__ void emptyKernel() {} + +/* + * Kernel which doesn't use cooperative groups and takes an argument + * and updates the value with 100 + */ +__global__ void argKernel(int *val) { *val = 100; } + +/* + * Kernel which uses cooperative groups and without any arguments + */ +__global__ void coopEmptykernel() { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + grid.sync(); +} + +/* + * Kernel which uses cooperative groups and takes arguments + * and performs below operations + * 1) All the threads in the block fills the element in the arr array + * based on the blockIdx.x + * 2) Wait for all the blocks completes it's operations + * 3) Fill each element in the output array with sum of elements in arr + */ +__global__ void coopFillArrayKernel(int *arr, int *output, int N) { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + + if (blockIdx.x == 0) + arr[0] = 10; + else if (blockIdx.x == 1) + arr[1] = 20; + else if (blockIdx.x == 2) + arr[2] = 30; + else if (blockIdx.x == 3) + arr[3] = 40; + else if (blockIdx.x == 4) + arr[4] = 50; + else if (blockIdx.x == 5) + arr[5] = 60; + else if (blockIdx.x == 6) + arr[6] = 70; + else if (blockIdx.x == 7) + arr[7] = 80; + else if (blockIdx.x == 8) + arr[8] = 90; + else if (blockIdx.x == 9) + arr[9] = 100; + + grid.sync(); + + for (int i = 0; i < N; i++) { + output[blockIdx.x] = output[blockIdx.x] + arr[i]; + } +} +} + diff --git a/projects/hip-tests/catch/unit/module/hipDrvLaunchKernelEx.cc b/projects/hip-tests/catch/unit/module/hipDrvLaunchKernelEx.cc new file mode 100644 index 0000000000..bf82a7a443 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipDrvLaunchKernelEx.cc @@ -0,0 +1,450 @@ +/* +Copyright (c) 2025 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 +#include +#include +#include + +static constexpr int N = 10; +constexpr auto CODE_OBJ_SINGLEARCH = "coopKernel.code"; +static constexpr auto kernel_name = "cooperativeKernelEx"; + +/** + * @addtogroup hipDrvLaunchKernelEx hipDrvLaunchKernelEx + * @{ + * @ingroup ModuleTest + * `hipError_t hipDrvLaunchKernelEx(const HIP_LAUNCH_CONFIG* config, + hipFunction_t f, void** params, void** extra)` + * Launches a HIP kernel using the driver API with the specified configuration. + */ + +/** + * Test Description + * ------------------------ + * This test verfies the different negative scenarios of hipDrvLaunchKernelEx + * Api Test source + * ------------------------ + * - catch/unit/module/hipDrvLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDrvLaunchKernelEx_NegTsts") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + int totalThreads = 64; + int blockSize = 16; + int numBlocks = (totalThreads + blockSize - 1) / blockSize; + + int *d_output = nullptr; + HIP_CHECK(hipMalloc(&d_output, totalThreads * sizeof(int))); + HIP_CHECK(hipMemset(d_output, 0, totalThreads * sizeof(int))); + + // Set up the HIP_LAUNCH_CONFIG structure + HIP_LAUNCH_CONFIG config = {}; + config.gridDimX = numBlocks; + config.gridDimY = 1; + config.gridDimZ = 1; + config.blockDimX = blockSize; + config.blockDimY = 1; + config.blockDimZ = 1; + config.sharedMemBytes = 0; + config.hStream = 0; // default stream + + // Set up a cooperative launch attribute + hipDrvLaunchAttribute attr; + attr.id = hipDrvLaunchAttributeCooperative; + // Zero out the padding bytes + memset(attr.pad, 0, sizeof(attr.pad)); + attr.value.cooperative = 1; + + config.attrs = &attr; + config.numAttrs = 1; + + // Kernel parameters: address of d_output and totalThreads. + void *kernelParams[] = {&d_output, &totalThreads}; + + hipModule_t module; + HIP_CHECK(hipModuleLoad(&module, CODE_OBJ_SINGLEARCH)); + + hipFunction_t function; + HIP_CHECK(hipModuleGetFunction(&function, module, kernel_name)); + + SECTION("Kernel config as nullptr") { + HIP_CHECK_ERROR(hipDrvLaunchKernelEx(nullptr, function, kernelParams, NULL), + hipErrorInvalidValue); + } + SECTION("Kernel function as nullptr") { + HIP_CHECK_ERROR(hipDrvLaunchKernelEx(&config, nullptr, kernelParams, NULL), + hipErrorInvalidResourceHandle); + } + SECTION("Kernel parameter as nullptr") { + HIP_CHECK_ERROR(hipDrvLaunchKernelEx(&config, function, nullptr, NULL), + hipErrorInvalidValue); + } + HIP_LAUNCH_CONFIG invalidConfig = {}; + invalidConfig.gridDimX = 0; + invalidConfig.gridDimY = 1; + invalidConfig.gridDimZ = 1; + invalidConfig.blockDimX = 0; + invalidConfig.blockDimY = 1; + invalidConfig.blockDimZ = 1; + invalidConfig.sharedMemBytes = 0; + invalidConfig.hStream = 0; // default stream + + // Set up a cooperative launch attribute + hipDrvLaunchAttribute invalidAttr; + invalidAttr.id = hipDrvLaunchAttributeCooperative; + // Zero out the padding bytes + memset(invalidAttr.pad, 0, sizeof(invalidAttr.pad)); + invalidAttr.value.cooperative = 1; + + invalidConfig.attrs = &invalidAttr; + invalidConfig.numAttrs = 1; + + SECTION("Invalid Kernel config") { + HIP_CHECK_ERROR( + hipDrvLaunchKernelEx(&invalidConfig, function, kernelParams, NULL), + hipErrorInvalidValue); + } +} + +bool runTestDrvLaunch(const char *testName, std::string kernelFunc, + int totalThreads, int blockSize, int flagValue) { + int numBlocks = (totalThreads + blockSize - 1) / blockSize; + + int *d_output = nullptr; + HIP_CHECK(hipMalloc(&d_output, totalThreads * sizeof(int))); + HIP_CHECK(hipMemset(d_output, 0, totalThreads * sizeof(int))); + + // Set up the HIP_LAUNCH_CONFIG structure + HIP_LAUNCH_CONFIG config = {}; + config.gridDimX = numBlocks; + config.gridDimY = 1; + config.gridDimZ = 1; + config.blockDimX = blockSize; + config.blockDimY = 1; + config.blockDimZ = 1; + config.sharedMemBytes = 0; + config.hStream = 0; // default stream + + // Set up a cooperative launch attribute + hipDrvLaunchAttribute attr; + attr.id = hipDrvLaunchAttributeCooperative; + // Zero out the padding bytes + memset(attr.pad, 0, sizeof(attr.pad)); + attr.value.cooperative = 1; + + config.attrs = &attr; + config.numAttrs = 1; + + // Kernel parameters: address of d_output and totalThreads. + void *kernelParams[] = {&d_output, &totalThreads}; + + hipModule_t module; + HIP_CHECK(hipModuleLoad(&module, CODE_OBJ_SINGLEARCH)); + + hipFunction_t function; + HIP_CHECK(hipModuleGetFunction(&function, module, kernelFunc.c_str())); + + // Launch the kernel using the driver API function. + hipError_t err = hipDrvLaunchKernelEx(&config, function, kernelParams, NULL); + if (err != hipSuccess) { + printf("%s test failed to launch kernel: error code %d\n", testName, err); + HIP_CHECK(hipFree(d_output)); + return false; + } + + HIP_CHECK(hipDeviceSynchronize()); + + int *h_output = (int *)malloc(totalThreads * sizeof(int)); + HIP_CHECK(hipMemcpy(h_output, d_output, totalThreads * sizeof(int), + hipMemcpyDeviceToHost)); + + // Verify results. + bool success = true; + if (h_output[0] != flagValue) { + printf("%s test failed: Expected flag %d at index 0, got %d\n", testName, + flagValue, h_output[0]); + success = false; + } + for (int i = 1; i < totalThreads; i++) { + int expectedValue = (flagValue == 1111) ? i : (i * 3); + if (h_output[i] != expectedValue) { + printf("%s test failed at index %d: Expected %d, got %d\n", testName, i, + expectedValue, h_output[i]); + success = false; + break; + } + } + + HIP_CHECK(hipFree(d_output)); + free(h_output); + return success; +} +/** + * Test Description + * ------------------------ + * This test verfies basic positive test of hipDrvLaunchKernelEx + * Api Test source + * ------------------------ + * - catch/unit/module/hipDrvLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDrvLaunchKernelEx_Functional") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + REQUIRE(runTestDrvLaunch("hipDrvLaunchKernelEx", kernel_name, 64, 16, 2222) == + true); +} + +/** + * Test Description + * ------------------------ + * - This testcase checks following scenarios(Uses kernels from the module) + * - 1) Launch Normal kernel which has no arguments using hipDrvLaunchKernelEx + * - 2) Launch Normal kernel which has arguments by using hipDrvLaunchKernelEx + * - 3) Launch Cooperative kernel which has no arguments by using + * - hipDrvLaunchKernelEx + * Test source + * ------------------------ + * - unit/module/hipDrvLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDrvLaunchKernelEx_With_Different_Kernels") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipModule_t module; + HIP_CHECK(hipModuleLoad(&module, "coopKernel.code")); + + HIP_LAUNCH_CONFIG config = {}; + config.gridDimX = 1; + config.gridDimY = 1; + config.gridDimZ = 1; + config.blockDimX = 1; + config.blockDimY = 1; + config.blockDimZ = 1; + config.sharedMemBytes = 0; + config.hStream = 0; + hipDrvLaunchAttribute attr; + attr.id = hipDrvLaunchAttributeCooperative; + // Zero out the padding bytes + memset(attr.pad, 0, sizeof(attr.pad)); + attr.value.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + SECTION("Normal kernel with no arguments") { + hipFunction_t simpleKernel; + HIP_CHECK(hipModuleGetFunction(&simpleKernel, module, "emptyKernel")); + + HIP_CHECK(hipDrvLaunchKernelEx(&config, simpleKernel, nullptr, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("Kernel with arguments using kernelParams") { + int *devMem = nullptr; + HIP_CHECK(hipMalloc(&devMem, sizeof(int))); + + void *kernel_args[1] = {&devMem}; + + hipFunction_t argKernel; + HIP_CHECK(hipModuleGetFunction(&argKernel, module, "argKernel")); + + HIP_CHECK(hipDrvLaunchKernelEx(&config, argKernel, kernel_args, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + + int result = 0; + HIP_CHECK(hipMemcpy(&result, devMem, sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 100); + } + + SECTION("Cooperative kernel with no arguments") { + hipFunction_t coopKernel; + HIP_CHECK(hipModuleGetFunction(&coopKernel, module, "coopEmptykernel")); + + HIP_CHECK(hipDrvLaunchKernelEx(&config, coopKernel, nullptr, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + HIP_CHECK(hipModuleUnload(module)); +} + +/** + * Test Description + * ------------------------ + * - This testcase launches a kernel(Uses kernels from the module) which + * - uses the cooperative groups andwith N number of blocks and one thread + * - in each block by using the hipDrvLaunchKernelEx and validates + * - the kernel output + * Test source + * ------------------------ + * - unit/module/hipDrvLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDrvLaunchKernelEx_With_CooperativeKernelWithArgs") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipModule_t module; + HIP_CHECK(hipModuleLoad(&module, "coopKernel.code")); + + HIP_LAUNCH_CONFIG config = {}; + config.gridDimX = N; + config.gridDimY = 1; + config.gridDimZ = 1; + config.blockDimX = 1; + config.blockDimY = 1; + config.blockDimZ = 1; + config.sharedMemBytes = 0; + config.hStream = 0; + hipDrvLaunchAttribute attr; + attr.id = hipDrvLaunchAttributeCooperative; + // Zero out the padding bytes + memset(attr.pad, 0, sizeof(attr.pad)); + attr.value.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + int hostMem[N]; + for (int i = 0; i < N; i++) { + hostMem[i] = 0; + } + + int *devMem1 = nullptr; + HIP_CHECK(hipMalloc(&devMem1, N * sizeof(int))); + HIP_CHECK(hipMemcpy(devMem1, hostMem, N * sizeof(int), hipMemcpyDefault)); + int *devMem2 = nullptr; + HIP_CHECK(hipMalloc(&devMem2, N * sizeof(int))); + HIP_CHECK(hipMemcpy(devMem2, hostMem, N * sizeof(int), hipMemcpyDefault)); + + int size = N; + void *kernel_args[3] = {&devMem1, &devMem2, &size}; + + hipFunction_t argKernel; + HIP_CHECK(hipModuleGetFunction(&argKernel, module, "coopFillArrayKernel")); + + HIP_CHECK(hipDrvLaunchKernelEx(&config, argKernel, kernel_args, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + + HIP_CHECK(hipMemcpy(hostMem, devMem2, N * sizeof(int), hipMemcpyDefault)); + for (int i = 0; i < N; i++) { + REQUIRE(hostMem[i] == 550); + } + + HIP_CHECK(hipFree(devMem1)); + HIP_CHECK(hipFree(devMem2)); + HIP_CHECK(hipModuleUnload(module)); +} + +/** + * Test Description + * ------------------------ + * - This testcase checks following scenarios(Uses kernels from the module) + * - 1) Launch a kernel with one block and maximum allowed threads in a block + * - in the x direction by using hipLaunchKernelEx and hipDrvLaunchKernelEx + * - 2) Launch a kernel with one block and maximum allowed threads in a block + * - in the y direction by using hipLaunchKernelEx and hipDrvLaunchKernelEx + * - 3) Launch a kernel with one block and maximum allowed threads in a block + * - in the z direction by using hipLaunchKernelEx and hipDrvLaunchKernelEx + * Test source + * ------------------------ + * - unit/module/hipDrvLaunchKernelEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.5 + */ +TEST_CASE("Unit_hipDrvLaunchKernelEx_With_MaxBlockDims") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + hipModule_t module; + HIP_CHECK(hipModuleLoad(&module, "coopKernel.code")); + + hipFunction_t kernel; + HIP_CHECK(hipModuleGetFunction(&kernel, module, "emptyKernel")); + + HIP_LAUNCH_CONFIG config = {}; + config.gridDimX = 1; + config.gridDimY = 1; + config.gridDimZ = 1; + config.blockDimX = 1; + config.blockDimY = 1; + config.blockDimZ = 1; + config.sharedMemBytes = 0; + config.hStream = 0; + hipDrvLaunchAttribute attr; + attr.id = hipDrvLaunchAttributeCooperative; + // Zero out the padding bytes + memset(attr.pad, 0, sizeof(attr.pad)); + attr.value.cooperative = 1; + config.attrs = &attr; + config.numAttrs = 1; + + SECTION("blockDim.x == maxBlockDimX") { + const unsigned int x = + GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); + config.blockDimX = x; + + HIP_CHECK(hipDrvLaunchKernelEx(&config, kernel, nullptr, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("blockDim.y == maxBlockDimY") { + const unsigned int y = + GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0); + config.blockDimY = y; + + HIP_CHECK(hipDrvLaunchKernelEx(&config, kernel, nullptr, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("blockDim.z == maxBlockDimZ") { + const unsigned int z = + GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0); + config.blockDimY = z; + + HIP_CHECK(hipDrvLaunchKernelEx(&config, kernel, nullptr, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + HIP_CHECK(hipModuleUnload(module)); +} +/** + * End doxygen group ModuleTest. + * @} + */