From 62da82d48cad26cb894f6de23e5af41c45c151b2 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 28 Apr 2016 13:43:32 -0500 Subject: [PATCH] Properly pair #if/#endif [ROCm/clr commit: c60614eca5354b46b0e10ff703cfb1cb085b172b] --- projects/clr/hipamd/include/hcc_detail/hip_ldg.h | 4 ++-- projects/clr/hipamd/include/hcc_detail/hip_runtime.h | 11 +++++------ 2 files changed, 7 insertions(+), 8 deletions(-) diff --git a/projects/clr/hipamd/include/hcc_detail/hip_ldg.h b/projects/clr/hipamd/include/hcc_detail/hip_ldg.h index 2a041cf9cb..71d39990ed 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_ldg.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_ldg.h @@ -100,9 +100,9 @@ __device__ double4 __ldg(const double4* ); #endif // __hcc_workweek__ -#endif // __HCC__ +#endif // USE_LDG -#endif +#endif // __HCC__ #endif // HIP_LDG_H diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime.h index b20fd1e7a9..30799bcd06 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime.h @@ -96,7 +96,7 @@ extern int HIP_TRACE_API; #define __HIP_ARCH_HAS_3DGRID__ (1) #define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0) -#endif +#endif /* Device feature flags */ //TODO-HCC this is currently ignored by HCC target of HIP @@ -298,7 +298,7 @@ extern const int warpSize; #define clock_t long long int __device__ long long int clock64(); __device__ clock_t clock(); -#if __cplusplus + //atomicAdd() __device__ int atomicAdd(int* address, int val); __device__ unsigned int atomicAdd(unsigned int* address, @@ -401,7 +401,6 @@ __device__ unsigned long long int __ballot( int input); // warp shuffle functions #ifdef __cplusplus - __device__ int __shfl(int input, int lane, int width=warpSize); __device__ int __shfl_up(int input, unsigned int lane_delta, int width=warpSize); __device__ int __shfl_down(int input, unsigned int lane_delta, int width=warpSize); @@ -448,6 +447,7 @@ __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)) @@ -484,12 +484,12 @@ __device__ float __dsqrt_rz(double x); #define hipGridDim_y (hc_get_num_groups(1)) #define hipGridDim_z (hc_get_num_groups(0)) -#endif +#endif // __hcc_workweek__ check #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) #define HIP_KERNEL_NAME(...) __VA_ARGS__ -#include + #ifdef __HCC_CPP__ hipStream_t ihipPreLaunchKernel(hipStream_t stream, grid_launch_parm *lp); void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp); @@ -573,4 +573,3 @@ do {\ #endif -#endif