HIP: Heterogenous-computing Interface for Portability
math_fwd.h
1 /*
2 Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved.
3 
4 Permission is hereby granted, free of charge, to any person obtaining a copy
5 of this software and associated documentation files (the "Software"), to deal
6 in the Software without restriction, including without limitation the rights
7 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
8 copies of the Software, and to permit persons to whom the Software is
9 furnished to do so, subject to the following conditions:
10 
11 The above copyright notice and this permission notice shall be included in
12 all copies or substantial portions of the Software.
13 
14 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
19 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
20 THE SOFTWARE.
21 */
22 
23 #pragma once
24 
25 #if !defined(__HIPCC_RTC__)
26 #include "host_defines.h"
27 #include "amd_hip_vector_types.h" // For Native_vec_
28 #endif
29 
30 #if defined(__cplusplus)
31  extern "C" {
32 #endif
33 
34 // DOT FUNCTIONS
35 #if defined(__clang__) && defined(__HIP__)
36 __device__
37 __attribute__((const))
38 int __ockl_sdot2(
39  HIP_vector_base<short, 2>::Native_vec_,
40  HIP_vector_base<short, 2>::Native_vec_,
41  int, bool);
42 
43 __device__
44 __attribute__((const))
45 unsigned int __ockl_udot2(
46  HIP_vector_base<unsigned short, 2>::Native_vec_,
47  HIP_vector_base<unsigned short, 2>::Native_vec_,
48  unsigned int, bool);
49 
50 __device__
51 __attribute__((const))
52 int __ockl_sdot4(
53  HIP_vector_base<char, 4>::Native_vec_,
54  HIP_vector_base<char, 4>::Native_vec_,
55  int, bool);
56 
57 __device__
58 __attribute__((const))
59 unsigned int __ockl_udot4(
60  HIP_vector_base<unsigned char, 4>::Native_vec_,
61  HIP_vector_base<unsigned char, 4>::Native_vec_,
62  unsigned int, bool);
63 
64 __device__
65 __attribute__((const))
66 int __ockl_sdot8(int, int, int, bool);
67 
68 __device__
69 __attribute__((const))
70 unsigned int __ockl_udot8(unsigned int, unsigned int, unsigned int, bool);
71 #endif
72 
73 #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
74 // BEGIN FLOAT
75 __device__
76 __attribute__((const))
77 float __ocml_acos_f32(float);
78 __device__
79 __attribute__((pure))
80 float __ocml_acosh_f32(float);
81 __device__
82 __attribute__((const))
83 float __ocml_asin_f32(float);
84 __device__
85 __attribute__((pure))
86 float __ocml_asinh_f32(float);
87 __device__
88 __attribute__((const))
89 float __ocml_atan2_f32(float, float);
90 __device__
91 __attribute__((const))
92 float __ocml_atan_f32(float);
93 __device__
94 __attribute__((pure))
95 float __ocml_atanh_f32(float);
96 __device__
97 __attribute__((pure))
98 float __ocml_cbrt_f32(float);
99 __device__
100 __attribute__((const))
101 float __ocml_ceil_f32(float);
102 __device__
103 __attribute__((const))
104 __device__
105 float __ocml_copysign_f32(float, float);
106 __device__
107 float __ocml_cos_f32(float);
108 __device__
109 float __ocml_native_cos_f32(float);
110 __device__
111 __attribute__((pure))
112 __device__
113 float __ocml_cosh_f32(float);
114 __device__
115 float __ocml_cospi_f32(float);
116 __device__
117 float __ocml_i0_f32(float);
118 __device__
119 float __ocml_i1_f32(float);
120 __device__
121 __attribute__((pure))
122 float __ocml_erfc_f32(float);
123 __device__
124 __attribute__((pure))
125 float __ocml_erfcinv_f32(float);
126 __device__
127 __attribute__((pure))
128 float __ocml_erfcx_f32(float);
129 __device__
130 __attribute__((pure))
131 float __ocml_erf_f32(float);
132 __device__
133 __attribute__((pure))
134 float __ocml_erfinv_f32(float);
135 __device__
136 __attribute__((pure))
137 float __ocml_exp10_f32(float);
138 __device__
139 __attribute__((pure))
140 float __ocml_native_exp10_f32(float);
141 __device__
142 __attribute__((pure))
143 float __ocml_exp2_f32(float);
144 __device__
145 __attribute__((pure))
146 float __ocml_exp_f32(float);
147 __device__
148 __attribute__((pure))
149 float __ocml_native_exp_f32(float);
150 __device__
151 __attribute__((pure))
152 float __ocml_expm1_f32(float);
153 __device__
154 __attribute__((const))
155 float __ocml_fabs_f32(float);
156 __device__
157 __attribute__((const))
158 float __ocml_fdim_f32(float, float);
159 __device__
160 __attribute__((const))
161 float __ocml_floor_f32(float);
162 __device__
163 __attribute__((const))
164 float __ocml_fma_f32(float, float, float);
165 __device__
166 __attribute__((const))
167 float __ocml_fmax_f32(float, float);
168 __device__
169 __attribute__((const))
170 float __ocml_fmin_f32(float, float);
171 __device__
172 __attribute__((const))
173 __device__
174 float __ocml_fmod_f32(float, float);
175 __device__
176 float __ocml_frexp_f32(float, __attribute__((address_space(5))) int*);
177 __device__
178 __attribute__((const))
179 float __ocml_hypot_f32(float, float);
180 __device__
181 __attribute__((const))
182 int __ocml_ilogb_f32(float);
183 __device__
184 __attribute__((const))
185 int __ocml_isfinite_f32(float);
186 __device__
187 __attribute__((const))
188 int __ocml_isinf_f32(float);
189 __device__
190 __attribute__((const))
191 int __ocml_isnan_f32(float);
192 __device__
193 float __ocml_j0_f32(float);
194 __device__
195 float __ocml_j1_f32(float);
196 __device__
197 __attribute__((const))
198 float __ocml_ldexp_f32(float, int);
199 __device__
200 float __ocml_lgamma_f32(float);
201 __device__
202 __attribute__((pure))
203 float __ocml_log10_f32(float);
204 __device__
205 __attribute__((pure))
206 float __ocml_native_log10_f32(float);
207 __device__
208 __attribute__((pure))
209 float __ocml_log1p_f32(float);
210 __device__
211 __attribute__((pure))
212 float __ocml_log2_f32(float);
213 __device__
214 __attribute__((pure))
215 float __ocml_native_log2_f32(float);
216 __device__
217 __attribute__((const))
218 float __ocml_logb_f32(float);
219 __device__
220 __attribute__((pure))
221 float __ocml_log_f32(float);
222 __device__
223 __attribute__((pure))
224 float __ocml_native_log_f32(float);
225 __device__
226 float __ocml_modf_f32(float, __attribute__((address_space(5))) float*);
227 __device__
228 __attribute__((const))
229 float __ocml_nearbyint_f32(float);
230 __device__
231 __attribute__((const))
232 float __ocml_nextafter_f32(float, float);
233 __device__
234 __attribute__((const))
235 float __ocml_len3_f32(float, float, float);
236 __device__
237 __attribute__((const))
238 float __ocml_len4_f32(float, float, float, float);
239 __device__
240 __attribute__((pure))
241 float __ocml_ncdf_f32(float);
242 __device__
243 __attribute__((pure))
244 float __ocml_ncdfinv_f32(float);
245 __device__
246 __attribute__((pure))
247 float __ocml_pow_f32(float, float);
248 __device__
249 __attribute__((pure))
250 float __ocml_pown_f32(float, int);
251 __device__
252 __attribute__((pure))
253 float __ocml_rcbrt_f32(float);
254 __device__
255 __attribute__((const))
256 float __ocml_remainder_f32(float, float);
257 __device__
258 float __ocml_remquo_f32(float, float, __attribute__((address_space(5))) int*);
259 __device__
260 __attribute__((const))
261 float __ocml_rhypot_f32(float, float);
262 __device__
263 __attribute__((const))
264 float __ocml_rint_f32(float);
265 __device__
266 __attribute__((const))
267 float __ocml_rlen3_f32(float, float, float);
268 __device__
269 __attribute__((const))
270 float __ocml_rlen4_f32(float, float, float, float);
271 __device__
272 __attribute__((const))
273 float __ocml_round_f32(float);
274 __device__
275 __attribute__((pure))
276 float __ocml_rsqrt_f32(float);
277 __device__
278 __attribute__((const))
279 float __ocml_scalb_f32(float, float);
280 __device__
281 __attribute__((const))
282 float __ocml_scalbn_f32(float, int);
283 __device__
284 __attribute__((const))
285 int __ocml_signbit_f32(float);
286 __device__
287 float __ocml_sincos_f32(float, __attribute__((address_space(5))) float*);
288 __device__
289 float __ocml_sincospi_f32(float, __attribute__((address_space(5))) float*);
290 __device__
291 float __ocml_sin_f32(float);
292 __device__
293 float __ocml_native_sin_f32(float);
294 __device__
295 __attribute__((pure))
296 float __ocml_sinh_f32(float);
297 __device__
298 float __ocml_sinpi_f32(float);
299 __device__
300 __attribute__((const))
301 float __ocml_sqrt_f32(float);
302 __device__
303 __attribute__((const))
304 float __ocml_native_sqrt_f32(float);
305 __device__
306 float __ocml_tan_f32(float);
307 __device__
308 __attribute__((pure))
309 float __ocml_tanh_f32(float);
310 __device__
311 float __ocml_tgamma_f32(float);
312 __device__
313 __attribute__((const))
314 float __ocml_trunc_f32(float);
315 __device__
316 float __ocml_y0_f32(float);
317 __device__
318 float __ocml_y1_f32(float);
319 
320 // BEGIN INTRINSICS
321 __device__
322 __attribute__((const))
323 float __ocml_add_rte_f32(float, float);
324 __device__
325 __attribute__((const))
326 float __ocml_add_rtn_f32(float, float);
327 __device__
328 __attribute__((const))
329 float __ocml_add_rtp_f32(float, float);
330 __device__
331 __attribute__((const))
332 float __ocml_add_rtz_f32(float, float);
333 __device__
334 __attribute__((const))
335 float __ocml_sub_rte_f32(float, float);
336 __device__
337 __attribute__((const))
338 float __ocml_sub_rtn_f32(float, float);
339 __device__
340 __attribute__((const))
341 float __ocml_sub_rtp_f32(float, float);
342 __device__
343 __attribute__((const))
344 float __ocml_sub_rtz_f32(float, float);
345 __device__
346 __attribute__((const))
347 float __ocml_mul_rte_f32(float, float);
348 __device__
349 __attribute__((const))
350 float __ocml_mul_rtn_f32(float, float);
351 __device__
352 __attribute__((const))
353 float __ocml_mul_rtp_f32(float, float);
354 __device__
355 __attribute__((const))
356 float __ocml_mul_rtz_f32(float, float);
357 __device__
358 __attribute__((const))
359 float __ocml_div_rte_f32(float, float);
360 __device__
361 __attribute__((const))
362 float __ocml_div_rtn_f32(float, float);
363 __device__
364 __attribute__((const))
365 float __ocml_div_rtp_f32(float, float);
366 __device__
367 __attribute__((const))
368 float __ocml_div_rtz_f32(float, float);
369 __device__
370 __attribute__((const))
371 float __ocml_sqrt_rte_f32(float);
372 __device__
373 __attribute__((const))
374 float __ocml_sqrt_rtn_f32(float);
375 __device__
376 __attribute__((const))
377 float __ocml_sqrt_rtp_f32(float);
378 __device__
379 __attribute__((const))
380 float __ocml_sqrt_rtz_f32(float);
381 __device__
382 __attribute__((const))
383 float __ocml_fma_rte_f32(float, float, float);
384 __device__
385 __attribute__((const))
386 float __ocml_fma_rtn_f32(float, float, float);
387 __device__
388 __attribute__((const))
389 float __ocml_fma_rtp_f32(float, float, float);
390 __device__
391 __attribute__((const))
392 float __ocml_fma_rtz_f32(float, float, float);
393 // END INTRINSICS
394 // END FLOAT
395 
396 // BEGIN DOUBLE
397 __device__
398 __attribute__((const))
399 double __ocml_acos_f64(double);
400 __device__
401 __attribute__((pure))
402 double __ocml_acosh_f64(double);
403 __device__
404 __attribute__((const))
405 double __ocml_asin_f64(double);
406 __device__
407 __attribute__((pure))
408 double __ocml_asinh_f64(double);
409 __device__
410 __attribute__((const))
411 double __ocml_atan2_f64(double, double);
412 __device__
413 __attribute__((const))
414 double __ocml_atan_f64(double);
415 __device__
416 __attribute__((pure))
417 double __ocml_atanh_f64(double);
418 __device__
419 __attribute__((pure))
420 double __ocml_cbrt_f64(double);
421 __device__
422 __attribute__((const))
423 double __ocml_ceil_f64(double);
424 __device__
425 __attribute__((const))
426 double __ocml_copysign_f64(double, double);
427 __device__
428 double __ocml_cos_f64(double);
429 __device__
430 __attribute__((pure))
431 double __ocml_cosh_f64(double);
432 __device__
433 double __ocml_cospi_f64(double);
434 __device__
435 double __ocml_i0_f64(double);
436 __device__
437 double __ocml_i1_f64(double);
438 __device__
439 __attribute__((pure))
440 double __ocml_erfc_f64(double);
441 __device__
442 __attribute__((pure))
443 double __ocml_erfcinv_f64(double);
444 __device__
445 __attribute__((pure))
446 double __ocml_erfcx_f64(double);
447 __device__
448 __attribute__((pure))
449 double __ocml_erf_f64(double);
450 __device__
451 __attribute__((pure))
452 double __ocml_erfinv_f64(double);
453 __device__
454 __attribute__((pure))
455 double __ocml_exp10_f64(double);
456 __device__
457 __attribute__((pure))
458 double __ocml_exp2_f64(double);
459 __device__
460 __attribute__((pure))
461 double __ocml_exp_f64(double);
462 __device__
463 __attribute__((pure))
464 double __ocml_expm1_f64(double);
465 __device__
466 __attribute__((const))
467 double __ocml_fabs_f64(double);
468 __device__
469 __attribute__((const))
470 double __ocml_fdim_f64(double, double);
471 __device__
472 __attribute__((const))
473 double __ocml_floor_f64(double);
474 __device__
475 __attribute__((const))
476 double __ocml_fma_f64(double, double, double);
477 __device__
478 __attribute__((const))
479 double __ocml_fmax_f64(double, double);
480 __device__
481 __attribute__((const))
482 double __ocml_fmin_f64(double, double);
483 __device__
484 __attribute__((const))
485 double __ocml_fmod_f64(double, double);
486 __device__
487 double __ocml_frexp_f64(double, __attribute__((address_space(5))) int*);
488 __device__
489 __attribute__((const))
490 double __ocml_hypot_f64(double, double);
491 __device__
492 __attribute__((const))
493 int __ocml_ilogb_f64(double);
494 __device__
495 __attribute__((const))
496 int __ocml_isfinite_f64(double);
497 __device__
498 __attribute__((const))
499 int __ocml_isinf_f64(double);
500 __device__
501 __attribute__((const))
502 int __ocml_isnan_f64(double);
503 __device__
504 double __ocml_j0_f64(double);
505 __device__
506 double __ocml_j1_f64(double);
507 __device__
508 __attribute__((const))
509 double __ocml_ldexp_f64(double, int);
510 __device__
511 double __ocml_lgamma_f64(double);
512 __device__
513 __attribute__((pure))
514 double __ocml_log10_f64(double);
515 __device__
516 __attribute__((pure))
517 double __ocml_log1p_f64(double);
518 __device__
519 __attribute__((pure))
520 double __ocml_log2_f64(double);
521 __device__
522 __attribute__((const))
523 double __ocml_logb_f64(double);
524 __device__
525 __attribute__((pure))
526 double __ocml_log_f64(double);
527 __device__
528 double __ocml_modf_f64(double, __attribute__((address_space(5))) double*);
529 __device__
530 __attribute__((const))
531 double __ocml_nearbyint_f64(double);
532 __device__
533 __attribute__((const))
534 double __ocml_nextafter_f64(double, double);
535 __device__
536 __attribute__((const))
537 double __ocml_len3_f64(double, double, double);
538 __device__
539 __attribute__((const))
540 double __ocml_len4_f64(double, double, double, double);
541 __device__
542 __attribute__((pure))
543 double __ocml_ncdf_f64(double);
544 __device__
545 __attribute__((pure))
546 double __ocml_ncdfinv_f64(double);
547 __device__
548 __attribute__((pure))
549 double __ocml_pow_f64(double, double);
550 __device__
551 __attribute__((pure))
552 double __ocml_pown_f64(double, int);
553 __device__
554 __attribute__((pure))
555 double __ocml_rcbrt_f64(double);
556 __device__
557 __attribute__((const))
558 double __ocml_remainder_f64(double, double);
559 __device__
560 double __ocml_remquo_f64(
561  double, double, __attribute__((address_space(5))) int*);
562 __device__
563 __attribute__((const))
564 double __ocml_rhypot_f64(double, double);
565 __device__
566 __attribute__((const))
567 double __ocml_rint_f64(double);
568 __device__
569 __attribute__((const))
570 double __ocml_rlen3_f64(double, double, double);
571 __device__
572 __attribute__((const))
573 double __ocml_rlen4_f64(double, double, double, double);
574 __device__
575 __attribute__((const))
576 double __ocml_round_f64(double);
577 __device__
578 __attribute__((pure))
579 double __ocml_rsqrt_f64(double);
580 __device__
581 __attribute__((const))
582 double __ocml_scalb_f64(double, double);
583 __device__
584 __attribute__((const))
585 double __ocml_scalbn_f64(double, int);
586 __device__
587 __attribute__((const))
588 int __ocml_signbit_f64(double);
589 __device__
590 double __ocml_sincos_f64(double, __attribute__((address_space(5))) double*);
591 __device__
592 double __ocml_sincospi_f64(double, __attribute__((address_space(5))) double*);
593 __device__
594 double __ocml_sin_f64(double);
595 __device__
596 __attribute__((pure))
597 double __ocml_sinh_f64(double);
598 __device__
599 double __ocml_sinpi_f64(double);
600 __device__
601 __attribute__((const))
602 double __ocml_sqrt_f64(double);
603 __device__
604 double __ocml_tan_f64(double);
605 __device__
606 __attribute__((pure))
607 double __ocml_tanh_f64(double);
608 __device__
609 double __ocml_tgamma_f64(double);
610 __device__
611 __attribute__((const))
612 double __ocml_trunc_f64(double);
613 __device__
614 double __ocml_y0_f64(double);
615 __device__
616 double __ocml_y1_f64(double);
617 
618 // BEGIN INTRINSICS
619 __device__
620 __attribute__((const))
621 double __ocml_add_rte_f64(double, double);
622 __device__
623 __attribute__((const))
624 double __ocml_add_rtn_f64(double, double);
625 __device__
626 __attribute__((const))
627 double __ocml_add_rtp_f64(double, double);
628 __device__
629 __attribute__((const))
630 double __ocml_add_rtz_f64(double, double);
631 __device__
632 __attribute__((const))
633 double __ocml_sub_rte_f64(double, double);
634 __device__
635 __attribute__((const))
636 double __ocml_sub_rtn_f64(double, double);
637 __device__
638 __attribute__((const))
639 double __ocml_sub_rtp_f64(double, double);
640 __device__
641 __attribute__((const))
642 double __ocml_sub_rtz_f64(double, double);
643 __device__
644 __attribute__((const))
645 double __ocml_mul_rte_f64(double, double);
646 __device__
647 __attribute__((const))
648 double __ocml_mul_rtn_f64(double, double);
649 __device__
650 __attribute__((const))
651 double __ocml_mul_rtp_f64(double, double);
652 __device__
653 __attribute__((const))
654 double __ocml_mul_rtz_f64(double, double);
655 __device__
656 __attribute__((const))
657 double __ocml_div_rte_f64(double, double);
658 __device__
659 __attribute__((const))
660 double __ocml_div_rtn_f64(double, double);
661 __device__
662 __attribute__((const))
663 double __ocml_div_rtp_f64(double, double);
664 __device__
665 __attribute__((const))
666 double __ocml_div_rtz_f64(double, double);
667 __device__
668 __attribute__((const))
669 double __ocml_sqrt_rte_f64(double);
670 __device__
671 __attribute__((const))
672 double __ocml_sqrt_rtn_f64(double);
673 __device__
674 __attribute__((const))
675 double __ocml_sqrt_rtp_f64(double);
676 __device__
677 __attribute__((const))
678 double __ocml_sqrt_rtz_f64(double);
679 __device__
680 __attribute__((const))
681 double __ocml_fma_rte_f64(double, double, double);
682 __device__
683 __attribute__((const))
684 double __ocml_fma_rtn_f64(double, double, double);
685 __device__
686 __attribute__((const))
687 double __ocml_fma_rtp_f64(double, double, double);
688 __device__
689 __attribute__((const))
690 double __ocml_fma_rtz_f64(double, double, double);
691 // END INTRINSICS
692 // END DOUBLE
693 
694 #endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
695 
696 #if defined(__cplusplus)
697  } // extern "C"
698 #endif
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Definition: hip_fp16_math_fwd.h:57