Merge pull request #18 from AMDComputeLibraries/grid_launch_index

Grid launch index fix

include/hcc_detail/hip_runtime.h

@@ -428,6 +428,27 @@ __device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
 /**
  * Kernel launching
  */
+
+#if __hcc_workweek__ >= 16123
+
+#define hipThreadIdx_x (amp_get_local_id(0))
+#define hipThreadIdx_y (amp_get_local_id(1))
+#define hipThreadIdx_z (amp_get_local_id(2))
+
+#define hipBlockIdx_x  (hc_get_group_id(0))
+#define hipBlockIdx_y  (hc_get_group_id(1))
+#define hipBlockIdx_z  (hc_get_group_id(2))
+
+#define hipBlockDim_x  (amp_get_local_size(0))
+#define hipBlockDim_y  (amp_get_local_size(1))
+#define hipBlockDim_z  (amp_get_local_size(2))
+
+#define hipGridDim_x   (hc_get_num_groups(0))
+#define hipGridDim_y   (hc_get_num_groups(1))
+#define hipGridDim_z   (hc_get_num_groups(2))
+
+#else
+
 #define hipThreadIdx_x (amp_get_local_id(2))
 #define hipThreadIdx_y (amp_get_local_id(1))
 #define hipThreadIdx_z (amp_get_local_id(0))
@@ -444,7 +465,7 @@ __device__ inline float __dsqrt_rz(double x) {return hc::fast_math::sqrt(x); };
 #define hipGridDim_y   (hc_get_num_groups(1))
 #define hipGridDim_z   (hc_get_num_groups(0))
 
-
+#endif
 
 
 #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)