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);
65 extern "C" __device__ __attribute__((const)) uint32_t __ockl_lane_u32();
66 extern "C" __device__ __attribute__((const))
int __ockl_grid_is_valid(
void);
67 extern "C" __device__ __attribute__((convergent))
void __ockl_grid_sync(
void);
68 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_num_grids(
void);
69 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_grid_rank(
void);
70 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_size(
void);
71 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank(
void);
72 extern "C" __device__ __attribute__((const))
int __ockl_multi_grid_is_valid(
void);
73 extern "C" __device__ __attribute__((convergent))
void __ockl_multi_grid_sync(
void);
75 extern "C" __device__
void __ockl_atomic_add_noret_f32(
float*,
float);
78 #define __local __attribute__((address_space(3))) 80 #ifdef __HIP_DEVICE_COMPILE__ 81 __device__
inline static __local
void* __to_local(
unsigned x) {
return (__local
void*)x; }
82 #endif //__HIP_DEVICE_COMPILE__ 84 #if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3) 86 extern "C" __device__
void __llvm_fence_acq_sg(
void);
87 extern "C" __device__
void __llvm_fence_acq_wg(
void);
88 extern "C" __device__
void __llvm_fence_acq_dev(
void);
89 extern "C" __device__
void __llvm_fence_acq_sys(
void);
91 extern "C" __device__
void __llvm_fence_rel_sg(
void);
92 extern "C" __device__
void __llvm_fence_rel_wg(
void);
93 extern "C" __device__
void __llvm_fence_rel_dev(
void);
94 extern "C" __device__
void __llvm_fence_rel_sys(
void);
96 extern "C" __device__
void __llvm_fence_ar_sg(
void);
97 extern "C" __device__
void __llvm_fence_ar_wg(
void);
98 extern "C" __device__
void __llvm_fence_ar_dev(
void);
99 extern "C" __device__
void __llvm_fence_ar_sys(
void);
102 extern "C" __device__
void __llvm_fence_sc_sg(
void);
103 extern "C" __device__
void __llvm_fence_sc_wg(
void);
104 extern "C" __device__
void __llvm_fence_sc_dev(
void);
105 extern "C" __device__
void __llvm_fence_sc_sys(
void);
108 #define __CLK_LOCAL_MEM_FENCE 0x01 109 typedef unsigned __cl_mem_fence_flags;
111 typedef enum __memory_scope {
112 __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
113 __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
114 __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
115 __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
116 __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
120 typedef enum __memory_order
122 __memory_order_relaxed = __ATOMIC_RELAXED,
123 __memory_order_acquire = __ATOMIC_ACQUIRE,
124 __memory_order_release = __ATOMIC_RELEASE,
125 __memory_order_acq_rel = __ATOMIC_ACQ_REL,
126 __memory_order_seq_cst = __ATOMIC_SEQ_CST
130 extern "C" __device__
void 131 __atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);