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