28 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_LIBRARY_DECLS_H 29 #define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_LIBRARY_DECLS_H 33 typedef unsigned char uchar;
34 typedef unsigned short ushort;
35 typedef unsigned int uint;
36 typedef unsigned long ulong;
37 typedef unsigned long long ullong;
39 extern "C" __device__ __attribute__((
const)) bool __ockl_wfany_i32(
int);
40 extern "C" __device__ __attribute__((const))
bool __ockl_wfall_i32(
int);
41 extern "C" __device__ uint __ockl_activelane_u32(
void);
43 extern "C" __device__ __attribute__((const)) uint __ockl_mul24_u32(uint, uint);
44 extern "C" __device__ __attribute__((const))
int __ockl_mul24_i32(
int,
int);
45 extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint);
46 extern "C" __device__ __attribute__((const))
int __ockl_mul_hi_i32(
int,
int);
47 extern "C" __device__ __attribute__((const)) uint __ockl_sad_u32(uint, uint, uint);
49 extern "C" __device__ __attribute__((const)) uchar __ockl_clz_u8(uchar);
50 extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort);
51 extern "C" __device__ __attribute__((const)) uint __ockl_clz_u32(uint);
52 extern "C" __device__ __attribute__((const)) ullong __ockl_clz_u64(ullong);
54 extern "C" __device__ __attribute__((const))
float __ocml_floor_f32(
float);
55 extern "C" __device__ __attribute__((const))
float __ocml_rint_f32(
float);
56 extern "C" __device__ __attribute__((const))
float __ocml_ceil_f32(
float);
57 extern "C" __device__ __attribute__((const))
float __ocml_trunc_f32(
float);
59 extern "C" __device__ __attribute__((const))
float __ocml_fmin_f32(
float,
float);
60 extern "C" __device__ __attribute__((const))
float __ocml_fmax_f32(
float,
float);
62 extern "C" __device__ __attribute__((convergent))
void __ockl_gws_init(uint nwm1, uint rid);
63 extern "C" __device__ __attribute__((convergent))
void __ockl_gws_barrier(uint nwm1, uint rid);
67 #define __local __attribute__((address_space(3))) 69 #ifdef __HIP_DEVICE_COMPILE__ 70 __device__
inline static __local
void* __to_local(
unsigned x) {
return (__local
void*)x; }
71 #endif //__HIP_DEVICE_COMPILE__ 73 #if defined(__HCC__) && (__hcc_minor__ < 3) 75 extern "C" __device__
void __llvm_fence_acq_sg(
void);
76 extern "C" __device__
void __llvm_fence_acq_wg(
void);
77 extern "C" __device__
void __llvm_fence_acq_dev(
void);
78 extern "C" __device__
void __llvm_fence_acq_sys(
void);
80 extern "C" __device__
void __llvm_fence_rel_sg(
void);
81 extern "C" __device__
void __llvm_fence_rel_wg(
void);
82 extern "C" __device__
void __llvm_fence_rel_dev(
void);
83 extern "C" __device__
void __llvm_fence_rel_sys(
void);
85 extern "C" __device__
void __llvm_fence_ar_sg(
void);
86 extern "C" __device__
void __llvm_fence_ar_wg(
void);
87 extern "C" __device__
void __llvm_fence_ar_dev(
void);
88 extern "C" __device__
void __llvm_fence_ar_sys(
void);
91 extern "C" __device__
void __llvm_fence_sc_sg(
void);
92 extern "C" __device__
void __llvm_fence_sc_wg(
void);
93 extern "C" __device__
void __llvm_fence_sc_dev(
void);
94 extern "C" __device__
void __llvm_fence_sc_sys(
void);
97 #define __CLK_LOCAL_MEM_FENCE 0x01 98 typedef unsigned __cl_mem_fence_flags;
100 typedef enum __memory_scope {
101 __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
102 __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
103 __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
104 __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
105 __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
109 typedef enum __memory_order
111 __memory_order_relaxed = __ATOMIC_RELAXED,
112 __memory_order_acquire = __ATOMIC_ACQUIRE,
113 __memory_order_release = __ATOMIC_RELEASE,
114 __memory_order_acq_rel = __ATOMIC_ACQ_REL,
115 __memory_order_seq_cst = __ATOMIC_SEQ_CST
119 extern "C" __device__
void 120 __atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);