From 303a8be7d38dfb0a1ce7c0f84df487a51045b851 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*. [ROCm/clr commit: ce9ca744dc8d4cae0c0e94065e6442c18c010df8] --- .../clr/hipamd/include/hip/hcc_detail/device_functions.h | 4 ++++ .../hipamd/include/hip/hcc_detail/device_library_decls.h | 6 +++++- 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index aae0706033..32509ffffd 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/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/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h b/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h index fba2d53e2e..a636c2c950 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_library_decls.h +++ b/projects/clr/hipamd/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