1 /*===---- __clang_hip_libdevice_declares.h - HIP device library decls -------===
2  *
3  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4  * See https://llvm.org/LICENSE.txt for license information.
5  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6  *
7  *===-----------------------------------------------------------------------===
8  */
9 
10 #ifndef __CLANG_HIP_LIBDEVICE_DECLARES_H__
11 #define __CLANG_HIP_LIBDEVICE_DECLARES_H__
12 
13 #ifdef __cplusplus
14 extern "C" {
15 #endif
16 
17 // BEGIN FLOAT
18 __device__ __attribute__((const)) float __ocml_acos_f32(float);
19 __device__ __attribute__((pure)) float __ocml_acosh_f32(float);
20 __device__ __attribute__((const)) float __ocml_asin_f32(float);
21 __device__ __attribute__((pure)) float __ocml_asinh_f32(float);
22 __device__ __attribute__((const)) float __ocml_atan2_f32(float, float);
23 __device__ __attribute__((const)) float __ocml_atan_f32(float);
24 __device__ __attribute__((pure)) float __ocml_atanh_f32(float);
25 __device__ __attribute__((pure)) float __ocml_cbrt_f32(float);
26 __device__ __attribute__((const)) float __ocml_ceil_f32(float);
27 __device__ __attribute__((const)) __device__ float __ocml_copysign_f32(float,
28                                                                        float);
29 __device__ float __ocml_cos_f32(float);
30 __device__ float __ocml_native_cos_f32(float);
31 __device__ __attribute__((pure)) __device__ float __ocml_cosh_f32(float);
32 __device__ float __ocml_cospi_f32(float);
33 __device__ float __ocml_i0_f32(float);
34 __device__ float __ocml_i1_f32(float);
35 __device__ __attribute__((pure)) float __ocml_erfc_f32(float);
36 __device__ __attribute__((pure)) float __ocml_erfcinv_f32(float);
37 __device__ __attribute__((pure)) float __ocml_erfcx_f32(float);
38 __device__ __attribute__((pure)) float __ocml_erf_f32(float);
39 __device__ __attribute__((pure)) float __ocml_erfinv_f32(float);
40 __device__ __attribute__((pure)) float __ocml_exp10_f32(float);
41 __device__ __attribute__((pure)) float __ocml_native_exp10_f32(float);
42 __device__ __attribute__((pure)) float __ocml_exp2_f32(float);
43 __device__ __attribute__((pure)) float __ocml_exp_f32(float);
44 __device__ __attribute__((pure)) float __ocml_native_exp_f32(float);
45 __device__ __attribute__((pure)) float __ocml_expm1_f32(float);
46 __device__ __attribute__((const)) float __ocml_fabs_f32(float);
47 __device__ __attribute__((const)) float __ocml_fdim_f32(float, float);
48 __device__ __attribute__((const)) float __ocml_floor_f32(float);
49 __device__ __attribute__((const)) float __ocml_fma_f32(float, float, float);
50 __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
51 __device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
52 __device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
53                                                                    float);
54 __device__ float __ocml_frexp_f32(float,
55                                   __attribute__((address_space(5))) int *);
56 __device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
57 __device__ __attribute__((const)) int __ocml_ilogb_f32(float);
58 __device__ __attribute__((const)) int __ocml_isfinite_f32(float);
59 __device__ __attribute__((const)) int __ocml_isinf_f32(float);
60 __device__ __attribute__((const)) int __ocml_isnan_f32(float);
61 __device__ float __ocml_j0_f32(float);
62 __device__ float __ocml_j1_f32(float);
63 __device__ __attribute__((const)) float __ocml_ldexp_f32(float, int);
64 __device__ float __ocml_lgamma_f32(float);
65 __device__ __attribute__((pure)) float __ocml_log10_f32(float);
66 __device__ __attribute__((pure)) float __ocml_native_log10_f32(float);
67 __device__ __attribute__((pure)) float __ocml_log1p_f32(float);
68 __device__ __attribute__((pure)) float __ocml_log2_f32(float);
69 __device__ __attribute__((pure)) float __ocml_native_log2_f32(float);
70 __device__ __attribute__((const)) float __ocml_logb_f32(float);
71 __device__ __attribute__((pure)) float __ocml_log_f32(float);
72 __device__ __attribute__((pure)) float __ocml_native_log_f32(float);
73 __device__ float __ocml_modf_f32(float,
74                                  __attribute__((address_space(5))) float *);
75 __device__ __attribute__((const)) float __ocml_nearbyint_f32(float);
76 __device__ __attribute__((const)) float __ocml_nextafter_f32(float, float);
77 __device__ __attribute__((const)) float __ocml_len3_f32(float, float, float);
78 __device__ __attribute__((const)) float __ocml_len4_f32(float, float, float,
79                                                         float);
80 __device__ __attribute__((pure)) float __ocml_ncdf_f32(float);
81 __device__ __attribute__((pure)) float __ocml_ncdfinv_f32(float);
82 __device__ __attribute__((pure)) float __ocml_pow_f32(float, float);
83 __device__ __attribute__((pure)) float __ocml_pown_f32(float, int);
84 __device__ __attribute__((pure)) float __ocml_rcbrt_f32(float);
85 __device__ __attribute__((const)) float __ocml_remainder_f32(float, float);
86 __device__ float __ocml_remquo_f32(float, float,
87                                    __attribute__((address_space(5))) int *);
88 __device__ __attribute__((const)) float __ocml_rhypot_f32(float, float);
89 __device__ __attribute__((const)) float __ocml_rint_f32(float);
90 __device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float);
91 __device__ __attribute__((const)) float __ocml_rlen4_f32(float, float, float,
92                                                          float);
93 __device__ __attribute__((const)) float __ocml_round_f32(float);
94 __device__ __attribute__((pure)) float __ocml_rsqrt_f32(float);
95 __device__ __attribute__((const)) float __ocml_scalb_f32(float, float);
96 __device__ __attribute__((const)) float __ocml_scalbn_f32(float, int);
97 __device__ __attribute__((const)) int __ocml_signbit_f32(float);
98 __device__ float __ocml_sincos_f32(float,
99                                    __attribute__((address_space(5))) float *);
100 __device__ float __ocml_sincospi_f32(float,
101                                      __attribute__((address_space(5))) float *);
102 __device__ float __ocml_sin_f32(float);
103 __device__ float __ocml_native_sin_f32(float);
104 __device__ __attribute__((pure)) float __ocml_sinh_f32(float);
105 __device__ float __ocml_sinpi_f32(float);
106 __device__ __attribute__((const)) float __ocml_sqrt_f32(float);
107 __device__ __attribute__((const)) float __ocml_native_sqrt_f32(float);
108 __device__ float __ocml_tan_f32(float);
109 __device__ __attribute__((pure)) float __ocml_tanh_f32(float);
110 __device__ float __ocml_tgamma_f32(float);
111 __device__ __attribute__((const)) float __ocml_trunc_f32(float);
112 __device__ float __ocml_y0_f32(float);
113 __device__ float __ocml_y1_f32(float);
114 
115 // BEGIN INTRINSICS
116 __device__ __attribute__((const)) float __ocml_add_rte_f32(float, float);
117 __device__ __attribute__((const)) float __ocml_add_rtn_f32(float, float);
118 __device__ __attribute__((const)) float __ocml_add_rtp_f32(float, float);
119 __device__ __attribute__((const)) float __ocml_add_rtz_f32(float, float);
120 __device__ __attribute__((const)) float __ocml_sub_rte_f32(float, float);
121 __device__ __attribute__((const)) float __ocml_sub_rtn_f32(float, float);
122 __device__ __attribute__((const)) float __ocml_sub_rtp_f32(float, float);
123 __device__ __attribute__((const)) float __ocml_sub_rtz_f32(float, float);
124 __device__ __attribute__((const)) float __ocml_mul_rte_f32(float, float);
125 __device__ __attribute__((const)) float __ocml_mul_rtn_f32(float, float);
126 __device__ __attribute__((const)) float __ocml_mul_rtp_f32(float, float);
127 __device__ __attribute__((const)) float __ocml_mul_rtz_f32(float, float);
128 __device__ __attribute__((const)) float __ocml_div_rte_f32(float, float);
129 __device__ __attribute__((const)) float __ocml_div_rtn_f32(float, float);
130 __device__ __attribute__((const)) float __ocml_div_rtp_f32(float, float);
131 __device__ __attribute__((const)) float __ocml_div_rtz_f32(float, float);
132 __device__ __attribute__((const)) float __ocml_sqrt_rte_f32(float);
133 __device__ __attribute__((const)) float __ocml_sqrt_rtn_f32(float);
134 __device__ __attribute__((const)) float __ocml_sqrt_rtp_f32(float);
135 __device__ __attribute__((const)) float __ocml_sqrt_rtz_f32(float);
136 __device__ __attribute__((const)) float __ocml_fma_rte_f32(float, float, float);
137 __device__ __attribute__((const)) float __ocml_fma_rtn_f32(float, float, float);
138 __device__ __attribute__((const)) float __ocml_fma_rtp_f32(float, float, float);
139 __device__ __attribute__((const)) float __ocml_fma_rtz_f32(float, float, float);
140 
141 __device__ __attribute__((const)) float
142 __llvm_amdgcn_cos_f32(float) __asm("llvm.amdgcn.cos.f32");
143 __device__ __attribute__((const)) float
144 __llvm_amdgcn_rcp_f32(float) __asm("llvm.amdgcn.rcp.f32");
145 __device__ __attribute__((const)) float
146 __llvm_amdgcn_rsq_f32(float) __asm("llvm.amdgcn.rsq.f32");
147 __device__ __attribute__((const)) float
148 __llvm_amdgcn_sin_f32(float) __asm("llvm.amdgcn.sin.f32");
149 // END INTRINSICS
150 // END FLOAT
151 
152 // BEGIN DOUBLE
153 __device__ __attribute__((const)) double __ocml_acos_f64(double);
154 __device__ __attribute__((pure)) double __ocml_acosh_f64(double);
155 __device__ __attribute__((const)) double __ocml_asin_f64(double);
156 __device__ __attribute__((pure)) double __ocml_asinh_f64(double);
157 __device__ __attribute__((const)) double __ocml_atan2_f64(double, double);
158 __device__ __attribute__((const)) double __ocml_atan_f64(double);
159 __device__ __attribute__((pure)) double __ocml_atanh_f64(double);
160 __device__ __attribute__((pure)) double __ocml_cbrt_f64(double);
161 __device__ __attribute__((const)) double __ocml_ceil_f64(double);
162 __device__ __attribute__((const)) double __ocml_copysign_f64(double, double);
163 __device__ double __ocml_cos_f64(double);
164 __device__ __attribute__((pure)) double __ocml_cosh_f64(double);
165 __device__ double __ocml_cospi_f64(double);
166 __device__ double __ocml_i0_f64(double);
167 __device__ double __ocml_i1_f64(double);
168 __device__ __attribute__((pure)) double __ocml_erfc_f64(double);
169 __device__ __attribute__((pure)) double __ocml_erfcinv_f64(double);
170 __device__ __attribute__((pure)) double __ocml_erfcx_f64(double);
171 __device__ __attribute__((pure)) double __ocml_erf_f64(double);
172 __device__ __attribute__((pure)) double __ocml_erfinv_f64(double);
173 __device__ __attribute__((pure)) double __ocml_exp10_f64(double);
174 __device__ __attribute__((pure)) double __ocml_exp2_f64(double);
175 __device__ __attribute__((pure)) double __ocml_exp_f64(double);
176 __device__ __attribute__((pure)) double __ocml_expm1_f64(double);
177 __device__ __attribute__((const)) double __ocml_fabs_f64(double);
178 __device__ __attribute__((const)) double __ocml_fdim_f64(double, double);
179 __device__ __attribute__((const)) double __ocml_floor_f64(double);
180 __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
181 __device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
182 __device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
183 __device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
184 __device__ double __ocml_frexp_f64(double,
185                                    __attribute__((address_space(5))) int *);
186 __device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
187 __device__ __attribute__((const)) int __ocml_ilogb_f64(double);
188 __device__ __attribute__((const)) int __ocml_isfinite_f64(double);
189 __device__ __attribute__((const)) int __ocml_isinf_f64(double);
190 __device__ __attribute__((const)) int __ocml_isnan_f64(double);
191 __device__ double __ocml_j0_f64(double);
192 __device__ double __ocml_j1_f64(double);
193 __device__ __attribute__((const)) double __ocml_ldexp_f64(double, int);
194 __device__ double __ocml_lgamma_f64(double);
195 __device__ __attribute__((pure)) double __ocml_log10_f64(double);
196 __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
197 __device__ __attribute__((pure)) double __ocml_log2_f64(double);
198 __device__ __attribute__((const)) double __ocml_logb_f64(double);
199 __device__ __attribute__((pure)) double __ocml_log_f64(double);
200 __device__ double __ocml_modf_f64(double,
201                                   __attribute__((address_space(5))) double *);
202 __device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
203 __device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
204 __device__ __attribute__((const)) double __ocml_len3_f64(double, double,
205                                                          double);
206 __device__ __attribute__((const)) double __ocml_len4_f64(double, double, double,
207                                                          double);
208 __device__ __attribute__((pure)) double __ocml_ncdf_f64(double);
209 __device__ __attribute__((pure)) double __ocml_ncdfinv_f64(double);
210 __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
211 __device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
212 __device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
213 __device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
214 __device__ double __ocml_remquo_f64(double, double,
215                                     __attribute__((address_space(5))) int *);
216 __device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
217 __device__ __attribute__((const)) double __ocml_rint_f64(double);
218 __device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
219                                                           double);
220 __device__ __attribute__((const)) double __ocml_rlen4_f64(double, double,
221                                                           double, double);
222 __device__ __attribute__((const)) double __ocml_round_f64(double);
223 __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
224 __device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
225 __device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
226 __device__ __attribute__((const)) int __ocml_signbit_f64(double);
227 __device__ double __ocml_sincos_f64(double,
228                                     __attribute__((address_space(5))) double *);
229 __device__ double
230 __ocml_sincospi_f64(double, __attribute__((address_space(5))) double *);
231 __device__ double __ocml_sin_f64(double);
232 __device__ __attribute__((pure)) double __ocml_sinh_f64(double);
233 __device__ double __ocml_sinpi_f64(double);
234 __device__ __attribute__((const)) double __ocml_sqrt_f64(double);
235 __device__ double __ocml_tan_f64(double);
236 __device__ __attribute__((pure)) double __ocml_tanh_f64(double);
237 __device__ double __ocml_tgamma_f64(double);
238 __device__ __attribute__((const)) double __ocml_trunc_f64(double);
239 __device__ double __ocml_y0_f64(double);
240 __device__ double __ocml_y1_f64(double);
241 
242 // BEGIN INTRINSICS
243 __device__ __attribute__((const)) double __ocml_add_rte_f64(double, double);
244 __device__ __attribute__((const)) double __ocml_add_rtn_f64(double, double);
245 __device__ __attribute__((const)) double __ocml_add_rtp_f64(double, double);
246 __device__ __attribute__((const)) double __ocml_add_rtz_f64(double, double);
247 __device__ __attribute__((const)) double __ocml_sub_rte_f64(double, double);
248 __device__ __attribute__((const)) double __ocml_sub_rtn_f64(double, double);
249 __device__ __attribute__((const)) double __ocml_sub_rtp_f64(double, double);
250 __device__ __attribute__((const)) double __ocml_sub_rtz_f64(double, double);
251 __device__ __attribute__((const)) double __ocml_mul_rte_f64(double, double);
252 __device__ __attribute__((const)) double __ocml_mul_rtn_f64(double, double);
253 __device__ __attribute__((const)) double __ocml_mul_rtp_f64(double, double);
254 __device__ __attribute__((const)) double __ocml_mul_rtz_f64(double, double);
255 __device__ __attribute__((const)) double __ocml_div_rte_f64(double, double);
256 __device__ __attribute__((const)) double __ocml_div_rtn_f64(double, double);
257 __device__ __attribute__((const)) double __ocml_div_rtp_f64(double, double);
258 __device__ __attribute__((const)) double __ocml_div_rtz_f64(double, double);
259 __device__ __attribute__((const)) double __ocml_sqrt_rte_f64(double);
260 __device__ __attribute__((const)) double __ocml_sqrt_rtn_f64(double);
261 __device__ __attribute__((const)) double __ocml_sqrt_rtp_f64(double);
262 __device__ __attribute__((const)) double __ocml_sqrt_rtz_f64(double);
263 __device__ __attribute__((const)) double __ocml_fma_rte_f64(double, double,
264                                                             double);
265 __device__ __attribute__((const)) double __ocml_fma_rtn_f64(double, double,
266                                                             double);
267 __device__ __attribute__((const)) double __ocml_fma_rtp_f64(double, double,
268                                                             double);
269 __device__ __attribute__((const)) double __ocml_fma_rtz_f64(double, double,
270                                                             double);
271 
272 __device__ __attribute__((const)) double
273 __llvm_amdgcn_rcp_f64(double) __asm("llvm.amdgcn.rcp.f64");
274 __device__ __attribute__((const)) double
275 __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64");
276 
277 __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
278 __device__ _Float16 __ocml_cos_f16(_Float16);
279 __device__ __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
280 __device__ __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
281 __device__ __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
282 __device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
283 __device__ __attribute__((const)) _Float16 __ocml_fma_f16(_Float16, _Float16,
284                                                           _Float16);
285 __device__ __attribute__((const)) _Float16 __ocml_fabs_f16(_Float16);
286 __device__ __attribute__((const)) int __ocml_isinf_f16(_Float16);
287 __device__ __attribute__((const)) int __ocml_isnan_f16(_Float16);
288 __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
289 __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
290 __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
291 __device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
292 __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
293 __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
294 __device__ _Float16 __ocml_sin_f16(_Float16);
295 __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
296 __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
297 __device__ __attribute__((pure)) _Float16 __ocml_pown_f16(_Float16, int);
298 
299 typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
300 typedef short __2i16 __attribute__((ext_vector_type(2)));
301 
302 __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b,
303                                                      float c, bool s);
304 __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
305 __device__ __attribute__((const)) __2f16 __ocml_fabs_2f16(__2f16);
306 __device__ __2f16 __ocml_cos_2f16(__2f16);
307 __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
308 __device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
309 __device__ __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
310 __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
311 __device__ __attribute__((const))
312 __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
313 __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
314 __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
315 __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
316 __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
317 __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
318 __device__ inline __2f16
__llvm_amdgcn_rcp_2f16(__2f16 __x)319 __llvm_amdgcn_rcp_2f16(__2f16 __x) // Not currently exposed by ROCDL.
320 {
321   return (__2f16)(__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y));
322 }
323 __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
324 __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
325 __device__ __2f16 __ocml_sin_2f16(__2f16);
326 __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
327 __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
328 __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
329 
330 #ifdef __cplusplus
331 } // extern "C"
332 #endif
333 
334 #endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__
335