Merge branch 'amd-master' into amd-develop

[ROCm/clr commit: d3a0de99a1]

projects/clr/hipamd/include/hcc_detail/hip_runtime.h

@@ -549,7 +549,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)
 
@@ -573,7 +573,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 */

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:

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)