From 58de525d41abb87d89979d2e20947ece43c1272e Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Mon, 20 Jun 2016 23:28:45 -0500
Subject: [PATCH] Switch to hc_* coordinate builtins (replace amp_)

Change-Id: I0a8871f0c9f047eb45a7391fd032100af2bbd4e0


[ROCm/hip commit: 100a744d926ee6b0ace439c0aa07b2803d17ded4]
---
 projects/hip/include/hcc_detail/hip_runtime.h | 34 ++++---------------
 1 file changed, 6 insertions(+), 28 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_runtime.h b/projects/hip/include/hcc_detail/hip_runtime.h
index 3d0bd7b1f2..066db4bbec 100644
--- a/projects/hip/include/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hcc_detail/hip_runtime.h
@@ -485,45 +485,23 @@ __device__ float __dsqrt_rz(double x);
  * Kernel launching
  */
 
-// Choose correct polarity of xyz/zyx ordering:
-#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 hipThreadIdx_x (hc_get_workitem_id(0))
+#define hipThreadIdx_y (hc_get_workitem_id(1))
+#define hipThreadIdx_z (hc_get_workitem_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 hipBlockDim_x  (hc_get_group_size(0))
+#define hipBlockDim_y  (hc_get_group_size(1))
+#define hipBlockDim_z  (hc_get_group_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))
-
-#define hipBlockIdx_x  (hc_get_group_id(2))
-#define hipBlockIdx_y  (hc_get_group_id(1))
-#define hipBlockIdx_z  (hc_get_group_id(0))
-
-#define hipBlockDim_x  (amp_get_local_size(2))
-#define hipBlockDim_y  (amp_get_local_size(1))
-#define hipBlockDim_z  (amp_get_local_size(0))
-
-#define hipGridDim_x   (hc_get_num_groups(2))
-#define hipGridDim_y   (hc_get_num_groups(1))
-#define hipGridDim_z   (hc_get_num_groups(0))
-
-#endif // __hcc_workweek__ check
-
 #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
 
 #define HIP_KERNEL_NAME(...) __VA_ARGS__