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