From 845c0a5b8a0bce523fdfc3af85f2bedf42fa3363 Mon Sep 17 00:00:00 2001 From: Ronak Chauhan Date: Thu, 25 Jun 2020 07:23:28 -0400 Subject: [PATCH] Support passing macros to hipLaunchKernelGGL This makes hipLaunchKernelGGL take a variable argument list, that will be expanded before being fed to hipLaunchKernelGGLInternal. This is different from 4dfef9a25b6e4ff323ed90ac72dac3b3e45e4b94. We try to accomodate the case when a kernel template has multiple type parameters. Change-Id: I87577d402c92b0f3b51e298f8293f4065e1f6de8 [ROCm/clr commit: 133251a6e1dc857d0bf795d812d5deec76021801] --- .../hipamd/include/hip/hcc_detail/hip_runtime.h | 6 ++++-- .../hipamd/include/hip/nvcc_detail/hip_runtime.h | 6 ++++-- .../hipamd/tests/src/kernel/hipLaunchParm.cpp | 16 ++++++++++++++++ 3 files changed, 24 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 13f6e6dd4c..b277ba038c 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -435,10 +435,12 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, hipLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream); } #else -#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ +#define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...) \ do { \ - kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__); \ + kernelName<<<(numBlocks), (numThreads), (memPerBlock), (streamId)>>>(__VA_ARGS__); \ } while (0) + +#define hipLaunchKernelGGL(kernelName, ...) hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__) #endif #include diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h index d159e116c0..c13540df54 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h @@ -31,11 +31,13 @@ THE SOFTWARE. typedef int hipLaunchParm; -#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ +#define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...) \ do { \ - kernelName<<>>(__VA_ARGS__); \ + kernelName<<>>(__VA_ARGS__); \ } while (0) +#define hipLaunchKernelGGL(kernelName, ...) hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__) + #define hipReadModeElementType cudaReadModeElementType #ifdef __CUDA_ARCH__ diff --git a/projects/clr/hipamd/tests/src/kernel/hipLaunchParm.cpp b/projects/clr/hipamd/tests/src/kernel/hipLaunchParm.cpp index 797bc84ced..ff67988ddb 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipLaunchParm.cpp +++ b/projects/clr/hipamd/tests/src/kernel/hipLaunchParm.cpp @@ -571,6 +571,10 @@ __global__ void hipLaunchKernelStructFunc21( __global__ void vAdd(float* a) {} +template +__global__ void myKernel(T1 a, T2 b) {} + + //--- // Some wrapper macro for testing: #define WRAP(...) __VA_ARGS__ @@ -913,6 +917,18 @@ int main() { hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), 1, 0, 0, Ad); hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), dim3(1), 0, 0, Ad); + // Test: Passing macro to hipLaunchKernelGGL +#define KERNEL_CONFIG dim3(1024), dim3(1), 0, 0 + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), KERNEL_CONFIG, Ad); + + // Test: Same thing with templates: + int a; + float b; + hipLaunchKernelGGL(HIP_KERNEL_NAME(myKernel), KERNEL_CONFIG, a, b); + +#define TYPE_PARAM_CONFIG int, float + hipLaunchKernelGGL(HIP_KERNEL_NAME(myKernel), KERNEL_CONFIG, a, b); + // Test: Passing hipLaunchKernelGGL inside another macro: float e0; MY_LAUNCH_MACRO(hipLaunchKernelGGL(vAdd, dim3(1024),