From 07de5cb3342bfbc8142020179e6f3680938eb270 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 4 Jul 2018 17:56:50 +0000 Subject: [PATCH] Workaround cast warning of smaller integer type for __to_local For now, guard the __to_local function for device compile only since a local pointer should be same size as unsigned int on GPU compile. Also change to void* instead of char*. --- include/hip/hcc_detail/device_functions.h | 4 ++++ include/hip/hcc_detail/device_library_decls.h | 6 +++++- 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index aae0706033..32509ffffd 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -734,6 +734,8 @@ int64_t __lanemask_lt() return ballot; } +#ifdef __HIP_DEVICE_COMPILE__ + __device__ inline void* __get_dynamicgroupbaseptr() @@ -748,6 +750,8 @@ void *__amdgcn_get_dynamicgroupbaseptr() { return __get_dynamicgroupbaseptr(); } +#endif // __HIP_DEVICE_COMPILE__ + #endif // __HCC_OR_HIP_CLANG__ #ifdef __HCC__ diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index fba2d53e2e..a636c2c950 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -45,7 +45,11 @@ extern "C" __device__ float __ocml_trunc_f32(float); // Introduce local address space #define __local __attribute__((address_space(3))) -__device__ inline static __local char* __to_local(unsigned x) { return (__local char*)x; } + +#ifdef __HIP_DEVICE_COMPILE__ +__device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; } +#endif //__HIP_DEVICE_COMPILE__ + extern "C" __device__ void* __local_to_generic(__local void* p); // __llvm_fence* functions from device-libs/irif/src/fence.ll