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_sadd_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);
77 extern "C" __device__ __attribute__((convergent))
int __ockl_wgred_add_i32(
int a);
78 extern "C" __device__ __attribute__((convergent))
int __ockl_wgred_and_i32(
int a);
79 extern "C" __device__ __attribute__((convergent))
int __ockl_wgred_or_i32(
int a);
83 #define __local __attribute__((address_space(3)))
85 #ifdef __HIP_DEVICE_COMPILE__
86 __device__
inline static __local
void* __to_local(
unsigned x) {
return (__local
void*)x; }
87 #endif //__HIP_DEVICE_COMPILE__
89 #if defined(__HCC__) && (__hcc_major__ < 3) && (__hcc_minor__ < 3)
91 extern "C" __device__
void __llvm_fence_acq_sg(
void);
92 extern "C" __device__
void __llvm_fence_acq_wg(
void);
93 extern "C" __device__
void __llvm_fence_acq_dev(
void);
94 extern "C" __device__
void __llvm_fence_acq_sys(
void);
96 extern "C" __device__
void __llvm_fence_rel_sg(
void);
97 extern "C" __device__
void __llvm_fence_rel_wg(
void);
98 extern "C" __device__
void __llvm_fence_rel_dev(
void);
99 extern "C" __device__
void __llvm_fence_rel_sys(
void);
101 extern "C" __device__
void __llvm_fence_ar_sg(
void);
102 extern "C" __device__
void __llvm_fence_ar_wg(
void);
103 extern "C" __device__
void __llvm_fence_ar_dev(
void);
104 extern "C" __device__
void __llvm_fence_ar_sys(
void);
107 extern "C" __device__
void __llvm_fence_sc_sg(
void);
108 extern "C" __device__
void __llvm_fence_sc_wg(
void);
109 extern "C" __device__
void __llvm_fence_sc_dev(
void);
110 extern "C" __device__
void __llvm_fence_sc_sys(
void);
113 #define __CLK_LOCAL_MEM_FENCE 0x01
114 typedef unsigned __cl_mem_fence_flags;
116 typedef enum __memory_scope {
117 __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
118 __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
119 __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
120 __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
121 __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
125 typedef enum __memory_order
127 __memory_order_relaxed = __ATOMIC_RELAXED,
128 __memory_order_acquire = __ATOMIC_ACQUIRE,
129 __memory_order_release = __ATOMIC_RELEASE,
130 __memory_order_acq_rel = __ATOMIC_ACQ_REL,
131 __memory_order_seq_cst = __ATOMIC_SEQ_CST
135 extern "C" __device__
void
136 __atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);