diff --git a/include/hcc_detail/hip_runtime.h b/include/hcc_detail/hip_runtime.h index b9a6981fee..79aca00f9a 100644 --- a/include/hcc_detail/hip_runtime.h +++ b/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)