From 275c5516b63a331acfe7539a61255659cad9eee9 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Fri, 17 Jun 2016 14:11:46 +0530
Subject: [PATCH] hipLaunchKernel supports empty kernel with no arguments

Change-Id: Iaa46ced7fe0e3d72734d02250dfb48b98bef615d
---
 hipamd/include/hcc_detail/hip_runtime.h  | 4 ++--
 hipamd/include/nvcc_detail/hip_runtime.h | 2 +-
 2 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h
index aa634cdf1e..51d4be40c3 100644
--- a/hipamd/include/hcc_detail/hip_runtime.h
+++ b/hipamd/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 */
diff --git a/hipamd/include/nvcc_detail/hip_runtime.h b/hipamd/include/nvcc_detail/hip_runtime.h
index 06c6ffb9b3..ecc1be8c11 100644
--- a/hipamd/include/nvcc_detail/hip_runtime.h
+++ b/hipamd/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)