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__ inline __attribute__((const)) float
__llvm_amdgcn_cos_f32(float __x)142 __llvm_amdgcn_cos_f32(float __x) {
143 return __builtin_amdgcn_cosf(__x);
144 }
145 __device__ inline __attribute__((const)) float
__llvm_amdgcn_rcp_f32(float __x)146 __llvm_amdgcn_rcp_f32(float __x) {
147 return __builtin_amdgcn_rcpf(__x);
148 }
149 __device__ inline __attribute__((const)) float
__llvm_amdgcn_rsq_f32(float __x)150 __llvm_amdgcn_rsq_f32(float __x) {
151 return __builtin_amdgcn_rsqf(__x);
152 }
153 __device__ inline __attribute__((const)) float
__llvm_amdgcn_sin_f32(float __x)154 __llvm_amdgcn_sin_f32(float __x) {
155 return __builtin_amdgcn_sinf(__x);
156 }
157 // END INTRINSICS
158 // END FLOAT
159
160 // BEGIN DOUBLE
161 __device__ __attribute__((const)) double __ocml_acos_f64(double);
162 __device__ __attribute__((pure)) double __ocml_acosh_f64(double);
163 __device__ __attribute__((const)) double __ocml_asin_f64(double);
164 __device__ __attribute__((pure)) double __ocml_asinh_f64(double);
165 __device__ __attribute__((const)) double __ocml_atan2_f64(double, double);
166 __device__ __attribute__((const)) double __ocml_atan_f64(double);
167 __device__ __attribute__((pure)) double __ocml_atanh_f64(double);
168 __device__ __attribute__((pure)) double __ocml_cbrt_f64(double);
169 __device__ __attribute__((const)) double __ocml_ceil_f64(double);
170 __device__ __attribute__((const)) double __ocml_copysign_f64(double, double);
171 __device__ double __ocml_cos_f64(double);
172 __device__ __attribute__((pure)) double __ocml_cosh_f64(double);
173 __device__ double __ocml_cospi_f64(double);
174 __device__ double __ocml_i0_f64(double);
175 __device__ double __ocml_i1_f64(double);
176 __device__ __attribute__((pure)) double __ocml_erfc_f64(double);
177 __device__ __attribute__((pure)) double __ocml_erfcinv_f64(double);
178 __device__ __attribute__((pure)) double __ocml_erfcx_f64(double);
179 __device__ __attribute__((pure)) double __ocml_erf_f64(double);
180 __device__ __attribute__((pure)) double __ocml_erfinv_f64(double);
181 __device__ __attribute__((pure)) double __ocml_exp10_f64(double);
182 __device__ __attribute__((pure)) double __ocml_exp2_f64(double);
183 __device__ __attribute__((pure)) double __ocml_exp_f64(double);
184 __device__ __attribute__((pure)) double __ocml_expm1_f64(double);
185 __device__ __attribute__((const)) double __ocml_fabs_f64(double);
186 __device__ __attribute__((const)) double __ocml_fdim_f64(double, double);
187 __device__ __attribute__((const)) double __ocml_floor_f64(double);
188 __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
189 __device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
190 __device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
191 __device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
192 __device__ double __ocml_frexp_f64(double,
193 __attribute__((address_space(5))) int *);
194 __device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
195 __device__ __attribute__((const)) int __ocml_ilogb_f64(double);
196 __device__ __attribute__((const)) int __ocml_isfinite_f64(double);
197 __device__ __attribute__((const)) int __ocml_isinf_f64(double);
198 __device__ __attribute__((const)) int __ocml_isnan_f64(double);
199 __device__ double __ocml_j0_f64(double);
200 __device__ double __ocml_j1_f64(double);
201 __device__ __attribute__((const)) double __ocml_ldexp_f64(double, int);
202 __device__ double __ocml_lgamma_f64(double);
203 __device__ __attribute__((pure)) double __ocml_log10_f64(double);
204 __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
205 __device__ __attribute__((pure)) double __ocml_log2_f64(double);
206 __device__ __attribute__((const)) double __ocml_logb_f64(double);
207 __device__ __attribute__((pure)) double __ocml_log_f64(double);
208 __device__ double __ocml_modf_f64(double,
209 __attribute__((address_space(5))) double *);
210 __device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
211 __device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
212 __device__ __attribute__((const)) double __ocml_len3_f64(double, double,
213 double);
214 __device__ __attribute__((const)) double __ocml_len4_f64(double, double, double,
215 double);
216 __device__ __attribute__((pure)) double __ocml_ncdf_f64(double);
217 __device__ __attribute__((pure)) double __ocml_ncdfinv_f64(double);
218 __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
219 __device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
220 __device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
221 __device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
222 __device__ double __ocml_remquo_f64(double, double,
223 __attribute__((address_space(5))) int *);
224 __device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
225 __device__ __attribute__((const)) double __ocml_rint_f64(double);
226 __device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
227 double);
228 __device__ __attribute__((const)) double __ocml_rlen4_f64(double, double,
229 double, double);
230 __device__ __attribute__((const)) double __ocml_round_f64(double);
231 __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
232 __device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
233 __device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
234 __device__ __attribute__((const)) int __ocml_signbit_f64(double);
235 __device__ double __ocml_sincos_f64(double,
236 __attribute__((address_space(5))) double *);
237 __device__ double
238 __ocml_sincospi_f64(double, __attribute__((address_space(5))) double *);
239 __device__ double __ocml_sin_f64(double);
240 __device__ __attribute__((pure)) double __ocml_sinh_f64(double);
241 __device__ double __ocml_sinpi_f64(double);
242 __device__ __attribute__((const)) double __ocml_sqrt_f64(double);
243 __device__ double __ocml_tan_f64(double);
244 __device__ __attribute__((pure)) double __ocml_tanh_f64(double);
245 __device__ double __ocml_tgamma_f64(double);
246 __device__ __attribute__((const)) double __ocml_trunc_f64(double);
247 __device__ double __ocml_y0_f64(double);
248 __device__ double __ocml_y1_f64(double);
249
250 // BEGIN INTRINSICS
251 __device__ __attribute__((const)) double __ocml_add_rte_f64(double, double);
252 __device__ __attribute__((const)) double __ocml_add_rtn_f64(double, double);
253 __device__ __attribute__((const)) double __ocml_add_rtp_f64(double, double);
254 __device__ __attribute__((const)) double __ocml_add_rtz_f64(double, double);
255 __device__ __attribute__((const)) double __ocml_sub_rte_f64(double, double);
256 __device__ __attribute__((const)) double __ocml_sub_rtn_f64(double, double);
257 __device__ __attribute__((const)) double __ocml_sub_rtp_f64(double, double);
258 __device__ __attribute__((const)) double __ocml_sub_rtz_f64(double, double);
259 __device__ __attribute__((const)) double __ocml_mul_rte_f64(double, double);
260 __device__ __attribute__((const)) double __ocml_mul_rtn_f64(double, double);
261 __device__ __attribute__((const)) double __ocml_mul_rtp_f64(double, double);
262 __device__ __attribute__((const)) double __ocml_mul_rtz_f64(double, double);
263 __device__ __attribute__((const)) double __ocml_div_rte_f64(double, double);
264 __device__ __attribute__((const)) double __ocml_div_rtn_f64(double, double);
265 __device__ __attribute__((const)) double __ocml_div_rtp_f64(double, double);
266 __device__ __attribute__((const)) double __ocml_div_rtz_f64(double, double);
267 __device__ __attribute__((const)) double __ocml_sqrt_rte_f64(double);
268 __device__ __attribute__((const)) double __ocml_sqrt_rtn_f64(double);
269 __device__ __attribute__((const)) double __ocml_sqrt_rtp_f64(double);
270 __device__ __attribute__((const)) double __ocml_sqrt_rtz_f64(double);
271 __device__ __attribute__((const)) double __ocml_fma_rte_f64(double, double,
272 double);
273 __device__ __attribute__((const)) double __ocml_fma_rtn_f64(double, double,
274 double);
275 __device__ __attribute__((const)) double __ocml_fma_rtp_f64(double, double,
276 double);
277 __device__ __attribute__((const)) double __ocml_fma_rtz_f64(double, double,
278 double);
279
280 __device__ inline __attribute__((const)) double
__llvm_amdgcn_rcp_f64(double __x)281 __llvm_amdgcn_rcp_f64(double __x) {
282 return __builtin_amdgcn_rcp(__x);
283 }
284 __device__ inline __attribute__((const)) double
__llvm_amdgcn_rsq_f64(double __x)285 __llvm_amdgcn_rsq_f64(double __x) {
286 return __builtin_amdgcn_rsq(__x);
287 }
288
289 __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
290 __device__ _Float16 __ocml_cos_f16(_Float16);
291 __device__ __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
292 __device__ __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
293 __device__ __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
294 __device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
295 __device__ __attribute__((const)) _Float16 __ocml_fma_f16(_Float16, _Float16,
296 _Float16);
297 __device__ __attribute__((const)) _Float16 __ocml_fabs_f16(_Float16);
298 __device__ __attribute__((const)) int __ocml_isinf_f16(_Float16);
299 __device__ __attribute__((const)) int __ocml_isnan_f16(_Float16);
300 __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
301 __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
302 __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
303 __device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
304 __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
305 __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
306 __device__ _Float16 __ocml_sin_f16(_Float16);
307 __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
308 __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
309 __device__ __attribute__((pure)) _Float16 __ocml_pown_f16(_Float16, int);
310
311 typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
312 typedef short __2i16 __attribute__((ext_vector_type(2)));
313
314 __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b,
315 float c, bool s);
316 __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
317 __device__ __attribute__((const)) __2f16 __ocml_fabs_2f16(__2f16);
318 __device__ __2f16 __ocml_cos_2f16(__2f16);
319 __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
320 __device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
321 __device__ __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
322 __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
323 __device__ __attribute__((const))
324 __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
325 __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
326 __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
327 __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
328 __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
329 __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
330 __device__ inline __2f16
__llvm_amdgcn_rcp_2f16(__2f16 __x)331 __llvm_amdgcn_rcp_2f16(__2f16 __x) // Not currently exposed by ROCDL.
332 {
333 return (__2f16)(__llvm_amdgcn_rcp_f16(__x.x), __llvm_amdgcn_rcp_f16(__x.y));
334 }
335 __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
336 __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
337 __device__ __2f16 __ocml_sin_2f16(__2f16);
338 __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
339 __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
340 __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
341
342 #ifdef __cplusplus
343 } // extern "C"
344 #endif
345
346 #endif // __CLANG_HIP_LIBDEVICE_DECLARES_H__
347