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),