Add workaround __local_to_generic

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__

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);