hipLaunchKernel supports empty kernel with no arguments

Change-Id: Iaa46ced7fe0e3d72734d02250dfb48b98bef615d

include/hcc_detail/hip_runtime.h

@@ -526,7 +526,7 @@ do {\
         fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:(%d,%d,%d) groupDim:(%d,%d,%d) groupMem:+%d stream=%p\n" KNRM, \
                 #_kernelName, lp.gridDim.x, lp.gridDim.y, lp.gridDim.z, lp.groupDim.x, lp.groupDim.y, lp.groupDim.z, lp.groupMemBytes, (void*)(_stream));\
     }\
-  _kernelName (lp, __VA_ARGS__);\
+  _kernelName (lp, ##__VA_ARGS__);\
   ihipPostLaunchKernel(trueStream, lp);\
 } while(0)
 
@@ -550,7 +550,7 @@ do {\
         fprintf(stderr, "==hip-api: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \
                 #_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\
     }\
-  _kernelName (lp, __VA_ARGS__);\
+  _kernelName (lp, ##__VA_ARGS__);\
   ihipPostLaunchKernel(trueStream, cf);\
 } while(0)
 /*end hipLaunchKernel */

include/nvcc_detail/hip_runtime.h

@@ -31,7 +31,7 @@ typedef int hipLaunchParm ;
 
 #define hipLaunchKernel(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
 do {\
-kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, __VA_ARGS__);\
+kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, ##__VA_ARGS__);\
 } while(0)