diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index 0b7dd81d67..1938170ce4 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -734,13 +734,15 @@ int64_t __lanemask_lt() return ballot; } +__device__ inline void* __local_to_generic(void* p) { return p; } + #ifdef __HIP_DEVICE_COMPILE__ __device__ inline void* __get_dynamicgroupbaseptr() { // Get group segment base pointer. - return (char*)__local_to_generic(__to_local(__llvm_amdgcn_groupstaticsize())); + return (char*)__local_to_generic((void*)__to_local(__llvm_amdgcn_groupstaticsize())); } #else __device__ diff --git a/hipamd/include/hip/hcc_detail/device_library_decls.h b/hipamd/include/hip/hcc_detail/device_library_decls.h index a7e81a1968..53ad7595fe 100644 --- a/hipamd/include/hip/hcc_detail/device_library_decls.h +++ b/hipamd/include/hip/hcc_detail/device_library_decls.h @@ -53,8 +53,6 @@ extern "C" __device__ float __ocml_fmax_f32(float, float); __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 extern "C" __device__ void __llvm_fence_acq_sg(void); extern "C" __device__ void __llvm_fence_acq_wg(void);