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*.
Bu işleme şunda yer alıyor:
@@ -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__
|
||||
|
||||
@@ -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
|
||||
|
||||
Yeni konuda referans
Bir kullanıcı engelle