VTK-m  2.1
Math.h
Go to the documentation of this file.
1 //============================================================================
2 // Copyright (c) Kitware, Inc.
3 // All rights reserved.
4 // See LICENSE.txt for details.
5 //
6 // This software is distributed WITHOUT ANY WARRANTY; without even
7 // the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
8 // PURPOSE. See the above copyright notice for more information.
9 //============================================================================
10 // **** DO NOT EDIT THIS FILE!!! ****
11 // This file is automatically generated by Math.h.in
12 
13 #ifndef vtk_m_Math_h
14 #define vtk_m_Math_h
15 
16 #include <vtkm/TypeTraits.h>
17 #include <vtkm/Types.h>
18 #include <vtkm/VecTraits.h>
19 
20 #include <limits> // must be found with or without CUDA.
21 #ifndef VTKM_CUDA
22 #include <cmath>
23 #include <cstring>
24 #include <limits.h>
25 #include <math.h>
26 #include <stdlib.h>
27 #endif // !VTKM_CUDA
28 
29 #if !defined(VTKM_CUDA_DEVICE_PASS)
30 #define VTKM_USE_STL
31 #include <algorithm>
32 #endif
33 
34 #ifdef VTKM_MSVC
35 #include <intrin.h> // For bitwise intrinsics (__popcnt, etc)
36 #include <vtkm/internal/Windows.h> // for types used by MSVC intrinsics.
37 #ifndef VTKM_CUDA
38 #include <math.h>
39 #endif // VTKM_CUDA
40 #endif // VTKM_MSVC
41 
42 #define VTKM_CUDA_MATH_FUNCTION_32(func) func##f
43 #define VTKM_CUDA_MATH_FUNCTION_64(func) func
44 
45 namespace vtkm
46 {
47 
48 //-----------------------------------------------------------------------------
49 namespace detail
50 {
51 template <typename T>
52 struct FloatingPointReturnType
53 {
54  using ctype = typename vtkm::VecTraits<T>::ComponentType;
55  using representable_as_float_type =
56  std::integral_constant<bool,
57  ((sizeof(ctype) < sizeof(float)) ||
58  std::is_same<ctype, vtkm::Float32>::value)>;
59  using Type = typename std::
60  conditional<representable_as_float_type::value, vtkm::Float32, vtkm::Float64>::type;
61 };
62 } // namespace detail
63 
66 template <typename T = vtkm::Float64>
67 static constexpr inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type TwoPi()
68 {
69  using FT = typename detail::FloatingPointReturnType<T>::Type;
70  return static_cast<FT>(6.28318530717958647692528676655900576);
71 }
72 
75 template <typename T = vtkm::Float64>
76 static constexpr inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi()
77 {
78  using FT = typename detail::FloatingPointReturnType<T>::Type;
79  return static_cast<FT>(3.14159265358979323846264338327950288);
80 }
81 
84 template <typename T = vtkm::Float64>
85 static constexpr inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_2()
86 {
87  using FT = typename detail::FloatingPointReturnType<T>::Type;
88  return static_cast<FT>(1.57079632679489661923132169163975144);
89 }
90 
93 template <typename T = vtkm::Float64>
94 static constexpr inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_3()
95 {
96  using FT = typename detail::FloatingPointReturnType<T>::Type;
97  return static_cast<FT>(1.04719755119659774615421446109316762);
98 }
99 
102 template <typename T = vtkm::Float64>
103 static constexpr inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_4()
104 {
105  using FT = typename detail::FloatingPointReturnType<T>::Type;
106  return static_cast<FT>(0.78539816339744830961566084581987572);
107 }
108 
111 template <typename T = vtkm::Float64>
112 static constexpr inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_180()
113 {
114  using FT = typename detail::FloatingPointReturnType<T>::Type;
115  return static_cast<FT>(0.01745329251994329547437168059786927);
116 }
117 
120 static constexpr inline VTKM_EXEC_CONT vtkm::Float32 TwoPif()
121 {
122  return TwoPi<vtkm::Float32>();
123 }
124 
127 static constexpr inline VTKM_EXEC_CONT vtkm::Float32 Pif()
128 {
129  return Pi<vtkm::Float32>();
130 }
131 
134 static constexpr inline VTKM_EXEC_CONT vtkm::Float32 Pi_2f()
135 {
136  return Pi_2<vtkm::Float32>();
137 }
138 
141 static constexpr inline VTKM_EXEC_CONT vtkm::Float32 Pi_3f()
142 {
143  return Pi_3<vtkm::Float32>();
144 }
145 
148 static constexpr inline VTKM_EXEC_CONT vtkm::Float32 Pi_4f()
149 {
150  return Pi_4<vtkm::Float32>();
151 }
152 
155 static constexpr inline VTKM_EXEC_CONT vtkm::Float32 Pi_180f()
156 {
157  return Pi_180<vtkm::Float32>();
158 }
159 
160 // clang-format off
161 
165 
167 {
168 #ifdef VTKM_CUDA
169  return VTKM_CUDA_MATH_FUNCTION_32(sin)(x);
170 #else
171  return std::sin(x);
172 #endif
173 }
174 
176 {
177 #ifdef VTKM_CUDA
178  return VTKM_CUDA_MATH_FUNCTION_64(sin)(x);
179 #else
180  return std::sin(x);
181 #endif
182 }
183 template <typename T>
184 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Sin(const T& x)
185 {
186  using RT = typename detail::FloatingPointReturnType<T>::Type;
187  return vtkm::Sin(static_cast<RT>(x));
188 }
189 template <typename T, vtkm::IdComponent N>
191  const vtkm::Vec<T, N>& x)
192 {
194  for (vtkm::IdComponent index = 0; index < N; index++)
195  {
196  result[index] = vtkm::Sin(x[index]);
197  }
198  return result;
199 }
200 template <typename T>
202  const vtkm::Vec<T, 4>& x)
203 {
205  vtkm::Sin(x[0]), vtkm::Sin(x[1]), vtkm::Sin(x[2]), vtkm::Sin(x[3]));
206 }
207 template <typename T>
209  const vtkm::Vec<T, 3>& x)
210 {
212  vtkm::Sin(x[0]), vtkm::Sin(x[1]), vtkm::Sin(x[2]));
213 }
214 template <typename T>
216  const vtkm::Vec<T, 2>& x)
217 {
219  vtkm::Sin(x[1]));
220 }
222 
226 
228 {
229 #ifdef VTKM_CUDA
230  return VTKM_CUDA_MATH_FUNCTION_32(cos)(x);
231 #else
232  return std::cos(x);
233 #endif
234 }
235 
237 {
238 #ifdef VTKM_CUDA
239  return VTKM_CUDA_MATH_FUNCTION_64(cos)(x);
240 #else
241  return std::cos(x);
242 #endif
243 }
244 template <typename T>
245 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Cos(const T& x)
246 {
247  using RT = typename detail::FloatingPointReturnType<T>::Type;
248  return vtkm::Cos(static_cast<RT>(x));
249 }
250 template <typename T, vtkm::IdComponent N>
252  const vtkm::Vec<T, N>& x)
253 {
255  for (vtkm::IdComponent index = 0; index < N; index++)
256  {
257  result[index] = vtkm::Cos(x[index]);
258  }
259  return result;
260 }
261 template <typename T>
263  const vtkm::Vec<T, 4>& x)
264 {
266  vtkm::Cos(x[0]), vtkm::Cos(x[1]), vtkm::Cos(x[2]), vtkm::Cos(x[3]));
267 }
268 template <typename T>
270  const vtkm::Vec<T, 3>& x)
271 {
273  vtkm::Cos(x[0]), vtkm::Cos(x[1]), vtkm::Cos(x[2]));
274 }
275 template <typename T>
277  const vtkm::Vec<T, 2>& x)
278 {
280  vtkm::Cos(x[1]));
281 }
283 
287 
289 {
290 #ifdef VTKM_CUDA
291  return VTKM_CUDA_MATH_FUNCTION_32(tan)(x);
292 #else
293  return std::tan(x);
294 #endif
295 }
296 
298 {
299 #ifdef VTKM_CUDA
300  return VTKM_CUDA_MATH_FUNCTION_64(tan)(x);
301 #else
302  return std::tan(x);
303 #endif
304 }
305 template <typename T>
306 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Tan(const T& x)
307 {
308  using RT = typename detail::FloatingPointReturnType<T>::Type;
309  return vtkm::Tan(static_cast<RT>(x));
310 }
311 template <typename T, vtkm::IdComponent N>
313  const vtkm::Vec<T, N>& x)
314 {
316  for (vtkm::IdComponent index = 0; index < N; index++)
317  {
318  result[index] = vtkm::Tan(x[index]);
319  }
320  return result;
321 }
322 template <typename T>
324  const vtkm::Vec<T, 4>& x)
325 {
327  vtkm::Tan(x[0]), vtkm::Tan(x[1]), vtkm::Tan(x[2]), vtkm::Tan(x[3]));
328 }
329 template <typename T>
331  const vtkm::Vec<T, 3>& x)
332 {
334  vtkm::Tan(x[0]), vtkm::Tan(x[1]), vtkm::Tan(x[2]));
335 }
336 template <typename T>
338  const vtkm::Vec<T, 2>& x)
339 {
341  vtkm::Tan(x[1]));
342 }
344 
348 
350 {
351 #ifdef VTKM_CUDA
352  return VTKM_CUDA_MATH_FUNCTION_32(asin)(x);
353 #else
354  return std::asin(x);
355 #endif
356 }
357 
359 {
360 #ifdef VTKM_CUDA
361  return VTKM_CUDA_MATH_FUNCTION_64(asin)(x);
362 #else
363  return std::asin(x);
364 #endif
365 }
366 template <typename T>
367 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ASin(const T& x)
368 {
369  using RT = typename detail::FloatingPointReturnType<T>::Type;
370  return vtkm::ASin(static_cast<RT>(x));
371 }
372 template <typename T, vtkm::IdComponent N>
374  const vtkm::Vec<T, N>& x)
375 {
377  for (vtkm::IdComponent index = 0; index < N; index++)
378  {
379  result[index] = vtkm::ASin(x[index]);
380  }
381  return result;
382 }
383 template <typename T>
385  const vtkm::Vec<T, 4>& x)
386 {
388  vtkm::ASin(x[0]), vtkm::ASin(x[1]), vtkm::ASin(x[2]), vtkm::ASin(x[3]));
389 }
390 template <typename T>
392  const vtkm::Vec<T, 3>& x)
393 {
395  vtkm::ASin(x[0]), vtkm::ASin(x[1]), vtkm::ASin(x[2]));
396 }
397 template <typename T>
399  const vtkm::Vec<T, 2>& x)
400 {
402  vtkm::ASin(x[1]));
403 }
405 
409 
411 {
412 #ifdef VTKM_CUDA
413  return VTKM_CUDA_MATH_FUNCTION_32(acos)(x);
414 #else
415  return std::acos(x);
416 #endif
417 }
418 
420 {
421 #ifdef VTKM_CUDA
422  return VTKM_CUDA_MATH_FUNCTION_64(acos)(x);
423 #else
424  return std::acos(x);
425 #endif
426 }
427 template <typename T>
428 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ACos(const T& x)
429 {
430  using RT = typename detail::FloatingPointReturnType<T>::Type;
431  return vtkm::ACos(static_cast<RT>(x));
432 }
433 template <typename T, vtkm::IdComponent N>
435  const vtkm::Vec<T, N>& x)
436 {
438  for (vtkm::IdComponent index = 0; index < N; index++)
439  {
440  result[index] = vtkm::ACos(x[index]);
441  }
442  return result;
443 }
444 template <typename T>
446  const vtkm::Vec<T, 4>& x)
447 {
449  vtkm::ACos(x[0]), vtkm::ACos(x[1]), vtkm::ACos(x[2]), vtkm::ACos(x[3]));
450 }
451 template <typename T>
453  const vtkm::Vec<T, 3>& x)
454 {
456  vtkm::ACos(x[0]), vtkm::ACos(x[1]), vtkm::ACos(x[2]));
457 }
458 template <typename T>
460  const vtkm::Vec<T, 2>& x)
461 {
463  vtkm::ACos(x[1]));
464 }
466 
470 
472 {
473 #ifdef VTKM_CUDA
474  return VTKM_CUDA_MATH_FUNCTION_32(atan)(x);
475 #else
476  return std::atan(x);
477 #endif
478 }
479 
481 {
482 #ifdef VTKM_CUDA
483  return VTKM_CUDA_MATH_FUNCTION_64(atan)(x);
484 #else
485  return std::atan(x);
486 #endif
487 }
488 template <typename T>
489 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ATan(const T& x)
490 {
491  using RT = typename detail::FloatingPointReturnType<T>::Type;
492  return vtkm::ATan(static_cast<RT>(x));
493 }
494 template <typename T, vtkm::IdComponent N>
496  const vtkm::Vec<T, N>& x)
497 {
499  for (vtkm::IdComponent index = 0; index < N; index++)
500  {
501  result[index] = vtkm::ATan(x[index]);
502  }
503  return result;
504 }
505 template <typename T>
507  const vtkm::Vec<T, 4>& x)
508 {
510  vtkm::ATan(x[0]), vtkm::ATan(x[1]), vtkm::ATan(x[2]), vtkm::ATan(x[3]));
511 }
512 template <typename T>
514  const vtkm::Vec<T, 3>& x)
515 {
517  vtkm::ATan(x[0]), vtkm::ATan(x[1]), vtkm::ATan(x[2]));
518 }
519 template <typename T>
521  const vtkm::Vec<T, 2>& x)
522 {
524  vtkm::ATan(x[1]));
525 }
527 
533 {
534 #ifdef VTKM_CUDA
535  return VTKM_CUDA_MATH_FUNCTION_32(atan2)(x, y);
536 #else
537  return std::atan2(x, y);
538 #endif
539 }
541 {
542 #ifdef VTKM_CUDA
543  return VTKM_CUDA_MATH_FUNCTION_64(atan2)(x, y);
544 #else
545  return std::atan2(x, y);
546 #endif
547 }
549 
553 
555 {
556 #ifdef VTKM_CUDA
557  return VTKM_CUDA_MATH_FUNCTION_32(sinh)(x);
558 #else
559  return std::sinh(x);
560 #endif
561 }
562 
564 {
565 #ifdef VTKM_CUDA
566  return VTKM_CUDA_MATH_FUNCTION_64(sinh)(x);
567 #else
568  return std::sinh(x);
569 #endif
570 }
571 template <typename T>
572 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type SinH(const T& x)
573 {
574  using RT = typename detail::FloatingPointReturnType<T>::Type;
575  return vtkm::SinH(static_cast<RT>(x));
576 }
577 template <typename T, vtkm::IdComponent N>
579  const vtkm::Vec<T, N>& x)
580 {
582  for (vtkm::IdComponent index = 0; index < N; index++)
583  {
584  result[index] = vtkm::SinH(x[index]);
585  }
586  return result;
587 }
588 template <typename T>
590  const vtkm::Vec<T, 4>& x)
591 {
593  vtkm::SinH(x[0]), vtkm::SinH(x[1]), vtkm::SinH(x[2]), vtkm::SinH(x[3]));
594 }
595 template <typename T>
597  const vtkm::Vec<T, 3>& x)
598 {
600  vtkm::SinH(x[0]), vtkm::SinH(x[1]), vtkm::SinH(x[2]));
601 }
602 template <typename T>
604  const vtkm::Vec<T, 2>& x)
605 {
607  vtkm::SinH(x[1]));
608 }
610 
614 
616 {
617 #ifdef VTKM_CUDA
618  return VTKM_CUDA_MATH_FUNCTION_32(cosh)(x);
619 #else
620  return std::cosh(x);
621 #endif
622 }
623 
625 {
626 #ifdef VTKM_CUDA
627  return VTKM_CUDA_MATH_FUNCTION_64(cosh)(x);
628 #else
629  return std::cosh(x);
630 #endif
631 }
632 template <typename T>
633 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type CosH(const T& x)
634 {
635  using RT = typename detail::FloatingPointReturnType<T>::Type;
636  return vtkm::CosH(static_cast<RT>(x));
637 }
638 template <typename T, vtkm::IdComponent N>
640  const vtkm::Vec<T, N>& x)
641 {
643  for (vtkm::IdComponent index = 0; index < N; index++)
644  {
645  result[index] = vtkm::CosH(x[index]);
646  }
647  return result;
648 }
649 template <typename T>
651  const vtkm::Vec<T, 4>& x)
652 {
654  vtkm::CosH(x[0]), vtkm::CosH(x[1]), vtkm::CosH(x[2]), vtkm::CosH(x[3]));
655 }
656 template <typename T>
658  const vtkm::Vec<T, 3>& x)
659 {
661  vtkm::CosH(x[0]), vtkm::CosH(x[1]), vtkm::CosH(x[2]));
662 }
663 template <typename T>
665  const vtkm::Vec<T, 2>& x)
666 {
668  vtkm::CosH(x[1]));
669 }
671 
675 
677 {
678 #ifdef VTKM_CUDA
679  return VTKM_CUDA_MATH_FUNCTION_32(tanh)(x);
680 #else
681  return std::tanh(x);
682 #endif
683 }
684 
686 {
687 #ifdef VTKM_CUDA
688  return VTKM_CUDA_MATH_FUNCTION_64(tanh)(x);
689 #else
690  return std::tanh(x);
691 #endif
692 }
693 template <typename T>
694 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type TanH(const T& x)
695 {
696  using RT = typename detail::FloatingPointReturnType<T>::Type;
697  return vtkm::TanH(static_cast<RT>(x));
698 }
699 template <typename T, vtkm::IdComponent N>
701  const vtkm::Vec<T, N>& x)
702 {
704  for (vtkm::IdComponent index = 0; index < N; index++)
705  {
706  result[index] = vtkm::TanH(x[index]);
707  }
708  return result;
709 }
710 template <typename T>
712  const vtkm::Vec<T, 4>& x)
713 {
715  vtkm::TanH(x[0]), vtkm::TanH(x[1]), vtkm::TanH(x[2]), vtkm::TanH(x[3]));
716 }
717 template <typename T>
719  const vtkm::Vec<T, 3>& x)
720 {
722  vtkm::TanH(x[0]), vtkm::TanH(x[1]), vtkm::TanH(x[2]));
723 }
724 template <typename T>
726  const vtkm::Vec<T, 2>& x)
727 {
729  vtkm::TanH(x[1]));
730 }
732 
736 
738 {
739 #ifdef VTKM_CUDA
740  return VTKM_CUDA_MATH_FUNCTION_32(asinh)(x);
741 #else
742  return std::asinh(x);
743 #endif
744 }
745 
747 {
748 #ifdef VTKM_CUDA
749  return VTKM_CUDA_MATH_FUNCTION_64(asinh)(x);
750 #else
751  return std::asinh(x);
752 #endif
753 }
754 template <typename T>
755 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ASinH(const T& x)
756 {
757  using RT = typename detail::FloatingPointReturnType<T>::Type;
758  return vtkm::ASinH(static_cast<RT>(x));
759 }
760 template <typename T, vtkm::IdComponent N>
762  const vtkm::Vec<T, N>& x)
763 {
765  for (vtkm::IdComponent index = 0; index < N; index++)
766  {
767  result[index] = vtkm::ASinH(x[index]);
768  }
769  return result;
770 }
771 template <typename T>
773  const vtkm::Vec<T, 4>& x)
774 {
776  vtkm::ASinH(x[0]), vtkm::ASinH(x[1]), vtkm::ASinH(x[2]), vtkm::ASinH(x[3]));
777 }
778 template <typename T>
780  const vtkm::Vec<T, 3>& x)
781 {
783  vtkm::ASinH(x[0]), vtkm::ASinH(x[1]), vtkm::ASinH(x[2]));
784 }
785 template <typename T>
787  const vtkm::Vec<T, 2>& x)
788 {
790  vtkm::ASinH(x[1]));
791 }
793 
797 
799 {
800 #ifdef VTKM_CUDA
801  return VTKM_CUDA_MATH_FUNCTION_32(acosh)(x);
802 #else
803  return std::acosh(x);
804 #endif
805 }
806 
808 {
809 #ifdef VTKM_CUDA
810  return VTKM_CUDA_MATH_FUNCTION_64(acosh)(x);
811 #else
812  return std::acosh(x);
813 #endif
814 }
815 template <typename T>
816 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ACosH(const T& x)
817 {
818  using RT = typename detail::FloatingPointReturnType<T>::Type;
819  return vtkm::ACosH(static_cast<RT>(x));
820 }
821 template <typename T, vtkm::IdComponent N>
823  const vtkm::Vec<T, N>& x)
824 {
826  for (vtkm::IdComponent index = 0; index < N; index++)
827  {
828  result[index] = vtkm::ACosH(x[index]);
829  }
830  return result;
831 }
832 template <typename T>
834  const vtkm::Vec<T, 4>& x)
835 {
837  vtkm::ACosH(x[0]), vtkm::ACosH(x[1]), vtkm::ACosH(x[2]), vtkm::ACosH(x[3]));
838 }
839 template <typename T>
841  const vtkm::Vec<T, 3>& x)
842 {
844  vtkm::ACosH(x[0]), vtkm::ACosH(x[1]), vtkm::ACosH(x[2]));
845 }
846 template <typename T>
848  const vtkm::Vec<T, 2>& x)
849 {
851  vtkm::ACosH(x[1]));
852 }
854 
858 
860 {
861 #ifdef VTKM_CUDA
862  return VTKM_CUDA_MATH_FUNCTION_32(atanh)(x);
863 #else
864  return std::atanh(x);
865 #endif
866 }
867 
869 {
870 #ifdef VTKM_CUDA
871  return VTKM_CUDA_MATH_FUNCTION_64(atanh)(x);
872 #else
873  return std::atanh(x);
874 #endif
875 }
876 template <typename T>
877 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ATanH(const T& x)
878 {
879  using RT = typename detail::FloatingPointReturnType<T>::Type;
880  return vtkm::ATanH(static_cast<RT>(x));
881 }
882 template <typename T, vtkm::IdComponent N>
884  const vtkm::Vec<T, N>& x)
885 {
887  for (vtkm::IdComponent index = 0; index < N; index++)
888  {
889  result[index] = vtkm::ATanH(x[index]);
890  }
891  return result;
892 }
893 template <typename T>
895  const vtkm::Vec<T, 4>& x)
896 {
898  vtkm::ATanH(x[0]), vtkm::ATanH(x[1]), vtkm::ATanH(x[2]), vtkm::ATanH(x[3]));
899 }
900 template <typename T>
902  const vtkm::Vec<T, 3>& x)
903 {
905  vtkm::ATanH(x[0]), vtkm::ATanH(x[1]), vtkm::ATanH(x[2]));
906 }
907 template <typename T>
909  const vtkm::Vec<T, 2>& x)
910 {
912  vtkm::ATanH(x[1]));
913 }
915 
916 //-----------------------------------------------------------------------------
921 {
922 #ifdef VTKM_CUDA
923  return VTKM_CUDA_MATH_FUNCTION_32(pow)(x, y);
924 #else
925  return std::pow(x, y);
926 #endif
927 }
929 {
930 #ifdef VTKM_CUDA
931  return VTKM_CUDA_MATH_FUNCTION_64(pow)(x, y);
932 #else
933  return std::pow(x, y);
934 #endif
935 }
937 
942 
944 {
945 #ifdef VTKM_CUDA
946  return VTKM_CUDA_MATH_FUNCTION_32(sqrt)(x);
947 #else
948  return std::sqrt(x);
949 #endif
950 }
951 
953 {
954 #ifdef VTKM_CUDA
955  return VTKM_CUDA_MATH_FUNCTION_64(sqrt)(x);
956 #else
957  return std::sqrt(x);
958 #endif
959 }
960 template <typename T>
961 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Sqrt(const T& x)
962 {
963  using RT = typename detail::FloatingPointReturnType<T>::Type;
964  return vtkm::Sqrt(static_cast<RT>(x));
965 }
966 template <typename T, vtkm::IdComponent N>
968  const vtkm::Vec<T, N>& x)
969 {
971  for (vtkm::IdComponent index = 0; index < N; index++)
972  {
973  result[index] = vtkm::Sqrt(x[index]);
974  }
975  return result;
976 }
977 template <typename T>
979  const vtkm::Vec<T, 4>& x)
980 {
982  vtkm::Sqrt(x[0]), vtkm::Sqrt(x[1]), vtkm::Sqrt(x[2]), vtkm::Sqrt(x[3]));
983 }
984 template <typename T>
986  const vtkm::Vec<T, 3>& x)
987 {
989  vtkm::Sqrt(x[0]), vtkm::Sqrt(x[1]), vtkm::Sqrt(x[2]));
990 }
991 template <typename T>
993  const vtkm::Vec<T, 2>& x)
994 {
996  vtkm::Sqrt(x[1]));
997 }
999 
1006 #ifdef VTKM_CUDA
1007 static inline VTKM_EXEC_CONT vtkm::Float32 RSqrt(vtkm::Float32 x)
1008 {
1009  return rsqrtf(x);
1010 }
1011 static inline VTKM_EXEC_CONT vtkm::Float64 RSqrt(vtkm::Float64 x)
1012 {
1013  return rsqrt(x);
1014 }
1015 template <typename T>
1016 static inline VTKM_EXEC_CONT vtkm::Float64 RSqrt(T x)
1017 {
1018  return rsqrt(static_cast<vtkm::Float64>(x));
1019 }
1020 #else // !VTKM_CUDA
1021 static inline VTKM_EXEC_CONT vtkm::Float32 RSqrt(vtkm::Float32 x)
1022 {
1023  return 1 / vtkm::Sqrt(x);
1024 }
1025 static inline VTKM_EXEC_CONT vtkm::Float64 RSqrt(vtkm::Float64 x)
1026 {
1027  return 1 / vtkm::Sqrt(x);
1028 }
1029 template <typename T>
1030 static inline VTKM_EXEC_CONT vtkm::Float64 RSqrt(T x)
1031 {
1032  return 1 / static_cast<vtkm::Float64>(x);
1033 }
1034 #endif // !VTKM_CUDA
1035 
1036 template <typename T, vtkm::IdComponent N>
1038  const vtkm::Vec<T, N>& x)
1039 {
1041  for (vtkm::IdComponent index = 0; index < N; index++)
1042  {
1043  result[index] = vtkm::RSqrt(x[index]);
1044  }
1045  return result;
1046 }
1047 template <typename T>
1049  const vtkm::Vec<T, 4>& x)
1050 {
1052  vtkm::RSqrt(x[0]), vtkm::RSqrt(x[1]), vtkm::RSqrt(x[2]), vtkm::RSqrt(x[3]));
1053 }
1054 template <typename T>
1056  const vtkm::Vec<T, 3>& x)
1057 {
1059  vtkm::RSqrt(x[0]), vtkm::RSqrt(x[1]), vtkm::RSqrt(x[2]));
1060 }
1061 template <typename T>
1063  const vtkm::Vec<T, 2>& x)
1064 {
1066  vtkm::RSqrt(x[1]));
1067 }
1069 
1073 
1075 {
1076 #ifdef VTKM_CUDA
1077  return VTKM_CUDA_MATH_FUNCTION_32(cbrt)(x);
1078 #else
1079  return std::cbrt(x);
1080 #endif
1081 }
1082 
1084 {
1085 #ifdef VTKM_CUDA
1086  return VTKM_CUDA_MATH_FUNCTION_64(cbrt)(x);
1087 #else
1088  return std::cbrt(x);
1089 #endif
1090 }
1091 template <typename T>
1092 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Cbrt(const T& x)
1093 {
1094  using RT = typename detail::FloatingPointReturnType<T>::Type;
1095  return vtkm::Cbrt(static_cast<RT>(x));
1096 }
1097 template <typename T, vtkm::IdComponent N>
1099  const vtkm::Vec<T, N>& x)
1100 {
1102  for (vtkm::IdComponent index = 0; index < N; index++)
1103  {
1104  result[index] = vtkm::Cbrt(x[index]);
1105  }
1106  return result;
1107 }
1108 template <typename T>
1110  const vtkm::Vec<T, 4>& x)
1111 {
1113  vtkm::Cbrt(x[0]), vtkm::Cbrt(x[1]), vtkm::Cbrt(x[2]), vtkm::Cbrt(x[3]));
1114 }
1115 template <typename T>
1117  const vtkm::Vec<T, 3>& x)
1118 {
1120  vtkm::Cbrt(x[0]), vtkm::Cbrt(x[1]), vtkm::Cbrt(x[2]));
1121 }
1122 template <typename T>
1124  const vtkm::Vec<T, 2>& x)
1125 {
1127  vtkm::Cbrt(x[1]));
1128 }
1130 
1137 #ifdef VTKM_CUDA
1138 static inline VTKM_EXEC_CONT vtkm::Float32 RCbrt(vtkm::Float32 x)
1139 {
1140  return rcbrtf(x);
1141 }
1142 static inline VTKM_EXEC_CONT vtkm::Float64 RCbrt(vtkm::Float64 x)
1143 {
1144  return rcbrt(x);
1145 }
1146 template <typename T>
1147 static inline VTKM_EXEC_CONT vtkm::Float64 RCbrt(T x)
1148 {
1149  return rcbrt(static_cast<vtkm::Float64>(x));
1150 }
1151 #else // !VTKM_CUDA
1152 static inline VTKM_EXEC_CONT vtkm::Float32 RCbrt(vtkm::Float32 x)
1153 {
1154  return 1 / vtkm::Cbrt(x);
1155 }
1156 static inline VTKM_EXEC_CONT vtkm::Float64 RCbrt(vtkm::Float64 x)
1157 {
1158  return 1 / vtkm::Cbrt(x);
1159 }
1160 template <typename T>
1161 static inline VTKM_EXEC_CONT vtkm::Float64 RCbrt(T x)
1162 {
1163  return 1 / vtkm::Cbrt(static_cast<vtkm::Float64>(x));
1164 }
1165 #endif // !VTKM_CUDA
1166 
1167 template <typename T, vtkm::IdComponent N>
1169  const vtkm::Vec<T, N>& x)
1170 {
1172  for (vtkm::IdComponent index = 0; index < N; index++)
1173  {
1174  result[index] = vtkm::RCbrt(x[index]);
1175  }
1176  return result;
1177 }
1178 template <typename T>
1180  const vtkm::Vec<T, 4>& x)
1181 {
1183  vtkm::RCbrt(x[0]), vtkm::RCbrt(x[1]), vtkm::RCbrt(x[2]), vtkm::RCbrt(x[3]));
1184 }
1185 template <typename T>
1187  const vtkm::Vec<T, 3>& x)
1188 {
1190  vtkm::RCbrt(x[0]), vtkm::RCbrt(x[1]), vtkm::RCbrt(x[2]));
1191 }
1192 template <typename T>
1194  const vtkm::Vec<T, 2>& x)
1195 {
1197  vtkm::RCbrt(x[1]));
1198 }
1200 
1204 
1206 {
1207 #ifdef VTKM_CUDA
1208  return VTKM_CUDA_MATH_FUNCTION_32(exp)(x);
1209 #else
1210  return std::exp(x);
1211 #endif
1212 }
1213 
1215 {
1216 #ifdef VTKM_CUDA
1217  return VTKM_CUDA_MATH_FUNCTION_64(exp)(x);
1218 #else
1219  return std::exp(x);
1220 #endif
1221 }
1222 template <typename T>
1223 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Exp(const T& x)
1224 {
1225  using RT = typename detail::FloatingPointReturnType<T>::Type;
1226  return vtkm::Exp(static_cast<RT>(x));
1227 }
1228 template <typename T, vtkm::IdComponent N>
1230  const vtkm::Vec<T, N>& x)
1231 {
1233  for (vtkm::IdComponent index = 0; index < N; index++)
1234  {
1235  result[index] = vtkm::Exp(x[index]);
1236  }
1237  return result;
1238 }
1239 template <typename T>
1241  const vtkm::Vec<T, 4>& x)
1242 {
1244  vtkm::Exp(x[0]), vtkm::Exp(x[1]), vtkm::Exp(x[2]), vtkm::Exp(x[3]));
1245 }
1246 template <typename T>
1248  const vtkm::Vec<T, 3>& x)
1249 {
1251  vtkm::Exp(x[0]), vtkm::Exp(x[1]), vtkm::Exp(x[2]));
1252 }
1253 template <typename T>
1255  const vtkm::Vec<T, 2>& x)
1256 {
1258  vtkm::Exp(x[1]));
1259 }
1261 
1265 
1267 {
1268 #ifdef VTKM_CUDA
1269  return VTKM_CUDA_MATH_FUNCTION_32(exp2)(x);
1270 #else
1271  return std::exp2(x);
1272 #endif
1273 }
1274 
1276 {
1277 #ifdef VTKM_CUDA
1278  return VTKM_CUDA_MATH_FUNCTION_64(exp2)(x);
1279 #else
1280  return std::exp2(x);
1281 #endif
1282 }
1283 template <typename T>
1284 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Exp2(const T& x)
1285 {
1286  using RT = typename detail::FloatingPointReturnType<T>::Type;
1287  return vtkm::Exp2(static_cast<RT>(x));
1288 }
1289 template <typename T, vtkm::IdComponent N>
1291  const vtkm::Vec<T, N>& x)
1292 {
1294  for (vtkm::IdComponent index = 0; index < N; index++)
1295  {
1296  result[index] = vtkm::Exp2(x[index]);
1297  }
1298  return result;
1299 }
1300 template <typename T>
1302  const vtkm::Vec<T, 4>& x)
1303 {
1305  vtkm::Exp2(x[0]), vtkm::Exp2(x[1]), vtkm::Exp2(x[2]), vtkm::Exp2(x[3]));
1306 }
1307 template <typename T>
1309  const vtkm::Vec<T, 3>& x)
1310 {
1312  vtkm::Exp2(x[0]), vtkm::Exp2(x[1]), vtkm::Exp2(x[2]));
1313 }
1314 template <typename T>
1316  const vtkm::Vec<T, 2>& x)
1317 {
1319  vtkm::Exp2(x[1]));
1320 }
1322 
1327 
1329 {
1330 #ifdef VTKM_CUDA
1331  return VTKM_CUDA_MATH_FUNCTION_32(expm1)(x);
1332 #else
1333  return std::expm1(x);
1334 #endif
1335 }
1336 
1338 {
1339 #ifdef VTKM_CUDA
1340  return VTKM_CUDA_MATH_FUNCTION_64(expm1)(x);
1341 #else
1342  return std::expm1(x);
1343 #endif
1344 }
1345 template <typename T>
1346 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ExpM1(const T& x)
1347 {
1348  using RT = typename detail::FloatingPointReturnType<T>::Type;
1349  return vtkm::ExpM1(static_cast<RT>(x));
1350 }
1351 template <typename T, vtkm::IdComponent N>
1353  const vtkm::Vec<T, N>& x)
1354 {
1356  for (vtkm::IdComponent index = 0; index < N; index++)
1357  {
1358  result[index] = vtkm::ExpM1(x[index]);
1359  }
1360  return result;
1361 }
1362 template <typename T>
1364  const vtkm::Vec<T, 4>& x)
1365 {
1367  vtkm::ExpM1(x[0]), vtkm::ExpM1(x[1]), vtkm::ExpM1(x[2]), vtkm::ExpM1(x[3]));
1368 }
1369 template <typename T>
1371  const vtkm::Vec<T, 3>& x)
1372 {
1374  vtkm::ExpM1(x[0]), vtkm::ExpM1(x[1]), vtkm::ExpM1(x[2]));
1375 }
1376 template <typename T>
1378  const vtkm::Vec<T, 2>& x)
1379 {
1381  vtkm::ExpM1(x[1]));
1382 }
1384 
1388 #ifdef VTKM_CUDA
1389 static inline VTKM_EXEC_CONT vtkm::Float32 Exp10(vtkm::Float32 x)
1390 {
1391  return exp10f(x);
1392 }
1393 static inline VTKM_EXEC_CONT vtkm::Float64 Exp10(vtkm::Float64 x)
1394 {
1395  return exp10(x);
1396 }
1397 template <typename T>
1398 static inline VTKM_EXEC_CONT vtkm::Float64 Exp10(T x)
1399 {
1400  return exp10(static_cast<vtkm::Float64>(x));
1401 }
1402 #else // !VTKM_CUDA
1403 static inline VTKM_EXEC_CONT vtkm::Float32 Exp10(vtkm::Float32 x)
1404 {
1405  return vtkm::Pow(10, x);
1406 }
1407 static inline VTKM_EXEC_CONT vtkm::Float64 Exp10(vtkm::Float64 x)
1408 {
1409  return vtkm::Pow(10, x);
1410 }
1411 template <typename T>
1412 static inline VTKM_EXEC_CONT vtkm::Float64 Exp10(T x)
1413 {
1414  return vtkm::Pow(10, static_cast<vtkm::Float64>(x));
1415 }
1416 #endif // !VTKM_CUDA
1417 
1418 template <typename T, vtkm::IdComponent N>
1420  const vtkm::Vec<T, N>& x)
1421 {
1423  for (vtkm::IdComponent index = 0; index < N; index++)
1424  {
1425  result[index] = vtkm::Exp10(x[index]);
1426  }
1427  return result;
1428 }
1429 template <typename T>
1431  const vtkm::Vec<T, 4>& x)
1432 {
1434  vtkm::Exp10(x[0]), vtkm::Exp10(x[1]), vtkm::Exp10(x[2]), vtkm::Exp10(x[3]));
1435 }
1436 template <typename T>
1438  const vtkm::Vec<T, 3>& x)
1439 {
1441  vtkm::Exp10(x[0]), vtkm::Exp10(x[1]), vtkm::Exp10(x[2]));
1442 }
1443 template <typename T>
1445  const vtkm::Vec<T, 2>& x)
1446 {
1448  vtkm::Exp10(x[1]));
1449 }
1451 
1455 
1457 {
1458 #ifdef VTKM_CUDA
1459  return VTKM_CUDA_MATH_FUNCTION_32(log)(x);
1460 #else
1461  return std::log(x);
1462 #endif
1463 }
1464 
1466 {
1467 #ifdef VTKM_CUDA
1468  return VTKM_CUDA_MATH_FUNCTION_64(log)(x);
1469 #else
1470  return std::log(x);
1471 #endif
1472 }
1473 template <typename T>
1474 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log(const T& x)
1475 {
1476  using RT = typename detail::FloatingPointReturnType<T>::Type;
1477  return vtkm::Log(static_cast<RT>(x));
1478 }
1479 template <typename T, vtkm::IdComponent N>
1481  const vtkm::Vec<T, N>& x)
1482 {
1484  for (vtkm::IdComponent index = 0; index < N; index++)
1485  {
1486  result[index] = vtkm::Log(x[index]);
1487  }
1488  return result;
1489 }
1490 template <typename T>
1492  const vtkm::Vec<T, 4>& x)
1493 {
1495  vtkm::Log(x[0]), vtkm::Log(x[1]), vtkm::Log(x[2]), vtkm::Log(x[3]));
1496 }
1497 template <typename T>
1499  const vtkm::Vec<T, 3>& x)
1500 {
1502  vtkm::Log(x[0]), vtkm::Log(x[1]), vtkm::Log(x[2]));
1503 }
1504 template <typename T>
1506  const vtkm::Vec<T, 2>& x)
1507 {
1509  vtkm::Log(x[1]));
1510 }
1512 
1516 
1518 {
1519 #ifdef VTKM_CUDA
1520  return VTKM_CUDA_MATH_FUNCTION_32(log2)(x);
1521 #else
1522  return std::log2(x);
1523 #endif
1524 }
1525 
1527 {
1528 #ifdef VTKM_CUDA
1529  return VTKM_CUDA_MATH_FUNCTION_64(log2)(x);
1530 #else
1531  return std::log2(x);
1532 #endif
1533 }
1534 template <typename T>
1535 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log2(const T& x)
1536 {
1537  using RT = typename detail::FloatingPointReturnType<T>::Type;
1538  return vtkm::Log2(static_cast<RT>(x));
1539 }
1540 template <typename T, vtkm::IdComponent N>
1542  const vtkm::Vec<T, N>& x)
1543 {
1545  for (vtkm::IdComponent index = 0; index < N; index++)
1546  {
1547  result[index] = vtkm::Log2(x[index]);
1548  }
1549  return result;
1550 }
1551 template <typename T>
1553  const vtkm::Vec<T, 4>& x)
1554 {
1556  vtkm::Log2(x[0]), vtkm::Log2(x[1]), vtkm::Log2(x[2]), vtkm::Log2(x[3]));
1557 }
1558 template <typename T>
1560  const vtkm::Vec<T, 3>& x)
1561 {
1563  vtkm::Log2(x[0]), vtkm::Log2(x[1]), vtkm::Log2(x[2]));
1564 }
1565 template <typename T>
1567  const vtkm::Vec<T, 2>& x)
1568 {
1570  vtkm::Log2(x[1]));
1571 }
1573 
1577 
1579 {
1580 #ifdef VTKM_CUDA
1581  return VTKM_CUDA_MATH_FUNCTION_32(log10)(x);
1582 #else
1583  return std::log10(x);
1584 #endif
1585 }
1586 
1588 {
1589 #ifdef VTKM_CUDA
1590  return VTKM_CUDA_MATH_FUNCTION_64(log10)(x);
1591 #else
1592  return std::log10(x);
1593 #endif
1594 }
1595 template <typename T>
1596 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log10(const T& x)
1597 {
1598  using RT = typename detail::FloatingPointReturnType<T>::Type;
1599  return vtkm::Log10(static_cast<RT>(x));
1600 }
1601 template <typename T, vtkm::IdComponent N>
1603  const vtkm::Vec<T, N>& x)
1604 {
1606  for (vtkm::IdComponent index = 0; index < N; index++)
1607  {
1608  result[index] = vtkm::Log10(x[index]);
1609  }
1610  return result;
1611 }
1612 template <typename T>
1614  const vtkm::Vec<T, 4>& x)
1615 {
1617  vtkm::Log10(x[0]), vtkm::Log10(x[1]), vtkm::Log10(x[2]), vtkm::Log10(x[3]));
1618 }
1619 template <typename T>
1621  const vtkm::Vec<T, 3>& x)
1622 {
1624  vtkm::Log10(x[0]), vtkm::Log10(x[1]), vtkm::Log10(x[2]));
1625 }
1626 template <typename T>
1628  const vtkm::Vec<T, 2>& x)
1629 {
1631  vtkm::Log10(x[1]));
1632 }
1634 
1638 
1640 {
1641 #ifdef VTKM_CUDA
1642  return VTKM_CUDA_MATH_FUNCTION_32(log1p)(x);
1643 #else
1644  return std::log1p(x);
1645 #endif
1646 }
1647 
1649 {
1650 #ifdef VTKM_CUDA
1651  return VTKM_CUDA_MATH_FUNCTION_64(log1p)(x);
1652 #else
1653  return std::log1p(x);
1654 #endif
1655 }
1656 template <typename T>
1657 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log1P(const T& x)
1658 {
1659  using RT = typename detail::FloatingPointReturnType<T>::Type;
1660  return vtkm::Log1P(static_cast<RT>(x));
1661 }
1662 template <typename T, vtkm::IdComponent N>
1664  const vtkm::Vec<T, N>& x)
1665 {
1667  for (vtkm::IdComponent index = 0; index < N; index++)
1668  {
1669  result[index] = vtkm::Log1P(x[index]);
1670  }
1671  return result;
1672 }
1673 template <typename T>
1675  const vtkm::Vec<T, 4>& x)
1676 {
1678  vtkm::Log1P(x[0]), vtkm::Log1P(x[1]), vtkm::Log1P(x[2]), vtkm::Log1P(x[3]));
1679 }
1680 template <typename T>
1682  const vtkm::Vec<T, 3>& x)
1683 {
1685  vtkm::Log1P(x[0]), vtkm::Log1P(x[1]), vtkm::Log1P(x[2]));
1686 }
1687 template <typename T>
1689  const vtkm::Vec<T, 2>& x)
1690 {
1692  vtkm::Log1P(x[1]));
1693 }
1695 
1696 //-----------------------------------------------------------------------------
1700 template <typename T>
1701 static inline VTKM_EXEC_CONT T Max(const T& x, const T& y);
1702 #ifdef VTKM_USE_STL
1704 {
1705  return (std::max)(x, y);
1706 }
1708 {
1709  return (std::max)(x, y);
1710 }
1711 #else // !VTKM_USE_STL
1713 {
1714 #ifdef VTKM_CUDA
1715  return VTKM_CUDA_MATH_FUNCTION_32(fmax)(x, y);
1716 #else
1717  return std::fmax(x, y);
1718 #endif
1719 }
1721 {
1722 #ifdef VTKM_CUDA
1723  return VTKM_CUDA_MATH_FUNCTION_64(fmax)(x, y);
1724 #else
1725  return std::fmax(x, y);
1726 #endif
1727 }
1728 #endif // !VTKM_USE_STL
1729 
1734 template <typename T>
1735 static inline VTKM_EXEC_CONT T Min(const T& x, const T& y);
1736 #if defined(VTKM_USE_STL) && !defined(VTKM_HIP)
1738 {
1739  return (std::min)(x, y);
1740 }
1742 {
1743  return (std::min)(x, y);
1744 }
1745 #else // !VTKM_USE_STL OR HIP
1747 {
1748 #ifdef VTKM_CUDA
1749  return VTKM_CUDA_MATH_FUNCTION_32(fmin)(x, y);
1750 #else
1751  return std::fmin(x, y);
1752 #endif
1753 }
1755 {
1756 #ifdef VTKM_CUDA
1757  return VTKM_CUDA_MATH_FUNCTION_64(fmin)(x, y);
1758 #else
1759  return std::fmin(x, y);
1760 #endif
1761 }
1762 #endif // !VTKM_USE_STL
1763 
1765 namespace detail
1766 {
1767 
1768 template <typename T>
1769 static inline VTKM_EXEC_CONT T Max(T x, T y, vtkm::TypeTraitsScalarTag)
1770 {
1771  return (x < y) ? y : x;
1772 }
1773 
1774 template <typename T>
1775 static inline VTKM_EXEC_CONT T Max(const T& x, const T& y, vtkm::TypeTraitsVectorTag)
1776 {
1777  using Traits = vtkm::VecTraits<T>;
1778  T result;
1779  for (vtkm::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++)
1780  {
1781  Traits::SetComponent(
1782  result, index, vtkm::Max(Traits::GetComponent(x, index), Traits::GetComponent(y, index)));
1783  }
1784  return result;
1785 }
1786 
1787 template <typename T>
1788 static inline VTKM_EXEC_CONT T Min(T x, T y, vtkm::TypeTraitsScalarTag)
1789 {
1790  return (x < y) ? x : y;
1791 }
1792 
1793 template <typename T>
1794 static inline VTKM_EXEC_CONT T Min(const T& x, const T& y, vtkm::TypeTraitsVectorTag)
1795 {
1796  using Traits = vtkm::VecTraits<T>;
1797  T result;
1798  for (vtkm::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++)
1799  {
1800  Traits::SetComponent(
1801  result, index, vtkm::Min(Traits::GetComponent(x, index), Traits::GetComponent(y, index)));
1802  }
1803  return result;
1804 }
1805 
1806 } // namespace detail
1807 
1810 template <typename T>
1811 static inline VTKM_EXEC_CONT T Max(const T& x, const T& y)
1812 {
1813  return detail::Max(x, y, typename vtkm::TypeTraits<T>::DimensionalityTag());
1814 }
1815 
1818 template <typename T>
1819 static inline VTKM_EXEC_CONT T Min(const T& x, const T& y)
1820 {
1821  return detail::Min(x, y, typename vtkm::TypeTraits<T>::DimensionalityTag());
1822 }
1823 
1828 {
1829  return x > lo ? (x < hi ? x : hi) : lo;
1830 }
1831 
1833 {
1834  return x > lo ? (x < hi ? x : hi) : lo;
1835 }
1837 
1838 //-----------------------------------------------------------------------------
1839 
1840 //#ifdef VTKM_CUDA
1841 #define VTKM_USE_IEEE_NONFINITE
1842 //#endif
1843 
1844 #ifdef VTKM_USE_IEEE_NONFINITE
1845 
1846 namespace detail
1847 {
1848 
1849 union IEEE754Bits32 {
1850  vtkm::UInt32 bits;
1851  vtkm::Float32 scalar;
1852 };
1853 #define VTKM_NAN_BITS_32 0x7FC00000U
1854 #define VTKM_INF_BITS_32 0x7F800000U
1855 #define VTKM_NEG_INF_BITS_32 0xFF800000U
1856 #define VTKM_EPSILON_32 1e-5f
1857 
1858 union IEEE754Bits64 {
1859  vtkm::UInt64 bits;
1860  vtkm::Float64 scalar;
1861 };
1862 #define VTKM_NAN_BITS_64 0x7FF8000000000000ULL
1863 #define VTKM_INF_BITS_64 0x7FF0000000000000ULL
1864 #define VTKM_NEG_INF_BITS_64 0xFFF0000000000000ULL
1865 #define VTKM_EPSILON_64 1e-9
1866 
1867 template <typename T>
1868 struct FloatLimits;
1869 
1870 template <>
1871 struct FloatLimits<vtkm::Float32>
1872 {
1873  using BitsType = vtkm::detail::IEEE754Bits32;
1874 
1876  static vtkm::Float32 Nan()
1877  {
1878  BitsType nan = { VTKM_NAN_BITS_32 };
1879  return nan.scalar;
1880  }
1881 
1883  static vtkm::Float32 Infinity()
1884  {
1885  BitsType inf = { VTKM_INF_BITS_32 };
1886  return inf.scalar;
1887  }
1888 
1890  static vtkm::Float32 NegativeInfinity()
1891  {
1892  BitsType neginf = { VTKM_NEG_INF_BITS_32 };
1893  return neginf.scalar;
1894  }
1895 
1897  static vtkm::Float32 Epsilon() { return VTKM_EPSILON_32; }
1898 };
1899 
1900 template <int N>
1901 struct FloatLimits<vtkm::Vec<vtkm::Float32, N>>
1902 {
1903  using BitsType = vtkm::detail::IEEE754Bits32;
1904 
1906  static vtkm::Vec<vtkm::Float32, N> Nan()
1907  {
1908  BitsType nan = { VTKM_NAN_BITS_32 };
1909  return vtkm::Vec<vtkm::Float32, N>(nan.scalar);
1910  }
1911 
1913  static vtkm::Vec<vtkm::Float32, N> Infinity()
1914  {
1915  BitsType inf = { VTKM_INF_BITS_32 };
1916  return vtkm::Vec<vtkm::Float32, N>(inf.scalar);
1917  }
1918 
1920  static vtkm::Vec<vtkm::Float32, N> NegativeInfinity()
1921  {
1922  BitsType neginf = { VTKM_NEG_INF_BITS_32 };
1923  return vtkm::Vec<vtkm::Float32, N>(neginf.scalar);
1924  }
1925 
1927  static vtkm::Vec<vtkm::Float32, N> Epsilon()
1928  {
1930  }
1931 };
1932 
1933 template <>
1934 struct FloatLimits<vtkm::Float64>
1935 {
1936  using BitsType = vtkm::detail::IEEE754Bits64;
1937 
1939  static vtkm::Float64 Nan()
1940  {
1941  BitsType nan = { VTKM_NAN_BITS_64 };
1942  return nan.scalar;
1943  }
1944 
1946  static vtkm::Float64 Infinity()
1947  {
1948  BitsType inf = { VTKM_INF_BITS_64 };
1949  return inf.scalar;
1950  }
1951 
1953  static vtkm::Float64 NegativeInfinity()
1954  {
1955  BitsType neginf = { VTKM_NEG_INF_BITS_64 };
1956  return neginf.scalar;
1957  }
1958 
1960  static vtkm::Float64 Epsilon() { return VTKM_EPSILON_64; }
1961 };
1962 
1963 template <int N>
1964 struct FloatLimits<vtkm::Vec<vtkm::Float64, N>>
1965 {
1966  using BitsType = vtkm::detail::IEEE754Bits64;
1967 
1969  static vtkm::Vec<vtkm::Float64, N> Nan()
1970  {
1971  BitsType nan = { VTKM_NAN_BITS_64 };
1972  return vtkm::Vec<vtkm::Float64, N>(nan.scalar);
1973  }
1974 
1976  static vtkm::Vec<vtkm::Float64, N> Infinity()
1977  {
1978  BitsType inf = { VTKM_INF_BITS_64 };
1979  return vtkm::Vec<vtkm::Float64, N>(inf.scalar);
1980  }
1981 
1983  static vtkm::Vec<vtkm::Float64, N> NegativeInfinity()
1984  {
1985  BitsType neginf = { VTKM_NEG_INF_BITS_64 };
1986  return vtkm::Vec<vtkm::Float64, N>(neginf.scalar);
1987  }
1988 
1990  static vtkm::Vec<vtkm::Float64, N> Epsilon()
1991  {
1993  }
1994 };
1995 
1996 #undef VTKM_NAN_BITS_32
1997 #undef VTKM_INF_BITS_32
1998 #undef VTKM_NEG_INF_BITS_32
1999 #undef VTKM_EPSILON_32
2000 #undef VTKM_NAN_BITS_64
2001 #undef VTKM_INF_BITS_64
2002 #undef VTKM_NEG_INF_BITS_64
2003 #undef VTKM_EPSILON_64
2004 
2005 } // namespace detail
2006 #endif //VTKM_USE_IEEE_NONFINITE
2007 
2015 #ifdef VTKM_USE_IEEE_NONFINITE
2016 template <typename T>
2017 static inline VTKM_EXEC_CONT T Nan()
2018 {
2019  return detail::FloatLimits<T>::Nan();
2020 }
2021 #else // !VTKM_USE_IEEE_NONFINITE
2022 template <typename T>
2023 static inline VTKM_EXEC_CONT T Nan()
2024 {
2025  return std::numeric_limits<T>::quiet_NaN();
2026 }
2027 #endif // !VTKM_USE_IEEE_NONFINITE
2028 static inline VTKM_EXEC_CONT vtkm::Float32 Nan32()
2029 {
2030  return vtkm::Nan<vtkm::Float32>();
2031 }
2032 static inline VTKM_EXEC_CONT vtkm::Float64 Nan64()
2033 {
2034  return vtkm::Nan<vtkm::Float64>();
2035 }
2037 
2045 #ifdef VTKM_USE_IEEE_NONFINITE
2046 template <typename T>
2047 static inline VTKM_EXEC_CONT T Infinity()
2048 {
2049  return detail::FloatLimits<T>::Infinity();
2050 }
2051 #else // !VTKM_USE_IEEE_NONFINITE
2052 template <typename T>
2053 static inline VTKM_EXEC_CONT T Infinity()
2054 {
2055  return std::numeric_limits<T>::infinity();
2056 }
2057 #endif // !VTKM_USE_IEEE_NONFINITE
2058 static inline VTKM_EXEC_CONT vtkm::Float32 Infinity32()
2059 {
2060  return vtkm::Infinity<vtkm::Float32>();
2061 }
2062 static inline VTKM_EXEC_CONT vtkm::Float64 Infinity64()
2063 {
2064  return vtkm::Infinity<vtkm::Float64>();
2065 }
2067 
2076 #ifdef VTKM_USE_IEEE_NONFINITE
2077 template <typename T>
2078 static inline VTKM_EXEC_CONT T NegativeInfinity()
2079 {
2080  return detail::FloatLimits<T>::NegativeInfinity();
2081 }
2082 #else // !VTKM_USE_IEEE_NONFINITE
2083 template <typename T>
2084 static inline VTKM_EXEC_CONT T NegativeInfinity()
2085 {
2086  return -std::numeric_limits<T>::infinity();
2087 }
2088 #endif // !VTKM_USE_IEEE_NONFINITE
2089 static inline VTKM_EXEC_CONT vtkm::Float32 NegativeInfinity32()
2090 {
2091  return vtkm::NegativeInfinity<vtkm::Float32>();
2092 }
2093 static inline VTKM_EXEC_CONT vtkm::Float64 NegativeInfinity64()
2094 {
2095  return vtkm::NegativeInfinity<vtkm::Float64>();
2096 }
2098 
2106 #ifdef VTKM_USE_IEEE_NONFINITE
2107 template <typename T>
2108 static inline VTKM_EXEC_CONT T Epsilon()
2109 {
2110  return detail::FloatLimits<T>::Epsilon();
2111 }
2112 #else // !VTKM_USE_IEEE_NONFINITE
2113 template <typename T>
2114 static inline VTKM_EXEC_CONT T Epsilon()
2115 {
2116  return std::numeric_limits<T>::epsilon();
2117 }
2118 #endif // !VTKM_USE_IEEE_NONFINITE
2119 static inline VTKM_EXEC_CONT vtkm::Float32 Epsilon32()
2120 {
2121  return vtkm::Epsilon<vtkm::Float32>();
2122 }
2123 static inline VTKM_EXEC_CONT vtkm::Float64 Epsilon64()
2124 {
2125  return vtkm::Epsilon<vtkm::Float64>();
2126 }
2128 
2129 
2130 //-----------------------------------------------------------------------------
2133 template <typename T>
2134 static inline VTKM_EXEC_CONT bool IsNan(T x)
2135 {
2136 #ifndef VTKM_CUDA
2137  using std::isnan;
2138 #endif
2139  return (isnan(x) != 0);
2140 }
2141 
2144 template <typename T>
2145 static inline VTKM_EXEC_CONT bool IsInf(T x)
2146 {
2147 #ifndef VTKM_CUDA
2148  using std::isinf;
2149 #endif
2150  return (isinf(x) != 0);
2151 }
2152 
2155 template <typename T>
2156 static inline VTKM_EXEC_CONT bool IsFinite(T x)
2157 {
2158 #ifndef VTKM_CUDA
2159  using std::isfinite;
2160 #endif
2161  return (isfinite(x) != 0);
2162 }
2163 
2164 //-----------------------------------------------------------------------------
2168 
2170 {
2171 #ifdef VTKM_CUDA
2172  return VTKM_CUDA_MATH_FUNCTION_32(ceil)(x);
2173 #else
2174  return std::ceil(x);
2175 #endif
2176 }
2177 
2179 {
2180 #ifdef VTKM_CUDA
2181  return VTKM_CUDA_MATH_FUNCTION_64(ceil)(x);
2182 #else
2183  return std::ceil(x);
2184 #endif
2185 }
2186 template <typename T>
2187 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Ceil(const T& x)
2188 {
2189  using RT = typename detail::FloatingPointReturnType<T>::Type;
2190  return vtkm::Ceil(static_cast<RT>(x));
2191 }
2192 template <typename T, vtkm::IdComponent N>
2194  const vtkm::Vec<T, N>& x)
2195 {
2197  for (vtkm::IdComponent index = 0; index < N; index++)
2198  {
2199  result[index] = vtkm::Ceil(x[index]);
2200  }
2201  return result;
2202 }
2203 template <typename T>
2205  const vtkm::Vec<T, 4>& x)
2206 {
2208  vtkm::Ceil(x[0]), vtkm::Ceil(x[1]), vtkm::Ceil(x[2]), vtkm::Ceil(x[3]));
2209 }
2210 template <typename T>
2212  const vtkm::Vec<T, 3>& x)
2213 {
2215  vtkm::Ceil(x[0]), vtkm::Ceil(x[1]), vtkm::Ceil(x[2]));
2216 }
2217 template <typename T>
2219  const vtkm::Vec<T, 2>& x)
2220 {
2222  vtkm::Ceil(x[1]));
2223 }
2225 
2229 
2231 {
2232 #ifdef VTKM_CUDA
2233  return VTKM_CUDA_MATH_FUNCTION_32(floor)(x);
2234 #else
2235  return std::floor(x);
2236 #endif
2237 }
2238 
2240 {
2241 #ifdef VTKM_CUDA
2242  return VTKM_CUDA_MATH_FUNCTION_64(floor)(x);
2243 #else
2244  return std::floor(x);
2245 #endif
2246 }
2247 template <typename T>
2248 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Floor(const T& x)
2249 {
2250  using RT = typename detail::FloatingPointReturnType<T>::Type;
2251  return vtkm::Floor(static_cast<RT>(x));
2252 }
2253 template <typename T, vtkm::IdComponent N>
2255  const vtkm::Vec<T, N>& x)
2256 {
2258  for (vtkm::IdComponent index = 0; index < N; index++)
2259  {
2260  result[index] = vtkm::Floor(x[index]);
2261  }
2262  return result;
2263 }
2264 template <typename T>
2266  const vtkm::Vec<T, 4>& x)
2267 {
2269  vtkm::Floor(x[0]), vtkm::Floor(x[1]), vtkm::Floor(x[2]), vtkm::Floor(x[3]));
2270 }
2271 template <typename T>
2273  const vtkm::Vec<T, 3>& x)
2274 {
2276  vtkm::Floor(x[0]), vtkm::Floor(x[1]), vtkm::Floor(x[2]));
2277 }
2278 template <typename T>
2280  const vtkm::Vec<T, 2>& x)
2281 {
2283  vtkm::Floor(x[1]));
2284 }
2286 
2290 
2292 {
2293 #ifdef VTKM_CUDA
2294  return VTKM_CUDA_MATH_FUNCTION_32(round)(x);
2295 #else
2296  return std::round(x);
2297 #endif
2298 }
2299 
2301 {
2302 #ifdef VTKM_CUDA
2303  return VTKM_CUDA_MATH_FUNCTION_64(round)(x);
2304 #else
2305  return std::round(x);
2306 #endif
2307 }
2308 template <typename T>
2309 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Round(const T& x)
2310 {
2311  using RT = typename detail::FloatingPointReturnType<T>::Type;
2312  return vtkm::Round(static_cast<RT>(x));
2313 }
2314 template <typename T, vtkm::IdComponent N>
2316  const vtkm::Vec<T, N>& x)
2317 {
2319  for (vtkm::IdComponent index = 0; index < N; index++)
2320  {
2321  result[index] = vtkm::Round(x[index]);
2322  }
2323  return result;
2324 }
2325 template <typename T>
2327  const vtkm::Vec<T, 4>& x)
2328 {
2330  vtkm::Round(x[0]), vtkm::Round(x[1]), vtkm::Round(x[2]), vtkm::Round(x[3]));
2331 }
2332 template <typename T>
2334  const vtkm::Vec<T, 3>& x)
2335 {
2337  vtkm::Round(x[0]), vtkm::Round(x[1]), vtkm::Round(x[2]));
2338 }
2339 template <typename T>
2341  const vtkm::Vec<T, 2>& x)
2342 {
2344  vtkm::Round(x[1]));
2345 }
2347 
2348 //-----------------------------------------------------------------------------
2356 {
2357 #ifdef VTKM_CUDA
2358  return VTKM_CUDA_MATH_FUNCTION_32(fmod)(x, y);
2359 #else
2360  return std::fmod(x, y);
2361 #endif
2362 }
2364 {
2365 #ifdef VTKM_CUDA
2366  return VTKM_CUDA_MATH_FUNCTION_64(fmod)(x, y);
2367 #else
2368  return std::fmod(x, y);
2369 #endif
2370 }
2372 
2380 #ifdef VTKM_MSVC
2381 template <typename T>
2382 static inline VTKM_EXEC_CONT T Remainder(T numerator, T denominator)
2383 {
2384  T quotient = vtkm::Round(numerator / denominator);
2385  return numerator - quotient * denominator;
2386 }
2387 #else // !VTKM_MSVC
2388 static inline VTKM_EXEC_CONT vtkm::Float32 Remainder(vtkm::Float32 x, vtkm::Float32 y)
2389 {
2390 #ifdef VTKM_CUDA
2391  return VTKM_CUDA_MATH_FUNCTION_32(remainder)(x, y);
2392 #else
2393  return std::remainder(x, y);
2394 #endif
2395 }
2396 static inline VTKM_EXEC_CONT vtkm::Float64 Remainder(vtkm::Float64 x, vtkm::Float64 y)
2397 {
2398 #ifdef VTKM_CUDA
2399  return VTKM_CUDA_MATH_FUNCTION_64(remainder)(x, y);
2400 #else
2401  return std::remainder(x, y);
2402 #endif
2403 }
2404 #endif // !VTKM_MSVC
2405 
2412 template <typename QType>
2413 static inline VTKM_EXEC_CONT vtkm::Float32 RemainderQuotient(vtkm::Float32 numerator,
2414  vtkm::Float32 denominator,
2415  QType& quotient)
2416 {
2417  int iQuotient;
2418  // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
2419 #if defined(VTKM_CUDA) || defined(VTKM_HIP)
2420  const vtkm::Float32 result =
2421  VTKM_CUDA_MATH_FUNCTION_32(remquo)(numerator, denominator, &iQuotient);
2422 #else
2423  const vtkm::Float32 result = std::remquo(numerator, denominator, &iQuotient);
2424 #endif
2425  quotient = static_cast<QType>(iQuotient);
2426  return result;
2427 }
2428 template <typename QType>
2429 static inline VTKM_EXEC_CONT vtkm::Float64 RemainderQuotient(vtkm::Float64 numerator,
2430  vtkm::Float64 denominator,
2431  QType& quotient)
2432 {
2433  int iQuotient;
2434 #ifdef VTKM_CUDA
2435  const vtkm::Float64 result =
2436  VTKM_CUDA_MATH_FUNCTION_64(remquo)(numerator, denominator, &iQuotient);
2437 #else
2438  const vtkm::Float64 result = std::remquo(numerator, denominator, &iQuotient);
2439 #endif
2440  quotient = static_cast<QType>(iQuotient);
2441  return result;
2442 }
2444 
2449 static inline VTKM_EXEC_CONT vtkm::Float32 ModF(vtkm::Float32 x, vtkm::Float32& integral)
2450 {
2451  // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
2452 #if defined(VTKM_CUDA) || defined(VTKM_HIP)
2453  return VTKM_CUDA_MATH_FUNCTION_32(modf)(x, &integral);
2454 #else
2455  return std::modf(x, &integral);
2456 #endif
2457 }
2458 static inline VTKM_EXEC_CONT vtkm::Float64 ModF(vtkm::Float64 x, vtkm::Float64& integral)
2459 {
2460 #if defined(VTKM_CUDA)
2461  return VTKM_CUDA_MATH_FUNCTION_64(modf)(x, &integral);
2462 #else
2463  return std::modf(x, &integral);
2464 #endif
2465 }
2467 
2468 //-----------------------------------------------------------------------------
2473 static inline VTKM_EXEC_CONT vtkm::Int32 Abs(vtkm::Int32 x)
2474 {
2475  return abs(x);
2476 }
2477 static inline VTKM_EXEC_CONT vtkm::Int64 Abs(vtkm::Int64 x)
2478 {
2479 #if VTKM_SIZE_LONG == 8
2480  return labs(x);
2481 #elif VTKM_SIZE_LONG_LONG == 8
2482  return llabs(x);
2483 #else
2484 #error Unknown size of Int64.
2485 #endif
2486 }
2487 static inline VTKM_EXEC_CONT vtkm::Float32 Abs(vtkm::Float32 x)
2488 {
2489 #ifdef VTKM_CUDA
2490  return VTKM_CUDA_MATH_FUNCTION_32(fabs)(x);
2491 #else
2492  return std::fabs(x);
2493 #endif
2494 }
2495 static inline VTKM_EXEC_CONT vtkm::Float64 Abs(vtkm::Float64 x)
2496 {
2497 #ifdef VTKM_CUDA
2498  return VTKM_CUDA_MATH_FUNCTION_64(fabs)(x);
2499 #else
2500  return std::fabs(x);
2501 #endif
2502 }
2503 template <typename T>
2504 static inline VTKM_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Abs(T x)
2505 {
2506 #ifdef VTKM_CUDA
2507  return VTKM_CUDA_MATH_FUNCTION_64(fabs)(static_cast<vtkm::Float64>(x));
2508 #else
2509  return std::fabs(static_cast<vtkm::Float64>(x));
2510 #endif
2511 }
2512 template <typename T, vtkm::IdComponent N>
2513 static inline VTKM_EXEC_CONT vtkm::Vec<T, N> Abs(const vtkm::Vec<T, N>& x)
2514 {
2515  vtkm::Vec<T, N> result;
2516  for (vtkm::IdComponent index = 0; index < N; index++)
2517  {
2518  result[index] = vtkm::Abs(x[index]);
2519  }
2520  return result;
2521 }
2522 template <typename T>
2523 static inline VTKM_EXEC_CONT vtkm::Vec<T, 4> Abs(const vtkm::Vec<T, 4>& x)
2524 {
2525  return vtkm::Vec<T, 4>(vtkm::Abs(x[0]), vtkm::Abs(x[1]), vtkm::Abs(x[2]), vtkm::Abs(x[3]));
2526 }
2527 template <typename T>
2528 static inline VTKM_EXEC_CONT vtkm::Vec<T, 3> Abs(const vtkm::Vec<T, 3>& x)
2529 {
2530  return vtkm::Vec<T, 3>(vtkm::Abs(x[0]), vtkm::Abs(x[1]), vtkm::Abs(x[2]));
2531 }
2532 template <typename T>
2533 static inline VTKM_EXEC_CONT vtkm::Vec<T, 2> Abs(const vtkm::Vec<T, 2>& x)
2534 {
2535  return vtkm::Vec<T, 2>(vtkm::Abs(x[0]), vtkm::Abs(x[1]));
2536 }
2538 
2542 static inline VTKM_EXEC_CONT vtkm::Int32 SignBit(vtkm::Float32 x)
2543 {
2544 #ifndef VTKM_CUDA
2545  using std::signbit;
2546 #endif
2547  return static_cast<vtkm::Int32>(signbit(x));
2548 }
2549 static inline VTKM_EXEC_CONT vtkm::Int32 SignBit(vtkm::Float64 x)
2550 {
2551 #ifndef VTKM_CUDA
2552  using std::signbit;
2553 #endif
2554  return static_cast<vtkm::Int32>(signbit(x));
2555 }
2557 
2561 static inline VTKM_EXEC_CONT bool IsNegative(vtkm::Float32 x)
2562 {
2563  return (vtkm::SignBit(x) != 0);
2564 }
2565 static inline VTKM_EXEC_CONT bool IsNegative(vtkm::Float64 x)
2566 {
2567  return (vtkm::SignBit(x) != 0);
2568 }
2570 
2575 static inline VTKM_EXEC_CONT vtkm::Float32 CopySign(vtkm::Float32 x, vtkm::Float32 y)
2576 {
2577 #ifdef VTKM_CUDA
2578  return VTKM_CUDA_MATH_FUNCTION_32(copysign)(x, y);
2579 #else
2580  return std::copysign(x, y);
2581 #endif
2582 }
2583 static inline VTKM_EXEC_CONT vtkm::Float64 CopySign(vtkm::Float64 x, vtkm::Float64 y)
2584 {
2585 #ifdef VTKM_CUDA
2586  return VTKM_CUDA_MATH_FUNCTION_64(copysign)(x, y);
2587 #else
2588  return std::copysign(x, y);
2589 #endif
2590 }
2591 
2592 template <typename T, vtkm::IdComponent N>
2593 static inline VTKM_EXEC_CONT vtkm::Vec<T, N> CopySign(const vtkm::Vec<T, N>& x,
2594  const vtkm::Vec<T, N>& y)
2595 {
2596  vtkm::Vec<T, N> result;
2597  for (vtkm::IdComponent index = 0; index < N; index++)
2598  {
2599  result[index] = vtkm::CopySign(x[index], y[index]);
2600  }
2601  return result;
2602 }
2604 
2609 {
2610  // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
2611 #if defined(VTKM_CUDA) || defined(VTKM_HIP)
2612  return VTKM_CUDA_MATH_FUNCTION_32(frexp)(x, exponent);
2613 #else
2614  return std::frexp(x, exponent);
2615 #endif
2616 }
2617 
2619 {
2620 #ifdef VTKM_CUDA
2621  return VTKM_CUDA_MATH_FUNCTION_64(frexp)(x, exponent);
2622 #else
2623  return std::frexp(x, exponent);
2624 #endif
2625 }
2627 
2629 {
2630 #ifdef VTKM_CUDA
2631  return VTKM_CUDA_MATH_FUNCTION_32(ldexp)(x, exponent);
2632 #else
2633  return std::ldexp(x, exponent);
2634 #endif
2635 }
2636 
2638 {
2639 #ifdef VTKM_CUDA
2640  return VTKM_CUDA_MATH_FUNCTION_64(ldexp)(x, exponent);
2641 #else
2642  return std::ldexp(x, exponent);
2643 #endif
2644 }
2645 
2647 // Float distance.
2648 // See: https://randomascii.wordpress.com/2012/01/23/stupid-float-tricks-2/ for why this works.
2649 
2650 namespace detail
2651 {
2652 
2653 inline VTKM_EXEC_CONT vtkm::UInt64 FloatDistancePositive(vtkm::Float64 x, vtkm::Float64 y)
2654 {
2655  VTKM_ASSERT(x >= 0);
2656  VTKM_ASSERT(y >= 0);
2657 
2658  // Note that:
2659  // int64_t xi = *reinterpret_cast<int64_t*>(&x);
2660  // int64_t yi = *reinterpret_cast<int64_t*>(&y);
2661  // also works (usually), but generates warnings because it is technically undefined behavior
2662  // according to the C++ standard.
2663  // Good option to have if we get compile errors off memcpy or don't want to #include <cstring> though.
2664  // At least on gcc, both versions generate the same assembly.
2665  vtkm::UInt64 xi;
2666  vtkm::UInt64 yi;
2667  memcpy(&xi, &x, sizeof(vtkm::UInt64));
2668  memcpy(&yi, &y, sizeof(vtkm::UInt64));
2669  if (yi > xi) {
2670  return yi - xi;
2671  }
2672  return xi - yi;
2673 }
2674 
2675 
2676 inline VTKM_EXEC_CONT vtkm::UInt64 FloatDistancePositive(vtkm::Float32 x, vtkm::Float32 y)
2677 {
2678  VTKM_ASSERT(x >= 0);
2679  VTKM_ASSERT(y >= 0);
2680 
2681  vtkm::UInt32 xi_32;
2682  vtkm::UInt32 yi_32;
2683  memcpy(&xi_32, &x, sizeof(vtkm::UInt32));
2684  memcpy(&yi_32, &y, sizeof(vtkm::UInt32));
2685  vtkm::UInt64 xi = xi_32;
2686  vtkm::UInt64 yi = yi_32;
2687  if (yi > xi) {
2688  return yi - xi;
2689  }
2690  return xi - yi;
2691 }
2692 
2693 } // namespace detail
2694 
2699 {
2700  static_assert(sizeof(vtkm::Float64) == sizeof(vtkm::UInt64), "vtkm::Float64 is incorrect size.");
2701  static_assert(std::numeric_limits<vtkm::Float64>::has_denorm == std::denorm_present, "FloatDistance presumes the floating-point type has subnormal numbers.");
2702 
2703  if (!vtkm::IsFinite(x) || !vtkm::IsFinite(y)) {
2704  return 0xFFFFFFFFFFFFFFFFL;
2705  }
2706 
2707  // Signed zero is the sworn enemy of this process.
2708  if (y == 0) {
2709  y = vtkm::Abs(y);
2710  }
2711  if (x == 0) {
2712  x = vtkm::Abs(x);
2713  }
2714 
2715  if ( (x < 0 && y >= 0) || (x >= 0 && y < 0) )
2716  {
2717  vtkm::UInt64 dx, dy;
2718  if (x < 0) {
2719  dy = detail::FloatDistancePositive(0.0, y);
2720  dx = detail::FloatDistancePositive(0.0, -x);
2721  }
2722  else {
2723  dy = detail::FloatDistancePositive(0.0, -y);
2724  dx = detail::FloatDistancePositive(0.0, x);
2725  }
2726 
2727  return dx + dy;
2728  }
2729 
2730  if (x < 0 && y < 0) {
2731  return detail::FloatDistancePositive(-x, -y);
2732  }
2733 
2734  return detail::FloatDistancePositive(x, y);
2735 }
2736 
2739 {
2740  static_assert(sizeof(vtkm::Float32) == sizeof(vtkm::Int32), "vtkm::Float32 is incorrect size.");
2741  static_assert(std::numeric_limits<vtkm::Float32>::has_denorm == std::denorm_present, "FloatDistance presumes the floating-point type has subnormal numbers.");
2742 
2743  if (!vtkm::IsFinite(x) || !vtkm::IsFinite(y)) {
2744  return 0xFFFFFFFFFFFFFFFFL;
2745  }
2746 
2747  if (y == 0) {
2748  y = vtkm::Abs(y);
2749  }
2750  if (x == 0) {
2751  x = vtkm::Abs(x);
2752  }
2753 
2754  if ( (x < 0 && y >= 0) || (x >= 0 && y < 0) )
2755  {
2756  vtkm::UInt64 dx, dy;
2757  if (x < 0) {
2758  dy = detail::FloatDistancePositive(0.0f, y);
2759  dx = detail::FloatDistancePositive(0.0f, -x);
2760  }
2761  else {
2762  dy = detail::FloatDistancePositive(0.0f, -y);
2763  dx = detail::FloatDistancePositive(0.0f, x);
2764  }
2765  return dx + dy;
2766  }
2767 
2768  if (x < 0 && y < 0) {
2769  return detail::FloatDistancePositive(-x, -y);
2770  }
2771 
2772  return detail::FloatDistancePositive(x, y);
2773 }
2774 
2775 // Computes ab - cd.
2776 // See: https://pharr.org/matt/blog/2019/11/03/difference-of-floats.html
2777 template<typename T>
2778 inline VTKM_EXEC_CONT T DifferenceOfProducts(T a, T b, T c, T d)
2779 {
2780  T cd = c * d;
2781  T err = std::fma(-c, d, cd);
2782  T dop = std::fma(a, b, -cd);
2783  return dop + err;
2784 }
2785 
2794 template<typename T>
2796 {
2797  if (a == 0)
2798  {
2799  if (b == 0)
2800  {
2801  if (c == 0)
2802  {
2803  // A degenerate case. All real numbers are roots; hopefully this arbitrary decision interacts gracefully with use.
2804  return vtkm::Vec<T,2>(0,0);
2805  }
2806  else
2807  {
2808  return vtkm::Vec<T,2>(vtkm::Nan<T>(), vtkm::Nan<T>());
2809  }
2810  }
2811  return vtkm::Vec<T,2>(-c/b, -c/b);
2812  }
2813  T delta = DifferenceOfProducts(b, b, 4*a, c);
2814  if (delta < 0)
2815  {
2816  return vtkm::Vec<T,2>(vtkm::Nan<T>(), vtkm::Nan<T>());
2817  }
2818 
2819  T q = -(b + vtkm::CopySign(vtkm::Sqrt(delta), b)) / 2;
2820  T r0 = q / a;
2821  T r1 = c / q;
2822  if (r0 < r1)
2823  {
2824  return vtkm::Vec<T,2>(r0, r1);
2825  }
2826  return vtkm::Vec<T,2>(r1, r0);
2827 }
2828 
2831 
2834 #ifdef VTKM_CUDA_DEVICE_PASS
2835 // Need to explicitly mark this as __device__ since __ffs is device only.
2836 inline __device__
2838 {
2839  // Output is [0,32], with ffs(0) == 0
2840  return __ffs(static_cast<int>(word));
2841 }
2842 #else // CUDA_DEVICE_PASS
2843 inline VTKM_EXEC_CONT
2845 {
2846 # if defined(VTKM_GCC) || defined(VTKM_CLANG)
2847 
2848  // Output is [0,32], with ffs(0) == 0
2849  return __builtin_ffs(static_cast<int>(word));
2850 
2851 # elif defined(VTKM_MSVC)
2852 
2853  // Output is [0, 31], check return code to see if bits are set:
2854  vtkm::UInt32 firstSet;
2855  return _BitScanForward(reinterpret_cast<DWORD*>(&firstSet), word) != 0
2856  ? static_cast<vtkm::Int32>(firstSet + 1) : 0;
2857 
2858 # elif defined(VTKM_ICC)
2859 
2860  // Output is [0, 31], undefined if word is 0.
2861  return word != 0 ? _bit_scan_forward(word) + 1 : 0;
2862 
2863 # else
2864 
2865  // Naive implementation:
2866  if (word == 0)
2867  {
2868  return 0;
2869  }
2870 
2871  vtkm::Int32 bit = 1;
2872  while ((word & 0x1) == 0)
2873  {
2874  word >>= 1;
2875  ++bit;
2876  }
2877  return bit;
2878 
2879 # endif
2880 }
2881 #endif // CUDA_DEVICE_PASS
2882 
2885 #ifdef VTKM_CUDA_DEVICE_PASS
2886 // Need to explicitly mark this as __device__ since __ffsll is device only.
2887 inline __device__
2889 {
2890 
2891  // Output is [0,64], with ffs(0) == 0
2892  return __ffsll(static_cast<long long int>(word));
2893 }
2894 #else // CUDA_DEVICE_PASS
2895 inline VTKM_EXEC_CONT
2897 {
2898 # if defined(VTKM_GCC) || defined(VTKM_CLANG) || defined(VTKM_ICC)
2899 
2900  // Output is [0,64], with ffs(0) == 0
2901  return __builtin_ffsll(static_cast<long long int>(word));
2902 
2903 # elif defined(VTKM_MSVC)
2904 
2905  // Output is [0, 63], check return code to see if bits are set:
2906  vtkm::UInt32 firstSet;
2907  return _BitScanForward64(reinterpret_cast<DWORD*>(&firstSet), word) != 0
2908  ? static_cast<vtkm::Int32>(firstSet + 1) : 0;
2909 
2910 # else
2911 
2912  // Naive implementation:
2913  if (word == 0)
2914  {
2915  return 0;
2916  }
2917 
2918  vtkm::Int32 bit = 1;
2919  while ((word & 0x1) == 0)
2920  {
2921  word >>= 1;
2922  ++bit;
2923  }
2924  return bit;
2925 
2926 # endif
2927 }
2928 #endif // CUDA_DEVICE_PASS
2929 
2931 #ifdef VTKM_CUDA_DEVICE_PASS
2932 // Need to explicitly mark this as __device__ since __popc is device only.
2933 inline __device__
2935 {
2936  return __popc(word);
2937 }
2938 #else // CUDA_DEVICE_PASS
2939 inline VTKM_EXEC_CONT
2941 {
2942 # if defined(VTKM_GCC) || defined(VTKM_CLANG)
2943 
2944  return __builtin_popcount(word);
2945 
2946 # elif defined(VTKM_MSVC) && !defined(_M_ARM64)
2947 
2948  return static_cast<vtkm::Int32>(__popcnt(word));
2949 
2950 # elif defined(VTKM_ICC)
2951 
2952  return _popcnt32(static_cast<int>(word));
2953 
2954 # else
2955 
2956  // Naive implementation:
2957  vtkm::Int32 bits = 0;
2958  while (word)
2959  {
2960  if (word & 0x1)
2961  {
2962  ++bits;
2963  }
2964  word >>= 1;
2965  }
2966  return bits;
2967 
2968 # endif
2969 }
2970 #endif // CUDA_DEVICE_PASS
2971 
2973 #ifdef VTKM_CUDA_DEVICE_PASS
2974 // Need to explicitly mark this as __device__ since __popcll is device only.
2975 inline __device__
2977 {
2978  return __popcll(word);
2979 }
2980 #else // CUDA_DEVICE_PASS
2981 inline VTKM_EXEC_CONT
2983 {
2984 # if defined(VTKM_GCC) || defined(VTKM_CLANG)
2985 
2986  return __builtin_popcountll(word);
2987 
2988 # elif defined(VTKM_MSVC) && !defined(_M_ARM64)
2989 
2990  return static_cast<vtkm::Int32>(__popcnt64(word));
2991 
2992 # elif defined(VTKM_ICC)
2993 
2994  return _popcnt64(static_cast<vtkm::Int64>(word));
2995 
2996 # else
2997 
2998  // Naive implementation:
2999  vtkm::Int32 bits = 0;
3000  while (word)
3001  {
3002  if (word & 0x1)
3003  {
3004  ++bits;
3005  }
3006  word >>= 1;
3007  }
3008  return bits;
3009 
3010 # endif
3011 }
3012 #endif // CUDA_DEVICE_PASS
3013 
3014 } // namespace vtkm
3015 // clang-format on
3016 
3017 #endif //vtk_m_Math_h
vtkm::CosH
vtkm::Float32 CosH(vtkm::Float32 x)
Definition: Math.h:615
vtkm::ACosH
vtkm::Float32 ACosH(vtkm::Float32 x)
Definition: Math.h:798
vtkm::TanH
vtkm::Float32 TanH(vtkm::Float32 x)
Definition: Math.h:676
VTKM_NAN_BITS_64
#define VTKM_NAN_BITS_64
Definition: Math.h:1862
vtkm::Floor
vtkm::Float32 Floor(vtkm::Float32 x)
Definition: Math.h:2230
vtkm
Groups connected points that have the same field value.
Definition: Atomic.h:19
VTKM_CUDA_MATH_FUNCTION_64
#define VTKM_CUDA_MATH_FUNCTION_64(func)
Definition: Math.h:43
vtkm::Sqrt
vtkm::Float32 Sqrt(vtkm::Float32 x)
Definition: Math.h:943
vtkm::Clamp
vtkm::Float32 Clamp(vtkm::Float32 x, vtkm::Float32 lo, vtkm::Float32 hi)
Definition: Math.h:1827
VTKM_INF_BITS_64
#define VTKM_INF_BITS_64
Definition: Math.h:1863
Types.h
VTKM_ASSERT
#define VTKM_ASSERT(condition)
Definition: Assert.h:43
VTKM_EPSILON_32
#define VTKM_EPSILON_32
Definition: Math.h:1856
VTKM_EXEC_CONT
#define VTKM_EXEC_CONT
Definition: ExportMacros.h:52
vtkm::IdComponent
vtkm::Int32 IdComponent
Base type to use to index small lists.
Definition: Types.h:194
vtkm::Log10
vtkm::Float32 Log10(vtkm::Float32 x)
Definition: Math.h:1578
vtkm::VecTraits::ComponentType
T ComponentType
Type of the components in the vector.
Definition: VecTraits.h:71
vtkm::Tan
vtkm::Float32 Tan(vtkm::Float32 x)
Definition: Math.h:288
VTKM_NEG_INF_BITS_64
#define VTKM_NEG_INF_BITS_64
Definition: Math.h:1864
vtkm::Log1P
vtkm::Float32 Log1P(vtkm::Float32 x)
Definition: Math.h:1639
vtkm::DifferenceOfProducts
T DifferenceOfProducts(T a, T b, T c, T d)
Definition: Math.h:2778
vtkm::Exp
vtkm::Float32 Exp(vtkm::Float32 x)
Definition: Math.h:1205
vtkm::Log2
vtkm::Float32 Log2(vtkm::Float32 x)
Definition: Math.h:1517
TypeTraits.h
vtkm::Frexp
vtkm::Float32 Frexp(vtkm::Float32 x, vtkm::Int32 *exponent)
Definition: Math.h:2608
vtkm::Ceil
vtkm::Float32 Ceil(vtkm::Float32 x)
Definition: Math.h:2169
vtkm::FindFirstSetBit
vtkm::Int32 FindFirstSetBit(vtkm::UInt32 word)
Bitwise operations.
Definition: Math.h:2844
VTKM_NAN_BITS_32
#define VTKM_NAN_BITS_32
Definition: Math.h:1853
vtkm::Vec< T, 4 >
Definition: Types.h:1131
vtkm::ACos
vtkm::Float32 ACos(vtkm::Float32 x)
Definition: Math.h:410
vtkm::Ldexp
vtkm::Float32 Ldexp(vtkm::Float32 x, vtkm::Int32 exponent)
Definition: Math.h:2628
vtkm::Round
vtkm::Float32 Round(vtkm::Float32 x)
Definition: Math.h:2291
vtkm::Vec< T, 2 >
Definition: Types.h:897
vtkm::Int64
signed long long Int64
Base type to use for 64-bit signed integer numbers.
Definition: Types.h:204
VTKM_INF_BITS_32
#define VTKM_INF_BITS_32
Definition: Math.h:1854
vtkm::TypeTraitsUnknownTag
Tag used to identify types that aren't Real, Integer, Scalar or Vector.
Definition: TypeTraits.h:20
vtkm::TypeTraitsScalarTag
Tag used to identify 0 dimensional types (scalars).
Definition: TypeTraits.h:44
vtkm::TypeTraitsVectorTag
Tag used to identify 1 dimensional types (vectors).
Definition: TypeTraits.h:51
vtkm::SinH
vtkm::Float32 SinH(vtkm::Float32 x)
Definition: Math.h:554
vtkm::Exp2
vtkm::Float32 Exp2(vtkm::Float32 x)
Definition: Math.h:1266
vtkm::Vec< T, 3 >
Definition: Types.h:1013
vtkm::Cbrt
vtkm::Float32 Cbrt(vtkm::Float32 x)
Definition: Math.h:1074
vtkm::Vec
A short fixed-length array.
Definition: Types.h:357
vtkm::UInt32
uint32_t UInt32
Base type to use for 32-bit unsigned integer numbers.
Definition: Types.h:185
vtkm::ASin
vtkm::Float32 ASin(vtkm::Float32 x)
Definition: Math.h:349
vtkm::Log
vtkm::Float32 Log(vtkm::Float32 x)
Definition: Math.h:1456
vtkm::Sin
vtkm::Float32 Sin(vtkm::Float32 x)
Definition: Math.h:166
vtkm::Float32
float Float32
Base type to use for 32-bit floating-point numbers.
Definition: Types.h:157
vtkm::UInt64
unsigned long long UInt64
Base type to use for 64-bit signed integer numbers.
Definition: Types.h:207
vtkm::Int32
int32_t Int32
Base type to use for 32-bit signed integer numbers.
Definition: Types.h:181
vtkm::Float64
double Float64
Base type to use for 64-bit floating-point numbers.
Definition: Types.h:161
vtkm::ExpM1
vtkm::Float32 ExpM1(vtkm::Float32 x)
Definition: Math.h:1328
vtkm::QuadraticRoots
vtkm::Vec< T, 2 > QuadraticRoots(T a, T b, T c)
Solves ax² + bx + c = 0.
Definition: Math.h:2795
VTKM_CUDA_MATH_FUNCTION_32
#define VTKM_CUDA_MATH_FUNCTION_32(func)
Definition: Math.h:42
VTKM_EPSILON_64
#define VTKM_EPSILON_64
Definition: Math.h:1865
vtkm::VecTraits
Traits that can be queried to treat any type as a Vec.
Definition: VecTraits.h:61
vtkm::ATan
vtkm::Float32 ATan(vtkm::Float32 x)
Definition: Math.h:471
VTKM_NEG_INF_BITS_32
#define VTKM_NEG_INF_BITS_32
Definition: Math.h:1855
vtkm::ASinH
vtkm::Float32 ASinH(vtkm::Float32 x)
Definition: Math.h:737
vtkm::Cos
vtkm::Float32 Cos(vtkm::Float32 x)
Definition: Math.h:227
vtkm::ATanH
vtkm::Float32 ATanH(vtkm::Float32 x)
Definition: Math.h:859
Windows.h
VecTraits.h
vtkm::CountSetBits
vtkm::Int32 CountSetBits(vtkm::UInt32 word)
Count the total number of bits set in word.
Definition: Math.h:2940
vtkm::FloatDistance
vtkm::UInt64 FloatDistance(vtkm::Float64 x, vtkm::Float64 y)
Computes the number of representables between two floating point numbers.
Definition: Math.h:2698