2019-04-15 23:24:21 +00:00
//============================================================================
2015-05-27 04:32:10 +00:00
// Copyright (c) Kitware, Inc.
// All rights reserved.
// See LICENSE.txt for details.
//
// This software is distributed WITHOUT ANY WARRANTY; without even
// the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
// PURPOSE. See the above copyright notice for more information.
2019-04-15 23:24:21 +00:00
//============================================================================
2015-05-27 04:32:10 +00:00
$ # This file uses the pyexpander macro processing utility to build the
$ # FunctionInterface facilities that use a variable number of arguments .
$ # Information , documentation , and downloads for pyexpander can be found at :
$ #
$ # http : //pyexpander.sourceforge.net/
$ #
$ # To build the source code , execute the following ( after installing
$ # pyexpander , of course ) :
$ #
$ # expander . py Math . h . in > Math . h
$ #
$ # Ignore the following comment . It is meant for the generated file .
// **** DO NOT EDIT THIS FILE!!! ****
2016-09-02 17:32:26 +00:00
// This file is automatically generated by Math.h.in
2015-05-27 04:32:10 +00:00
# ifndef vtk_m_Math_h
# define vtk_m_Math_h
2015-07-22 17:27:44 +00:00
# include <vtkm/TypeTraits.h>
2017-05-26 00:50:45 +00:00
# include <vtkm/Types.h>
2015-07-22 17:27:44 +00:00
# include <vtkm/VecTraits.h>
2015-05-27 04:32:10 +00:00
2020-06-24 14:52:59 +00:00
# include <limits> // must be found with or without CUDA.
2015-05-27 04:32:10 +00:00
# ifndef VTKM_CUDA
2017-05-26 00:50:45 +00:00
# include <cmath>
2020-06-24 14:52:59 +00:00
# include <cstring>
2015-06-23 23:35:55 +00:00
# include <limits.h>
2015-05-27 04:32:10 +00:00
# include <math.h>
2015-06-24 21:40:18 +00:00
# include <stdlib.h>
2015-06-23 23:35:55 +00:00
# endif // !VTKM_CUDA
2015-05-27 04:32:10 +00:00
2019-01-24 16:08:37 +00:00
# if !defined(VTKM_CUDA_DEVICE_PASS)
2016-09-02 17:32:26 +00:00
# define VTKM_USE_STL
2015-12-02 21:06:46 +00:00
# include <algorithm>
# endif
2019-03-05 16:46:32 +00:00
# ifdef VTKM_MSVC
2019-04-23 20:28:35 +00:00
# include <intrin.h> // For bitwise intrinsics (__popcnt, etc)
2019-03-05 16:46:32 +00:00
# include <vtkm/internal/Windows.h> // for types used by MSVC intrinsics.
# ifndef VTKM_CUDA
2016-09-02 17:32:26 +00:00
# include <math.h>
2019-03-05 16:46:32 +00:00
# endif // VTKM_CUDA
# endif // VTKM_MSVC
2017-05-26 00:50:45 +00:00
# define VTKM_CUDA_MATH_FUNCTION_32(func) func##f
2016-09-02 17:32:26 +00:00
# define VTKM_CUDA_MATH_FUNCTION_64(func) func
2015-05-27 04:32:10 +00:00
$ py (
2017-01-26 20:30:28 +00:00
def unary_function ( name , type , returntype , cuda_expression , std_expression , template_header , static_keywords ) :
return ' ' ' { 5 }
2017-05-26 00:50:45 +00:00
{ 6 } VTKM_EXEC_CONT { 2 } { 0 } ( { 1 } x )
{ {
2016-09-02 17:32:26 +00:00
# ifdef VTKM_CUDA
2015-06-24 21:40:18 +00:00
return { 3 } ;
2016-09-02 17:32:26 +00:00
# else
return { 4 } ;
# endif
2015-05-27 04:32:10 +00:00
} }
2017-01-26 20:30:28 +00:00
' ' ' . format ( name , type , returntype , cuda_expression , std_expression , template_header , static_keywords )
2017-01-24 21:47:51 +00:00
2015-05-27 04:32:10 +00:00
def unary_Vec_function ( vtkmname ) :
2017-05-26 00:50:45 +00:00
return ' ' ' template < typename T , vtkm : : IdComponent N >
static inline VTKM_EXEC_CONT vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , N > { 0 } (
const vtkm : : Vec < T , N > & x )
{ {
vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , N > result ;
2015-05-27 04:32:10 +00:00
for ( vtkm : : IdComponent index = 0 ; index < N ; index + + )
{ {
result [ index ] = vtkm : : { 0 } ( x [ index ] ) ;
} }
return result ;
} }
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , 4 > { 0 } (
const vtkm : : Vec < T , 4 > & x )
{ {
return vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , 4 > (
vtkm : : { 0 } ( x [ 0 ] ) , vtkm : : { 0 } ( x [ 1 ] ) , vtkm : : { 0 } ( x [ 2 ] ) , vtkm : : { 0 } ( x [ 3 ] ) ) ;
2015-05-27 04:32:10 +00:00
} }
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , 3 > { 0 } (
const vtkm : : Vec < T , 3 > & x )
{ {
return vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , 3 > (
vtkm : : { 0 } ( x [ 0 ] ) , vtkm : : { 0 } ( x [ 1 ] ) , vtkm : : { 0 } ( x [ 2 ] ) ) ;
2015-05-27 04:32:10 +00:00
} }
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , 2 > { 0 } (
const vtkm : : Vec < T , 2 > & x )
{ {
return vtkm : : Vec < typename detail : : FloatingPointReturnType < T > : : Type , 2 > ( vtkm : : { 0 } ( x [ 0 ] ) ,
vtkm : : { 0 } ( x [ 1 ] ) ) ;
2015-05-27 04:32:10 +00:00
} }
' ' ' . format ( vtkmname )
2018-01-04 18:19:39 +00:00
def unary_math_function_no_vec ( vtkmname , sysname ) :
general_type = ' ' ' template < typename T >
static inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type { 0 } ( const T & x )
{ {
using RT = typename detail : : FloatingPointReturnType < T > : : Type ;
return vtkm : : { 0 } ( static_cast < RT > ( x ) ) ;
} }
' ' ' . format ( vtkmname )
specialization_types = unary_function ( vtkmname ,
' vtkm : : Float32 ' ,
' vtkm : : Float32 ' ,
' VTKM_CUDA_MATH_FUNCTION_32 ( ' + sysname + ' ) ( x ) ' ,
' std : : ' + sysname + ' ( x ) ' ,
' ' ,
' inline ' )
specialization_types + = unary_function ( vtkmname ,
' vtkm : : Float64 ' ,
' vtkm : : Float64 ' ,
' VTKM_CUDA_MATH_FUNCTION_64 ( ' + sysname + ' ) ( x ) ' ,
' std : : ' + sysname + ' ( x ) ' ,
' ' ,
' inline ' )
return specialization_types + general_type
2015-06-22 22:24:18 +00:00
def unary_math_function ( vtkmname , sysname ) :
return unary_math_function_no_vec ( vtkmname , sysname ) + \
2015-05-27 04:32:10 +00:00
unary_Vec_function ( vtkmname )
2015-06-22 22:24:18 +00:00
2015-06-24 21:40:18 +00:00
def unary_template_function_no_vec ( vtkmname ,
expression ,
returntype = None ,
preexpression = ' ' ) :
2017-05-26 00:50:45 +00:00
return ' ' ' static inline VTKM_EXEC_CONT { 2 } { 0 } ( vtkm : : Float32 x )
{ {
2015-06-24 21:40:18 +00:00
{ 3 } return { 1 } ;
2015-06-22 22:24:18 +00:00
} }
2015-06-24 21:40:18 +00:00
' ' ' . format ( vtkmname ,
expression ,
' vtkm : : Float32 ' if returntype = = None else returntype ,
preexpression ) + \
2017-05-26 00:50:45 +00:00
' ' ' static inline VTKM_EXEC_CONT { 2 } { 0 } ( vtkm : : Float64 x )
{ {
2015-06-24 21:40:18 +00:00
{ 3 } return { 1 } ;
2015-06-22 22:24:18 +00:00
} }
2015-06-24 21:40:18 +00:00
' ' ' . format ( vtkmname ,
expression ,
' vtkm : : Float64 ' if returntype = = None else returntype ,
preexpression )
2015-06-22 22:24:18 +00:00
2016-09-02 17:32:26 +00:00
def binary_function ( name , type , cuda_expression , std_expression ) :
2017-05-26 00:50:45 +00:00
return ' ' ' static inline VTKM_EXEC_CONT { 1 } { 0 } ( { 1 } x , { 1 } y )
{ {
2016-09-02 17:32:26 +00:00
# ifdef VTKM_CUDA
2015-06-22 22:24:18 +00:00
return { 2 } ;
2016-09-02 17:32:26 +00:00
# else
return { 3 } ;
# endif
2015-06-22 22:24:18 +00:00
} }
2016-09-02 17:32:26 +00:00
' ' ' . format ( name , type , cuda_expression , std_expression )
2015-06-22 22:24:18 +00:00
def binary_math_function ( vtkmname , sysname ) :
return binary_function ( vtkmname ,
' vtkm : : Float32 ' ,
2017-05-26 00:50:45 +00:00
' VTKM_CUDA_MATH_FUNCTION_32 ( ' + sysname + ' ) ( x , y ) ' ,
' std : : ' + sysname + ' ( x , y ) ' ) + \
2015-06-22 22:24:18 +00:00
binary_function ( vtkmname ,
' vtkm : : Float64 ' ,
2017-05-26 00:50:45 +00:00
' VTKM_CUDA_MATH_FUNCTION_64 ( ' + sysname + ' ) ( x , y ) ' ,
' std : : ' + sysname + ' ( x , y ) ' )
2015-06-22 23:11:23 +00:00
def binary_template_function ( vtkmname , expression ) :
2017-05-26 00:50:45 +00:00
return ' ' ' static inline VTKM_EXEC_CONT vtkm : : Float32 { 0 } ( vtkm : : Float32 x , vtkm : : Float32 y )
{ {
2015-06-22 23:11:23 +00:00
return { 1 } ;
} }
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 { 0 } ( vtkm : : Float64 x , vtkm : : Float64 y )
{ {
2015-06-22 23:11:23 +00:00
return { 1 } ;
} }
' ' ' . format ( vtkmname , expression )
2018-05-15 21:29:31 +00:00
2017-05-26 00:50:45 +00:00
) \
$ extend ( unary_math_function ) \
$ extend ( unary_math_function_no_vec ) \
$ extend ( unary_Vec_function ) \
$ extend ( unary_template_function_no_vec ) \
$ extend ( binary_math_function ) \
$ extend ( binary_template_function ) \
\
namespace vtkm
{
2015-05-27 04:32:10 +00:00
2015-06-24 22:27:07 +00:00
//-----------------------------------------------------------------------------
2018-05-15 21:29:31 +00:00
namespace detail
{
template < typename T >
struct FloatingPointReturnType
{
using ctype = typename vtkm : : VecTraits < T > : : ComponentType ;
2024-02-08 13:38:40 +00:00
using representable_as_float_type =
std : : integral_constant < bool ,
( ( sizeof ( ctype ) < sizeof ( float ) ) | |
std : : is_same < ctype , vtkm : : Float32 > : : value ) > ;
using Type = typename std : :
conditional < representable_as_float_type : : value , vtkm : : Float32 , vtkm : : Float64 > : : type ;
2018-05-15 21:29:31 +00:00
} ;
} // namespace detail
2024-02-08 13:38:40 +00:00
/// Returns the constant 2 times Pi.
2018-05-15 21:29:31 +00:00
///
2024-02-08 13:38:40 +00:00
template < typename T = vtkm : : Float64 >
static constexpr inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type TwoPi ( )
2018-05-15 21:29:31 +00:00
{
2024-02-08 13:38:40 +00:00
using FT = typename detail : : FloatingPointReturnType < T > : : Type ;
return static_cast < FT > ( 6.28318530717958647692528676655900576 ) ;
2018-05-15 21:29:31 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi.
2018-05-15 21:29:31 +00:00
///
2024-02-08 13:38:40 +00:00
template < typename T = vtkm : : Float64 >
static constexpr inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type Pi ( )
2018-05-15 21:29:31 +00:00
{
2024-02-08 13:38:40 +00:00
using FT = typename detail : : FloatingPointReturnType < T > : : Type ;
return static_cast < FT > ( 3.14159265358979323846264338327950288 ) ;
2018-05-15 21:29:31 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi halves.
2018-05-15 21:29:31 +00:00
///
2024-02-08 13:38:40 +00:00
template < typename T = vtkm : : Float64 >
static constexpr inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type Pi_2 ( )
2018-05-15 21:29:31 +00:00
{
2024-02-08 13:38:40 +00:00
using FT = typename detail : : FloatingPointReturnType < T > : : Type ;
return static_cast < FT > ( 1.57079632679489661923132169163975144 ) ;
2018-05-15 21:29:31 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi thirds.
2018-05-15 21:29:31 +00:00
///
2024-02-08 13:38:40 +00:00
template < typename T = vtkm : : Float64 >
static constexpr inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type Pi_3 ( )
2018-05-15 21:29:31 +00:00
{
2024-02-08 13:38:40 +00:00
using FT = typename detail : : FloatingPointReturnType < T > : : Type ;
return static_cast < FT > ( 1.04719755119659774615421446109316762 ) ;
2018-05-15 21:29:31 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi fourths.
2018-05-15 21:29:31 +00:00
///
2024-02-08 13:38:40 +00:00
template < typename T = vtkm : : Float64 >
static constexpr inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type Pi_4 ( )
2018-05-15 21:29:31 +00:00
{
2024-02-08 13:38:40 +00:00
using FT = typename detail : : FloatingPointReturnType < T > : : Type ;
return static_cast < FT > ( 0.78539816339744830961566084581987572 ) ;
2018-05-15 21:29:31 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi one hundred and eightieth.
2015-09-02 17:06:02 +00:00
///
2024-02-08 13:38:40 +00:00
template < typename T = vtkm : : Float64 >
static constexpr inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type Pi_180 ( )
2018-05-15 21:29:31 +00:00
{
2024-02-08 13:38:40 +00:00
using FT = typename detail : : FloatingPointReturnType < T > : : Type ;
return static_cast < FT > ( 0.01745329251994329547437168059786927 ) ;
2018-05-15 21:29:31 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant 2 times Pi in float32.
2018-05-15 21:29:31 +00:00
///
2024-02-08 13:38:40 +00:00
static constexpr inline VTKM_EXEC_CONT vtkm : : Float32 TwoPif ( )
2015-09-02 17:06:02 +00:00
{
2024-02-08 13:38:40 +00:00
return TwoPi < vtkm : : Float32 > ( ) ;
2015-09-02 17:06:02 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi in float32.
2015-06-24 22:27:07 +00:00
///
2024-02-08 13:38:40 +00:00
static constexpr inline VTKM_EXEC_CONT vtkm : : Float32 Pif ( )
2015-06-24 22:27:07 +00:00
{
2024-02-08 13:38:40 +00:00
return Pi < vtkm : : Float32 > ( ) ;
2015-06-24 22:27:07 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi halves in float32.
2015-09-02 17:06:02 +00:00
///
2024-02-08 13:38:40 +00:00
static constexpr inline VTKM_EXEC_CONT vtkm : : Float32 Pi_2f ( )
2015-09-02 17:06:02 +00:00
{
2024-02-08 13:38:40 +00:00
return Pi_2 < vtkm : : Float32 > ( ) ;
2015-09-02 17:06:02 +00:00
}
2018-05-15 21:29:31 +00:00
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi thirds in float32.
2015-09-02 17:06:02 +00:00
///
2024-02-08 13:38:40 +00:00
static constexpr inline VTKM_EXEC_CONT vtkm : : Float32 Pi_3f ( )
2015-09-02 17:06:02 +00:00
{
2024-02-08 13:38:40 +00:00
return Pi_3 < vtkm : : Float32 > ( ) ;
2015-09-02 17:06:02 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi fourths in float32.
2015-09-02 17:06:02 +00:00
///
2024-02-08 13:38:40 +00:00
static constexpr inline VTKM_EXEC_CONT vtkm : : Float32 Pi_4f ( )
2015-09-02 17:06:02 +00:00
{
2024-02-08 13:38:40 +00:00
return Pi_4 < vtkm : : Float32 > ( ) ;
2015-09-02 17:06:02 +00:00
}
2024-02-08 13:38:40 +00:00
/// Returns the constant Pi one hundred and eightieth in float32.
2018-05-15 21:29:31 +00:00
///
2024-02-08 13:38:40 +00:00
static constexpr inline VTKM_EXEC_CONT vtkm : : Float32 Pi_180f ( )
2017-05-26 00:50:45 +00:00
{
2024-02-08 13:38:40 +00:00
return Pi_180 < vtkm : : Float32 > ( ) ;
2018-05-15 21:29:31 +00:00
}
2024-02-08 13:38:40 +00:00
// clang-format off
2017-01-26 20:30:28 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the sine of \p x.
///
$ unary_math_function ( ' Sin ' , ' sin ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the cosine of \p x.
///
$ unary_math_function ( ' Cos ' , ' cos ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the tangent of \p x.
///
$ unary_math_function ( ' Tan ' , ' tan ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the arc sine of \p x.
///
$ unary_math_function ( ' ASin ' , ' asin ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the arc cosine of \p x.
///
$ unary_math_function ( ' ACos ' , ' acos ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the arc tangent of \p x.
///
$ unary_math_function ( ' ATan ' , ' atan ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the arc tangent of \p x / \p y using the signs of both arguments
/// to determine the quadrant of the return value.
///
$ binary_math_function ( ' ATan2 ' , ' atan2 ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the hyperbolic sine of \p x.
///
$ unary_math_function ( ' SinH ' , ' sinh ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the hyperbolic cosine of \p x.
///
$ unary_math_function ( ' CosH ' , ' cosh ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the hyperbolic tangent of \p x.
///
$ unary_math_function ( ' TanH ' , ' tanh ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the hyperbolic arc sine of \p x.
///
$ unary_math_function_no_vec ( ' ASinH ' , ' asinh ' ) \
$ #
$ unary_Vec_function ( ' ASinH ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the hyperbolic arc cosine of \p x.
///
$ unary_math_function_no_vec ( ' ACosH ' , ' acosh ' ) \
$ #
$ unary_Vec_function ( ' ACosH ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 22:27:07 +00:00
/// Compute the hyperbolic arc tangent of \p x.
///
$ unary_math_function_no_vec ( ' ATanH ' , ' atanh ' ) \
$ #
$ unary_Vec_function ( ' ATanH ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 22:27:07 +00:00
2015-06-23 23:35:55 +00:00
//-----------------------------------------------------------------------------
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 22:24:18 +00:00
/// Computes \p x raised to the power of \p y.
///
2015-06-24 21:40:18 +00:00
$ binary_math_function ( ' Pow ' , ' pow ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-05-27 04:32:10 +00:00
/// Compute the square root of \p x.
2024-02-11 22:05:57 +00:00
/// On some hardware it is faster to find the reciprocal square root, so `RSqrt`
/// should be used if you actually plan to divide by the square root.
2015-06-24 21:40:18 +00:00
$ unary_math_function ( ' Sqrt ' , ' sqrt ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-05-27 04:32:10 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 22:24:18 +00:00
/// Compute the reciprocal square root of \p x. The result of this function is
/// equivalent to <tt>1/Sqrt(x)</tt>. However, on some devices it is faster to
/// compute the reciprocal square root than the regular square root. Thus, you
/// should use this function whenever dividing by the square root.
///
# ifdef VTKM_CUDA
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 RSqrt ( vtkm : : Float32 x )
{
2016-09-02 17:32:26 +00:00
return rsqrtf ( x ) ;
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 RSqrt ( vtkm : : Float64 x )
{
2016-09-02 17:32:26 +00:00
return rsqrt ( x ) ;
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Float64 RSqrt ( T x )
{
2017-01-26 20:30:28 +00:00
return rsqrt ( static_cast < vtkm : : Float64 > ( x ) ) ;
}
2017-05-26 00:50:45 +00:00
# else // !VTKM_CUDA
static inline VTKM_EXEC_CONT vtkm : : Float32 RSqrt ( vtkm : : Float32 x )
{
return 1 / vtkm : : Sqrt ( x ) ;
2016-09-02 17:32:26 +00:00
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 RSqrt ( vtkm : : Float64 x )
{
return 1 / vtkm : : Sqrt ( x ) ;
2016-09-02 17:32:26 +00:00
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Float64 RSqrt ( T x )
{
return 1 / static_cast < vtkm : : Float64 > ( x ) ;
2017-01-26 20:30:28 +00:00
}
2015-06-22 22:24:18 +00:00
# endif // !VTKM_CUDA
2016-09-02 17:32:26 +00:00
2015-06-24 21:40:18 +00:00
$ unary_Vec_function ( ' RSqrt ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 22:24:18 +00:00
/// Compute the cube root of \p x.
///
2015-06-24 21:40:18 +00:00
$ unary_math_function_no_vec ( ' Cbrt ' , ' cbrt ' ) \
$ #
$ unary_Vec_function ( ' Cbrt ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 22:24:18 +00:00
/// Compute the reciprocal cube root of \p x. The result of this function is
/// equivalent to <tt>1/Cbrt(x)</tt>. However, on some devices it is faster to
/// compute the reciprocal cube root than the regular cube root. Thus, you
/// should use this function whenever dividing by the cube root.
///
# ifdef VTKM_CUDA
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 RCbrt ( vtkm : : Float32 x )
{
2016-09-02 17:32:26 +00:00
return rcbrtf ( x ) ;
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 RCbrt ( vtkm : : Float64 x )
{
2016-09-02 17:32:26 +00:00
return rcbrt ( x ) ;
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Float64 RCbrt ( T x )
{
2017-01-26 20:30:28 +00:00
return rcbrt ( static_cast < vtkm : : Float64 > ( x ) ) ;
}
2017-05-26 00:50:45 +00:00
# else // !VTKM_CUDA
static inline VTKM_EXEC_CONT vtkm : : Float32 RCbrt ( vtkm : : Float32 x )
{
return 1 / vtkm : : Cbrt ( x ) ;
2016-09-02 17:32:26 +00:00
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 RCbrt ( vtkm : : Float64 x )
{
return 1 / vtkm : : Cbrt ( x ) ;
2016-09-02 17:32:26 +00:00
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Float64 RCbrt ( T x )
{
return 1 / vtkm : : Cbrt ( static_cast < vtkm : : Float64 > ( x ) ) ;
2017-01-26 20:30:28 +00:00
}
2015-06-22 22:24:18 +00:00
# endif // !VTKM_CUDA
2016-09-02 17:32:26 +00:00
2015-06-24 21:40:18 +00:00
$ unary_Vec_function ( ' RCbrt ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
/// Computes e^x, the base-e exponential of `x`.
2015-06-22 22:24:18 +00:00
///
2015-06-24 21:40:18 +00:00
$ unary_math_function ( ' Exp ' , ' exp ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
/// Computes 2^x, the base-2 exponential of `x`.
2015-06-22 22:24:18 +00:00
///
2015-06-24 21:40:18 +00:00
$ unary_math_function_no_vec ( ' Exp2 ' , ' exp2 ' ) \
$ #
$ unary_Vec_function ( ' Exp2 ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
/// Computes (e^x) - 1, the of base-e exponental of `x` then minus 1. The
/// accuracy of this function is good even for very small values of `x`.
2015-06-22 22:24:18 +00:00
///
2015-06-24 21:40:18 +00:00
$ unary_math_function_no_vec ( ' ExpM1 ' , ' expm1 ' ) \
$ #
$ unary_Vec_function ( ' ExpM1 ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
/// Computes 10^x, the base-10 exponential of `x`.
2015-06-22 22:24:18 +00:00
///
# ifdef VTKM_CUDA
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 Exp10 ( vtkm : : Float32 x )
{
2016-09-02 17:32:26 +00:00
return exp10f ( x ) ;
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 Exp10 ( vtkm : : Float64 x )
{
2016-09-02 17:32:26 +00:00
return exp10 ( x ) ;
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Float64 Exp10 ( T x )
{
2017-01-26 20:30:28 +00:00
return exp10 ( static_cast < vtkm : : Float64 > ( x ) ) ;
}
2017-05-26 00:50:45 +00:00
# else // !VTKM_CUDA
static inline VTKM_EXEC_CONT vtkm : : Float32 Exp10 ( vtkm : : Float32 x )
{
return vtkm : : Pow ( 10 , x ) ;
2016-09-02 17:32:26 +00:00
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 Exp10 ( vtkm : : Float64 x )
{
return vtkm : : Pow ( 10 , x ) ;
2016-09-02 17:32:26 +00:00
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Float64 Exp10 ( T x )
{
return vtkm : : Pow ( 10 , static_cast < vtkm : : Float64 > ( x ) ) ;
2017-01-26 20:30:28 +00:00
}
2015-06-22 22:24:18 +00:00
# endif // !VTKM_CUDA
2016-09-02 17:32:26 +00:00
2015-06-24 21:40:18 +00:00
$ unary_Vec_function ( ' Exp10 ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 22:24:18 +00:00
/// Computes the natural logarithm of \p x.
///
2015-06-24 21:40:18 +00:00
$ unary_math_function ( ' Log ' , ' log ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 22:24:18 +00:00
/// Computes the logarithm base 2 of \p x.
///
2015-06-24 21:40:18 +00:00
$ unary_math_function_no_vec ( ' Log2 ' , ' log2 ' ) \
$ #
$ unary_Vec_function ( ' Log2 ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 22:24:18 +00:00
/// Computes the logarithm base 10 of \p x.
///
2015-06-24 21:40:18 +00:00
$ unary_math_function ( ' Log10 ' , ' log10 ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
/// Computes the value of log(1+x). This method is more accurate for very small values of x
/// than the `vtkm::Log` function.
2015-06-24 21:40:18 +00:00
$ unary_math_function_no_vec ( ' Log1P ' , ' log1p ' ) \
$ #
$ unary_Vec_function ( ' Log1P ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 22:24:18 +00:00
2015-06-22 23:11:23 +00:00
//-----------------------------------------------------------------------------
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 23:11:23 +00:00
/// Returns \p x or \p y, whichever is larger.
///
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Max ( const T & x , const T & y ) ;
2016-09-02 17:32:26 +00:00
# ifdef VTKM_USE_STL
2015-06-24 21:40:18 +00:00
$ binary_template_function ( ' Max ' , ' ( std : : max ) ( x , y ) ' ) \
$ #
2016-09-02 17:32:26 +00:00
# else // !VTKM_USE_STL
2015-06-24 21:40:18 +00:00
$ binary_math_function ( ' Max ' , ' fmax ' ) \
$ #
2016-09-02 17:32:26 +00:00
# endif // !VTKM_USE_STL
2024-02-11 22:05:57 +00:00
///@}
2015-06-22 23:11:23 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-22 23:11:23 +00:00
/// Returns \p x or \p y, whichever is smaller.
///
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Min ( const T & x , const T & y ) ;
2020-12-28 18:29:32 +00:00
# if defined(VTKM_USE_STL) && !defined(VTKM_HIP)
2015-06-24 21:40:18 +00:00
$ binary_template_function ( ' Min ' , ' ( std : : min ) ( x , y ) ' ) \
$ #
2020-12-28 18:29:32 +00:00
# else // !VTKM_USE_STL OR HIP
2015-06-24 21:40:18 +00:00
$ binary_math_function ( ' Min ' , ' fmin ' ) \
$ #
2016-09-02 17:32:26 +00:00
# endif // !VTKM_USE_STL
2024-02-11 22:05:57 +00:00
///@}
2015-06-23 23:35:55 +00:00
2017-05-26 00:50:45 +00:00
namespace detail
{
2015-07-22 17:27:44 +00:00
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Max ( T x , T y , vtkm : : TypeTraitsScalarTag )
2015-07-22 17:27:44 +00:00
{
return ( x < y ) ? y : x ;
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Max ( const T & x , const T & y , vtkm : : TypeTraitsVectorTag )
2015-07-22 17:27:44 +00:00
{
2017-06-23 18:50:34 +00:00
using Traits = vtkm : : VecTraits < T > ;
2015-07-22 17:27:44 +00:00
T result ;
for ( vtkm : : IdComponent index = 0 ; index < Traits : : NUM_COMPONENTS ; index + + )
{
2017-05-31 15:37:29 +00:00
Traits : : SetComponent (
result , index , vtkm : : Max ( Traits : : GetComponent ( x , index ) , Traits : : GetComponent ( y , index ) ) ) ;
2015-07-22 17:27:44 +00:00
}
return result ;
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Min ( T x , T y , vtkm : : TypeTraitsScalarTag )
2015-07-22 17:27:44 +00:00
{
return ( x < y ) ? x : y ;
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Min ( const T & x , const T & y , vtkm : : TypeTraitsVectorTag )
2015-07-22 17:27:44 +00:00
{
2017-06-23 18:50:34 +00:00
using Traits = vtkm : : VecTraits < T > ;
2015-07-22 17:27:44 +00:00
T result ;
for ( vtkm : : IdComponent index = 0 ; index < Traits : : NUM_COMPONENTS ; index + + )
{
2017-05-31 15:37:29 +00:00
Traits : : SetComponent (
result , index , vtkm : : Min ( Traits : : GetComponent ( x , index ) , Traits : : GetComponent ( y , index ) ) ) ;
2015-07-22 17:27:44 +00:00
}
return result ;
}
} // namespace detail
/// Returns \p x or \p y, whichever is larger.
///
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Max ( const T & x , const T & y )
{
2015-07-22 17:27:44 +00:00
return detail : : Max ( x , y , typename vtkm : : TypeTraits < T > : : DimensionalityTag ( ) ) ;
}
/// Returns \p x or \p y, whichever is smaller.
///
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Min ( const T & x , const T & y )
{
2015-07-22 17:27:44 +00:00
return detail : : Min ( x , y , typename vtkm : : TypeTraits < T > : : DimensionalityTag ( ) ) ;
}
2024-02-11 22:05:57 +00:00
///@{
2018-06-08 04:27:30 +00:00
/// Clamp \p x to the given range.
///
inline VTKM_EXEC_CONT vtkm : : Float32 Clamp ( vtkm : : Float32 x , vtkm : : Float32 lo , vtkm : : Float32 hi )
{
return x > lo ? ( x < hi ? x : hi ) : lo ;
}
inline VTKM_EXEC_CONT vtkm : : Float64 Clamp ( vtkm : : Float64 x , vtkm : : Float64 lo , vtkm : : Float64 hi )
{
return x > lo ? ( x < hi ? x : hi ) : lo ;
}
2024-02-11 22:05:57 +00:00
///@}
2015-06-23 23:35:55 +00:00
//-----------------------------------------------------------------------------
2015-06-25 18:20:29 +00:00
//#ifdef VTKM_CUDA
2015-06-23 23:35:55 +00:00
# define VTKM_USE_IEEE_NONFINITE
2015-06-25 18:20:29 +00:00
//#endif
2015-06-23 23:35:55 +00:00
# ifdef VTKM_USE_IEEE_NONFINITE
2017-05-26 00:50:45 +00:00
namespace detail
{
2015-06-23 23:35:55 +00:00
union IEEE754Bits32 {
vtkm : : UInt32 bits ;
vtkm : : Float32 scalar ;
} ;
2017-05-26 00:50:45 +00:00
# define VTKM_NAN_BITS_32 0x7FC00000U
# define VTKM_INF_BITS_32 0x7F800000U
# define VTKM_NEG_INF_BITS_32 0xFF800000U
# define VTKM_EPSILON_32 1e-5f
2015-06-23 23:35:55 +00:00
union IEEE754Bits64 {
vtkm : : UInt64 bits ;
vtkm : : Float64 scalar ;
} ;
2017-05-26 00:50:45 +00:00
# define VTKM_NAN_BITS_64 0x7FF8000000000000ULL
# define VTKM_INF_BITS_64 0x7FF0000000000000ULL
# define VTKM_NEG_INF_BITS_64 0xFFF0000000000000ULL
# define VTKM_EPSILON_64 1e-9
2015-06-23 23:35:55 +00:00
2017-05-26 00:50:45 +00:00
template < typename T >
struct FloatLimits ;
2015-06-23 23:35:55 +00:00
2017-05-26 00:50:45 +00:00
template < >
2015-06-25 18:20:29 +00:00
struct FloatLimits < vtkm : : Float32 >
2015-06-23 23:35:55 +00:00
{
2017-06-23 18:50:34 +00:00
using BitsType = vtkm : : detail : : IEEE754Bits32 ;
2015-06-25 18:20:29 +00:00
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float32 Nan ( )
{
BitsType nan = { VTKM_NAN_BITS_32 } ;
2015-06-25 18:20:29 +00:00
return nan . scalar ;
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float32 Infinity ( )
{
BitsType inf = { VTKM_INF_BITS_32 } ;
2015-06-25 18:20:29 +00:00
return inf . scalar ;
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float32 NegativeInfinity ( )
{
BitsType neginf = { VTKM_NEG_INF_BITS_32 } ;
2015-06-25 18:20:29 +00:00
return neginf . scalar ;
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float32 Epsilon ( ) { return VTKM_EPSILON_32 ; }
2015-06-23 23:35:55 +00:00
} ;
2017-05-26 00:50:45 +00:00
template < int N >
struct FloatLimits < vtkm : : Vec < vtkm : : Float32 , N > >
2016-11-11 22:04:32 +00:00
{
2017-06-23 18:50:34 +00:00
using BitsType = vtkm : : detail : : IEEE754Bits32 ;
2016-11-11 22:04:32 +00:00
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float32 , N > Nan ( )
{
BitsType nan = { VTKM_NAN_BITS_32 } ;
return vtkm : : Vec < vtkm : : Float32 , N > ( nan . scalar ) ;
2016-11-11 22:04:32 +00:00
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float32 , N > Infinity ( )
{
BitsType inf = { VTKM_INF_BITS_32 } ;
return vtkm : : Vec < vtkm : : Float32 , N > ( inf . scalar ) ;
2016-11-11 22:04:32 +00:00
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float32 , N > NegativeInfinity ( )
{
BitsType neginf = { VTKM_NEG_INF_BITS_32 } ;
return vtkm : : Vec < vtkm : : Float32 , N > ( neginf . scalar ) ;
2016-11-11 22:04:32 +00:00
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float32 , N > Epsilon ( )
{
return vtkm : : Vec < vtkm : : Float32 , N > ( VTKM_EPSILON_32 ) ;
2016-11-11 22:04:32 +00:00
}
} ;
2017-05-26 00:50:45 +00:00
template < >
2015-06-25 18:20:29 +00:00
struct FloatLimits < vtkm : : Float64 >
2015-06-23 23:35:55 +00:00
{
2017-06-23 18:50:34 +00:00
using BitsType = vtkm : : detail : : IEEE754Bits64 ;
2015-06-25 18:20:29 +00:00
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float64 Nan ( )
{
BitsType nan = { VTKM_NAN_BITS_64 } ;
2015-06-25 18:20:29 +00:00
return nan . scalar ;
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float64 Infinity ( )
{
BitsType inf = { VTKM_INF_BITS_64 } ;
2015-06-25 18:20:29 +00:00
return inf . scalar ;
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float64 NegativeInfinity ( )
{
BitsType neginf = { VTKM_NEG_INF_BITS_64 } ;
2015-06-25 18:20:29 +00:00
return neginf . scalar ;
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Float64 Epsilon ( ) { return VTKM_EPSILON_64 ; }
2015-06-23 23:35:55 +00:00
} ;
2017-05-26 00:50:45 +00:00
template < int N >
struct FloatLimits < vtkm : : Vec < vtkm : : Float64 , N > >
2016-11-11 22:04:32 +00:00
{
2017-06-23 18:50:34 +00:00
using BitsType = vtkm : : detail : : IEEE754Bits64 ;
2016-11-11 22:04:32 +00:00
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float64 , N > Nan ( )
{
BitsType nan = { VTKM_NAN_BITS_64 } ;
return vtkm : : Vec < vtkm : : Float64 , N > ( nan . scalar ) ;
2016-11-11 22:04:32 +00:00
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float64 , N > Infinity ( )
{
BitsType inf = { VTKM_INF_BITS_64 } ;
return vtkm : : Vec < vtkm : : Float64 , N > ( inf . scalar ) ;
2016-11-11 22:04:32 +00:00
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float64 , N > NegativeInfinity ( )
{
BitsType neginf = { VTKM_NEG_INF_BITS_64 } ;
return vtkm : : Vec < vtkm : : Float64 , N > ( neginf . scalar ) ;
2016-11-11 22:04:32 +00:00
}
2016-10-19 22:42:58 +00:00
VTKM_EXEC_CONT
2017-05-26 00:50:45 +00:00
static vtkm : : Vec < vtkm : : Float64 , N > Epsilon ( )
{
return vtkm : : Vec < vtkm : : Float64 , N > ( VTKM_EPSILON_64 ) ;
2016-11-11 22:04:32 +00:00
}
} ;
2015-06-23 23:35:55 +00:00
# undef VTKM_NAN_BITS_32
# undef VTKM_INF_BITS_32
# undef VTKM_NEG_INF_BITS_32
# undef VTKM_EPSILON_32
# undef VTKM_NAN_BITS_64
# undef VTKM_INF_BITS_64
# undef VTKM_NEG_INF_BITS_64
# undef VTKM_EPSILON_64
} // namespace detail
2024-02-11 22:05:57 +00:00
/// Returns the representation for infinity. The result is greater than any other
/// number except another infinity or NaN. When comparing two infinities or infinity
/// to NaN, neither is greater than, less than, nor equal to the other. The
/// `Infinity` method is templated to specify either a 32 or 64 bit floating point
/// number. The convenience methods `Infinity32` and`Infinity64` are non-templated
/// versions that return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Nan ( )
2015-06-23 23:35:55 +00:00
{
2015-06-25 18:20:29 +00:00
return detail : : FloatLimits < T > : : Nan ( ) ;
2015-06-23 23:35:55 +00:00
}
2024-02-11 22:05:57 +00:00
/// Returns the representation for infinity. The result is greater than any other
/// number except another infinity or NaN. When comparing two infinities or infinity
/// to NaN, neither is greater than, less than, nor equal to the other. The
/// `Infinity` method is templated to specify either a 32 or 64 bit floating point
/// number. The convenience methods `Infinity32` and`Infinity64` are non-templated
/// versions that return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Infinity ( )
2015-06-23 23:35:55 +00:00
{
2015-06-25 18:20:29 +00:00
return detail : : FloatLimits < T > : : Infinity ( ) ;
2015-06-23 23:35:55 +00:00
}
2024-02-11 22:05:57 +00:00
/// Returns the representation for negative infinity. The result is less than any
/// other number except another negative infinity or NaN. When comparing two
/// negative infinities or negative infinity to NaN, neither is greater than, less
/// than, nor equal to the other. The `NegativeInfinity` method is templated to
/// specify either a 32 or 64 bit floating point number. The convenience methods
/// `NegativeInfinity32` and`NegativeInfinity64` are non-templated versions that
/// return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T NegativeInfinity ( )
2015-06-23 23:35:55 +00:00
{
2015-06-25 18:20:29 +00:00
return detail : : FloatLimits < T > : : NegativeInfinity ( ) ;
2015-06-23 23:35:55 +00:00
}
2024-02-11 22:05:57 +00:00
/// Returns the difference between 1 and the least value greater than 1 that
/// is representable by a floating point number. Epsilon is useful for specifying
/// the tolerance one should have when considering numerical error. The `Epsilon`
/// method is templated to specify either a 32 or 64 bit floating point number. The
/// convenience methods `Epsilon32` and`Epsilon64` are non-templated versions that
/// return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Epsilon ( )
2015-06-23 23:35:55 +00:00
{
2015-06-25 18:20:29 +00:00
return detail : : FloatLimits < T > : : Epsilon ( ) ;
2015-06-23 23:35:55 +00:00
}
2017-05-26 00:50:45 +00:00
# else // !VTKM_USE_IEEE_NONFINITE
2015-06-23 23:35:55 +00:00
2024-02-11 22:05:57 +00:00
/// Returns the representation for infinity. The result is greater than any other
/// number except another infinity or NaN. When comparing two infinities or infinity
/// to NaN, neither is greater than, less than, nor equal to the other. The
/// `Infinity` method is templated to specify either a 32 or 64 bit floating point
/// number. The convenience methods `Infinity32` and`Infinity64` are non-templated
/// versions that return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Nan ( )
2015-06-23 23:35:55 +00:00
{
return std : : numeric_limits < T > : : quiet_NaN ( ) ;
}
2024-02-11 22:05:57 +00:00
/// Returns the representation for infinity. The result is greater than any other
/// number except another infinity or NaN. When comparing two infinities or infinity
/// to NaN, neither is greater than, less than, nor equal to the other. The
/// `Infinity` method is templated to specify either a 32 or 64 bit floating point
/// number. The convenience methods `Infinity32` and`Infinity64` are non-templated
/// versions that return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Infinity ( )
2015-06-23 23:35:55 +00:00
{
return std : : numeric_limits < T > : : infinity ( ) ;
}
2024-02-11 22:05:57 +00:00
/// Returns the representation for negative infinity. The result is less than any
/// other number except another negative infinity or NaN. When comparing two
/// negative infinities or negative infinity to NaN, neither is greater than, less
/// than, nor equal to the other. The `NegativeInfinity` method is templated to
/// specify either a 32 or 64 bit floating point number. The convenience methods
/// `NegativeInfinity32` and`NegativeInfinity64` are non-templated versions that
/// return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T NegativeInfinity ( )
2015-06-23 23:35:55 +00:00
{
return - std : : numeric_limits < T > : : infinity ( ) ;
}
2024-02-11 22:05:57 +00:00
/// Returns the difference between 1 and the least value greater than 1 that
/// is representable by a floating point number. Epsilon is useful for specifying
/// the tolerance one should have when considering numerical error. The `Epsilon`
/// method is templated to specify either a 32 or 64 bit floating point number. The
/// convenience methods `Epsilon32` and`Epsilon64` are non-templated versions that
/// return the precision for a particular precision.
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Epsilon ( )
2015-06-23 23:35:55 +00:00
{
return std : : numeric_limits < T > : : epsilon ( ) ;
}
# endif // !VTKM_USE_IEEE_NONFINITE
2024-02-11 22:05:57 +00:00
/// @copydoc Nan
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 Nan32 ( )
{
2015-06-23 23:35:55 +00:00
return vtkm : : Nan < vtkm : : Float32 > ( ) ;
}
2024-02-11 22:05:57 +00:00
/// @copydoc Nan
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 Nan64 ( )
{
2015-06-23 23:35:55 +00:00
return vtkm : : Nan < vtkm : : Float64 > ( ) ;
}
2024-02-11 22:05:57 +00:00
/// @copydoc Infinity
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 Infinity32 ( )
{
2015-06-23 23:35:55 +00:00
return vtkm : : Infinity < vtkm : : Float32 > ( ) ;
}
2024-02-11 22:05:57 +00:00
/// @copydoc Infinity
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 Infinity64 ( )
{
2015-06-23 23:35:55 +00:00
return vtkm : : Infinity < vtkm : : Float64 > ( ) ;
}
2024-02-11 22:05:57 +00:00
/// @copydoc NegativeInfinity
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 NegativeInfinity32 ( )
{
2015-06-23 23:35:55 +00:00
return vtkm : : NegativeInfinity < vtkm : : Float32 > ( ) ;
}
2024-02-11 22:05:57 +00:00
/// @copydoc NegativeInfinity
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 NegativeInfinity64 ( )
{
2015-06-23 23:35:55 +00:00
return vtkm : : NegativeInfinity < vtkm : : Float64 > ( ) ;
}
2024-02-11 22:05:57 +00:00
/// @copydoc Epsilon
2016-10-19 22:42:58 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 Epsilon32 ( )
2015-07-02 17:20:21 +00:00
{
return vtkm : : Epsilon < vtkm : : Float32 > ( ) ;
}
2024-02-11 22:05:57 +00:00
/// @copydoc Epsilon
2016-10-19 22:42:58 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 Epsilon64 ( )
2015-07-02 17:20:21 +00:00
{
return vtkm : : Epsilon < vtkm : : Float64 > ( ) ;
}
2015-06-23 23:35:55 +00:00
//-----------------------------------------------------------------------------
/// Returns true if \p x is not a number.
///
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT bool IsNan ( T x )
2015-06-23 23:35:55 +00:00
{
2016-12-15 17:40:17 +00:00
# ifndef VTKM_CUDA
2016-09-02 17:32:26 +00:00
using std : : isnan ;
2016-12-15 17:40:17 +00:00
# endif
2015-06-23 23:35:55 +00:00
return ( isnan ( x ) ! = 0 ) ;
}
/// Returns true if \p x is positive or negative infinity.
///
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT bool IsInf ( T x )
2015-06-23 23:35:55 +00:00
{
2016-12-15 17:40:17 +00:00
# ifndef VTKM_CUDA
2016-09-02 17:32:26 +00:00
using std : : isinf ;
2016-12-15 17:40:17 +00:00
# endif
2015-06-23 23:35:55 +00:00
return ( isinf ( x ) ! = 0 ) ;
}
/// Returns true if \p x is a normal number (not NaN or infinite).
///
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT bool IsFinite ( T x )
2015-06-23 23:35:55 +00:00
{
2016-12-15 17:40:17 +00:00
# ifndef VTKM_CUDA
2016-09-02 17:32:26 +00:00
using std : : isfinite ;
2016-12-15 17:40:17 +00:00
# endif
2015-06-23 23:35:55 +00:00
return ( isfinite ( x ) ! = 0 ) ;
}
2015-06-30 14:50:02 +00:00
//-----------------------------------------------------------------------------
2024-02-11 22:05:57 +00:00
///@{
2015-06-30 14:50:02 +00:00
/// Round \p x to the smallest integer value not less than x.
///
$ unary_math_function ( ' Ceil ' , ' ceil ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-30 14:50:02 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-30 14:50:02 +00:00
/// Round \p x to the largest integer value not greater than x.
///
$ unary_math_function ( ' Floor ' , ' floor ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-30 14:50:02 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-30 14:50:02 +00:00
/// Round \p x to the nearest integral value.
///
$ unary_math_function_no_vec ( ' Round ' , ' round ' ) \
$ #
$ unary_Vec_function ( ' Round ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-30 14:50:02 +00:00
2015-06-23 23:35:55 +00:00
//-----------------------------------------------------------------------------
2024-02-11 22:05:57 +00:00
///@{
2015-06-23 23:35:55 +00:00
/// Computes the remainder on division of 2 floating point numbers. The return
/// value is \p numerator - n \p denominator, where n is the quotient of \p
/// numerator divided by \p denominator rounded towards zero to an integer. For
/// example, <tt>FMod(6.5, 2.3)</tt> returns 1.9, which is 6.5 - 2*2.3.
///
2015-06-24 21:40:18 +00:00
$ binary_math_function ( ' FMod ' , ' fmod ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-23 23:35:55 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-23 23:35:55 +00:00
/// Computes the remainder on division of 2 floating point numbers. The return
/// value is \p numerator - n \p denominator, where n is the quotient of \p
/// numerator divided by \p denominator rounded towards the nearest integer
/// (instead of toward zero like FMod). For example, <tt>FMod(6.5, 2.3)</tt>
/// returns -0.4, which is 6.5 - 3*2.3.
///
2015-06-30 14:50:02 +00:00
# ifdef VTKM_MSVC
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT T Remainder ( T numerator , T denominator )
2015-06-30 14:50:02 +00:00
{
2017-05-26 00:50:45 +00:00
T quotient = vtkm : : Round ( numerator / denominator ) ;
return numerator - quotient * denominator ;
2015-06-30 14:50:02 +00:00
}
# else // !VTKM_MSVC
2015-06-24 21:40:18 +00:00
$ binary_math_function ( ' Remainder ' , ' remainder ' ) \
2015-06-30 14:50:02 +00:00
$ #
# endif // !VTKM_MSVC
2024-02-11 22:05:57 +00:00
///@}
2015-06-23 23:35:55 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-23 23:35:55 +00:00
/// Returns the remainder on division of 2 floating point numbers just like
/// Remainder. In addition, this function also returns the \c quotient used to
/// get that remainder.
///
2017-05-26 00:50:45 +00:00
template < typename QType >
static inline VTKM_EXEC_CONT vtkm : : Float32 RemainderQuotient ( vtkm : : Float32 numerator ,
vtkm : : Float32 denominator ,
QType & quotient )
2015-06-23 23:35:55 +00:00
{
int iQuotient ;
2020-10-07 15:42:49 +00:00
// See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
# if defined(VTKM_CUDA) || defined(VTKM_HIP)
2018-04-20 19:23:25 +00:00
const vtkm : : Float32 result =
2018-01-04 22:14:40 +00:00
VTKM_CUDA_MATH_FUNCTION_32 ( remquo ) ( numerator , denominator , & iQuotient ) ;
# else
const vtkm : : Float32 result = std : : remquo ( numerator , denominator , & iQuotient ) ;
# endif
2018-04-20 19:23:25 +00:00
quotient = static_cast < QType > ( iQuotient ) ;
2015-06-23 23:35:55 +00:00
return result ;
}
2017-05-26 00:50:45 +00:00
template < typename QType >
static inline VTKM_EXEC_CONT vtkm : : Float64 RemainderQuotient ( vtkm : : Float64 numerator ,
vtkm : : Float64 denominator ,
QType & quotient )
2015-06-23 23:35:55 +00:00
{
int iQuotient ;
2018-01-04 22:14:40 +00:00
# ifdef VTKM_CUDA
const vtkm : : Float64 result =
VTKM_CUDA_MATH_FUNCTION_64 ( remquo ) ( numerator , denominator , & iQuotient ) ;
# else
const vtkm : : Float64 result = std : : remquo ( numerator , denominator , & iQuotient ) ;
# endif
2018-04-20 19:23:25 +00:00
quotient = static_cast < QType > ( iQuotient ) ;
2015-06-23 23:35:55 +00:00
return result ;
}
2024-02-11 22:05:57 +00:00
///@}
2015-06-23 23:35:55 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-23 23:35:55 +00:00
/// Gets the integral and fractional parts of \c x. The return value is the
/// fractional part and \c integral is set to the integral part.
///
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 ModF ( vtkm : : Float32 x , vtkm : : Float32 & integral )
2015-06-23 23:35:55 +00:00
{
2020-10-07 15:42:49 +00:00
// See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
# if defined(VTKM_CUDA) || defined(VTKM_HIP)
2020-10-07 15:40:41 +00:00
return VTKM_CUDA_MATH_FUNCTION_32 ( modf ) ( x , & integral ) ;
# else
2016-09-02 17:32:26 +00:00
return std : : modf ( x , & integral ) ;
2020-10-07 15:40:41 +00:00
# endif
2015-06-23 23:35:55 +00:00
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 ModF ( vtkm : : Float64 x , vtkm : : Float64 & integral )
2015-06-23 23:35:55 +00:00
{
2020-10-07 15:40:41 +00:00
# if defined(VTKM_CUDA)
return VTKM_CUDA_MATH_FUNCTION_64 ( modf ) ( x , & integral ) ;
# else
2016-09-02 17:32:26 +00:00
return std : : modf ( x , & integral ) ;
2020-10-07 15:40:41 +00:00
# endif
2015-06-23 23:35:55 +00:00
}
2024-02-11 22:05:57 +00:00
///@}
2015-06-23 23:35:55 +00:00
2015-06-24 21:40:18 +00:00
//-----------------------------------------------------------------------------
2024-02-11 22:05:57 +00:00
///@{
2015-07-22 15:51:14 +00:00
/// Return the absolute value of \p x. That is, return \p x if it is positive or
2015-06-24 21:40:18 +00:00
/// \p -x if it is negative.
///
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Int32 Abs ( vtkm : : Int32 x )
2015-06-24 21:40:18 +00:00
{
return abs ( x ) ;
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Int64 Abs ( vtkm : : Int64 x )
2015-06-24 21:40:18 +00:00
{
# if VTKM_SIZE_LONG == 8
return labs ( x ) ;
# elif VTKM_SIZE_LONG_LONG == 8
return llabs ( x ) ;
# else
# error Unknown size of Int64.
# endif
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float32 Abs ( vtkm : : Float32 x )
{
2017-02-06 22:50:13 +00:00
# ifdef VTKM_CUDA
return VTKM_CUDA_MATH_FUNCTION_32 ( fabs ) ( x ) ;
# else
return std : : fabs ( x ) ;
# endif
}
2017-05-26 00:50:45 +00:00
static inline VTKM_EXEC_CONT vtkm : : Float64 Abs ( vtkm : : Float64 x )
{
2017-02-06 22:50:13 +00:00
# ifdef VTKM_CUDA
return VTKM_CUDA_MATH_FUNCTION_64 ( fabs ) ( x ) ;
# else
return std : : fabs ( x ) ;
# endif
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT typename detail : : FloatingPointReturnType < T > : : Type Abs ( T x )
{
2017-02-06 22:50:13 +00:00
# ifdef VTKM_CUDA
return VTKM_CUDA_MATH_FUNCTION_64 ( fabs ) ( static_cast < vtkm : : Float64 > ( x ) ) ;
# else
return std : : fabs ( static_cast < vtkm : : Float64 > ( x ) ) ;
# endif
}
2017-05-26 00:50:45 +00:00
template < typename T , vtkm : : IdComponent N >
static inline VTKM_EXEC_CONT vtkm : : Vec < T , N > Abs ( const vtkm : : Vec < T , N > & x )
{
vtkm : : Vec < T , N > result ;
2017-02-06 22:50:13 +00:00
for ( vtkm : : IdComponent index = 0 ; index < N ; index + + )
{
result [ index ] = vtkm : : Abs ( x [ index ] ) ;
}
return result ;
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Vec < T , 4 > Abs ( const vtkm : : Vec < T , 4 > & x )
{
return vtkm : : Vec < T , 4 > ( vtkm : : Abs ( x [ 0 ] ) , vtkm : : Abs ( x [ 1 ] ) , vtkm : : Abs ( x [ 2 ] ) , vtkm : : Abs ( x [ 3 ] ) ) ;
2017-02-06 22:50:13 +00:00
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Vec < T , 3 > Abs ( const vtkm : : Vec < T , 3 > & x )
{
return vtkm : : Vec < T , 3 > ( vtkm : : Abs ( x [ 0 ] ) , vtkm : : Abs ( x [ 1 ] ) , vtkm : : Abs ( x [ 2 ] ) ) ;
2017-02-06 22:50:13 +00:00
}
2017-05-26 00:50:45 +00:00
template < typename T >
static inline VTKM_EXEC_CONT vtkm : : Vec < T , 2 > Abs ( const vtkm : : Vec < T , 2 > & x )
{
return vtkm : : Vec < T , 2 > ( vtkm : : Abs ( x [ 0 ] ) , vtkm : : Abs ( x [ 1 ] ) ) ;
2017-02-06 22:50:13 +00:00
}
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 21:40:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-07-22 15:51:14 +00:00
/// Returns a nonzero value if \p x is negative.
2015-06-24 21:40:18 +00:00
///
$ unary_template_function_no_vec ( ' SignBit ' ,
' static_cast < vtkm : : Int32 > ( signbit ( x ) ) ' ,
' vtkm : : Int32 ' ,
2016-11-30 20:07:54 +00:00
' ' ' # ifndef VTKM_CUDA
2016-09-02 17:32:26 +00:00
using std : : signbit ;
2016-11-30 20:07:54 +00:00
# endif
2015-06-24 21:40:18 +00:00
' ' ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 21:40:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 21:40:18 +00:00
/// Returns true if \p x is less than zero, false otherwise.
///
$ unary_template_function_no_vec ( ' IsNegative ' , ' ( vtkm : : SignBit ( x ) ! = 0 ) ' , ' bool ' ) \
2024-02-11 22:05:57 +00:00
///@}
2015-06-24 21:40:18 +00:00
2024-02-11 22:05:57 +00:00
///@{
2015-06-24 21:40:18 +00:00
/// Copies the sign of \p y onto \p x. If \p y is positive, returns Abs(\p x).
/// If \p y is negative, returns -Abs(\p x).
///
$ binary_math_function ( ' CopySign ' , ' copysign ' ) \
$ #
2016-09-02 17:32:26 +00:00
2017-05-26 00:50:45 +00:00
template < typename T , vtkm : : IdComponent N >
static inline VTKM_EXEC_CONT vtkm : : Vec < T , N > CopySign ( const vtkm : : Vec < T , N > & x ,
const vtkm : : Vec < T , N > & y )
2015-06-24 21:40:18 +00:00
{
2017-05-26 00:50:45 +00:00
vtkm : : Vec < T , N > result ;
2015-06-24 21:40:18 +00:00
for ( vtkm : : IdComponent index = 0 ; index < N ; index + + )
{
result [ index ] = vtkm : : CopySign ( x [ index ] , y [ index ] ) ;
}
return result ;
}
2024-02-11 22:05:57 +00:00
///@}
2015-06-23 23:35:55 +00:00
2024-02-11 22:05:57 +00:00
///@{
2018-08-27 16:00:46 +00:00
/// Decompose floating poing value
///
inline VTKM_EXEC_CONT vtkm : : Float32 Frexp ( vtkm : : Float32 x , vtkm : : Int32 * exponent )
{
2020-10-22 16:57:42 +00:00
// See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
# if defined(VTKM_CUDA) || defined(VTKM_HIP)
2018-08-27 16:00:46 +00:00
return VTKM_CUDA_MATH_FUNCTION_32 ( frexp ) ( x , exponent ) ;
# else
return std : : frexp ( x , exponent ) ;
# endif
}
inline VTKM_EXEC_CONT vtkm : : Float64 Frexp ( vtkm : : Float64 x , vtkm : : Int32 * exponent )
{
# ifdef VTKM_CUDA
return VTKM_CUDA_MATH_FUNCTION_64 ( frexp ) ( x , exponent ) ;
# else
return std : : frexp ( x , exponent ) ;
# endif
}
2024-02-11 22:05:57 +00:00
///@}
2018-08-27 16:00:46 +00:00
inline VTKM_EXEC_CONT vtkm : : Float32 Ldexp ( vtkm : : Float32 x , vtkm : : Int32 exponent )
{
# ifdef VTKM_CUDA
return VTKM_CUDA_MATH_FUNCTION_32 ( ldexp ) ( x , exponent ) ;
# else
return std : : ldexp ( x , exponent ) ;
# endif
}
inline VTKM_EXEC_CONT vtkm : : Float64 Ldexp ( vtkm : : Float64 x , vtkm : : Int32 exponent )
{
# ifdef VTKM_CUDA
return VTKM_CUDA_MATH_FUNCTION_64 ( ldexp ) ( x , exponent ) ;
# else
return std : : ldexp ( x , exponent ) ;
# endif
}
2021-07-20 21:35:23 +00:00
///////////////
// Float distance.
2020-06-24 14:52:59 +00:00
// See: https://randomascii.wordpress.com/2012/01/23/stupid-float-tricks-2/ for why this works.
2021-07-20 21:35:23 +00:00
namespace detail
{
inline VTKM_EXEC_CONT vtkm : : UInt64 FloatDistancePositive ( vtkm : : Float64 x , vtkm : : Float64 y )
{
VTKM_ASSERT ( x > = 0 ) ;
VTKM_ASSERT ( y > = 0 ) ;
// Note that:
// int64_t xi = *reinterpret_cast<int64_t*>(&x);
// int64_t yi = *reinterpret_cast<int64_t*>(&y);
// also works (usually), but generates warnings because it is technically undefined behavior
// according to the C++ standard.
// Good option to have if we get compile errors off memcpy or don't want to #include <cstring> though.
// At least on gcc, both versions generate the same assembly.
vtkm : : UInt64 xi ;
vtkm : : UInt64 yi ;
memcpy ( & xi , & x , sizeof ( vtkm : : UInt64 ) ) ;
memcpy ( & yi , & y , sizeof ( vtkm : : UInt64 ) ) ;
if ( yi > xi ) {
return yi - xi ;
}
return xi - yi ;
}
inline VTKM_EXEC_CONT vtkm : : UInt64 FloatDistancePositive ( vtkm : : Float32 x , vtkm : : Float32 y )
{
VTKM_ASSERT ( x > = 0 ) ;
VTKM_ASSERT ( y > = 0 ) ;
vtkm : : UInt32 xi_32 ;
vtkm : : UInt32 yi_32 ;
memcpy ( & xi_32 , & x , sizeof ( vtkm : : UInt32 ) ) ;
memcpy ( & yi_32 , & y , sizeof ( vtkm : : UInt32 ) ) ;
vtkm : : UInt64 xi = xi_32 ;
vtkm : : UInt64 yi = yi_32 ;
if ( yi > xi ) {
return yi - xi ;
}
return xi - yi ;
}
} // namespace detail
2024-02-11 22:05:57 +00:00
/// Computes the number of representables between two floating point numbers. This function
/// is non-negative and symmetric in its arguments. If either argument is non-finite, the
/// value returned is the maximum value allowed by 64-bit unsigned integers: 2^64-1.
2020-06-24 14:52:59 +00:00
inline VTKM_EXEC_CONT vtkm : : UInt64 FloatDistance ( vtkm : : Float64 x , vtkm : : Float64 y )
{
static_assert ( sizeof ( vtkm : : Float64 ) = = sizeof ( vtkm : : UInt64 ) , " vtkm::Float64 is incorrect size. " ) ;
static_assert ( std : : numeric_limits < vtkm : : Float64 > : : has_denorm = = std : : denorm_present , " FloatDistance presumes the floating-point type has subnormal numbers. " ) ;
if ( ! vtkm : : IsFinite ( x ) | | ! vtkm : : IsFinite ( y ) ) {
return 0xFFFFFFFFFFFFFFFFL ;
}
// Signed zero is the sworn enemy of this process.
if ( y = = 0 ) {
y = vtkm : : Abs ( y ) ;
}
if ( x = = 0 ) {
x = vtkm : : Abs ( x ) ;
}
if ( ( x < 0 & & y > = 0 ) | | ( x > = 0 & & y < 0 ) )
{
vtkm : : UInt64 dx , dy ;
if ( x < 0 ) {
2021-07-20 21:35:23 +00:00
dy = detail : : FloatDistancePositive ( 0.0 , y ) ;
dx = detail : : FloatDistancePositive ( 0.0 , - x ) ;
2020-06-24 14:52:59 +00:00
}
else {
2021-07-20 21:35:23 +00:00
dy = detail : : FloatDistancePositive ( 0.0 , - y ) ;
dx = detail : : FloatDistancePositive ( 0.0 , x ) ;
2020-06-24 14:52:59 +00:00
}
return dx + dy ;
}
if ( x < 0 & & y < 0 ) {
2021-07-20 21:35:23 +00:00
return detail : : FloatDistancePositive ( - x , - y ) ;
2020-06-24 14:52:59 +00:00
}
2021-07-20 21:35:23 +00:00
return detail : : FloatDistancePositive ( x , y ) ;
2020-06-24 14:52:59 +00:00
}
2024-02-11 22:05:57 +00:00
/// @copydoc FloatDistance
2020-06-24 14:52:59 +00:00
inline VTKM_EXEC_CONT vtkm : : UInt64 FloatDistance ( vtkm : : Float32 x , vtkm : : Float32 y )
{
static_assert ( sizeof ( vtkm : : Float32 ) = = sizeof ( vtkm : : Int32 ) , " vtkm::Float32 is incorrect size. " ) ;
static_assert ( std : : numeric_limits < vtkm : : Float32 > : : has_denorm = = std : : denorm_present , " FloatDistance presumes the floating-point type has subnormal numbers. " ) ;
if ( ! vtkm : : IsFinite ( x ) | | ! vtkm : : IsFinite ( y ) ) {
return 0xFFFFFFFFFFFFFFFFL ;
}
if ( y = = 0 ) {
y = vtkm : : Abs ( y ) ;
}
if ( x = = 0 ) {
x = vtkm : : Abs ( x ) ;
}
if ( ( x < 0 & & y > = 0 ) | | ( x > = 0 & & y < 0 ) )
{
vtkm : : UInt64 dx , dy ;
if ( x < 0 ) {
2021-07-20 21:35:23 +00:00
dy = detail : : FloatDistancePositive ( 0.0f , y ) ;
dx = detail : : FloatDistancePositive ( 0.0f , - x ) ;
2020-06-24 14:52:59 +00:00
}
else {
2021-07-20 21:35:23 +00:00
dy = detail : : FloatDistancePositive ( 0.0f , - y ) ;
dx = detail : : FloatDistancePositive ( 0.0f , x ) ;
2020-06-24 14:52:59 +00:00
}
return dx + dy ;
}
if ( x < 0 & & y < 0 ) {
2021-07-20 21:35:23 +00:00
return detail : : FloatDistancePositive ( - x , - y ) ;
2020-06-24 14:52:59 +00:00
}
2021-07-20 21:35:23 +00:00
return detail : : FloatDistancePositive ( x , y ) ;
2020-06-24 14:52:59 +00:00
}
2021-04-05 19:20:56 +00:00
// Computes ab - cd.
// See: https://pharr.org/matt/blog/2019/11/03/difference-of-floats.html
template < typename T >
inline VTKM_EXEC_CONT T DifferenceOfProducts ( T a , T b , T c , T d )
{
T cd = c * d ;
T err = std : : fma ( - c , d , cd ) ;
T dop = std : : fma ( a , b , - cd ) ;
return dop + err ;
}
2024-02-11 22:05:57 +00:00
/// @brief Solves ax² + bx + c = 0.
///
/// Only returns the real roots.
/// If there are real roots, the first element of the pair is less than or equal to the second.
/// If there are no real roots, both elements are NaNs.
/// If VTK-m is compiled with FMA support, each root is accurate to 3 ulps; otherwise
/// the discriminant is prone to catastrophic subtractive cancellation and no accuracy
/// guarantees can be provided.
2021-04-08 12:50:27 +00:00
template < typename T >
2021-04-08 18:45:59 +00:00
inline VTKM_EXEC_CONT vtkm : : Vec < T , 2 > QuadraticRoots ( T a , T b , T c )
2021-04-08 12:50:27 +00:00
{
if ( a = = 0 )
{
if ( b = = 0 )
{
if ( c = = 0 )
{
// A degenerate case. All real numbers are roots; hopefully this arbitrary decision interacts gracefully with use.
2021-04-08 18:45:59 +00:00
return vtkm : : Vec < T , 2 > ( 0 , 0 ) ;
2021-04-08 12:50:27 +00:00
}
else
{
2021-04-08 18:45:59 +00:00
return vtkm : : Vec < T , 2 > ( vtkm : : Nan < T > ( ) , vtkm : : Nan < T > ( ) ) ;
2021-04-08 12:50:27 +00:00
}
}
2021-04-08 18:45:59 +00:00
return vtkm : : Vec < T , 2 > ( - c / b , - c / b ) ;
2021-04-08 12:50:27 +00:00
}
T delta = DifferenceOfProducts ( b , b , 4 * a , c ) ;
if ( delta < 0 )
{
2021-04-08 18:45:59 +00:00
return vtkm : : Vec < T , 2 > ( vtkm : : Nan < T > ( ) , vtkm : : Nan < T > ( ) ) ;
2021-04-08 12:50:27 +00:00
}
T q = - ( b + vtkm : : CopySign ( vtkm : : Sqrt ( delta ) , b ) ) / 2 ;
T r0 = q / a ;
T r1 = c / q ;
if ( r0 < r1 )
{
2021-04-08 18:45:59 +00:00
return vtkm : : Vec < T , 2 > ( r0 , r1 ) ;
2021-04-08 12:50:27 +00:00
}
2021-04-08 18:45:59 +00:00
return vtkm : : Vec < T , 2 > ( r1 , r0 ) ;
2021-04-08 12:50:27 +00:00
}
2019-03-05 16:46:32 +00:00
/// Bitwise operations
///
/// Find the first set bit in @a word, and return its position (1-32). If no
/// bits are set, returns 0.
# ifdef VTKM_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __ffs is device only.
inline __device__
vtkm : : Int32 FindFirstSetBit ( vtkm : : UInt32 word )
{
// Output is [0,32], with ffs(0) == 0
return __ffs ( static_cast < int > ( word ) ) ;
}
# else // CUDA_DEVICE_PASS
inline VTKM_EXEC_CONT
vtkm : : Int32 FindFirstSetBit ( vtkm : : UInt32 word )
{
# if defined(VTKM_GCC) || defined(VTKM_CLANG)
// Output is [0,32], with ffs(0) == 0
return __builtin_ffs ( static_cast < int > ( word ) ) ;
# elif defined(VTKM_MSVC)
// Output is [0, 31], check return code to see if bits are set:
vtkm : : UInt32 firstSet ;
return _BitScanForward ( reinterpret_cast < DWORD * > ( & firstSet ) , word ) ! = 0
? static_cast < vtkm : : Int32 > ( firstSet + 1 ) : 0 ;
# elif defined(VTKM_ICC)
// Output is [0, 31], undefined if word is 0.
return word ! = 0 ? _bit_scan_forward ( word ) + 1 : 0 ;
# else
// Naive implementation:
if ( word = = 0 )
{
return 0 ;
}
vtkm : : Int32 bit = 1 ;
while ( ( word & 0x1 ) = = 0 )
{
word > > = 1 ;
+ + bit ;
}
return bit ;
# endif
}
# endif // CUDA_DEVICE_PASS
/// Find the first set bit in @a word, and return its position (1-64). If no
/// bits are set, returns 0.
# ifdef VTKM_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __ffsll is device only.
inline __device__
vtkm : : Int32 FindFirstSetBit ( vtkm : : UInt64 word )
{
// Output is [0,64], with ffs(0) == 0
return __ffsll ( static_cast < long long int > ( word ) ) ;
}
# else // CUDA_DEVICE_PASS
inline VTKM_EXEC_CONT
vtkm : : Int32 FindFirstSetBit ( vtkm : : UInt64 word )
{
2019-05-08 19:48:58 +00:00
# if defined(VTKM_GCC) || defined(VTKM_CLANG) || defined(VTKM_ICC)
2019-03-05 16:46:32 +00:00
// Output is [0,64], with ffs(0) == 0
return __builtin_ffsll ( static_cast < long long int > ( word ) ) ;
2019-05-08 19:48:58 +00:00
# elif defined(VTKM_MSVC)
2019-03-05 16:46:32 +00:00
// Output is [0, 63], check return code to see if bits are set:
vtkm : : UInt32 firstSet ;
return _BitScanForward64 ( reinterpret_cast < DWORD * > ( & firstSet ) , word ) ! = 0
? static_cast < vtkm : : Int32 > ( firstSet + 1 ) : 0 ;
# else
// Naive implementation:
if ( word = = 0 )
{
return 0 ;
}
vtkm : : Int32 bit = 1 ;
while ( ( word & 0x1 ) = = 0 )
{
word > > = 1 ;
+ + bit ;
}
return bit ;
# endif
}
# endif // CUDA_DEVICE_PASS
/// Count the total number of bits set in @a word.
# ifdef VTKM_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __popc is device only.
inline __device__
vtkm : : Int32 CountSetBits ( vtkm : : UInt32 word )
{
return __popc ( word ) ;
}
# else // CUDA_DEVICE_PASS
inline VTKM_EXEC_CONT
vtkm : : Int32 CountSetBits ( vtkm : : UInt32 word )
{
# if defined(VTKM_GCC) || defined(VTKM_CLANG)
return __builtin_popcount ( word ) ;
2022-11-28 00:46:13 +00:00
# elif defined(VTKM_MSVC) && !defined(_M_ARM64)
2019-03-05 16:46:32 +00:00
return static_cast < vtkm : : Int32 > ( __popcnt ( word ) ) ;
# elif defined(VTKM_ICC)
return _popcnt32 ( static_cast < int > ( word ) ) ;
# else
// Naive implementation:
vtkm : : Int32 bits = 0 ;
while ( word )
{
if ( word & 0x1 )
{
+ + bits ;
}
word > > = 1 ;
}
return bits ;
# endif
}
# endif // CUDA_DEVICE_PASS
/// Count the total number of bits set in @a word.
# ifdef VTKM_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __popcll is device only.
inline __device__
vtkm : : Int32 CountSetBits ( vtkm : : UInt64 word )
{
return __popcll ( word ) ;
}
# else // CUDA_DEVICE_PASS
inline VTKM_EXEC_CONT
vtkm : : Int32 CountSetBits ( vtkm : : UInt64 word )
{
# if defined(VTKM_GCC) || defined(VTKM_CLANG)
return __builtin_popcountll ( word ) ;
2022-11-28 00:46:13 +00:00
# elif defined(VTKM_MSVC) && !defined(_M_ARM64)
2019-03-05 16:46:32 +00:00
return static_cast < vtkm : : Int32 > ( __popcnt64 ( word ) ) ;
# elif defined(VTKM_ICC)
return _popcnt64 ( static_cast < vtkm : : Int64 > ( word ) ) ;
# else
// Naive implementation:
vtkm : : Int32 bits = 0 ;
while ( word )
{
if ( word & 0x1 )
{
+ + bits ;
}
word > > = 1 ;
}
return bits ;
# endif
}
# endif // CUDA_DEVICE_PASS
2015-05-27 04:32:10 +00:00
} // namespace vtkm
2018-01-09 20:58:59 +00:00
// clang-format on
2015-05-27 04:32:10 +00:00
# endif //vtk_m_Math_h