28 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_LIBRARY_DECLS_H
29 #define HIP_INCLUDE_HIP_AMD_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__
90 #define __CLK_LOCAL_MEM_FENCE 0x01
91 typedef unsigned __cl_mem_fence_flags;
93 typedef enum __memory_scope {
94 __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM,
95 __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP,
96 __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE,
97 __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES,
98 __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP
102 typedef enum __memory_order
104 __memory_order_relaxed = __ATOMIC_RELAXED,
105 __memory_order_acquire = __ATOMIC_ACQUIRE,
106 __memory_order_release = __ATOMIC_RELEASE,
107 __memory_order_acq_rel = __ATOMIC_ACQ_REL,
108 __memory_order_seq_cst = __ATOMIC_SEQ_CST
112 extern "C" __device__
void
113 __atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope);