Add HCC compatibility mode

bin/hipcc

@@ -57,6 +57,7 @@ $HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hi
 $HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'};
 $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'};
 $DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'};
+$HIP_CLANG_HCC_COMPAT_MODE=$ENV{'HIP_CLANG_HCC_COMPAT_MODE'}; # HCC compatibility mode
 
 #---
 # Read .hipInfo
@@ -158,6 +159,10 @@ if ($HIP_PLATFORM eq "clang") {
 
     $HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH";
     $HIPLDFLAGS .= " --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_LIB_PATH -Wl,--rpath=$HIP_LIB_PATH -lhip_hcc";
+    if ($HIP_CLANG_HCC_COMPAT_MODE) {
+        ## Allow __fp16 as function parameter and return type.
+        $HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1";
+    }
 } elsif ($HIP_PLATFORM eq "hcc") {
     $HIP_INCLUDE_PATH = "$HIP_PATH/include";
     $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";

include/hip/hcc_detail/hip_runtime.h

@@ -438,6 +438,26 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
 #define hipGridDim_y gridDim.y
 #define hipGridDim_z gridDim.z
 
+#if __HIP_HCC_COMPAT_MODE__
+// Define HCC work item functions in terms of HIP builtin variables.
+#pragma push_macro("__DEFINE_HCC_FUNC")
+#define __DEFINE_HCC_FUNC(hc_fun,hip_var) \
+inline __device__ __attribute__((always_inline)) uint hc_get_##hc_fun(uint i) { \
+  if (i==0) \
+    return hip_var.x; \
+  else if(i==1) \
+    return hip_var.y; \
+  else \
+    return hip_var.z; \
+}
+
+__DEFINE_HCC_FUNC(workitem_id, threadIdx)
+__DEFINE_HCC_FUNC(group_id, blockIdx)
+__DEFINE_HCC_FUNC(group_size, blockDim)
+__DEFINE_HCC_FUNC(num_groups, gridDim)
+#pragma pop_macro("__DEFINE_HCC_FUNC")
+#endif
+
 // Support std::complex.
 #pragma push_macro("__CUDA__")
 #define __CUDA__