/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/math_v2.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/math_v2.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/utility/math_v2.hpp Source File
math_v2.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #ifndef __HIP_DEVICE_COMPILE__
7 #include <cmath>
8 #endif
9 
10 #include "ck/utility/data_type.hpp"
11 #include "ck/utility/type.hpp"
13 
14 namespace ck {
15 namespace math {
16 
17 #if CK_WORKAROUND_SWDEV_383542
18 extern "C" __device__ float __ocml_native_recip_f32(float);
19 #endif
20 
21 // math functions for the host, some are implemented by calling C++ std functions
22 #if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC)
23 static inline __host__ float abs(float x) { return std::abs(x); };
24 
25 static inline __host__ double abs(double x) { return std::abs(x); };
26 
27 static inline __host__ int8_t abs(int8_t x)
28 {
29  int8_t sgn = x >> (8 - 1);
30 
31  return (x ^ sgn) - sgn;
32 };
33 
34 static inline __host__ int32_t abs(int32_t x)
35 {
36  int32_t sgn = x >> (32 - 1);
37 
38  return (x ^ sgn) - sgn;
39 };
40 
41 static inline __host__ half_t abs(half_t x)
42 {
43  uint16_t xx = ck::bit_cast<uint16_t>(x);
44 
45  uint16_t abs_xx = xx & 0x7fff;
46 
47  half_t abs_x = ck::bit_cast<half_t>(abs_xx);
48 
49  return abs_x;
50 };
51 
52 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
53 static inline __host__ int4_t abs(int4_t x)
54 {
55  int4_t sgn = x >> (4 - 1);
56  return (x ^ sgn) - sgn;
57 }
58 #endif
59 
60 static inline __host__ bool isnan(float x) { return std::isnan(x); };
61 
62 static inline __host__ bool isnan(double x) { return std::isnan(x); };
63 
64 static inline __host__ bool isnan(int8_t x)
65 {
66  (void)x;
67  return false;
68 };
69 
70 static inline __host__ bool isnan(int32_t x)
71 {
72  (void)x;
73  return false;
74 };
75 
76 static inline __host__ bool isnan(half_t x)
77 {
78  uint16_t xx = ck::bit_cast<uint16_t>(x);
79 
80  return (xx & 0x7FFF) > 0x7C00;
81 };
82 
83 static inline __host__ bool isnan(f8_t x) { return ck::fp8_is_nan(x); };
84 
85 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
86 static inline __host__ bool isnan(int4_t x)
87 {
88  (void)x;
89  return false;
90 };
91 #endif
92 
93 static inline __host__ half_t sqrt(half_t x)
94 {
95  return static_cast<half_t>(std::sqrt(static_cast<float>(x)));
96 };
97 
98 static inline __host__ float sqrt(float x) { return std::sqrt(x); };
99 
100 static inline __host__ double sqrt(double x) { return std::sqrt(x); };
101 
102 template <typename T>
103 inline __host__ T tanh(T x)
104 {
105  return ck::type_convert<T>(std::tanhf(ck::type_convert<float>(x)));
106 };
107 
108 template <>
109 inline __host__ float tanh<float>(float x)
110 {
111  return std::tanhf(x);
112 };
113 
114 template <>
115 inline __host__ double tanh<double>(double x)
116 {
117  return std::tanh(x);
118 };
119 
120 template <typename T>
121 inline __host__ T acos(T x)
122 {
123  return ck::type_convert<T>(std::acosf(ck::type_convert<float>(x)));
124 };
125 
126 template <>
127 inline __host__ float acos<float>(float x)
128 {
129  return std::acosf(x);
130 };
131 
132 template <>
133 inline __host__ double acos<double>(double x)
134 {
135  return std::acos(x);
136 };
137 
138 template <typename T>
139 inline __host__ T neg(T x)
140 {
141  return ck::type_convert<T>(-(ck::type_convert<float>(x)));
142 };
143 
144 template <>
145 inline __host__ float neg<float>(float x)
146 {
147  return -x;
148 };
149 
150 template <>
151 inline __host__ double neg<double>(double x)
152 {
153  return -x;
154 };
155 
156 template <>
157 inline __host__ int32_t neg<int32_t>(int32_t x)
158 {
159  return -x;
160 };
161 
162 template <>
163 inline __host__ int8_t neg<int8_t>(int8_t x)
164 {
165  return -x;
166 };
167 
168 template <typename T>
169 inline __host__ T atan(T x)
170 {
171  return ck::type_convert<T>(std::atanf(ck::type_convert<float>(x)));
172 };
173 
174 template <>
175 inline __host__ float atan<float>(float x)
176 {
177  return std::atanf(x);
178 };
179 
180 template <>
181 inline __host__ double atan<double>(double x)
182 {
183  return std::atan(x);
184 };
185 
186 template <typename T>
187 inline __host__ T sin(T x)
188 {
189  return ck::type_convert<T>(std::sinf(ck::type_convert<float>(x)));
190 };
191 
192 template <>
193 inline __host__ float sin<float>(float x)
194 {
195  return std::sinf(x);
196 };
197 
198 template <>
199 inline __host__ double sin<double>(double x)
200 {
201  return std::sin(x);
202 };
203 
204 template <typename T>
205 inline __host__ T asin(T x)
206 {
207  return ck::type_convert<T>(std::asinf(ck::type_convert<float>(x)));
208 };
209 
210 template <>
211 inline __host__ float asin<float>(float x)
212 {
213  return std::asinf(x);
214 };
215 
216 template <>
217 inline __host__ double asin<double>(double x)
218 {
219  return std::asin(x);
220 };
221 
222 template <typename T>
223 inline __host__ T asinh(T x)
224 {
225  return ck::type_convert<T>(std::asinhf(ck::type_convert<float>(x)));
226 };
227 
228 template <>
229 inline __host__ float asinh<float>(float x)
230 {
231  return std::asinhf(x);
232 };
233 
234 template <>
235 inline __host__ double asinh<double>(double x)
236 {
237  return std::asinh(x);
238 };
239 
240 template <typename T>
241 inline __host__ T cos(T x)
242 {
243  return ck::type_convert<T>(std::cosf(ck::type_convert<float>(x)));
244 };
245 
246 template <>
247 inline __host__ float cos<float>(float x)
248 {
249  return std::cosf(x);
250 };
251 
252 template <>
253 inline __host__ double cos<double>(double x)
254 {
255  return std::cos(x);
256 };
257 
258 template <typename T>
259 inline __host__ T acosh(T x)
260 {
261  return ck::type_convert<T>(std::acoshf(ck::type_convert<float>(x)));
262 };
263 
264 template <>
265 inline __host__ float acosh<float>(float x)
266 {
267  return std::acoshf(x);
268 };
269 
270 template <>
271 inline __host__ double acosh<double>(double x)
272 {
273  return std::acosh(x);
274 };
275 
276 template <typename T>
277 inline __host__ T tan(T x)
278 {
279  return ck::type_convert<T>(std::tanf(ck::type_convert<float>(x)));
280 };
281 
282 template <>
283 inline __host__ float tan<float>(float x)
284 {
285  return std::tanf(x);
286 };
287 
288 template <>
289 inline __host__ double tan<double>(double x)
290 {
291  return std::tan(x);
292 };
293 
294 template <typename T>
295 inline __host__ T atanh(T x)
296 {
297  return ck::type_convert<T>(std::atanhf(ck::type_convert<float>(x)));
298 };
299 
300 template <>
301 inline __host__ float atanh<float>(float x)
302 {
303  return std::atanhf(x);
304 };
305 
306 template <>
307 inline __host__ double atanh<double>(double x)
308 {
309  return std::atanh(x);
310 };
311 
312 template <typename T>
313 inline __host__ T sinh(T x)
314 {
315  return ck::type_convert<T>(std::sinhf(ck::type_convert<float>(x)));
316 };
317 
318 template <>
319 inline __host__ float sinh<float>(float x)
320 {
321  return std::sinhf(x);
322 };
323 
324 template <>
325 inline __host__ double sinh<double>(double x)
326 {
327  return std::sinh(x);
328 };
329 
330 template <typename T>
331 inline __host__ T ceil(T x)
332 {
333  return ck::type_convert<T>(std::ceilf(ck::type_convert<float>(x)));
334 };
335 
336 template <>
337 inline __host__ float ceil<float>(float x)
338 {
339  return std::ceilf(x);
340 };
341 
342 template <>
343 inline __host__ double ceil<double>(double x)
344 {
345  return std::ceil(x);
346 };
347 
348 template <typename T>
349 inline __host__ T cosh(T x)
350 {
351  return ck::type_convert<T>(std::coshf(ck::type_convert<float>(x)));
352 };
353 
354 template <>
355 inline __host__ float cosh<float>(float x)
356 {
357  return std::coshf(x);
358 };
359 
360 template <>
361 inline __host__ double cosh<double>(double x)
362 {
363  return std::cosh(x);
364 };
365 
366 template <typename T>
367 inline __host__ T floor(T x)
368 {
369  return ck::type_convert<T>(std::floorf(ck::type_convert<float>(x)));
370 };
371 
372 template <>
373 inline __host__ float floor<float>(float x)
374 {
375  return std::floorf(x);
376 };
377 
378 template <>
379 inline __host__ double floor<double>(double x)
380 {
381  return std::floor(x);
382 };
383 
384 template <typename T>
385 inline __host__ T rcp(T x)
386 {
387  return ck::type_convert<T>(1.f / ck::type_convert<float>(x));
388 };
389 
390 template <typename T>
391 inline __host__ T exp(T x)
392 {
393  return ck::type_convert<T>(std::expf(ck::type_convert<float>(x)));
394 }
395 
396 template <>
397 inline __host__ float exp<float>(float x)
398 {
399  return std::expf(x);
400 }
401 
402 template <>
403 inline __host__ double exp<double>(double x)
404 {
405  return std::exp(x);
406 }
407 
408 template <typename T>
409 inline __host__ T log(T x)
410 {
411  return ck::type_convert<T>(std::logf(ck::type_convert<float>(x)));
412 }
413 
414 template <>
415 inline __host__ float log<float>(float x)
416 {
417  return std::logf(x);
418 }
419 
420 template <>
421 inline __host__ double log<double>(double x)
422 {
423  return std::log(x);
424 }
425 
426 template <typename T>
427 inline __host__ T pow(T x, T gamma)
428 {
429  return ck::type_convert<T>(
430  std::powf(ck::type_convert<float>(x), ck::type_convert<float>(gamma)));
431 }
432 
433 template <>
434 inline __host__ float pow<float>(float x, float gamma)
435 {
436  return std::powf(x, gamma);
437 }
438 
439 template <>
440 inline __host__ double pow<double>(double x, double gamma)
441 {
442  return std::pow(x, gamma);
443 }
444 
445 template <typename T>
446 inline __host__ T expm1(T x)
447 {
448  return ck::type_convert<T>(std::expm1f(ck::type_convert<float>(x)));
449 }
450 
451 template <>
452 inline __host__ float expm1<float>(float x)
453 {
454  return std::expm1f(x);
455 }
456 
457 template <>
458 inline __host__ double expm1<double>(double x)
459 {
460  return std::expm1(x);
461 }
462 #endif
463 // math functions for the HIP kernel, some are implemented by calling hip builtin functions
464 
465 static inline __device__ float abs(float x) { return ::abs(x); };
466 
467 static inline __device__ double abs(double x) { return ::abs(x); };
468 
469 static inline __device__ int8_t abs(int8_t x)
470 {
471  int8_t sgn = x >> (8 - 1);
472 
473  return (x ^ sgn) - sgn;
474 };
475 
476 static inline __device__ int32_t abs(int32_t x)
477 {
478  int32_t sgn = x >> (32 - 1);
479 
480  return (x ^ sgn) - sgn;
481 };
482 
483 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
484 static inline __device__ int4_t abs(int4_t x)
485 {
486  int4_t sgn = x >> (4 - 1);
487 
488  return (x ^ sgn) - sgn;
489 };
490 #endif
491 
492 static inline __device__ half_t abs(half_t x)
493 {
494  uint16_t xx = ck::bit_cast<uint16_t>(x);
495 
496  uint16_t abs_xx = xx & 0x7fff;
497 
498  half_t abs_x = ck::bit_cast<half_t>(abs_xx);
499 
500  return abs_x;
501 };
502 
503 static inline __device__ bool isnan(float x) { return ::isnan(x); };
504 
505 static inline __device__ bool isnan(double x) { return ::isnan(x); };
506 
507 static inline __device__ bool isnan(int8_t x)
508 {
509  (void)x;
510  return false;
511 };
512 
513 static inline __device__ bool isnan(int32_t x)
514 {
515  (void)x;
516  return false;
517 };
518 
519 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
520 static inline __device__ bool isnan(int4_t x)
521 {
522  (void)x;
523  return false;
524 };
525 #endif
526 
527 static inline __device__ bool isnan(half_t x)
528 {
529  uint16_t xx = ck::bit_cast<uint16_t>(x);
530 
531  return (xx & 0x7FFF) > 0x7C00;
532 };
533 
534 static inline __device__ bool isnan(f8_t x) { return ck::fp8_is_nan(x); };
535 
536 static inline __device__ half_t sqrt(half_t x)
537 {
538  return static_cast<half_t>(__builtin_amdgcn_sqrtf(static_cast<float>(x)));
539 };
540 
541 static inline __device__ float sqrt(float x) { return __builtin_amdgcn_sqrtf(x); };
542 
543 static inline __device__ double sqrt(double x) { return __builtin_amdgcn_sqrt(x); };
544 
545 template <typename T>
546 inline __device__ T tanh(T x)
547 {
548  return ck::type_convert<T>(::tanhf(ck::type_convert<float>(x)));
549 };
550 
551 template <>
552 inline __device__ float tanh<float>(float x)
553 {
554  return ::tanhf(x);
555 };
556 
557 template <>
558 inline __device__ double tanh<double>(double x)
559 {
560  return ::tanh(x);
561 };
562 
563 template <typename T>
564 inline __device__ T acos(T x)
565 {
566  return ck::type_convert<T>(::acosf(ck::type_convert<float>(x)));
567 };
568 
569 template <>
570 inline __device__ float acos<float>(float x)
571 {
572  return ::acosf(x);
573 };
574 
575 template <>
576 inline __device__ double acos<double>(double x)
577 {
578  return ::acos(x);
579 };
580 
581 template <typename T>
582 inline __device__ T neg(T x)
583 {
584  return ck::type_convert<T>(-(ck::type_convert<float>(x)));
585 };
586 
587 template <>
588 inline __device__ float neg<float>(float x)
589 {
590  return -x;
591 };
592 
593 template <>
594 inline __device__ double neg<double>(double x)
595 {
596  return -x;
597 };
598 
599 template <>
600 inline __device__ int32_t neg<int32_t>(int32_t x)
601 {
602  return -x;
603 };
604 
605 template <>
606 inline __device__ int8_t neg<int8_t>(int8_t x)
607 {
608  return -x;
609 };
610 
611 template <>
612 inline __device__ half_t neg<half_t>(half_t x)
613 {
614  return __hneg(static_cast<__half>(x));
615 };
616 
617 template <typename T>
618 inline __device__ T atan(T x)
619 {
620  return ck::type_convert<T>(::atanf(ck::type_convert<float>(x)));
621 };
622 
623 template <>
624 inline __device__ float atan<float>(float x)
625 {
626  return ::atanf(x);
627 };
628 
629 template <>
630 inline __device__ double atan<double>(double x)
631 {
632  return ::atan(x);
633 };
634 
635 template <typename T>
636 inline __device__ T sin(T x)
637 {
638  return ck::type_convert<T>(::sinf(ck::type_convert<float>(x)));
639 };
640 
641 template <>
642 inline __device__ float sin<float>(float x)
643 {
644  return ::sinf(x);
645 };
646 
647 template <>
648 inline __device__ double sin<double>(double x)
649 {
650  return ::sin(x);
651 };
652 
653 template <>
654 inline __device__ half_t sin<half_t>(half_t x)
655 {
656  return hsin(static_cast<__half>(x));
657 };
658 
659 template <typename T>
660 inline __device__ T asin(T x)
661 {
662  return ck::type_convert<T>(::asinf(ck::type_convert<float>(x)));
663 };
664 
665 template <>
666 inline __device__ float asin<float>(float x)
667 {
668  return ::asinf(x);
669 };
670 
671 template <>
672 inline __device__ double asin<double>(double x)
673 {
674  return ::asin(x);
675 };
676 
677 template <typename T>
678 inline __device__ T asinh(T x)
679 {
680  return ck::type_convert<T>(::asinhf(ck::type_convert<float>(x)));
681 };
682 
683 template <>
684 inline __device__ float asinh<float>(float x)
685 {
686  return ::asinhf(x);
687 };
688 
689 template <>
690 inline __device__ double asinh<double>(double x)
691 {
692  return ::asinh(x);
693 };
694 
695 template <typename T>
696 inline __device__ T acosh(T x)
697 {
698  return ck::type_convert<T>(::acoshf(ck::type_convert<float>(x)));
699 };
700 
701 template <>
702 inline __device__ float acosh<float>(float x)
703 {
704  return ::acoshf(x);
705 };
706 
707 template <>
708 inline __device__ double acosh<double>(double x)
709 {
710  return ::acosh(x);
711 };
712 
713 template <typename T>
714 inline __device__ T tan(T x)
715 {
716  return ck::type_convert<T>(::tanf(ck::type_convert<float>(x)));
717 };
718 
719 template <>
720 inline __device__ float tan<float>(float x)
721 {
722  return ::tanf(x);
723 };
724 
725 template <>
726 inline __device__ double tan<double>(double x)
727 {
728  return ::tan(x);
729 };
730 
731 template <typename T>
732 inline __device__ T atanh(T x)
733 {
734  return ck::type_convert<T>(::atanhf(ck::type_convert<float>(x)));
735 };
736 
737 template <>
738 inline __device__ float atanh<float>(float x)
739 {
740  return ::atanhf(x);
741 };
742 
743 template <>
744 inline __device__ double atanh<double>(double x)
745 {
746  return ::atanh(x);
747 };
748 
749 template <typename T>
750 inline __device__ T sinh(T x)
751 {
752  return ck::type_convert<T>(::sinhf(ck::type_convert<float>(x)));
753 };
754 
755 template <>
756 inline __device__ float sinh<float>(float x)
757 {
758  return ::sinhf(x);
759 };
760 
761 template <>
762 inline __device__ double sinh<double>(double x)
763 {
764  return ::sinh(x);
765 };
766 
767 template <typename T>
768 inline __device__ T ceil(T x)
769 {
770  return ck::type_convert<T>(::ceilf(ck::type_convert<float>(x)));
771 };
772 
773 template <>
774 inline __device__ float ceil<float>(float x)
775 {
776  return ::ceilf(x);
777 };
778 
779 template <>
780 inline __device__ double ceil<double>(double x)
781 {
782  return ::ceil(x);
783 };
784 
785 template <>
786 inline __device__ half_t ceil<half_t>(half_t x)
787 {
788  return hceil(static_cast<__half>(x));
789 };
790 
791 template <typename T>
792 inline __device__ T cosh(T x)
793 {
794  return ck::type_convert<T>(::coshf(ck::type_convert<float>(x)));
795 };
796 
797 template <>
798 inline __device__ float cosh<float>(float x)
799 {
800  return ::coshf(x);
801 };
802 
803 template <>
804 inline __device__ double cosh<double>(double x)
805 {
806  return ::cosh(x);
807 };
808 
809 template <typename T>
810 inline __device__ T floor(T x)
811 {
812  return ck::type_convert<T>(::floorf(ck::type_convert<float>(x)));
813 };
814 
815 template <>
816 inline __device__ float floor<float>(float x)
817 {
818  return ::floorf(x);
819 };
820 
821 template <>
822 inline __device__ double floor<double>(double x)
823 {
824  return ::floor(x);
825 };
826 
827 template <>
828 inline __device__ half_t floor<half_t>(half_t x)
829 {
830  return hfloor(static_cast<__half>(x));
831 };
832 
833 template <typename T>
834 inline __device__ T rcp(T x)
835 {
836 #if !CK_WORKAROUND_SWDEV_383542
837  return __frcp_rn(x);
838 #else
839  return __ocml_native_recip_f32(x);
840 #endif
841 };
842 
843 template <typename T>
844 inline __device__ T exp(T x)
845 {
846  return ck::type_convert<T>(__ocml_exp_f32(ck::type_convert<float>(x)));
847 };
848 
849 template <>
850 inline __device__ half_t exp<half_t>(half_t x)
851 {
852  return hexp(static_cast<__half>(x));
853 };
854 
855 template <>
856 inline __device__ float exp<float>(float x)
857 {
858  return __ocml_exp_f32(x);
859 };
860 
861 template <>
862 inline __device__ double exp<double>(double x)
863 {
864  return exp(x);
865 };
866 
867 template <typename T>
868 inline __device__ T log(T x)
869 {
870  return ck::type_convert<T>(__logf(ck::type_convert<float>(x)));
871 };
872 
873 template <>
874 inline __device__ half_t log<half_t>(half_t x)
875 {
876  return hlog(static_cast<__half>(x));
877 };
878 
879 template <>
880 inline __device__ float log<float>(float x)
881 {
882  return __logf(x);
883 };
884 
885 template <>
886 inline __device__ double log<double>(double x)
887 {
888  return log(x);
889 };
890 
891 template <typename T>
892 inline __device__ T pow(T x, T gamma)
893 {
894  return ck::type_convert<T>(powf(ck::type_convert<float>(x), ck::type_convert<float>(gamma)));
895 };
896 
897 template <>
898 inline __device__ float pow<float>(float x, float gamma)
899 {
900  return powf(x, gamma);
901 };
902 
903 template <>
904 inline __device__ double pow<double>(double x, double gamma)
905 {
906  return pow(x, gamma);
907 };
908 
909 template <typename T>
910 inline __device__ T expm1(T x)
911 {
912  return ck::type_convert<T>(expm1f(ck::type_convert<float>(x)));
913 };
914 
915 template <>
916 inline __device__ float expm1<float>(float x)
917 {
918  return expm1f(x);
919 };
920 
921 template <>
922 inline __device__ double expm1<double>(double x)
923 {
924  return expm1(x);
925 };
926 
927 template <typename T>
928 inline __device__ T cos(T x)
929 {
930  return ck::type_convert<T>(cosf(ck::type_convert<float>(x)));
931 };
932 
933 template <>
934 inline __device__ float cos<float>(float x)
935 {
936  return cosf(x);
937 };
938 
939 template <>
940 inline __device__ double cos<double>(double x)
941 {
942  return cos(x);
943 };
944 
945 } // namespace math
946 } // namespace ck
__host__ float cos< float >(float x)
Definition: math_v2.hpp:247
__device__ T asin(T x)
Definition: math_v2.hpp:660
__host__ T log(T x)
Definition: math_v2.hpp:409
__host__ int8_t neg< int8_t >(int8_t x)
Definition: math_v2.hpp:163
__host__ T cosh(T x)
Definition: math_v2.hpp:349
__host__ double asin< double >(double x)
Definition: math_v2.hpp:217
__host__ T exp(T x)
Definition: math_v2.hpp:391
__host__ double atan< double >(double x)
Definition: math_v2.hpp:181
__host__ double sinh< double >(double x)
Definition: math_v2.hpp:325
__host__ float floor< float >(float x)
Definition: math_v2.hpp:373
__host__ double ceil< double >(double x)
Definition: math_v2.hpp:343
__host__ float acosh< float >(float x)
Definition: math_v2.hpp:265
__host__ float log< float >(float x)
Definition: math_v2.hpp:415
__host__ T rcp(T x)
Definition: math_v2.hpp:385
__host__ double expm1< double >(double x)
Definition: math_v2.hpp:458
__host__ float asinh< float >(float x)
Definition: math_v2.hpp:229
__host__ T tan(T x)
Definition: math_v2.hpp:277
__device__ T sinh(T x)
Definition: math_v2.hpp:750
__host__ double log< double >(double x)
Definition: math_v2.hpp:421
__device__ half_t floor< half_t >(half_t x)
Definition: math_v2.hpp:828
__host__ T sin(T x)
Definition: math_v2.hpp:187
__host__ float asin< float >(float x)
Definition: math_v2.hpp:211
__host__ T expm1(T x)
Definition: math_v2.hpp:446
__host__ double exp< double >(double x)
Definition: math_v2.hpp:403
__device__ half_t log< half_t >(half_t x)
Definition: math_v2.hpp:874
__host__ double tan< double >(double x)
Definition: math_v2.hpp:289
__device__ T asinh(T x)
Definition: math_v2.hpp:678
__device__ half_t sin< half_t >(half_t x)
Definition: math_v2.hpp:654
__device__ T acos(T x)
Definition: math_v2.hpp:564
__host__ T atan(T x)
Definition: math_v2.hpp:169
__host__ float sin< float >(float x)
Definition: math_v2.hpp:193
__host__ float pow< float >(float x, float gamma)
Definition: math_v2.hpp:434
__host__ T acos(T x)
Definition: math_v2.hpp:121
__host__ double acos< double >(double x)
Definition: math_v2.hpp:133
__device__ T expm1(T x)
Definition: math_v2.hpp:910
__host__ double acosh< double >(double x)
Definition: math_v2.hpp:271
__host__ double cos< double >(double x)
Definition: math_v2.hpp:253
__host__ double atanh< double >(double x)
Definition: math_v2.hpp:307
__device__ T log(T x)
Definition: math_v2.hpp:868
__host__ T ceil(T x)
Definition: math_v2.hpp:331
__host__ T asin(T x)
Definition: math_v2.hpp:205
__host__ T pow(T x, T gamma)
Definition: math_v2.hpp:427
__host__ T neg(T x)
Definition: math_v2.hpp:139
__device__ T cos(T x)
Definition: math_v2.hpp:928
__host__ double pow< double >(double x, double gamma)
Definition: math_v2.hpp:440
__device__ T sin(T x)
Definition: math_v2.hpp:636
__device__ half_t neg< half_t >(half_t x)
Definition: math_v2.hpp:612
__host__ int32_t neg< int32_t >(int32_t x)
Definition: math_v2.hpp:157
__host__ float tan< float >(float x)
Definition: math_v2.hpp:283
__host__ float expm1< float >(float x)
Definition: math_v2.hpp:452
__device__ T atanh(T x)
Definition: math_v2.hpp:732
__device__ T ceil(T x)
Definition: math_v2.hpp:768
__host__ double floor< double >(double x)
Definition: math_v2.hpp:379
__host__ float acos< float >(float x)
Definition: math_v2.hpp:127
__device__ T exp(T x)
Definition: math_v2.hpp:844
__host__ float exp< float >(float x)
Definition: math_v2.hpp:397
__host__ float atanh< float >(float x)
Definition: math_v2.hpp:301
__host__ float neg< float >(float x)
Definition: math_v2.hpp:145
__host__ double asinh< double >(double x)
Definition: math_v2.hpp:235
__host__ T cos(T x)
Definition: math_v2.hpp:241
__host__ T floor(T x)
Definition: math_v2.hpp:367
__host__ double cosh< double >(double x)
Definition: math_v2.hpp:361
__device__ T cosh(T x)
Definition: math_v2.hpp:792
__host__ float tanh< float >(float x)
Definition: math_v2.hpp:109
__device__ T tan(T x)
Definition: math_v2.hpp:714
__host__ float atan< float >(float x)
Definition: math_v2.hpp:175
__host__ float cosh< float >(float x)
Definition: math_v2.hpp:355
__host__ T asinh(T x)
Definition: math_v2.hpp:223
__host__ float sinh< float >(float x)
Definition: math_v2.hpp:319
__device__ T acosh(T x)
Definition: math_v2.hpp:696
__host__ float ceil< float >(float x)
Definition: math_v2.hpp:337
__device__ half_t exp< half_t >(half_t x)
Definition: math_v2.hpp:850
__host__ T atanh(T x)
Definition: math_v2.hpp:295
__host__ double neg< double >(double x)
Definition: math_v2.hpp:151
__host__ T tanh(T x)
Definition: math_v2.hpp:103
__host__ double tanh< double >(double x)
Definition: math_v2.hpp:115
__device__ T tanh(T x)
Definition: math_v2.hpp:546
__host__ double sin< double >(double x)
Definition: math_v2.hpp:199
__device__ T pow(T x, T gamma)
Definition: math_v2.hpp:892
__host__ T acosh(T x)
Definition: math_v2.hpp:259
__device__ half_t ceil< half_t >(half_t x)
Definition: math_v2.hpp:786
__device__ T floor(T x)
Definition: math_v2.hpp:810
__device__ T atan(T x)
Definition: math_v2.hpp:618
__host__ T sinh(T x)
Definition: math_v2.hpp:313
Definition: ck.hpp:267
f8_fnuz_t f8_t
Definition: amd_ck_fp8.hpp:1737
_Float16 half_t
Definition: data_type.hpp:30
_BitInt(4) int4_t
Definition: data_type.hpp:31
unsigned short uint16_t
Definition: stdint.h:125
signed int int32_t
Definition: stdint.h:123
signed char int8_t
Definition: stdint.h:121