From 475fd56684dc2a37e61633abab20ea4514499cb5 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 16 Jun 2016 16:18:35 +0530
Subject: [PATCH 1/2] Add compiler annotations to make_ functions

Change-Id: Ic16cfe402dc5cb387b1b63ad311ce8b57ea07361


[ROCm/clr commit: bdfd943b4a518114715808a7b54d52e41ee298d3]
---
 projects/clr/hipamd/include/hcc_detail/hip_vector_types.h | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/projects/clr/hipamd/include/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hcc_detail/hip_vector_types.h
index 0d1ad97663..1537b0fb70 100644
--- a/projects/clr/hipamd/include/hcc_detail/hip_vector_types.h
+++ b/projects/clr/hipamd/include/hcc_detail/hip_vector_types.h
@@ -185,10 +185,10 @@ __HIP_DEVICE__ double4 make_double4(double, double, double, double );
 /*
 ///---
 // Inline functions for creating vector types from basic types
-#define ONE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x) { VT t; t.x = x; return t; };
-#define TWO_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x, T y) { VT t; t.x=x; t.y=y; return t; };
-#define THREE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x, T y, T z) { VT t; t.x=x; t.y=y; t.z=z; return t; };
-#define FOUR_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x, T y, T z, T w) { VT t; t.x=x; t.y=y; t.z=z; t.w=w; return t; };
+#define ONE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x) { VT t; t.x = x; return t; };
+#define TWO_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x, T y) { VT t; t.x=x; t.y=y; return t; };
+#define THREE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x, T y, T z) { VT t; t.x=x; t.y=y; t.z=z; return t; };
+#define FOUR_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT [[hc]] [[cpu]] (T x, T y, T z, T w) { VT t; t.x=x; t.y=y; t.z=z; t.w=w; return t; };
 
 
 //signed:

From 75ecf38bb7521cc75fa59c051267c5af12deb882 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 2/2] hipLaunchKernel supports empty kernel with no arguments

Change-Id: Iaa46ced7fe0e3d72734d02250dfb48b98bef615d


[ROCm/clr commit: 275c5516b63a331acfe7539a61255659cad9eee9]
---
 projects/clr/hipamd/include/hcc_detail/hip_runtime.h  | 4 ++--
 projects/clr/hipamd/include/nvcc_detail/hip_runtime.h | 2 +-
 2 files changed, 3 insertions(+), 3 deletions(-)

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