Support passing macros to hipLaunchKernelGGL

This makes hipLaunchKernelGGL take a variable argument list, that will be
expanded before being fed to hipLaunchKernelGGLInternal.

Change-Id: Id76e2bf91acd5d68f56a24fc39f219f2eeb06d33


[ROCm/clr commit: fb6bb04f03]

projects/clr/hipamd/tests/src/kernel/hipLaunchParm.cpp

@@ -571,6 +571,10 @@ __global__ void hipLaunchKernelStructFunc21(
 
 __global__ void vAdd(float* a) {}
 
+template<class T1, class T2>
+__global__ void myKernel(T1 a, T2 b) {}
+
+
 //---
 // Some wrapper macro for testing:
 #define WRAP(...) __VA_ARGS__
@@ -913,6 +917,19 @@ 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:
+    // Additional parenthesis required case of templates with multiple type paramters.
+    int a;
+    float b;
+    hipLaunchKernelGGL(HIP_KERNEL_NAME((myKernel<int, float>)), KERNEL_CONFIG, a, b);
+
+#define TYPE_PARAM_CONFIG int, float
+    hipLaunchKernelGGL(HIP_KERNEL_NAME((myKernel<TYPE_PARAM_CONFIG>)), KERNEL_CONFIG, a, b);
+
     // Test: Passing hipLaunchKernelGGL inside another macro:
     float e0;
     MY_LAUNCH_MACRO(hipLaunchKernelGGL(vAdd, dim3(1024),