From e3e31f1486b690839f69bb14575fd54d529c868e Mon Sep 17 00:00:00 2001 From: "Betigeri, Sourabh" Date: Mon, 7 Apr 2025 12:54:01 +0900 Subject: [PATCH] SWDEV-523281 - Adds hipLaunchKernelEx C++ wrapper (#44) [ROCm/hip commit: 3998c11287490b3f086ec06170f65ae0199afc71] --- projects/hip/include/hip/hip_runtime_api.h | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index c835e27658..1eaff984ae 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -9805,6 +9805,28 @@ static inline hipError_t hipMallocFromPoolAsync( hipStream_t stream) { return hipMallocFromPoolAsync(reinterpret_cast(dev_ptr), size, mem_pool, stream); } +/** + * @brief Launches a HIP kernel using the specified configuration. + * @ingroup Execution + * + * This function dispatches the provided kernel with the given launch configuration and forwards the + * kernel arguments. + * + * @param [in] config Pointer to the kernel launch configuration structure. + * @param [in] kernel Pointer to the device kernel function to be launched. + * @param [in] args Variadic list of arguments to be passed to the kernel. + * + * @returns #hipSuccess if the kernel is launched successfully, otherwise an appropriate error code. + */ +template +static inline __host__ hipError_t hipLaunchKernelEx(const hipLaunchConfig_t* config, + void (*kernel)(KernelArgs...), + Params&&... args) { + return [&](KernelArgs... convertedArgs) { + void* pArgs[] = {&convertedArgs...}; + return ::hipLaunchKernelExC(config, reinterpret_cast(kernel), pArgs); + }(std::forward(args)...); +} /** * @} */