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 961717879d.

We try to accomodate the case when a kernel template has multiple
type parameters.

Change-Id: I87577d402c92b0f3b51e298f8293f4065e1f6de8

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 <hip/hip_runtime_api.h>

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<<<numblocks, numthreads, memperblock, streamId>>>(__VA_ARGS__);                 \
+        kernelName<<<numBlocks, numThreads, memPerBlock, streamId>>>(__VA_ARGS__);                 \
     } while (0)
 
+#define hipLaunchKernelGGL(kernelName, ...)  hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__)
+
 #define hipReadModeElementType cudaReadModeElementType
 
 #ifdef __CUDA_ARCH__