31 #ifndef HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_LIBRARY_DECLS_H
32 #define HIP_INCLUDE_HIP_AMD_DETAIL_DEVICE_LIBRARY_DECLS_H
34 #if !defined(__HIPCC_RTC__)
38 typedef unsigned char uchar;
39 typedef unsigned short ushort;
40 typedef unsigned int uint;
41 typedef unsigned long ulong;
42 typedef unsigned long long ullong;
44 extern "C" __device__
__attribute__((
const)) bool __ockl_wfany_i32(
int);
45 extern "C" __device__ __attribute__((const))
bool __ockl_wfall_i32(
int);
46 extern "C" __device__ uint __ockl_activelane_u32(
void);
48 extern "C" __device__ __attribute__((const)) uint __ockl_mul24_u32(uint, uint);
49 extern "C" __device__ __attribute__((const))
int __ockl_mul24_i32(
int,
int);
50 extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint);
51 extern "C" __device__ __attribute__((const))
int __ockl_mul_hi_i32(
int,
int);
52 extern "C" __device__ __attribute__((const)) uint __ockl_sadd_u32(uint, uint, uint);
54 extern "C" __device__ __attribute__((const)) uchar __ockl_clz_u8(uchar);
55 extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort);
56 extern "C" __device__ __attribute__((const)) uint __ockl_clz_u32(uint);
57 extern "C" __device__ __attribute__((const)) uint64_t __ockl_clz_u64(uint64_t);
59 extern "C" __device__ __attribute__((const))
float __ocml_floor_f32(
float);
60 extern "C" __device__ __attribute__((const))
float __ocml_rint_f32(
float);
61 extern "C" __device__ __attribute__((const))
float __ocml_ceil_f32(
float);
62 extern "C" __device__ __attribute__((const))
float __ocml_trunc_f32(
float);
64 extern "C" __device__ __attribute__((const))
float __ocml_fmin_f32(
float,
float);
65 extern "C" __device__ __attribute__((const))
float __ocml_fmax_f32(
float,
float);
67 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtn_f32_f64(
double);
68 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtp_f32_f64(
double);
69 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtz_f32_f64(
double);
71 extern "C" __device__ __attribute__((const)) _Float16 __ocml_cvtrtn_f16_f32(
float);
72 extern "C" __device__ __attribute__((const)) _Float16 __ocml_cvtrtp_f16_f32(
float);
73 extern "C" __device__ __attribute__((const)) _Float16 __ocml_cvtrtz_f16_f32(
float);
75 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtn_f32_s32(
int);
76 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtp_f32_s32(
int);
77 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtz_f32_s32(
int);
78 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtn_f32_u32(uint32_t);
79 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtp_f32_u32(uint32_t);
80 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtz_f32_u32(uint32_t);
81 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtn_f32_s64(int64_t);
82 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtp_f32_s64(int64_t);
83 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtz_f32_s64(int64_t);
84 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtn_f32_u64(uint64_t);
85 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtp_f32_u64(uint64_t);
86 extern "C" __device__ __attribute__((const))
float __ocml_cvtrtz_f32_u64(uint64_t);
87 extern "C" __device__ __attribute__((const))
double __ocml_cvtrtn_f64_s64(int64_t);
88 extern "C" __device__ __attribute__((const))
double __ocml_cvtrtp_f64_s64(int64_t);
89 extern "C" __device__ __attribute__((const))
double __ocml_cvtrtz_f64_s64(int64_t);
90 extern "C" __device__ __attribute__((const))
double __ocml_cvtrtn_f64_u64(uint64_t);
91 extern "C" __device__ __attribute__((const))
double __ocml_cvtrtp_f64_u64(uint64_t);
92 extern "C" __device__ __attribute__((const))
double __ocml_cvtrtz_f64_u64(uint64_t);
94 extern "C" __device__ __attribute__((convergent))
void __ockl_gws_init(uint nwm1, uint rid);
95 extern "C" __device__ __attribute__((convergent))
void __ockl_gws_barrier(uint nwm1, uint rid);
97 extern "C" __device__ __attribute__((const)) uint32_t __ockl_lane_u32();
98 extern "C" __device__ __attribute__((const))
int __ockl_grid_is_valid(
void);
99 extern "C" __device__ __attribute__((convergent))
void __ockl_grid_sync(
void);
100 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_num_grids(
void);
101 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_grid_rank(
void);
102 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_size(
void);
103 extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank(
void);
104 extern "C" __device__ __attribute__((const))
int __ockl_multi_grid_is_valid(
void);
105 extern "C" __device__ __attribute__((convergent))
void __ockl_multi_grid_sync(
void);
107 extern "C" __device__
void __ockl_atomic_add_noret_f32(
float*,
float);
109 extern "C" __device__ __attribute__((convergent))
int __ockl_wgred_add_i32(
int a);
110 extern "C" __device__ __attribute__((convergent))
int __ockl_wgred_and_i32(
int a);
111 extern "C" __device__ __attribute__((convergent))
int __ockl_wgred_or_i32(
int a);
113 extern "C" __device__ uint64_t __ockl_fprintf_stderr_begin();
114 extern "C" __device__ uint64_t __ockl_fprintf_append_args(uint64_t msg_desc, uint32_t num_args,
115 uint64_t value0, uint64_t value1,
116 uint64_t value2, uint64_t value3,
117 uint64_t value4, uint64_t value5,
118 uint64_t value6, uint32_t is_last);
119 extern "C" __device__ uint64_t __ockl_fprintf_append_string_n(uint64_t msg_desc, const
char* data,
120 uint64_t length, uint32_t is_last);
123 #define __local __attribute__((address_space(3)))
125 #ifdef __HIP_DEVICE_COMPILE__
126 __device__
inline static __local
void* __to_local(
unsigned x) {
return (__local
void*)x; }
130 #define __CLK_LOCAL_MEM_FENCE 0x01
131 typedef unsigned __cl_mem_fence_flags;
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Definition: hip_fp16_math_fwd.h:57