Math.h

Go to the documentation of this file.

//============================================================================
//  The contents of this file are covered by the Viskores license. See
//  LICENSE.txt for details.
//
//  By contributing to this file, all contributors agree to the Developer
//  Certificate of Origin Version 1.1 (DCO 1.1) as stated in DCO.txt.
//============================================================================
 
//============================================================================
//  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.
//============================================================================
// **** DO NOT EDIT THIS FILE!!! ****
// This file is automatically generated by Math.h.in
 
#ifndef viskores_Math_h
#define viskores_Math_h
 
#include <viskores/TypeTraits.h>
#include <viskores/Types.h>
#include <viskores/VecTraits.h>
 
#include <limits> // must be found with or without CUDA.
#ifndef VISKORES_CUDA
#include <cmath>
#include <cstring>
#include <limits.h>
#include <math.h>
#include <stdlib.h>
#endif // !VISKORES_CUDA
 
#if !defined(VISKORES_CUDA_DEVICE_PASS)
#define VISKORES_USE_STL
#include <algorithm>
#endif
 
#ifdef VISKORES_MSVC
#include <intrin.h>                // For bitwise intrinsics (__popcnt, etc)
#include <viskores/internal/Windows.h> // for types used by MSVC intrinsics.
#ifndef VISKORES_CUDA
#include <math.h>
#endif // VISKORES_CUDA
#endif // VISKORES_MSVC
 
#define VISKORES_CUDA_MATH_FUNCTION_32(func) func##f
#define VISKORES_CUDA_MATH_FUNCTION_64(func) func
 
namespace viskores
{
 
//-----------------------------------------------------------------------------
namespace detail
{
template <typename T>
struct FloatingPointReturnType
{
  using ctype = typename viskores::VecTraits<T>::ComponentType;
  using representable_as_float_type =
    std::integral_constant<bool,
                           ((sizeof(ctype) < sizeof(float)) ||
                            std::is_same<ctype, viskores::Float32>::value)>;
  using Type = typename std::
    conditional<representable_as_float_type::value, viskores::Float32, viskores::Float64>::type;
};
} // namespace detail
 
template <typename T = viskores::Float64>
static constexpr inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type TwoPi()
{
  using FT = typename detail::FloatingPointReturnType<T>::Type;
  return static_cast<FT>(6.28318530717958647692528676655900576);
}
 
template <typename T = viskores::Float64>
static constexpr inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi()
{
  using FT = typename detail::FloatingPointReturnType<T>::Type;
  return static_cast<FT>(3.14159265358979323846264338327950288);
}
 
template <typename T = viskores::Float64>
static constexpr inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_2()
{
  using FT = typename detail::FloatingPointReturnType<T>::Type;
  return static_cast<FT>(1.57079632679489661923132169163975144);
}
 
template <typename T = viskores::Float64>
static constexpr inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_3()
{
  using FT = typename detail::FloatingPointReturnType<T>::Type;
  return static_cast<FT>(1.04719755119659774615421446109316762);
}
 
template <typename T = viskores::Float64>
static constexpr inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_4()
{
  using FT = typename detail::FloatingPointReturnType<T>::Type;
  return static_cast<FT>(0.78539816339744830961566084581987572);
}
 
template <typename T = viskores::Float64>
static constexpr inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Pi_180()
{
  using FT = typename detail::FloatingPointReturnType<T>::Type;
  return static_cast<FT>(0.01745329251994329547437168059786927);
}
 
static constexpr inline VISKORES_EXEC_CONT viskores::Float32 TwoPif()
{
  return TwoPi<viskores::Float32>();
}
 
static constexpr inline VISKORES_EXEC_CONT viskores::Float32 Pif()
{
  return Pi<viskores::Float32>();
}
 
static constexpr inline VISKORES_EXEC_CONT viskores::Float32 Pi_2f()
{
  return Pi_2<viskores::Float32>();
}
 
static constexpr inline VISKORES_EXEC_CONT viskores::Float32 Pi_3f()
{
  return Pi_3<viskores::Float32>();
}
 
static constexpr inline VISKORES_EXEC_CONT viskores::Float32 Pi_4f()
{
  return Pi_4<viskores::Float32>();
}
 
static constexpr inline VISKORES_EXEC_CONT viskores::Float32 Pi_180f()
{
  return Pi_180<viskores::Float32>();
}
 
// clang-format off
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Sin(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(sin)(x);
#else
  return std::sin(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Sin(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(sin)(x);
#else
  return std::sin(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Sin(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Sin(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Sin(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Sin(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Sin(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Sin(x[0]), viskores::Sin(x[1]), viskores::Sin(x[2]), viskores::Sin(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Sin(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Sin(x[0]), viskores::Sin(x[1]), viskores::Sin(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Sin(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Sin(x[0]),
                                                                         viskores::Sin(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Cos(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(cos)(x);
#else
  return std::cos(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Cos(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(cos)(x);
#else
  return std::cos(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Cos(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Cos(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Cos(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Cos(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Cos(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Cos(x[0]), viskores::Cos(x[1]), viskores::Cos(x[2]), viskores::Cos(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Cos(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Cos(x[0]), viskores::Cos(x[1]), viskores::Cos(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Cos(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Cos(x[0]),
                                                                         viskores::Cos(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Tan(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(tan)(x);
#else
  return std::tan(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Tan(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(tan)(x);
#else
  return std::tan(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Tan(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Tan(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Tan(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Tan(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Tan(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Tan(x[0]), viskores::Tan(x[1]), viskores::Tan(x[2]), viskores::Tan(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Tan(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Tan(x[0]), viskores::Tan(x[1]), viskores::Tan(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Tan(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Tan(x[0]),
                                                                         viskores::Tan(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 ASin(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(asin)(x);
#else
  return std::asin(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 ASin(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(asin)(x);
#else
  return std::asin(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ASin(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::ASin(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> ASin(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::ASin(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> ASin(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::ASin(x[0]), viskores::ASin(x[1]), viskores::ASin(x[2]), viskores::ASin(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> ASin(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::ASin(x[0]), viskores::ASin(x[1]), viskores::ASin(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> ASin(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::ASin(x[0]),
                                                                         viskores::ASin(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 ACos(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(acos)(x);
#else
  return std::acos(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 ACos(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(acos)(x);
#else
  return std::acos(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ACos(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::ACos(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> ACos(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::ACos(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> ACos(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::ACos(x[0]), viskores::ACos(x[1]), viskores::ACos(x[2]), viskores::ACos(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> ACos(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::ACos(x[0]), viskores::ACos(x[1]), viskores::ACos(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> ACos(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::ACos(x[0]),
                                                                         viskores::ACos(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 ATan(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(atan)(x);
#else
  return std::atan(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 ATan(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(atan)(x);
#else
  return std::atan(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ATan(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::ATan(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> ATan(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::ATan(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> ATan(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::ATan(x[0]), viskores::ATan(x[1]), viskores::ATan(x[2]), viskores::ATan(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> ATan(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::ATan(x[0]), viskores::ATan(x[1]), viskores::ATan(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> ATan(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::ATan(x[0]),
                                                                         viskores::ATan(x[1]));
}
 
static inline VISKORES_EXEC_CONT viskores::Float32 ATan2(viskores::Float32 x, viskores::Float32 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(atan2)(x, y);
#else
  return std::atan2(x, y);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 ATan2(viskores::Float64 x, viskores::Float64 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(atan2)(x, y);
#else
  return std::atan2(x, y);
#endif
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 SinH(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(sinh)(x);
#else
  return std::sinh(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 SinH(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(sinh)(x);
#else
  return std::sinh(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type SinH(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::SinH(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> SinH(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::SinH(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> SinH(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::SinH(x[0]), viskores::SinH(x[1]), viskores::SinH(x[2]), viskores::SinH(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> SinH(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::SinH(x[0]), viskores::SinH(x[1]), viskores::SinH(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> SinH(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::SinH(x[0]),
                                                                         viskores::SinH(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 CosH(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(cosh)(x);
#else
  return std::cosh(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 CosH(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(cosh)(x);
#else
  return std::cosh(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type CosH(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::CosH(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> CosH(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::CosH(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> CosH(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::CosH(x[0]), viskores::CosH(x[1]), viskores::CosH(x[2]), viskores::CosH(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> CosH(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::CosH(x[0]), viskores::CosH(x[1]), viskores::CosH(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> CosH(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::CosH(x[0]),
                                                                         viskores::CosH(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 TanH(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(tanh)(x);
#else
  return std::tanh(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 TanH(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(tanh)(x);
#else
  return std::tanh(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type TanH(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::TanH(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> TanH(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::TanH(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> TanH(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::TanH(x[0]), viskores::TanH(x[1]), viskores::TanH(x[2]), viskores::TanH(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> TanH(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::TanH(x[0]), viskores::TanH(x[1]), viskores::TanH(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> TanH(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::TanH(x[0]),
                                                                         viskores::TanH(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 ASinH(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(asinh)(x);
#else
  return std::asinh(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 ASinH(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(asinh)(x);
#else
  return std::asinh(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ASinH(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::ASinH(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> ASinH(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::ASinH(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> ASinH(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::ASinH(x[0]), viskores::ASinH(x[1]), viskores::ASinH(x[2]), viskores::ASinH(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> ASinH(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::ASinH(x[0]), viskores::ASinH(x[1]), viskores::ASinH(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> ASinH(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::ASinH(x[0]),
                                                                         viskores::ASinH(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 ACosH(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(acosh)(x);
#else
  return std::acosh(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 ACosH(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(acosh)(x);
#else
  return std::acosh(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ACosH(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::ACosH(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> ACosH(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::ACosH(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> ACosH(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::ACosH(x[0]), viskores::ACosH(x[1]), viskores::ACosH(x[2]), viskores::ACosH(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> ACosH(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::ACosH(x[0]), viskores::ACosH(x[1]), viskores::ACosH(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> ACosH(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::ACosH(x[0]),
                                                                         viskores::ACosH(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 ATanH(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(atanh)(x);
#else
  return std::atanh(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 ATanH(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(atanh)(x);
#else
  return std::atanh(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ATanH(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::ATanH(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> ATanH(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::ATanH(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> ATanH(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::ATanH(x[0]), viskores::ATanH(x[1]), viskores::ATanH(x[2]), viskores::ATanH(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> ATanH(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::ATanH(x[0]), viskores::ATanH(x[1]), viskores::ATanH(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> ATanH(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::ATanH(x[0]),
                                                                         viskores::ATanH(x[1]));
}
 
//-----------------------------------------------------------------------------
static inline VISKORES_EXEC_CONT viskores::Float32 Pow(viskores::Float32 x, viskores::Float32 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(pow)(x, y);
#else
  return std::pow(x, y);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 Pow(viskores::Float64 x, viskores::Float64 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(pow)(x, y);
#else
  return std::pow(x, y);
#endif
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Sqrt(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(sqrt)(x);
#else
  return std::sqrt(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Sqrt(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(sqrt)(x);
#else
  return std::sqrt(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Sqrt(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Sqrt(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Sqrt(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Sqrt(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Sqrt(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Sqrt(x[0]), viskores::Sqrt(x[1]), viskores::Sqrt(x[2]), viskores::Sqrt(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Sqrt(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Sqrt(x[0]), viskores::Sqrt(x[1]), viskores::Sqrt(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Sqrt(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Sqrt(x[0]),
                                                                         viskores::Sqrt(x[1]));
}
 
#ifdef VISKORES_CUDA
static inline VISKORES_EXEC_CONT viskores::Float32 RSqrt(viskores::Float32 x)
{
  return rsqrtf(x);
}
static inline VISKORES_EXEC_CONT viskores::Float64 RSqrt(viskores::Float64 x)
{
  return rsqrt(x);
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Float64 RSqrt(T x)
{
  return rsqrt(static_cast<viskores::Float64>(x));
}
#else  // !VISKORES_CUDA
static inline VISKORES_EXEC_CONT viskores::Float32 RSqrt(viskores::Float32 x)
{
  return 1 / viskores::Sqrt(x);
}
static inline VISKORES_EXEC_CONT viskores::Float64 RSqrt(viskores::Float64 x)
{
  return 1 / viskores::Sqrt(x);
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Float64 RSqrt(T x)
{
  return 1 / static_cast<viskores::Float64>(x);
}
#endif // !VISKORES_CUDA
 
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> RSqrt(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::RSqrt(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> RSqrt(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::RSqrt(x[0]), viskores::RSqrt(x[1]), viskores::RSqrt(x[2]), viskores::RSqrt(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> RSqrt(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::RSqrt(x[0]), viskores::RSqrt(x[1]), viskores::RSqrt(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> RSqrt(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::RSqrt(x[0]),
                                                                         viskores::RSqrt(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Cbrt(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(cbrt)(x);
#else
  return std::cbrt(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Cbrt(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(cbrt)(x);
#else
  return std::cbrt(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Cbrt(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Cbrt(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Cbrt(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Cbrt(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Cbrt(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Cbrt(x[0]), viskores::Cbrt(x[1]), viskores::Cbrt(x[2]), viskores::Cbrt(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Cbrt(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Cbrt(x[0]), viskores::Cbrt(x[1]), viskores::Cbrt(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Cbrt(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Cbrt(x[0]),
                                                                         viskores::Cbrt(x[1]));
}
 
#ifdef VISKORES_CUDA
static inline VISKORES_EXEC_CONT viskores::Float32 RCbrt(viskores::Float32 x)
{
  return rcbrtf(x);
}
static inline VISKORES_EXEC_CONT viskores::Float64 RCbrt(viskores::Float64 x)
{
  return rcbrt(x);
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Float64 RCbrt(T x)
{
  return rcbrt(static_cast<viskores::Float64>(x));
}
#else  // !VISKORES_CUDA
static inline VISKORES_EXEC_CONT viskores::Float32 RCbrt(viskores::Float32 x)
{
  return 1 / viskores::Cbrt(x);
}
static inline VISKORES_EXEC_CONT viskores::Float64 RCbrt(viskores::Float64 x)
{
  return 1 / viskores::Cbrt(x);
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Float64 RCbrt(T x)
{
  return 1 / viskores::Cbrt(static_cast<viskores::Float64>(x));
}
#endif // !VISKORES_CUDA
 
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> RCbrt(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::RCbrt(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> RCbrt(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::RCbrt(x[0]), viskores::RCbrt(x[1]), viskores::RCbrt(x[2]), viskores::RCbrt(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> RCbrt(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::RCbrt(x[0]), viskores::RCbrt(x[1]), viskores::RCbrt(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> RCbrt(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::RCbrt(x[0]),
                                                                         viskores::RCbrt(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Exp(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(exp)(x);
#else
  return std::exp(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Exp(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(exp)(x);
#else
  return std::exp(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Exp(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Exp(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Exp(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Exp(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Exp(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Exp(x[0]), viskores::Exp(x[1]), viskores::Exp(x[2]), viskores::Exp(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Exp(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Exp(x[0]), viskores::Exp(x[1]), viskores::Exp(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Exp(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Exp(x[0]),
                                                                         viskores::Exp(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Exp2(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(exp2)(x);
#else
  return std::exp2(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Exp2(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(exp2)(x);
#else
  return std::exp2(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Exp2(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Exp2(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Exp2(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Exp2(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Exp2(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Exp2(x[0]), viskores::Exp2(x[1]), viskores::Exp2(x[2]), viskores::Exp2(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Exp2(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Exp2(x[0]), viskores::Exp2(x[1]), viskores::Exp2(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Exp2(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Exp2(x[0]),
                                                                         viskores::Exp2(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 ExpM1(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(expm1)(x);
#else
  return std::expm1(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 ExpM1(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(expm1)(x);
#else
  return std::expm1(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type ExpM1(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::ExpM1(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> ExpM1(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::ExpM1(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> ExpM1(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::ExpM1(x[0]), viskores::ExpM1(x[1]), viskores::ExpM1(x[2]), viskores::ExpM1(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> ExpM1(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::ExpM1(x[0]), viskores::ExpM1(x[1]), viskores::ExpM1(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> ExpM1(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::ExpM1(x[0]),
                                                                         viskores::ExpM1(x[1]));
}
 
#ifdef VISKORES_CUDA
static inline VISKORES_EXEC_CONT viskores::Float32 Exp10(viskores::Float32 x)
{
  return exp10f(x);
}
static inline VISKORES_EXEC_CONT viskores::Float64 Exp10(viskores::Float64 x)
{
  return exp10(x);
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Float64 Exp10(T x)
{
  return exp10(static_cast<viskores::Float64>(x));
}
#else  // !VISKORES_CUDA
static inline VISKORES_EXEC_CONT viskores::Float32 Exp10(viskores::Float32 x)
{
  return viskores::Pow(10, x);
}
static inline VISKORES_EXEC_CONT viskores::Float64 Exp10(viskores::Float64 x)
{
  return viskores::Pow(10, x);
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Float64 Exp10(T x)
{
  return viskores::Pow(10, static_cast<viskores::Float64>(x));
}
#endif // !VISKORES_CUDA
 
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Exp10(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Exp10(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Exp10(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Exp10(x[0]), viskores::Exp10(x[1]), viskores::Exp10(x[2]), viskores::Exp10(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Exp10(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Exp10(x[0]), viskores::Exp10(x[1]), viskores::Exp10(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Exp10(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Exp10(x[0]),
                                                                         viskores::Exp10(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Log(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(log)(x);
#else
  return std::log(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Log(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(log)(x);
#else
  return std::log(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Log(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Log(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Log(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Log(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Log(x[0]), viskores::Log(x[1]), viskores::Log(x[2]), viskores::Log(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Log(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Log(x[0]), viskores::Log(x[1]), viskores::Log(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Log(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Log(x[0]),
                                                                         viskores::Log(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Log2(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(log2)(x);
#else
  return std::log2(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Log2(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(log2)(x);
#else
  return std::log2(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log2(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Log2(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Log2(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Log2(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Log2(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Log2(x[0]), viskores::Log2(x[1]), viskores::Log2(x[2]), viskores::Log2(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Log2(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Log2(x[0]), viskores::Log2(x[1]), viskores::Log2(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Log2(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Log2(x[0]),
                                                                         viskores::Log2(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Log10(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(log10)(x);
#else
  return std::log10(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Log10(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(log10)(x);
#else
  return std::log10(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log10(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Log10(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Log10(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Log10(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Log10(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Log10(x[0]), viskores::Log10(x[1]), viskores::Log10(x[2]), viskores::Log10(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Log10(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Log10(x[0]), viskores::Log10(x[1]), viskores::Log10(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Log10(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Log10(x[0]),
                                                                         viskores::Log10(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Log1P(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(log1p)(x);
#else
  return std::log1p(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Log1P(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(log1p)(x);
#else
  return std::log1p(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Log1P(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Log1P(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Log1P(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Log1P(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Log1P(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Log1P(x[0]), viskores::Log1P(x[1]), viskores::Log1P(x[2]), viskores::Log1P(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Log1P(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Log1P(x[0]), viskores::Log1P(x[1]), viskores::Log1P(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Log1P(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Log1P(x[0]),
                                                                         viskores::Log1P(x[1]));
}
 
//-----------------------------------------------------------------------------
template <typename T>
static inline VISKORES_EXEC_CONT T Max(const T& x, const T& y);
#ifdef VISKORES_USE_STL
static inline VISKORES_EXEC_CONT viskores::Float32 Max(viskores::Float32 x, viskores::Float32 y)
{
  return (std::max)(x, y);
}
static inline VISKORES_EXEC_CONT viskores::Float64 Max(viskores::Float64 x, viskores::Float64 y)
{
  return (std::max)(x, y);
}
#else // !VISKORES_USE_STL
static inline VISKORES_EXEC_CONT viskores::Float32 Max(viskores::Float32 x, viskores::Float32 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(fmax)(x, y);
#else
  return std::fmax(x, y);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 Max(viskores::Float64 x, viskores::Float64 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(fmax)(x, y);
#else
  return std::fmax(x, y);
#endif
}
#endif // !VISKORES_USE_STL
 
template <typename T>
static inline VISKORES_EXEC_CONT T Min(const T& x, const T& y);
#if defined(VISKORES_USE_STL) && !defined(VISKORES_HIP)
static inline VISKORES_EXEC_CONT viskores::Float32 Min(viskores::Float32 x, viskores::Float32 y)
{
  return (std::min)(x, y);
}
static inline VISKORES_EXEC_CONT viskores::Float64 Min(viskores::Float64 x, viskores::Float64 y)
{
  return (std::min)(x, y);
}
#else // !VISKORES_USE_STL OR HIP
static inline VISKORES_EXEC_CONT viskores::Float32 Min(viskores::Float32 x, viskores::Float32 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(fmin)(x, y);
#else
  return std::fmin(x, y);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 Min(viskores::Float64 x, viskores::Float64 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(fmin)(x, y);
#else
  return std::fmin(x, y);
#endif
}
#endif // !VISKORES_USE_STL
 
namespace detail
{
 
template <typename T>
static inline VISKORES_EXEC_CONT T Max(T x, T y, viskores::TypeTraitsScalarTag)
{
  return (x < y) ? y : x;
}
 
template <typename T>
static inline VISKORES_EXEC_CONT T Max(const T& x, const T& y, viskores::TypeTraitsVectorTag)
{
  using Traits = viskores::VecTraits<T>;
  T result;
  for (viskores::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++)
  {
    Traits::SetComponent(
      result, index, viskores::Max(Traits::GetComponent(x, index), Traits::GetComponent(y, index)));
  }
  return result;
}
 
template <typename T>
static inline VISKORES_EXEC_CONT T Min(T x, T y, viskores::TypeTraitsScalarTag)
{
  return (x < y) ? x : y;
}
 
template <typename T>
static inline VISKORES_EXEC_CONT T Min(const T& x, const T& y, viskores::TypeTraitsVectorTag)
{
  using Traits = viskores::VecTraits<T>;
  T result;
  for (viskores::IdComponent index = 0; index < Traits::NUM_COMPONENTS; index++)
  {
    Traits::SetComponent(
      result, index, viskores::Min(Traits::GetComponent(x, index), Traits::GetComponent(y, index)));
  }
  return result;
}
 
} // namespace detail
 
template <typename T>
static inline VISKORES_EXEC_CONT T Max(const T& x, const T& y)
{
  return detail::Max(x, y, typename viskores::TypeTraits<T>::DimensionalityTag());
}
 
template <typename T>
static inline VISKORES_EXEC_CONT T Min(const T& x, const T& y)
{
  return detail::Min(x, y, typename viskores::TypeTraits<T>::DimensionalityTag());
}
 
inline VISKORES_EXEC_CONT viskores::Float32 Clamp(viskores::Float32 x, viskores::Float32 lo, viskores::Float32 hi)
{
  return x > lo ? (x < hi ? x : hi) : lo;
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Clamp(viskores::Float64 x, viskores::Float64 lo, viskores::Float64 hi)
{
  return x > lo ? (x < hi ? x : hi) : lo;
}
 
//-----------------------------------------------------------------------------
 
//#ifdef VISKORES_CUDA
#define VISKORES_USE_IEEE_NONFINITE
//#endif
 
#ifdef VISKORES_USE_IEEE_NONFINITE
 
namespace detail
{
 
union IEEE754Bits32 {
  viskores::UInt32 bits;
  viskores::Float32 scalar;
};
#define VISKORES_NAN_BITS_32 0x7FC00000U
#define VISKORES_INF_BITS_32 0x7F800000U
#define VISKORES_NEG_INF_BITS_32 0xFF800000U
#define VISKORES_EPSILON_32 1e-5f
 
union IEEE754Bits64 {
  viskores::UInt64 bits;
  viskores::Float64 scalar;
};
#define VISKORES_NAN_BITS_64 0x7FF8000000000000ULL
#define VISKORES_INF_BITS_64 0x7FF0000000000000ULL
#define VISKORES_NEG_INF_BITS_64 0xFFF0000000000000ULL
#define VISKORES_EPSILON_64 1e-9
 
template <typename T>
struct FloatLimits;
 
template <>
struct FloatLimits<viskores::Float32>
{
  using BitsType = viskores::detail::IEEE754Bits32;
 
  VISKORES_EXEC_CONT
  static viskores::Float32 Nan()
  {
    BitsType nan = { VISKORES_NAN_BITS_32 };
    return nan.scalar;
  }
 
  VISKORES_EXEC_CONT
  static viskores::Float32 Infinity()
  {
    BitsType inf = { VISKORES_INF_BITS_32 };
    return inf.scalar;
  }
 
  VISKORES_EXEC_CONT
  static viskores::Float32 NegativeInfinity()
  {
    BitsType neginf = { VISKORES_NEG_INF_BITS_32 };
    return neginf.scalar;
  }
 
  VISKORES_EXEC_CONT
  static viskores::Float32 Epsilon() { return VISKORES_EPSILON_32; }
};
 
template <int N>
struct FloatLimits<viskores::Vec<viskores::Float32, N>>
{
  using BitsType = viskores::detail::IEEE754Bits32;
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float32, N> Nan()
  {
    BitsType nan = { VISKORES_NAN_BITS_32 };
    return viskores::Vec<viskores::Float32, N>(nan.scalar);
  }
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float32, N> Infinity()
  {
    BitsType inf = { VISKORES_INF_BITS_32 };
    return viskores::Vec<viskores::Float32, N>(inf.scalar);
  }
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float32, N> NegativeInfinity()
  {
    BitsType neginf = { VISKORES_NEG_INF_BITS_32 };
    return viskores::Vec<viskores::Float32, N>(neginf.scalar);
  }
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float32, N> Epsilon()
  {
    return viskores::Vec<viskores::Float32, N>(VISKORES_EPSILON_32);
  }
};
 
template <>
struct FloatLimits<viskores::Float64>
{
  using BitsType = viskores::detail::IEEE754Bits64;
 
  VISKORES_EXEC_CONT
  static viskores::Float64 Nan()
  {
    BitsType nan = { VISKORES_NAN_BITS_64 };
    return nan.scalar;
  }
 
  VISKORES_EXEC_CONT
  static viskores::Float64 Infinity()
  {
    BitsType inf = { VISKORES_INF_BITS_64 };
    return inf.scalar;
  }
 
  VISKORES_EXEC_CONT
  static viskores::Float64 NegativeInfinity()
  {
    BitsType neginf = { VISKORES_NEG_INF_BITS_64 };
    return neginf.scalar;
  }
 
  VISKORES_EXEC_CONT
  static viskores::Float64 Epsilon() { return VISKORES_EPSILON_64; }
};
 
template <int N>
struct FloatLimits<viskores::Vec<viskores::Float64, N>>
{
  using BitsType = viskores::detail::IEEE754Bits64;
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float64, N> Nan()
  {
    BitsType nan = { VISKORES_NAN_BITS_64 };
    return viskores::Vec<viskores::Float64, N>(nan.scalar);
  }
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float64, N> Infinity()
  {
    BitsType inf = { VISKORES_INF_BITS_64 };
    return viskores::Vec<viskores::Float64, N>(inf.scalar);
  }
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float64, N> NegativeInfinity()
  {
    BitsType neginf = { VISKORES_NEG_INF_BITS_64 };
    return viskores::Vec<viskores::Float64, N>(neginf.scalar);
  }
 
  VISKORES_EXEC_CONT
  static viskores::Vec<viskores::Float64, N> Epsilon()
  {
    return viskores::Vec<viskores::Float64, N>(VISKORES_EPSILON_64);
  }
};
 
#undef VISKORES_NAN_BITS_32
#undef VISKORES_INF_BITS_32
#undef VISKORES_NEG_INF_BITS_32
#undef VISKORES_EPSILON_32
#undef VISKORES_NAN_BITS_64
#undef VISKORES_INF_BITS_64
#undef VISKORES_NEG_INF_BITS_64
#undef VISKORES_EPSILON_64
 
} // namespace detail
#endif //VISKORES_USE_IEEE_NONFINITE
 
#ifdef VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T Nan()
{
  return detail::FloatLimits<T>::Nan();
}
#else  // !VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T Nan()
{
  return std::numeric_limits<T>::quiet_NaN();
}
#endif // !VISKORES_USE_IEEE_NONFINITE
static inline VISKORES_EXEC_CONT viskores::Float32 Nan32()
{
  return viskores::Nan<viskores::Float32>();
}
static inline VISKORES_EXEC_CONT viskores::Float64 Nan64()
{
  return viskores::Nan<viskores::Float64>();
}
 
#ifdef VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T Infinity()
{
  return detail::FloatLimits<T>::Infinity();
}
#else  // !VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T Infinity()
{
  return std::numeric_limits<T>::infinity();
}
#endif // !VISKORES_USE_IEEE_NONFINITE
static inline VISKORES_EXEC_CONT viskores::Float32 Infinity32()
{
  return viskores::Infinity<viskores::Float32>();
}
static inline VISKORES_EXEC_CONT viskores::Float64 Infinity64()
{
  return viskores::Infinity<viskores::Float64>();
}
 
#ifdef VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T NegativeInfinity()
{
  return detail::FloatLimits<T>::NegativeInfinity();
}
#else  // !VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T NegativeInfinity()
{
  return -std::numeric_limits<T>::infinity();
}
#endif // !VISKORES_USE_IEEE_NONFINITE
static inline VISKORES_EXEC_CONT viskores::Float32 NegativeInfinity32()
{
  return viskores::NegativeInfinity<viskores::Float32>();
}
static inline VISKORES_EXEC_CONT viskores::Float64 NegativeInfinity64()
{
  return viskores::NegativeInfinity<viskores::Float64>();
}
 
#ifdef VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T Epsilon()
{
  return detail::FloatLimits<T>::Epsilon();
}
#else  // !VISKORES_USE_IEEE_NONFINITE
template <typename T>
static inline VISKORES_EXEC_CONT T Epsilon()
{
  return std::numeric_limits<T>::epsilon();
}
#endif // !VISKORES_USE_IEEE_NONFINITE
static inline VISKORES_EXEC_CONT viskores::Float32 Epsilon32()
{
  return viskores::Epsilon<viskores::Float32>();
}
static inline VISKORES_EXEC_CONT viskores::Float64 Epsilon64()
{
  return viskores::Epsilon<viskores::Float64>();
}
 
 
//-----------------------------------------------------------------------------
template <typename T>
static inline VISKORES_EXEC_CONT bool IsNan(T x)
{
#ifndef VISKORES_CUDA
  using std::isnan;
#endif
  return (isnan(x) != 0);
}
 
template <typename T>
static inline VISKORES_EXEC_CONT bool IsInf(T x)
{
#ifndef VISKORES_CUDA
  using std::isinf;
#endif
  return (isinf(x) != 0);
}
 
template <typename T>
static inline VISKORES_EXEC_CONT bool IsFinite(T x)
{
#ifndef VISKORES_CUDA
  using std::isfinite;
#endif
  return (isfinite(x) != 0);
}
 
//-----------------------------------------------------------------------------
 
inline VISKORES_EXEC_CONT viskores::Float32 Ceil(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(ceil)(x);
#else
  return std::ceil(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Ceil(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(ceil)(x);
#else
  return std::ceil(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Ceil(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Ceil(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Ceil(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Ceil(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Ceil(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Ceil(x[0]), viskores::Ceil(x[1]), viskores::Ceil(x[2]), viskores::Ceil(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Ceil(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Ceil(x[0]), viskores::Ceil(x[1]), viskores::Ceil(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Ceil(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Ceil(x[0]),
                                                                         viskores::Ceil(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Floor(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(floor)(x);
#else
  return std::floor(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Floor(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(floor)(x);
#else
  return std::floor(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Floor(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Floor(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Floor(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Floor(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Floor(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Floor(x[0]), viskores::Floor(x[1]), viskores::Floor(x[2]), viskores::Floor(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Floor(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Floor(x[0]), viskores::Floor(x[1]), viskores::Floor(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Floor(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Floor(x[0]),
                                                                         viskores::Floor(x[1]));
}
 
 
inline VISKORES_EXEC_CONT viskores::Float32 Round(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(round)(x);
#else
  return std::round(x);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Round(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(round)(x);
#else
  return std::round(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Round(const T& x)
{
  using RT = typename detail::FloatingPointReturnType<T>::Type;
  return viskores::Round(static_cast<RT>(x));
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> Round(
  const viskores::Vec<T, N>& x)
{
  viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Round(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4> Round(
  const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 4>(
    viskores::Round(x[0]), viskores::Round(x[1]), viskores::Round(x[2]), viskores::Round(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3> Round(
  const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 3>(
    viskores::Round(x[0]), viskores::Round(x[1]), viskores::Round(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2> Round(
  const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<typename detail::FloatingPointReturnType<T>::Type, 2>(viskores::Round(x[0]),
                                                                         viskores::Round(x[1]));
}
 
//-----------------------------------------------------------------------------
static inline VISKORES_EXEC_CONT viskores::Float32 FMod(viskores::Float32 x, viskores::Float32 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(fmod)(x, y);
#else
  return std::fmod(x, y);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 FMod(viskores::Float64 x, viskores::Float64 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(fmod)(x, y);
#else
  return std::fmod(x, y);
#endif
}
 
#ifdef VISKORES_MSVC
template <typename T>
static inline VISKORES_EXEC_CONT T Remainder(T numerator, T denominator)
{
  T quotient = viskores::Round(numerator / denominator);
  return numerator - quotient * denominator;
}
#else // !VISKORES_MSVC
static inline VISKORES_EXEC_CONT viskores::Float32 Remainder(viskores::Float32 x, viskores::Float32 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(remainder)(x, y);
#else
  return std::remainder(x, y);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 Remainder(viskores::Float64 x, viskores::Float64 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(remainder)(x, y);
#else
  return std::remainder(x, y);
#endif
}
#endif // !VISKORES_MSVC
 
template <typename QType>
static inline VISKORES_EXEC_CONT viskores::Float32 RemainderQuotient(viskores::Float32 numerator,
                                                             viskores::Float32 denominator,
                                                             QType& quotient)
{
  int iQuotient;
  // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
#if defined(VISKORES_CUDA) || defined(VISKORES_HIP)
  const viskores::Float32 result =
    VISKORES_CUDA_MATH_FUNCTION_32(remquo)(numerator, denominator, &iQuotient);
#else
  const viskores::Float32 result = std::remquo(numerator, denominator, &iQuotient);
#endif
  quotient = static_cast<QType>(iQuotient);
  return result;
}
template <typename QType>
static inline VISKORES_EXEC_CONT viskores::Float64 RemainderQuotient(viskores::Float64 numerator,
                                                             viskores::Float64 denominator,
                                                             QType& quotient)
{
  int iQuotient;
#ifdef VISKORES_CUDA
  const viskores::Float64 result =
    VISKORES_CUDA_MATH_FUNCTION_64(remquo)(numerator, denominator, &iQuotient);
#else
  const viskores::Float64 result = std::remquo(numerator, denominator, &iQuotient);
#endif
  quotient = static_cast<QType>(iQuotient);
  return result;
}
 
static inline VISKORES_EXEC_CONT viskores::Float32 ModF(viskores::Float32 x, viskores::Float32& integral)
{
  // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
#if defined(VISKORES_CUDA) || defined(VISKORES_HIP)
  return VISKORES_CUDA_MATH_FUNCTION_32(modf)(x, &integral);
#else
  return std::modf(x, &integral);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 ModF(viskores::Float64 x, viskores::Float64& integral)
{
#if defined(VISKORES_CUDA)
  return VISKORES_CUDA_MATH_FUNCTION_64(modf)(x, &integral);
#else
  return std::modf(x, &integral);
#endif
}
 
//-----------------------------------------------------------------------------
static inline VISKORES_EXEC_CONT viskores::Int32 Abs(viskores::Int32 x)
{
  return abs(x);
}
static inline VISKORES_EXEC_CONT viskores::Int64 Abs(viskores::Int64 x)
{
#if VISKORES_SIZE_LONG == 8
  return labs(x);
#elif VISKORES_SIZE_LONG_LONG == 8
  return llabs(x);
#else
#error Unknown size of Int64.
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float32 Abs(viskores::Float32 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(fabs)(x);
#else
  return std::fabs(x);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 Abs(viskores::Float64 x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(fabs)(x);
#else
  return std::fabs(x);
#endif
}
template <typename T>
static inline VISKORES_EXEC_CONT typename detail::FloatingPointReturnType<T>::Type Abs(T x)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(fabs)(static_cast<viskores::Float64>(x));
#else
  return std::fabs(static_cast<viskores::Float64>(x));
#endif
}
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<T, N> Abs(const viskores::Vec<T, N>& x)
{
  viskores::Vec<T, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::Abs(x[index]);
  }
  return result;
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<T, 4> Abs(const viskores::Vec<T, 4>& x)
{
  return viskores::Vec<T, 4>(viskores::Abs(x[0]), viskores::Abs(x[1]), viskores::Abs(x[2]), viskores::Abs(x[3]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<T, 3> Abs(const viskores::Vec<T, 3>& x)
{
  return viskores::Vec<T, 3>(viskores::Abs(x[0]), viskores::Abs(x[1]), viskores::Abs(x[2]));
}
template <typename T>
static inline VISKORES_EXEC_CONT viskores::Vec<T, 2> Abs(const viskores::Vec<T, 2>& x)
{
  return viskores::Vec<T, 2>(viskores::Abs(x[0]), viskores::Abs(x[1]));
}
 
static inline VISKORES_EXEC_CONT viskores::Int32 SignBit(viskores::Float32 x)
{
#ifndef VISKORES_CUDA
  using std::signbit;
#endif
  return static_cast<viskores::Int32>(signbit(x));
}
static inline VISKORES_EXEC_CONT viskores::Int32 SignBit(viskores::Float64 x)
{
#ifndef VISKORES_CUDA
  using std::signbit;
#endif
  return static_cast<viskores::Int32>(signbit(x));
}
 
static inline VISKORES_EXEC_CONT bool IsNegative(viskores::Float32 x)
{
  return (viskores::SignBit(x) != 0);
}
static inline VISKORES_EXEC_CONT bool IsNegative(viskores::Float64 x)
{
  return (viskores::SignBit(x) != 0);
}
 
static inline VISKORES_EXEC_CONT viskores::Float32 CopySign(viskores::Float32 x, viskores::Float32 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(copysign)(x, y);
#else
  return std::copysign(x, y);
#endif
}
static inline VISKORES_EXEC_CONT viskores::Float64 CopySign(viskores::Float64 x, viskores::Float64 y)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(copysign)(x, y);
#else
  return std::copysign(x, y);
#endif
}
 
template <typename T, viskores::IdComponent N>
static inline VISKORES_EXEC_CONT viskores::Vec<T, N> CopySign(const viskores::Vec<T, N>& x,
                                                      const viskores::Vec<T, N>& y)
{
  viskores::Vec<T, N> result;
  for (viskores::IdComponent index = 0; index < N; index++)
  {
    result[index] = viskores::CopySign(x[index], y[index]);
  }
  return result;
}
 
inline VISKORES_EXEC_CONT viskores::Float32 Frexp(viskores::Float32 x, viskores::Int32 *exponent)
{
  // See: https://github.com/ROCm-Developer-Tools/HIP/issues/2169
#if defined(VISKORES_CUDA) || defined(VISKORES_HIP)
  return VISKORES_CUDA_MATH_FUNCTION_32(frexp)(x, exponent);
#else
  return std::frexp(x, exponent);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Frexp(viskores::Float64 x, viskores::Int32 *exponent)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(frexp)(x, exponent);
#else
  return std::frexp(x, exponent);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float32 Ldexp(viskores::Float32 x, viskores::Int32 exponent)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_32(ldexp)(x, exponent);
#else
  return std::ldexp(x, exponent);
#endif
}
 
inline VISKORES_EXEC_CONT viskores::Float64 Ldexp(viskores::Float64 x, viskores::Int32 exponent)
{
#ifdef VISKORES_CUDA
  return VISKORES_CUDA_MATH_FUNCTION_64(ldexp)(x, exponent);
#else
  return std::ldexp(x, exponent);
#endif
}
 
// Float distance.
// See: https://randomascii.wordpress.com/2012/01/23/stupid-float-tricks-2/ for why this works.
 
namespace detail
{
 
inline VISKORES_EXEC_CONT viskores::UInt64 FloatDistancePositive(viskores::Float64 x, viskores::Float64 y)
{
  VISKORES_ASSERT(x >= 0);
  VISKORES_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.
  viskores::UInt64 xi;
  viskores::UInt64 yi;
  memcpy(&xi, &x, sizeof(viskores::UInt64));
  memcpy(&yi, &y, sizeof(viskores::UInt64));
  if (yi > xi) {
    return yi - xi;
  }
  return xi - yi;
}
 
 
inline VISKORES_EXEC_CONT viskores::UInt64 FloatDistancePositive(viskores::Float32 x, viskores::Float32 y)
{
  VISKORES_ASSERT(x >= 0);
  VISKORES_ASSERT(y >= 0);
 
  viskores::UInt32 xi_32;
  viskores::UInt32 yi_32;
  memcpy(&xi_32, &x, sizeof(viskores::UInt32));
  memcpy(&yi_32, &y, sizeof(viskores::UInt32));
  viskores::UInt64 xi = xi_32;
  viskores::UInt64 yi = yi_32;
  if (yi > xi) {
    return yi - xi;
  }
  return xi - yi;
}
 
} // namespace detail
 
inline VISKORES_EXEC_CONT viskores::UInt64 FloatDistance(viskores::Float64 x, viskores::Float64 y)
{
  static_assert(sizeof(viskores::Float64) == sizeof(viskores::UInt64), "viskores::Float64 is incorrect size.");
  static_assert(std::numeric_limits<viskores::Float64>::has_denorm == std::denorm_present, "FloatDistance presumes the floating-point type has subnormal numbers.");
 
  if (!viskores::IsFinite(x) || !viskores::IsFinite(y)) {
    return 0xFFFFFFFFFFFFFFFFL;
  }
 
  // Signed zero is the sworn enemy of this process.
  if (y == 0) {
    y = viskores::Abs(y);
  }
  if (x == 0) {
    x = viskores::Abs(x);
  }
 
  if ( (x < 0 && y >= 0) || (x >= 0 && y < 0) )
  {
    viskores::UInt64 dx, dy;
    if (x < 0) {
      dy = detail::FloatDistancePositive(0.0, y);
      dx = detail::FloatDistancePositive(0.0, -x);
    }
    else {
      dy = detail::FloatDistancePositive(0.0, -y);
      dx = detail::FloatDistancePositive(0.0, x);
    }
 
    return dx + dy;
  }
 
  if (x < 0 && y < 0) {
    return detail::FloatDistancePositive(-x, -y);
  }
 
  return detail::FloatDistancePositive(x, y);
}
 
inline VISKORES_EXEC_CONT viskores::UInt64 FloatDistance(viskores::Float32 x, viskores::Float32 y)
{
  static_assert(sizeof(viskores::Float32) == sizeof(viskores::Int32), "viskores::Float32 is incorrect size.");
  static_assert(std::numeric_limits<viskores::Float32>::has_denorm == std::denorm_present, "FloatDistance presumes the floating-point type has subnormal numbers.");
 
  if (!viskores::IsFinite(x) || !viskores::IsFinite(y)) {
    return 0xFFFFFFFFFFFFFFFFL;
  }
 
  if (y == 0) {
    y = viskores::Abs(y);
  }
  if (x == 0) {
    x = viskores::Abs(x);
  }
 
  if ( (x < 0 && y >= 0) || (x >= 0 && y < 0) )
  {
    viskores::UInt64 dx, dy;
    if (x < 0) {
      dy = detail::FloatDistancePositive(0.0f, y);
      dx = detail::FloatDistancePositive(0.0f, -x);
    }
    else {
      dy = detail::FloatDistancePositive(0.0f, -y);
      dx = detail::FloatDistancePositive(0.0f, x);
    }
    return dx + dy;
  }
 
  if (x < 0 && y < 0) {
    return detail::FloatDistancePositive(-x, -y);
  }
 
  return detail::FloatDistancePositive(x, y);
}
 
// Computes ab - cd.
// See: https://pharr.org/matt/blog/2019/11/03/difference-of-floats.html
template<typename T>
inline VISKORES_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;
}
 
template<typename T>
inline VISKORES_EXEC_CONT viskores::Vec<T, 2> QuadraticRoots(T a, T b, T c)
{
  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.
        return viskores::Vec<T,2>(0,0);
      }
      else
      {
        return viskores::Vec<T,2>(viskores::Nan<T>(), viskores::Nan<T>());
      }
    }
    return viskores::Vec<T,2>(-c/b, -c/b);
  }
  T delta = DifferenceOfProducts(b, b, 4*a, c);
  if (delta < 0)
  {
    return viskores::Vec<T,2>(viskores::Nan<T>(), viskores::Nan<T>());
  }
 
  T q = -(b + viskores::CopySign(viskores::Sqrt(delta), b)) / 2;
  T r0 = q / a;
  T r1 = c / q;
  if (r0 < r1)
  {
    return viskores::Vec<T,2>(r0, r1);
  }
  return viskores::Vec<T,2>(r1, r0);
}
 
 
#ifdef VISKORES_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __ffs is device only.
inline __device__
viskores::Int32 FindFirstSetBit(viskores::UInt32 word)
{
  // Output is [0,32], with ffs(0) == 0
  return __ffs(static_cast<int>(word));
}
#else // CUDA_DEVICE_PASS
inline VISKORES_EXEC_CONT
viskores::Int32 FindFirstSetBit(viskores::UInt32 word)
{
#  if defined(VISKORES_GCC) || defined(VISKORES_CLANG)
 
  // Output is [0,32], with ffs(0) == 0
  return __builtin_ffs(static_cast<int>(word));
 
#  elif defined(VISKORES_MSVC)
 
  // Output is [0, 31], check return code to see if bits are set:
  viskores::UInt32 firstSet;
  return _BitScanForward(reinterpret_cast<DWORD*>(&firstSet), word) != 0
         ? static_cast<viskores::Int32>(firstSet + 1) : 0;
 
#  elif defined(VISKORES_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;
  }
 
  viskores::Int32 bit = 1;
  while ((word & 0x1) == 0)
  {
    word >>= 1;
    ++bit;
  }
  return bit;
 
#  endif
}
#endif // CUDA_DEVICE_PASS
 
#ifdef VISKORES_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __ffsll is device only.
inline __device__
viskores::Int32 FindFirstSetBit(viskores::UInt64 word)
{
 
  // Output is [0,64], with ffs(0) == 0
  return __ffsll(static_cast<long long int>(word));
}
#else // CUDA_DEVICE_PASS
inline VISKORES_EXEC_CONT
viskores::Int32 FindFirstSetBit(viskores::UInt64 word)
{
#  if defined(VISKORES_GCC) || defined(VISKORES_CLANG) || defined(VISKORES_ICC)
 
  // Output is [0,64], with ffs(0) == 0
  return __builtin_ffsll(static_cast<long long int>(word));
 
#  elif defined(VISKORES_MSVC)
 
  // Output is [0, 63], check return code to see if bits are set:
  viskores::UInt32 firstSet;
  return _BitScanForward64(reinterpret_cast<DWORD*>(&firstSet), word) != 0
         ? static_cast<viskores::Int32>(firstSet + 1) : 0;
 
#  else
 
  // Naive implementation:
  if (word == 0)
  {
    return 0;
  }
 
  viskores::Int32 bit = 1;
  while ((word & 0x1) == 0)
  {
    word >>= 1;
    ++bit;
  }
  return bit;
 
#  endif
}
#endif // CUDA_DEVICE_PASS
 
#ifdef VISKORES_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __popc is device only.
inline __device__
viskores::Int32 CountSetBits(viskores::UInt32 word)
{
  return __popc(word);
}
#else // CUDA_DEVICE_PASS
inline VISKORES_EXEC_CONT
viskores::Int32 CountSetBits(viskores::UInt32 word)
{
#  if defined(VISKORES_GCC) || defined(VISKORES_CLANG)
 
  return __builtin_popcount(word);
 
#  elif defined(VISKORES_MSVC) && !defined(_M_ARM64)
 
  return static_cast<viskores::Int32>(__popcnt(word));
 
#  elif defined(VISKORES_ICC)
 
  return _popcnt32(static_cast<int>(word));
 
#  else
 
  // Naive implementation:
  viskores::Int32 bits = 0;
  while (word)
  {
    if (word & 0x1)
    {
      ++bits;
    }
    word >>= 1;
  }
  return bits;
 
#  endif
}
#endif // CUDA_DEVICE_PASS
 
#ifdef VISKORES_CUDA_DEVICE_PASS
// Need to explicitly mark this as __device__ since __popcll is device only.
inline __device__
viskores::Int32 CountSetBits(viskores::UInt64 word)
{
  return __popcll(word);
}
#else // CUDA_DEVICE_PASS
inline VISKORES_EXEC_CONT
viskores::Int32 CountSetBits(viskores::UInt64 word)
{
#  if defined(VISKORES_GCC) || defined(VISKORES_CLANG)
 
  return __builtin_popcountll(word);
 
#  elif defined(VISKORES_MSVC) && !defined(_M_ARM64)
 
  return static_cast<viskores::Int32>(__popcnt64(word));
 
#  elif defined(VISKORES_ICC)
 
  return _popcnt64(static_cast<viskores::Int64>(word));
 
#  else
 
  // Naive implementation:
  viskores::Int32 bits = 0;
  while (word)
  {
    if (word & 0x1)
    {
      ++bits;
    }
    word >>= 1;
  }
  return bits;
 
#  endif
}
#endif // CUDA_DEVICE_PASS
 
} // namespace viskores
// clang-format on
 
#endif //viskores_Math_h