From 1836891e1e02ccd1ae7afaa5e4157b93286cd214 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Fri, 8 Nov 2024 17:50:13 +0100 Subject: [PATCH] Move definition of floating point trait functions to their own file --- libcudacxx/include/cuda/std/__cccl/builtin.h | 48 +- libcudacxx/include/cuda/std/__cmath/nvbf16.h | 40 -- libcudacxx/include/cuda/std/__cmath/nvfp16.h | 40 -- libcudacxx/include/cuda/std/__cmath/traits.h | 469 ++++++++++++++++++ .../cuda/std/detail/libcxx/include/cmath | 113 +---- .../cuda/std/detail/libcxx/include/complex | 201 ++++---- .../std/numerics/c.math/fp_traits.pass.cpp | 410 +++++++++++++++ 7 files changed, 1033 insertions(+), 288 deletions(-) create mode 100644 libcudacxx/include/cuda/std/__cmath/traits.h create mode 100644 libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp diff --git a/libcudacxx/include/cuda/std/__cccl/builtin.h b/libcudacxx/include/cuda/std/__cccl/builtin.h index b047e106e66..762c49c485f 100644 --- a/libcudacxx/include/cuda/std/__cccl/builtin.h +++ b/libcudacxx/include/cuda/std/__cccl/builtin.h @@ -210,6 +210,33 @@ # undef _CCCL_BUILTIN_IS_CONSTANT_EVALUATED #endif // _CCCL_STD_VER < 2014 && _CCCL_CUDA_COMPILER_NVCC +#if _CCCL_CHECK_BUILTIN(builtin_isfinite) || defined(_CCCL_COMPILER_GCC) || defined(_CCCL_COMPILER_NVRTC) +# define _CCCL_BUILTIN_ISFINITE(...) __builtin_isfinite(__VA_ARGS__) +#endif // _CCCL_CHECK_BUILTIN(isfinite) + +// Below 11.7 nvcc treats the builtin as a host only function +#if _CCCL_CUDACC_BELOW(11, 7) +# undef _CCCL_BUILTIN_ISFINITE +#endif // _CCCL_CUDACC_BELOW(11, 7) + +#if _CCCL_CHECK_BUILTIN(builtin_isinf) || defined(_CCCL_COMPILER_GCC) +# define _CCCL_BUILTIN_ISINF(...) __builtin_isinf(__VA_ARGS__) +#endif // _CCCL_CHECK_BUILTIN(isinf) + +// Below 11.7 nvcc treats the builtin as a host only function +#if _CCCL_CUDACC_BELOW(11, 7) +# undef _CCCL_BUILTIN_ISINF +#endif // _CCCL_CUDACC_BELOW(11, 7) + +#if _CCCL_CHECK_BUILTIN(builtin_isnan) || defined(_CCCL_COMPILER_GCC) +# define _CCCL_BUILTIN_ISNAN(...) __builtin_isnan(__VA_ARGS__) +#endif // _CCCL_CHECK_BUILTIN(isnan) + +// Below 11.7 nvcc treats the builtin as a host only function +#if _CCCL_CUDACC_BELOW(11, 7) +# undef _CCCL_BUILTIN_ISNAN +#endif // _CCCL_CUDACC_BELOW(11, 7) + #if (_CCCL_CHECK_BUILTIN(builtin_launder) || (defined(_CCCL_COMPILER_GCC) && _CCCL_GCC_VERSION >= 70000)) # define _CCCL_BUILTIN_LAUNDER(...) __builtin_launder(__VA_ARGS__) #endif // _CCCL_CHECK_BUILTIN(builtin_launder) && gcc >= 7 @@ -321,6 +348,15 @@ # define _CCCL_BUILTIN_OPERATOR_NEW(...) __builtin_operator_new(__VA_ARGS__) #endif // _CCCL_CHECK_BUILTIN(__builtin_operator_new) && _CCCL_CHECK_BUILTIN(__builtin_operator_delete) +#if _CCCL_CHECK_BUILTIN(builtin_signbit) || defined(_CCCL_COMPILER_GCC) +# define _CCCL_BUILTIN_SIGNBIT(...) __builtin_signbit(__VA_ARGS__) +#endif // _CCCL_CHECK_BUILTIN(builtin_signbit) + +// Below 11.7 nvcc treats the builtin as a host only function +#if _CCCL_CUDACC_BELOW(11, 7) +# undef _CCCL_BUILTIN_SIGNBIT +#endif // _CCCL_CUDACC_BELOW(11, 7) + #if _CCCL_HAS_BUILTIN(__decay) && defined(_CCCL_CUDA_COMPILER_CLANG) # define _CCCL_BUILTIN_DECAY(...) __decay(__VA_ARGS__) #endif // _CCCL_HAS_BUILTIN(__decay) && clang-cuda @@ -593,18 +629,6 @@ # define _CCCL_BUILTIN_IS_VOLATILE(...) __is_volatile(__VA_ARGS__) #endif // _CCCL_HAS_BUILTIN(__is_volatile) -#if _CCCL_CHECK_BUILTIN(isfinite) -# define _CCCL_BUILTIN_ISFINITE(...) __builtin_isfinite(__VA_ARGS__) -#endif // _CCCL_CHECK_BUILTIN(isfinite) - -#if _CCCL_CHECK_BUILTIN(isinf) -# define _CCCL_BUILTIN_ISINF(...) __builtin_isinf(__VA_ARGS__) -#endif // _CCCL_CHECK_BUILTIN(isinf) - -#if _CCCL_CHECK_BUILTIN(isnan) -# define _CCCL_BUILTIN_ISNAN(...) __builtin_isnan(__VA_ARGS__) -#endif // _CCCL_CHECK_BUILTIN(isnan) - #if _CCCL_CHECK_BUILTIN(make_integer_seq) || (defined(_CCCL_COMPILER_MSVC) && _CCCL_MSVC_VERSION >= 1923) # define _CCCL_BUILTIN_MAKE_INTEGER_SEQ(...) __make_integer_seq<__VA_ARGS__> #endif // _CCCL_CHECK_BUILTIN(make_integer_seq) diff --git a/libcudacxx/include/cuda/std/__cmath/nvbf16.h b/libcudacxx/include/cuda/std/__cmath/nvbf16.h index 8c7385ad087..8f116968f8b 100644 --- a/libcudacxx/include/cuda/std/__cmath/nvbf16.h +++ b/libcudacxx/include/cuda/std/__cmath/nvbf16.h @@ -76,46 +76,6 @@ _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 sqrt(__nv_bfloat16 __x) } // floating point helper -_LIBCUDACXX_HIDE_FROM_ABI bool signbit(__nv_bfloat16 __v) -{ - return ::signbit(::__bfloat162float(__v)); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isnan(__nv_bfloat16 __x) noexcept -{ - return ::__hisnan(__x); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool isnan(__nv_bfloat16 __v) -{ - return _CUDA_VSTD::__constexpr_isnan(__v); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isinf(__nv_bfloat16 __x) noexcept -{ -# if _CCCL_STD_VER >= 2020 && _CCCL_CUDACC_BELOW(12, 3) - // this is a workaround for nvbug 4362808 - return !::__hisnan(__x) && ::__hisnan(__x - __x); -# else // ^^^ C++20 && below 12.3 ^^^ / vvv C++17 or 12.3+ vvv - return ::__hisinf(__x) != 0; -# endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3) -} - -_LIBCUDACXX_HIDE_FROM_ABI bool isinf(__nv_bfloat16 __v) -{ - return _CUDA_VSTD::__constexpr_isinf(__v); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isfinite(__nv_bfloat16 __x) noexcept -{ - return !_CUDA_VSTD::__constexpr_isnan(__x) && !_CUDA_VSTD::__constexpr_isinf(__x); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool isfinite(__nv_bfloat16 __v) -{ - return _CUDA_VSTD::__constexpr_isfinite(__v); -} - _LIBCUDACXX_HIDE_FROM_ABI __nv_bfloat16 __constexpr_copysign(__nv_bfloat16 __x, __nv_bfloat16 __y) noexcept { return __float2bfloat16(::copysignf(__bfloat162float(__x), __bfloat162float(__y))); diff --git a/libcudacxx/include/cuda/std/__cmath/nvfp16.h b/libcudacxx/include/cuda/std/__cmath/nvfp16.h index f7317e7cf48..dbcaebbb4ef 100644 --- a/libcudacxx/include/cuda/std/__cmath/nvfp16.h +++ b/libcudacxx/include/cuda/std/__cmath/nvfp16.h @@ -141,46 +141,6 @@ _LIBCUDACXX_HIDE_FROM_ABI __half sqrt(__half __x) } // floating point helper -_LIBCUDACXX_HIDE_FROM_ABI bool signbit(__half __v) -{ - return ::signbit(::__half2float(__v)); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isnan(__half __x) noexcept -{ - return ::__hisnan(__x); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool isnan(__half __v) -{ - return _CUDA_VSTD::__constexpr_isnan(__v); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isinf(__half __x) noexcept -{ -# if _CCCL_STD_VER >= 2020 && _CCCL_CUDACC_BELOW(12, 3) - // this is a workaround for nvbug 4362808 - return !::__hisnan(__x) && ::__hisnan(__x - __x); -# else // ^^^ C++20 && below 12.3 ^^^ / vvv C++17 or 12.3+ vvv - return ::__hisinf(__x) != 0; -# endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_BELOW(12, 3) -} - -_LIBCUDACXX_HIDE_FROM_ABI bool isinf(__half __v) -{ - return _CUDA_VSTD::__constexpr_isinf(__v); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool __constexpr_isfinite(__half __x) noexcept -{ - return !_CUDA_VSTD::__constexpr_isnan(__x) && !_CUDA_VSTD::__constexpr_isinf(__x); -} - -_LIBCUDACXX_HIDE_FROM_ABI bool isfinite(__half __v) -{ - return _CUDA_VSTD::__constexpr_isfinite(__v); -} - _LIBCUDACXX_HIDE_FROM_ABI __half __constexpr_copysign(__half __x, __half __y) noexcept { return __float2half(::copysignf(__half2float(__x), __half2float(__y))); diff --git a/libcudacxx/include/cuda/std/__cmath/traits.h b/libcudacxx/include/cuda/std/__cmath/traits.h new file mode 100644 index 00000000000..789a09e68af --- /dev/null +++ b/libcudacxx/include/cuda/std/__cmath/traits.h @@ -0,0 +1,469 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCUDACXX___CMATH_TRAITS_H +#define _LIBCUDACXX___CMATH_TRAITS_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include + +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_STD + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool signbit(float __x) noexcept +{ +#if defined(_CCCL_BUILTIN_SIGNBIT) + return __builtin_signbit(__x); +#else // ^^^ _CCCL_BUILTIN_SIGNBIT ^^^ / vvv !_CCCL_BUILTIN_SIGNBIT vvv + return ::signbit(__x); +#endif // !_CCCL_BUILTIN_SIGNBIT +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool signbit(double __x) noexcept +{ +#if defined(_CCCL_BUILTIN_SIGNBIT) + return __builtin_signbit(__x); +#else // ^^^ _CCCL_BUILTIN_SIGNBIT ^^^ / vvv !_CCCL_BUILTIN_SIGNBIT vvv + return ::signbit(__x); +#endif // !_CCCL_BUILTIN_SIGNBIT +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool signbit(long double __x) noexcept +{ +# if defined(_CCCL_BUILTIN_SIGNBIT) + return __builtin_signbit(__x); +# else // ^^^ _CCCL_BUILTIN_SIGNBIT ^^^ / vvv !_CCCL_BUILTIN_SIGNBIT vvv + return ::signbit(__x); +# endif // !_CCCL_BUILTIN_SIGNBIT +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool signbit(__half __x) noexcept +{ + return _CUDA_VSTD::signbit(__half2float(__x)); +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool signbit(__nv_bfloat16 __x) noexcept +{ + return _CUDA_VSTD::signbit(__bfloat162float(__x)); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool signbit(_A1 __x) noexcept +{ + return __x < 0; +} + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool signbit(_A1) noexcept +{ + return false; +} + +// isfinite + +#if defined(_CCCL_BUILTIN_ISFINITE) || (defined(_CCCL_BUILTIN_ISINF) && defined(_CCCL_BUILTIN_ISNAN)) +# define _CCCL_CONSTEXPR_ISFINITE constexpr +#else // ^^^ _CCCL_BUILTIN_ISFINITE ^^^ / vvv !_CCCL_BUILTIN_ISFINITE vvv +# define _CCCL_CONSTEXPR_ISFINITE +#endif // !_CCCL_BUILTIN_ISFINITE + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isfinite(_A1) noexcept +{ + return true; +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISFINITE bool isfinite(float __x) noexcept +{ +#if defined(_CCCL_BUILTIN_ISFINITE) + return _CCCL_BUILTIN_ISFINITE(__x); +#elif _CCCL_CUDACC_BELOW(11, 8) + return !::__isinf(__x) && !::__isnan(__x); +#else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isfinite(__x); +#endif // !_CCCL_CUDACC_BELOW(11, 8) +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISFINITE bool isfinite(double __x) noexcept +{ +#if defined(_CCCL_BUILTIN_ISFINITE) + return _CCCL_BUILTIN_ISFINITE(__x); +#elif _CCCL_CUDACC_BELOW(11, 8) + return !::__isinf(__x) && !::__isnan(__x); +#else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isfinite(__x); +#endif // !_CCCL_CUDACC_BELOW(11, 8) +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISFINITE bool isfinite(long double __x) noexcept +{ +# if defined(_CCCL_BUILTIN_ISFINITE) + return _CCCL_BUILTIN_ISFINITE(__x); +# elif _CCCL_CUDACC_BELOW(11, 8) + return !::__isinf(__x) && !::__isnan(__x); +# else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isfinite(__x); +# endif // !_CCCL_CUDACC_BELOW(11, 8) +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isfinite(__half __x) noexcept +{ + return !::__hisnan(__x) && !::__hisinf(__x); +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isfinite(__nv_bfloat16 __x) noexcept +{ + return !::__hisnan(__x) && !::__hisinf(__x); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +// isinf + +#if defined(_CCCL_BUILTIN_ISINF) +# define _CCCL_CONSTEXPR_ISINF constexpr +#else // ^^^ _CCCL_BUILTIN_ISINF ^^^ / vvv !_CCCL_BUILTIN_ISINF vvv +# define _CCCL_CONSTEXPR_ISINF +#endif // !_CCCL_BUILTIN_ISINF + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isinf(_A1) noexcept +{ + return false; +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISINF bool isinf(float __x) noexcept +{ +#if defined(_CCCL_BUILTIN_ISINF) + return _CCCL_BUILTIN_ISINF(__x); +#elif _CCCL_CUDACC_BELOW(11, 8) + return ::__isinf(__x); +#else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isinf(__x); +#endif // !_CCCL_CUDACC_BELOW(11, 8) +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISINF bool isinf(double __x) noexcept +{ +#if defined(_CCCL_BUILTIN_ISINF) + return _CCCL_BUILTIN_ISINF(__x); +#elif _CCCL_CUDACC_BELOW(11, 8) + return ::__isinf(__x); +#else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isinf(__x); +#endif // !_CCCL_CUDACC_BELOW(11, 8) +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISINF bool isinf(long double __x) noexcept +{ +# if defined(_CCCL_BUILTIN_ISINF) + return _CCCL_BUILTIN_ISINF(__x); +# elif _CCCL_CUDACC_BELOW(11, 8) + return ::__isinf(__x); +# else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isinf(__x); +# endif // !_CCCL_CUDACC_BELOW(11, 8) +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isinf(__half __x) noexcept +{ +# if _CCCL_STD_VER >= 2020 && _CCCL_CUDACC_BELOW(12, 3) + // this is a workaround for nvbug 4362808 + return !::__hisnan(__x) && ::__hisnan(__x - __x); +# else // ^^^ C++20 && below 12.3 ^^^ / vvv C++17 or 12.3+ vvv + return ::__hisinf(__x) != 0; +# endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_VER < 1203000 +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isinf(__nv_bfloat16 __x) noexcept +{ +# if _CCCL_STD_VER >= 2020 && _CCCL_CUDACC_BELOW(12, 3) + // this is a workaround for nvbug 4362808 + return !::__hisnan(__x) && ::__hisnan(__x - __x); +# else // ^^^ C++20 && below 12.3 ^^^ / vvv C++17 or 12.3+ vvv + return ::__hisinf(__x) != 0; +# endif // _CCCL_STD_VER <= 2017 || _CCCL_CUDACC_VER < 1203000 +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +// isnan + +#if defined(_CCCL_BUILTIN_ISNAN) +# define _CCCL_CONSTEXPR_ISNAN constexpr +#else // ^^^ _CCCL_BUILTIN_ISNAN ^^^ / vvv !_CCCL_BUILTIN_ISNAN vvv +# define _CCCL_CONSTEXPR_ISNAN +#endif // !_CCCL_BUILTIN_ISNAN + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr bool isnan(_A1) noexcept +{ + return false; +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISNAN bool isnan(float __x) noexcept +{ +#if defined(_CCCL_BUILTIN_ISNAN) + return _CCCL_BUILTIN_ISNAN(__x); +#elif _CCCL_CUDACC_BELOW(11, 8) + return ::__isnan(__x); +#else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isnan(__x); +#endif // !_CCCL_CUDACC_BELOW(11, 8) +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISNAN bool isnan(double __x) noexcept +{ +#if defined(_CCCL_BUILTIN_ISNAN) + return _CCCL_BUILTIN_ISNAN(__x); +#elif _CCCL_CUDACC_BELOW(11, 8) + return ::__isnan(__x); +#else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isnan(__x); +#endif // !_CCCL_CUDACC_BELOW(11, 8) +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_ISNAN bool isnan(long double __x) noexcept +{ +# if defined(_CCCL_BUILTIN_ISNAN) + return _CCCL_BUILTIN_ISNAN(__x); +# elif _CCCL_CUDACC_BELOW(11, 8) + return ::__isnan(__x); +# else // ^^^ _CCCL_CUDACC_BELOW(11, 8) ^^^ / vvv !_CCCL_CUDACC_BELOW(11, 8) vvv + return ::isnan(__x); +# endif // !_CCCL_CUDACC_BELOW(11, 8) +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnan(__half __x) noexcept +{ + return ::__hisnan(__x); +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnan(__nv_bfloat16 __x) noexcept +{ + return ::__hisnan(__x); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +// isnormal + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnormal(_A1 __x) noexcept +{ + return __x != 0; +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnormal(float __x) noexcept +{ + return __x != 0 && _CUDA_VSTD::isfinite(__x); +} + +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnormal(double __x) noexcept +{ + return __x != 0 && _CUDA_VSTD::isfinite(__x); +} + +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnormal(long double __x) noexcept +{ + return __x != 0 && _CUDA_VSTD::isfinite(__x); +} +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + +#if defined(_LIBCUDACXX_HAS_NVFP16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnormal(__half __x) noexcept +{ + return _CUDA_VSTD::isnormal(__half2float(__x)); +} +#endif // _LIBCUDACXX_HAS_NVFP16 + +#if defined(_LIBCUDACXX_HAS_NVBF16) +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isnormal(__nv_bfloat16 __x) noexcept +{ + return _CUDA_VSTD::isnormal(__bfloat162float(__x)); +} +#endif // _LIBCUDACXX_HAS_NVBF16 + +// isgreater + +template +struct __is_extended_arithmetic +{ + static constexpr bool value = _CCCL_TRAIT(is_arithmetic, _Tp) || _CCCL_TRAIT(__is_extended_floating_point, _Tp); +}; + +#if !defined(_CCCL_NO_VARIABLE_TEMPLATES) +template +_CCCL_INLINE_VAR constexpr bool __is_extended_arithmetic_v = + is_arithmetic_v<_Tp> || __is_extended_floating_point_v<_Tp>; +#endif // !_CCCL_NO_INLINE_VARIABLES + +template = 0> +_CCCL_NODISCARD _CCCL_DEVICE _CCCL_HIDE_FROM_ABI bool __device_isgreater(_A1 __x, _A1 __y) noexcept +{ + if (_CUDA_VSTD::isnan(__x) || _CUDA_VSTD::isnan(__y)) + { + return false; + } + return __x > __y; +} + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isgreater(_A1 __x, _A2 __y) noexcept +{ + using type = __promote_t<_A1, _A2>; + NV_IF_ELSE_TARGET(NV_IS_HOST, + (return ::isgreater((type) __x, (type) __y);), + (return _CUDA_VSTD::__device_isgreater((type) __x, (type) __y);)) +} + +// isgreaterequal + +template = 0> +_CCCL_NODISCARD _CCCL_DEVICE _CCCL_HIDE_FROM_ABI bool __device_isgreaterequal(_A1 __x, _A1 __y) noexcept +{ + if (_CUDA_VSTD::isnan(__x) || _CUDA_VSTD::isnan(__y)) + { + return false; + } + return __x >= __y; +} + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isgreaterequal(_A1 __x, _A2 __y) noexcept +{ + using type = __promote_t<_A1, _A2>; + NV_IF_ELSE_TARGET(NV_IS_HOST, + (return ::isgreaterequal((type) __x, (type) __y);), + (return _CUDA_VSTD::__device_isgreaterequal((type) __x, (type) __y);)) +} + +// isless + +template = 0> +_CCCL_NODISCARD _CCCL_DEVICE _CCCL_HIDE_FROM_ABI bool __device_isless(_A1 __x, _A1 __y) noexcept +{ + if (_CUDA_VSTD::isnan(__x) || _CUDA_VSTD::isnan(__y)) + { + return false; + } + return __x < __y; +} + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isless(_A1 __x, _A2 __y) noexcept +{ + using type = __promote_t<_A1, _A2>; + NV_IF_ELSE_TARGET(NV_IS_HOST, + (return ::isless((type) __x, (type) __y);), + (return _CUDA_VSTD::__device_isless((type) __x, (type) __y);)) +} + +// islessequal + +template = 0> +_CCCL_NODISCARD _CCCL_DEVICE _CCCL_HIDE_FROM_ABI bool __device_islessequal(_A1 __x, _A1 __y) noexcept +{ + if (_CUDA_VSTD::isnan(__x) || _CUDA_VSTD::isnan(__y)) + { + return false; + } + return __x <= __y; +} + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool islessequal(_A1 __x, _A2 __y) noexcept +{ + using type = __promote_t<_A1, _A2>; + NV_IF_ELSE_TARGET(NV_IS_HOST, + (return ::islessequal((type) __x, (type) __y);), + (return _CUDA_VSTD::__device_islessequal((type) __x, (type) __y);)) +} + +// islessgreater + +template = 0> +_CCCL_NODISCARD _CCCL_DEVICE _CCCL_HIDE_FROM_ABI bool __device_islessgreater(_A1 __x, _A1 __y) noexcept +{ + if (_CUDA_VSTD::isnan(__x) || _CUDA_VSTD::isnan(__y)) + { + return false; + } + return __x < __y || __x > __y; +} + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool islessgreater(_A1 __x, _A2 __y) noexcept +{ + using type = __promote_t<_A1, _A2>; + NV_IF_ELSE_TARGET(NV_IS_HOST, + (return ::islessgreater((type) __x, (type) __y);), + (return _CUDA_VSTD::__device_islessgreater((type) __x, (type) __y);)) +} + +// isunordered + +template = 0> +_CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI bool isunordered(_A1 __x, _A2 __y) noexcept +{ + using type = __promote_t<_A1, _A2>; + return _CUDA_VSTD::isnan((type) __x) || _CUDA_VSTD::isnan((type) __y); +} + +_LIBCUDACXX_END_NAMESPACE_STD + +#endif // _LIBCUDACXX___CMATH_TRAITS_H diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/cmath b/libcudacxx/include/cuda/std/detail/libcxx/include/cmath index d35fd9e7945..c228c614ca1 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/cmath +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/cmath @@ -321,6 +321,7 @@ long double truncl(long double x); #include #include #include +#include #include #include #include @@ -343,11 +344,6 @@ _CCCL_PUSH_MACROS _LIBCUDACXX_BEGIN_NAMESPACE_STD -using ::isfinite; -using ::isinf; -using ::isnan; -using ::signbit; - using ::acos; using ::acosf; using ::asin; @@ -399,14 +395,6 @@ using ::abs; #ifndef _CCCL_COMPILER_NVRTC using ::fpclassify; -using ::isgreater; -using ::isgreaterequal; -using ::isless; -using ::islessequal; -using ::islessgreater; -using ::isnormal; - -using ::isunordered; using ::double_t; using ::float_t; @@ -595,69 +583,6 @@ hypot(_A1 __lcpp_x, _A2 __lcpp_y, _A3 __lcpp_z) noexcept # define _CCCL_CONSTEXPR_CXX14_COMPLEX #endif // _LIBCUDACXX_HAS_NO_CONSTEXPR_COMPLEX_OPERATIONS -template -_LIBCUDACXX_HIDE_FROM_ABI constexpr enable_if_t::value, bool> -__constexpr_isnan(_A1 __lcpp_x) noexcept -{ -#if _CCCL_CUDACC_BELOW(11, 8) - return __isnan(__lcpp_x); -#elif defined(_CCCL_BUILTIN_ISNAN) - // nvcc at times has issues determining the type of __lcpp_x - return _CCCL_BUILTIN_ISNAN(static_cast(__lcpp_x)); -#else // ^^^ _CCCL_BUILTIN_ISNAN ^^^ / vvv !_CCCL_BUILTIN_ISNAN vvv - return ::isnan(__lcpp_x); -#endif // !_CCCL_BUILTIN_ISNAN -} - -template -_LIBCUDACXX_HIDE_FROM_ABI constexpr enable_if_t::value, bool> -__constexpr_isnan(_A1 __lcpp_x) noexcept -{ - return ::isnan(__lcpp_x); -} - -template -_LIBCUDACXX_HIDE_FROM_ABI constexpr enable_if_t::value, bool> -__constexpr_isinf(_A1 __lcpp_x) noexcept -{ -#if _CCCL_CUDACC_BELOW(11, 8) - return __isinf(__lcpp_x); -#elif defined(_CCCL_BUILTIN_ISINF) - // nvcc at times has issues determining the type of __lcpp_x - return __builtin_isinf(static_cast(__lcpp_x)); -#else // ^^^ _CCCL_BUILTIN_ISINF ^^^ / vvv !_CCCL_BUILTIN_ISINF vvv - return ::isinf(__lcpp_x); -#endif // !_CCCL_BUILTIN_ISINF -} - -template -_LIBCUDACXX_HIDE_FROM_ABI constexpr enable_if_t::value, bool> -__constexpr_isinf(_A1 __lcpp_x) noexcept -{ - return ::isinf(__lcpp_x); -} - -template -_LIBCUDACXX_HIDE_FROM_ABI constexpr enable_if_t::value, bool> -__constexpr_isfinite(_A1 __lcpp_x) noexcept -{ -#if _CCCL_CUDACC_BELOW(11, 8) - return !__isinf(__lcpp_x) && !__isnan(__lcpp_x); -#elif defined(_CCCL_BUILTIN_ISFINITE) - // nvcc at times has issues determining the type of __lcpp_x - return __builtin_isfinite(static_cast(__lcpp_x)); -#else // ^^^ _CCCL_BUILTIN_ISFINITE ^^^ / vvv !_CCCL_BUILTIN_ISFINITE vvv - return ::isfinite(__lcpp_x); -#endif // !_CCCL_BUILTIN_ISFINITE -} - -template -_LIBCUDACXX_HIDE_FROM_ABI constexpr enable_if_t::value, bool> -__constexpr_isfinite(_A1 __lcpp_x) noexcept -{ - return isfinite(__lcpp_x); -} - #if defined(_CCCL_COMPILER_MSVC) || defined(_CCCL_COMPILER_NVRTC) template _LIBCUDACXX_HIDE_FROM_ABI _A1 __constexpr_copysign(_A1 __x, _A1 __y) noexcept @@ -675,11 +600,13 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 double __constexpr_copysign(doub return __builtin_copysign(__x, __y); } +# if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 long double __constexpr_copysign(long double __x, long double __y) noexcept { return __builtin_copysignl(__x, __y); } +# endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE template _LIBCUDACXX_HIDE_FROM_ABI @@ -709,10 +636,12 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 double __constexpr_fabs(double _ return __builtin_fabs(__x); } +# if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 long double __constexpr_fabs(long double __x) noexcept { return __builtin_fabsl(__x); } +# endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE template ::value, int> = 0> _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 double __constexpr_fabs(_Tp __x) noexcept @@ -737,11 +666,11 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX float __constexpr_fmax(f if (_CCCL_BUILTIN_IS_CONSTANT_EVALUATED()) # endif // defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) { - if (_CUDA_VSTD::__constexpr_isnan(__x)) + if (_CUDA_VSTD::isnan(__x)) { return __y; } - if (_CUDA_VSTD::__constexpr_isnan(__y)) + if (_CUDA_VSTD::isnan(__y)) { return __x; } @@ -760,11 +689,11 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX double __constexpr_fmax( if (_CCCL_BUILTIN_IS_CONSTANT_EVALUATED()) # endif // defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) { - if (_CUDA_VSTD::__constexpr_isnan(__x)) + if (_CUDA_VSTD::isnan(__x)) { return __y; } - if (_CUDA_VSTD::__constexpr_isnan(__y)) + if (_CUDA_VSTD::isnan(__y)) { return __x; } @@ -774,29 +703,31 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX double __constexpr_fmax( return __builtin_fmax(__x, __y); } +# if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX long double __constexpr_fmax(long double __x, long double __y) noexcept { -# if defined(_CCCL_BUILTIN_IS_CONSTANT_EVALUATED) && !defined(_LIBCUDACXX_HAS_NO_CONSTEXPR_COMPLEX_OPERATIONS) -# if defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) +# if defined(_CCCL_BUILTIN_IS_CONSTANT_EVALUATED) && !defined(_LIBCUDACXX_HAS_NO_CONSTEXPR_COMPLEX_OPERATIONS) +# if defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) if (false) -# else // defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) +# else // defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) if (_CCCL_BUILTIN_IS_CONSTANT_EVALUATED()) -# endif // defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) +# endif // defined(_CCCL_COMPILER_ICC) && _NV_ISEMPTY(_CCCL_CONSTEXPR_CXX14_COMPLEX) { - if (_CUDA_VSTD::__constexpr_isnan(__x)) + if (_CUDA_VSTD::isnan(__x)) { return __y; } - if (_CUDA_VSTD::__constexpr_isnan(__y)) + if (_CUDA_VSTD::isnan(__y)) { return __x; } return __x < __y ? __y : __x; } -# endif // defined(_CCCL_BUILTIN_IS_CONSTANT_EVALUATED) +# endif // defined(_CCCL_BUILTIN_IS_CONSTANT_EVALUATED) return __builtin_fmax(__x, __y); } +# endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE template ::value && is_arithmetic<_Up>::value, int> = 0> _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX __promote_t<_Tp, _Up> __constexpr_fmax(_Tp __x, _Up __y) noexcept @@ -825,12 +756,12 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX _Tp __constexpr_logb(_Tp return -numeric_limits<_Tp>::infinity(); } - if (_CUDA_VSTD::__constexpr_isinf(__x)) + if (_CUDA_VSTD::isinf(__x)) { return numeric_limits<_Tp>::infinity(); } - if (_CUDA_VSTD::__constexpr_isnan(__x)) + if (_CUDA_VSTD::isnan(__x)) { return numeric_limits<_Tp>::quiet_NaN(); } @@ -887,7 +818,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX _Tp __constexpr_scalbn(_ return __x; } - if (_CUDA_VSTD::__constexpr_isinf(__x)) + if (_CUDA_VSTD::isinf(__x)) { return __x; } @@ -897,7 +828,7 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14_COMPLEX _Tp __constexpr_scalbn(_ return __x; } - if (_CUDA_VSTD::__constexpr_isnan(__x)) + if (_CUDA_VSTD::isnan(__x)) { return numeric_limits<_Tp>::quiet_NaN(); } diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/complex b/libcudacxx/include/cuda/std/detail/libcxx/include/complex index 7a24039ba2b..47110ee2365 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/complex +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/complex @@ -507,16 +507,14 @@ operator*(const complex<_Tp>& __z, const complex<_Tp>& __w) { bool __z_zero = __a == _Tp(0) && __b == _Tp(0); bool __w_zero = __c == _Tp(0) && __d == _Tp(0); - bool __z_inf = _CUDA_VSTD::__constexpr_isinf(__a) || _CUDA_VSTD::__constexpr_isinf(__b); - bool __w_inf = _CUDA_VSTD::__constexpr_isinf(__c) || _CUDA_VSTD::__constexpr_isinf(__d); + bool __z_inf = _CUDA_VSTD::isinf(__a) || _CUDA_VSTD::isinf(__b); + bool __w_inf = _CUDA_VSTD::isinf(__c) || _CUDA_VSTD::isinf(__d); bool __z_nan = !__z_inf - && ((_CUDA_VSTD::__constexpr_isnan(__a) && _CUDA_VSTD::__constexpr_isnan(__b)) - || (_CUDA_VSTD::__constexpr_isnan(__a) && __b == _Tp(0)) - || (__a == _Tp(0) && _CUDA_VSTD::__constexpr_isnan(__b))); + && ((_CUDA_VSTD::isnan(__a) && _CUDA_VSTD::isnan(__b)) || (_CUDA_VSTD::isnan(__a) && __b == _Tp(0)) + || (__a == _Tp(0) && _CUDA_VSTD::isnan(__b))); bool __w_nan = !__w_inf - && ((_CUDA_VSTD::__constexpr_isnan(__c) && _CUDA_VSTD::__constexpr_isnan(__d)) - || (_CUDA_VSTD::__constexpr_isnan(__c) && __d == _Tp(0)) - || (__c == _Tp(0) && _CUDA_VSTD::__constexpr_isnan(__d))); + && ((_CUDA_VSTD::isnan(__c) && _CUDA_VSTD::isnan(__d)) || (_CUDA_VSTD::isnan(__c) && __d == _Tp(0)) + || (__c == _Tp(0) && _CUDA_VSTD::isnan(__d))); if (__z_nan || __w_nan) { return complex<_Tp>(_Tp(numeric_limits<_Tp>::quiet_NaN()), _Tp(0)); @@ -529,10 +527,8 @@ operator*(const complex<_Tp>& __z, const complex<_Tp>& __w) } return complex<_Tp>(_Tp(numeric_limits<_Tp>::infinity()), _Tp(numeric_limits<_Tp>::infinity())); } - bool __z_nonzero_nan = - !__z_inf && !__z_nan && (_CUDA_VSTD::__constexpr_isnan(__a) || _CUDA_VSTD::__constexpr_isnan(__b)); - bool __w_nonzero_nan = - !__w_inf && !__w_nan && (_CUDA_VSTD::__constexpr_isnan(__c) || _CUDA_VSTD::__constexpr_isnan(__d)); + bool __z_nonzero_nan = !__z_inf && !__z_nan && (_CUDA_VSTD::isnan(__a) || _CUDA_VSTD::isnan(__b)); + bool __w_nonzero_nan = !__w_inf && !__w_nan && (_CUDA_VSTD::isnan(__c) || _CUDA_VSTD::isnan(__d)); if (__z_nonzero_nan || __w_nonzero_nan) { return complex<_Tp>(_Tp(numeric_limits<_Tp>::quiet_NaN()), _Tp(0)); @@ -545,54 +541,54 @@ operator*(const complex<_Tp>& __z, const complex<_Tp>& __w) _Tp __x = __partials.__ac - __partials.__bd; _Tp __y = __partials.__ad + __partials.__bc; #ifndef LIBCUDACXX_ENABLE_SIMPLIFIED_COMPLEX_MULTIPLICATION - if (_CUDA_VSTD::__constexpr_isnan(__x) && _CUDA_VSTD::__constexpr_isnan(__y)) + if (_CUDA_VSTD::isnan(__x) && _CUDA_VSTD::isnan(__y)) { bool __recalc = false; - if (_CUDA_VSTD::__constexpr_isinf(__a) || _CUDA_VSTD::__constexpr_isinf(__b)) + if (_CUDA_VSTD::isinf(__a) || _CUDA_VSTD::isinf(__b)) { - __a = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__a) ? _Tp(1) : _Tp(0), __a); - __b = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__b) ? _Tp(1) : _Tp(0), __b); - if (_CUDA_VSTD::__constexpr_isnan(__c)) + __a = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__a) ? _Tp(1) : _Tp(0), __a); + __b = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__b) ? _Tp(1) : _Tp(0), __b); + if (_CUDA_VSTD::isnan(__c)) { __c = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __c); } - if (_CUDA_VSTD::__constexpr_isnan(__d)) + if (_CUDA_VSTD::isnan(__d)) { __d = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __d); } __recalc = true; } - if (_CUDA_VSTD::__constexpr_isinf(__c) || _CUDA_VSTD::__constexpr_isinf(__d)) + if (_CUDA_VSTD::isinf(__c) || _CUDA_VSTD::isinf(__d)) { - __c = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__c) ? _Tp(1) : _Tp(0), __c); - __d = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__d) ? _Tp(1) : _Tp(0), __d); - if (_CUDA_VSTD::__constexpr_isnan(__a)) + __c = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__c) ? _Tp(1) : _Tp(0), __c); + __d = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__d) ? _Tp(1) : _Tp(0), __d); + if (_CUDA_VSTD::isnan(__a)) { __a = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __a); } - if (_CUDA_VSTD::__constexpr_isnan(__b)) + if (_CUDA_VSTD::isnan(__b)) { __b = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __b); } __recalc = true; } if (!__recalc - && (_CUDA_VSTD::__constexpr_isinf(__partials.__ac) || _CUDA_VSTD::__constexpr_isinf(__partials.__bd) - || _CUDA_VSTD::__constexpr_isinf(__partials.__ad) || _CUDA_VSTD::__constexpr_isinf(__partials.__bc))) + && (_CUDA_VSTD::isinf(__partials.__ac) || _CUDA_VSTD::isinf(__partials.__bd) + || _CUDA_VSTD::isinf(__partials.__ad) || _CUDA_VSTD::isinf(__partials.__bc))) { - if (_CUDA_VSTD::__constexpr_isnan(__a)) + if (_CUDA_VSTD::isnan(__a)) { __a = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __a); } - if (_CUDA_VSTD::__constexpr_isnan(__b)) + if (_CUDA_VSTD::isnan(__b)) { __b = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __b); } - if (_CUDA_VSTD::__constexpr_isnan(__c)) + if (_CUDA_VSTD::isnan(__c)) { __c = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __c); } - if (_CUDA_VSTD::__constexpr_isnan(__d)) + if (_CUDA_VSTD::isnan(__d)) { __d = _CUDA_VSTD::__constexpr_copysign(_Tp(0), __d); } @@ -637,7 +633,7 @@ operator/(const complex<_Tp>& __z, const complex<_Tp>& __w) _Tp __d = __w.imag(); _Tp __logbw = _CUDA_VSTD::__constexpr_logb( _CUDA_VSTD::__constexpr_fmax(_CUDA_VSTD::__constexpr_fabs(__c), _CUDA_VSTD::__constexpr_fabs(__d))); - if (_CUDA_VSTD::__constexpr_isfinite(__logbw)) + if (_CUDA_VSTD::isfinite(__logbw)) { __ilogbw = static_cast(__logbw); __c = _CUDA_VSTD::__constexpr_scalbn(__c, -__ilogbw); @@ -650,24 +646,20 @@ operator/(const complex<_Tp>& __z, const complex<_Tp>& __w) { bool __z_zero = __a == _Tp(0) && __b == _Tp(0); bool __w_zero = __c == _Tp(0) && __d == _Tp(0); - bool __z_inf = _CUDA_VSTD::__constexpr_isinf(__a) || _CUDA_VSTD::__constexpr_isinf(__b); - bool __w_inf = _CUDA_VSTD::__constexpr_isinf(__c) || _CUDA_VSTD::__constexpr_isinf(__d); + bool __z_inf = _CUDA_VSTD::isinf(__a) || _CUDA_VSTD::isinf(__b); + bool __w_inf = _CUDA_VSTD::isinf(__c) || _CUDA_VSTD::isinf(__d); bool __z_nan = !__z_inf - && ((_CUDA_VSTD::__constexpr_isnan(__a) && _CUDA_VSTD::__constexpr_isnan(__b)) - || (_CUDA_VSTD::__constexpr_isnan(__a) && __b == _Tp(0)) - || (__a == _Tp(0) && _CUDA_VSTD::__constexpr_isnan(__b))); + && ((_CUDA_VSTD::isnan(__a) && _CUDA_VSTD::isnan(__b)) || (_CUDA_VSTD::isnan(__a) && __b == _Tp(0)) + || (__a == _Tp(0) && _CUDA_VSTD::isnan(__b))); bool __w_nan = !__w_inf - && ((_CUDA_VSTD::__constexpr_isnan(__c) && _CUDA_VSTD::__constexpr_isnan(__d)) - || (_CUDA_VSTD::__constexpr_isnan(__c) && __d == _Tp(0)) - || (__c == _Tp(0) && _CUDA_VSTD::__constexpr_isnan(__d))); + && ((_CUDA_VSTD::isnan(__c) && _CUDA_VSTD::isnan(__d)) || (_CUDA_VSTD::isnan(__c) && __d == _Tp(0)) + || (__c == _Tp(0) && _CUDA_VSTD::isnan(__d))); if ((__z_nan || __w_nan) || (__z_inf && __w_inf)) { return complex<_Tp>(_Tp(numeric_limits<_Tp>::quiet_NaN()), _Tp(0)); } - bool __z_nonzero_nan = - !__z_inf && !__z_nan && (_CUDA_VSTD::__constexpr_isnan(__a) || _CUDA_VSTD::__constexpr_isnan(__b)); - bool __w_nonzero_nan = - !__w_inf && !__w_nan && (_CUDA_VSTD::__constexpr_isnan(__c) || _CUDA_VSTD::__constexpr_isnan(__d)); + bool __z_nonzero_nan = !__z_inf && !__z_nan && (_CUDA_VSTD::isnan(__a) || _CUDA_VSTD::isnan(__b)); + bool __w_nonzero_nan = !__w_inf && !__w_nan && (_CUDA_VSTD::isnan(__c) || _CUDA_VSTD::isnan(__d)); if (__z_nonzero_nan || __w_nonzero_nan) { if (__w_zero) @@ -702,26 +694,25 @@ operator/(const complex<_Tp>& __z, const complex<_Tp>& __w) _Tp __x = _CUDA_VSTD::__constexpr_scalbn((__partials.__ac + __partials.__bd) / __denom, -__ilogbw); _Tp __y = _CUDA_VSTD::__constexpr_scalbn((__partials.__bc - __partials.__ad) / __denom, -__ilogbw); #ifndef LIBCUDACXX_ENABLE_SIMPLIFIED_COMPLEX_DIVISION - if (_CUDA_VSTD::__constexpr_isnan(__x) && _CUDA_VSTD::__constexpr_isnan(__y)) + if (_CUDA_VSTD::isnan(__x) && _CUDA_VSTD::isnan(__y)) { - if ((__denom == _Tp(0)) && (!_CUDA_VSTD::__constexpr_isnan(__a) || !_CUDA_VSTD::__constexpr_isnan(__b))) + if ((__denom == _Tp(0)) && (!_CUDA_VSTD::isnan(__a) || !_CUDA_VSTD::isnan(__b))) { __x = _CUDA_VSTD::__constexpr_copysign(_Tp(INFINITY), __c) * __a; __y = _CUDA_VSTD::__constexpr_copysign(_Tp(INFINITY), __c) * __b; } - else if ((_CUDA_VSTD::__constexpr_isinf(__a) || _CUDA_VSTD::__constexpr_isinf(__b)) - && _CUDA_VSTD::__constexpr_isfinite(__c) && _CUDA_VSTD::__constexpr_isfinite(__d)) + else if ((_CUDA_VSTD::isinf(__a) || _CUDA_VSTD::isinf(__b)) && _CUDA_VSTD::isfinite(__c) + && _CUDA_VSTD::isfinite(__d)) { - __a = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__a) ? _Tp(1) : _Tp(0), __a); - __b = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__b) ? _Tp(1) : _Tp(0), __b); + __a = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__a) ? _Tp(1) : _Tp(0), __a); + __b = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__b) ? _Tp(1) : _Tp(0), __b); __x = _Tp(INFINITY) * (__a * __c + __b * __d); __y = _Tp(INFINITY) * (__b * __c - __a * __d); } - else if (_CUDA_VSTD::__constexpr_isinf(__logbw) && __logbw > _Tp(0) && _CUDA_VSTD::__constexpr_isfinite(__a) - && _CUDA_VSTD::__constexpr_isfinite(__b)) + else if (_CUDA_VSTD::isinf(__logbw) && __logbw > _Tp(0) && _CUDA_VSTD::isfinite(__a) && _CUDA_VSTD::isfinite(__b)) { - __c = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__c) ? _Tp(1) : _Tp(0), __c); - __d = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::__constexpr_isinf(__d) ? _Tp(1) : _Tp(0), __d); + __c = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__c) ? _Tp(1) : _Tp(0), __c); + __d = _CUDA_VSTD::__constexpr_copysign(_CUDA_VSTD::isinf(__d) ? _Tp(1) : _Tp(0), __d); __x = _Tp(0) * (__a * __c + __b * __d); __y = _Tp(0) * (__b * __c - __a * __d); } @@ -922,11 +913,11 @@ _LIBCUDACXX_HIDE_FROM_ABI enable_if_t::value, float> arg(_Tp template _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 _Tp norm(const complex<_Tp>& __c) { - if (_CUDA_VSTD::__constexpr_isinf(__c.real())) + if (_CUDA_VSTD::isinf(__c.real())) { return _CUDA_VSTD::abs(__c.real()); } - if (_CUDA_VSTD::__constexpr_isinf(__c.imag())) + if (_CUDA_VSTD::isinf(__c.imag())) { return _CUDA_VSTD::abs(__c.imag()); } @@ -959,7 +950,7 @@ template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> proj(const complex<_Tp>& __c) { complex<_Tp> __r = __c; - if (_CUDA_VSTD::__constexpr_isinf(__c.real()) || _CUDA_VSTD::__constexpr_isinf(__c.imag())) + if (_CUDA_VSTD::isinf(__c.real()) || _CUDA_VSTD::isinf(__c.imag())) { __r = complex<_Tp>(INFINITY, __constexpr_copysign(_Tp(0), __c.imag())); } @@ -969,7 +960,7 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> proj(const complex<_Tp>& __c) template _LIBCUDACXX_HIDE_FROM_ABI enable_if_t<__is_complex_float<_Tp>::value, __libcpp_complex_complex_type<_Tp>> proj(_Tp __re) { - if (_CUDA_VSTD::__constexpr_isinf(__re)) + if (_CUDA_VSTD::isinf(__re)) { __re = _CUDA_VSTD::abs(__re); } @@ -987,33 +978,33 @@ _LIBCUDACXX_HIDE_FROM_ABI enable_if_t::value, __libcpp_complex_ template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> polar(const _Tp& __rho, const _Tp& __theta = _Tp()) { - if (_CUDA_VSTD::__constexpr_isnan(__rho) || _CUDA_VSTD::signbit(__rho)) + if (_CUDA_VSTD::isnan(__rho) || _CUDA_VSTD::signbit(__rho)) { return complex<_Tp>(_Tp(NAN), _Tp(NAN)); } - if (_CUDA_VSTD::__constexpr_isnan(__theta)) + if (_CUDA_VSTD::isnan(__theta)) { - if (_CUDA_VSTD::__constexpr_isinf(__rho)) + if (_CUDA_VSTD::isinf(__rho)) { return complex<_Tp>(__rho, __theta); } return complex<_Tp>(__theta, __theta); } - if (_CUDA_VSTD::__constexpr_isinf(__theta)) + if (_CUDA_VSTD::isinf(__theta)) { - if (_CUDA_VSTD::__constexpr_isinf(__rho)) + if (_CUDA_VSTD::isinf(__rho)) { return complex<_Tp>(__rho, _Tp(NAN)); } return complex<_Tp>(_Tp(NAN), _Tp(NAN)); } _Tp __x = __rho * _CUDA_VSTD::cos(__theta); - if (_CUDA_VSTD::__constexpr_isnan(__x)) + if (_CUDA_VSTD::isnan(__x)) { __x = 0; } _Tp __y = __rho * _CUDA_VSTD::sin(__theta); - if (_CUDA_VSTD::__constexpr_isnan(__y)) + if (_CUDA_VSTD::isnan(__y)) { __y = 0; } @@ -1041,18 +1032,18 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> log10(const complex<_Tp>& __x) template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> sqrt(const complex<_Tp>& __x) { - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(_Tp(INFINITY), __x.imag()); } - if (_CUDA_VSTD::__constexpr_isinf(__x.real())) + if (_CUDA_VSTD::isinf(__x.real())) { if (__x.real() > _Tp(0)) { - return complex<_Tp>( - __x.real(), _CUDA_VSTD::__constexpr_isnan(__x.imag()) ? __x.imag() : __constexpr_copysign(_Tp(0), __x.imag())); + return complex<_Tp>(__x.real(), + _CUDA_VSTD::isnan(__x.imag()) ? __x.imag() : __constexpr_copysign(_Tp(0), __x.imag())); } - return complex<_Tp>(_CUDA_VSTD::__constexpr_isnan(__x.imag()) ? __x.imag() : _Tp(0), + return complex<_Tp>(_CUDA_VSTD::isnan(__x.imag()) ? __x.imag() : _Tp(0), __constexpr_copysign(__x.real(), __x.imag())); } return _CUDA_VSTD::polar(_CUDA_VSTD::sqrt(_CUDA_VSTD::abs(__x)), _CUDA_VSTD::arg(__x) / _Tp(2)); @@ -1068,18 +1059,18 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> exp(const complex<_Tp>& __x) { return complex<_Tp>(_CUDA_VSTD::exp(__x.real()), __constexpr_copysign(_Tp(0), __x.imag())); } - if (_CUDA_VSTD::__constexpr_isinf(__x.real())) + if (_CUDA_VSTD::isinf(__x.real())) { if (__x.real() < _Tp(0)) { - if (!_CUDA_VSTD::__constexpr_isfinite(__i)) + if (!_CUDA_VSTD::isfinite(__i)) { __i = _Tp(1); } } - else if (__i == _Tp(0) || !_CUDA_VSTD::__constexpr_isfinite(__i)) + else if (__i == _Tp(0) || !_CUDA_VSTD::isfinite(__i)) { - if (_CUDA_VSTD::__constexpr_isinf(__i)) + if (_CUDA_VSTD::isinf(__i)) { __i = _Tp(NAN); } @@ -1138,21 +1129,21 @@ template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> asinh(const complex<_Tp>& __x) { const _Tp __pi(static_cast<_Tp>(atan2(+0., -0.))); - if (_CUDA_VSTD::__constexpr_isinf(__x.real())) + if (_CUDA_VSTD::isinf(__x.real())) { - if (_CUDA_VSTD::__constexpr_isnan(__x.imag())) + if (_CUDA_VSTD::isnan(__x.imag())) { return __x; } - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(__x.real(), __constexpr_copysign(__pi * _Tp(0.25), __x.imag())); } return complex<_Tp>(__x.real(), __constexpr_copysign(_Tp(0), __x.imag())); } - if (_CUDA_VSTD::__constexpr_isnan(__x.real())) + if (_CUDA_VSTD::isnan(__x.real())) { - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(__x.imag(), __x.real()); } @@ -1162,7 +1153,7 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> asinh(const complex<_Tp>& __x) } return complex<_Tp>(__x.real(), __x.real()); } - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(__constexpr_copysign(__x.imag(), __x.real()), __constexpr_copysign(__pi / _Tp(2), __x.imag())); } @@ -1176,13 +1167,13 @@ template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> acosh(const complex<_Tp>& __x) { const _Tp __pi(static_cast<_Tp>(atan2(+0., -0.))); - if (_CUDA_VSTD::__constexpr_isinf(__x.real())) + if (_CUDA_VSTD::isinf(__x.real())) { - if (_CUDA_VSTD::__constexpr_isnan(__x.imag())) + if (_CUDA_VSTD::isnan(__x.imag())) { return complex<_Tp>(_CUDA_VSTD::abs(__x.real()), __x.imag()); } - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { if (__x.real() > _Tp(0)) { @@ -1199,15 +1190,15 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> acosh(const complex<_Tp>& __x) } return complex<_Tp>(__x.real(), __constexpr_copysign(_Tp(0), __x.imag())); } - if (_CUDA_VSTD::__constexpr_isnan(__x.real())) + if (_CUDA_VSTD::isnan(__x.real())) { - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(_CUDA_VSTD::abs(__x.imag()), __x.real()); } return complex<_Tp>(__x.real(), __x.real()); } - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(_CUDA_VSTD::abs(__x.imag()), __constexpr_copysign(__pi / _Tp(2), __x.imag())); } @@ -1221,23 +1212,23 @@ template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> atanh(const complex<_Tp>& __x) { const _Tp __pi(static_cast<_Tp>(atan2(+0., -0.))); - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(__constexpr_copysign(_Tp(0), __x.real()), __constexpr_copysign(__pi / _Tp(2), __x.imag())); } - if (_CUDA_VSTD::__constexpr_isnan(__x.imag())) + if (_CUDA_VSTD::isnan(__x.imag())) { - if (_CUDA_VSTD::__constexpr_isinf(__x.real()) || __x.real() == _Tp(0)) + if (_CUDA_VSTD::isinf(__x.real()) || __x.real() == _Tp(0)) { return complex<_Tp>(__constexpr_copysign(_Tp(0), __x.real()), __x.imag()); } return complex<_Tp>(__x.imag(), __x.imag()); } - if (_CUDA_VSTD::__constexpr_isnan(__x.real())) + if (_CUDA_VSTD::isnan(__x.real())) { return complex<_Tp>(__x.real(), __x.real()); } - if (_CUDA_VSTD::__constexpr_isinf(__x.real())) + if (_CUDA_VSTD::isinf(__x.real())) { return complex<_Tp>(__constexpr_copysign(_Tp(0), __x.real()), __constexpr_copysign(__pi / _Tp(2), __x.imag())); } @@ -1254,15 +1245,15 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> atanh(const complex<_Tp>& __x) template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> sinh(const complex<_Tp>& __x) { - if (_CUDA_VSTD::__constexpr_isinf(__x.real()) && !_CUDA_VSTD::__constexpr_isfinite(__x.imag())) + if (_CUDA_VSTD::isinf(__x.real()) && !_CUDA_VSTD::isfinite(__x.imag())) { return complex<_Tp>(__x.real(), _Tp(NAN)); } - if (__x.real() == _Tp(0) && !_CUDA_VSTD::__constexpr_isfinite(__x.imag())) + if (__x.real() == _Tp(0) && !_CUDA_VSTD::isfinite(__x.imag())) { return complex<_Tp>(__x.real(), _Tp(NAN)); } - if (__x.imag() == _Tp(0) && !_CUDA_VSTD::__constexpr_isfinite(__x.real())) + if (__x.imag() == _Tp(0) && !_CUDA_VSTD::isfinite(__x.real())) { return __x; } @@ -1275,11 +1266,11 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> sinh(const complex<_Tp>& __x) template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> cosh(const complex<_Tp>& __x) { - if (_CUDA_VSTD::__constexpr_isinf(__x.real()) && !_CUDA_VSTD::__constexpr_isfinite(__x.imag())) + if (_CUDA_VSTD::isinf(__x.real()) && !_CUDA_VSTD::isfinite(__x.imag())) { return complex<_Tp>(_CUDA_VSTD::abs(__x.real()), _Tp(NAN)); } - if (__x.real() == _Tp(0) && !_CUDA_VSTD::__constexpr_isfinite(__x.imag())) + if (__x.real() == _Tp(0) && !_CUDA_VSTD::isfinite(__x.imag())) { return complex<_Tp>(_Tp(NAN), __x.real()); } @@ -1287,7 +1278,7 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> cosh(const complex<_Tp>& __x) { return complex<_Tp>(_Tp(1), __x.imag()); } - if (__x.imag() == _Tp(0) && !_CUDA_VSTD::__constexpr_isfinite(__x.real())) + if (__x.imag() == _Tp(0) && !_CUDA_VSTD::isfinite(__x.real())) { return complex<_Tp>(_CUDA_VSTD::abs(__x.real()), __x.imag()); } @@ -1300,16 +1291,16 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> cosh(const complex<_Tp>& __x) template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> tanh(const complex<_Tp>& __x) { - if (_CUDA_VSTD::__constexpr_isinf(__x.real())) + if (_CUDA_VSTD::isinf(__x.real())) { - if (!_CUDA_VSTD::__constexpr_isfinite(__x.imag())) + if (!_CUDA_VSTD::isfinite(__x.imag())) { return complex<_Tp>(__constexpr_copysign(_Tp(1), __x.real()), _Tp(0)); } return complex<_Tp>(__constexpr_copysign(_Tp(1), __x.real()), __constexpr_copysign(_Tp(0), _CUDA_VSTD::sin(_Tp(2) * __x.imag()))); } - if (_CUDA_VSTD::__constexpr_isnan(__x.real()) && __x.imag() == _Tp(0)) + if (_CUDA_VSTD::isnan(__x.real()) && __x.imag() == _Tp(0)) { return __x; } @@ -1317,7 +1308,7 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> tanh(const complex<_Tp>& __x) _Tp __2i(_Tp(2) * __x.imag()); _Tp __d(_CUDA_VSTD::cosh(__2r) + _CUDA_VSTD::cos(__2i)); _Tp __2rsh(_CUDA_VSTD::sinh(__2r)); - if (_CUDA_VSTD::__constexpr_isinf(__2rsh) && _CUDA_VSTD::__constexpr_isinf(__d)) + if (_CUDA_VSTD::isinf(__2rsh) && _CUDA_VSTD::isinf(__d)) { return complex<_Tp>(__2rsh > _Tp(0) ? _Tp(1) : _Tp(-1), __2i > _Tp(0) ? _Tp(0) : _Tp(-0.)); } @@ -1339,13 +1330,13 @@ template _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> acos(const complex<_Tp>& __x) { const _Tp __pi(static_cast<_Tp>(atan2(+0., -0.))); - if (_CUDA_VSTD::__constexpr_isinf(__x.real())) + if (_CUDA_VSTD::isinf(__x.real())) { - if (_CUDA_VSTD::__constexpr_isnan(__x.imag())) + if (_CUDA_VSTD::isnan(__x.imag())) { return complex<_Tp>(__x.imag(), __x.real()); } - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { if (__x.real() < _Tp(0)) { @@ -1359,15 +1350,15 @@ _LIBCUDACXX_HIDE_FROM_ABI complex<_Tp> acos(const complex<_Tp>& __x) } return complex<_Tp>(_Tp(0), _CUDA_VSTD::signbit(__x.imag()) ? __x.real() : -__x.real()); } - if (_CUDA_VSTD::__constexpr_isnan(__x.real())) + if (_CUDA_VSTD::isnan(__x.real())) { - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(__x.real(), -__x.imag()); } return complex<_Tp>(__x.real(), __x.real()); } - if (_CUDA_VSTD::__constexpr_isinf(__x.imag())) + if (_CUDA_VSTD::isinf(__x.imag())) { return complex<_Tp>(__pi / _Tp(2), -__x.imag()); } diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp new file mode 100644 index 00000000000..a0f288d1c5a --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp @@ -0,0 +1,410 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// + +#include +#include +#include +#include + +#include "fp_compare.h" +#include "test_macros.h" + +template +__host__ __device__ void test_signbit(float val) +{ + static_assert((cuda::std::is_same::value), ""); + assert(cuda::std::signbit(T(val)) == false); + assert(cuda::std::signbit(T(-1.0)) == true); + assert(cuda::std::signbit(T(0.0)) == false); +} + +__host__ __device__ void test_signbit(float val) +{ + test_signbit(val); + test_signbit(val); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + test_signbit(val); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + test_signbit<__half>(val); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + test_signbit<__nv_bfloat16>(val); +#endif // _LIBCUDACXX_HAS_NVBF16 + + assert(cuda::std::signbit(0u) == false); + assert(cuda::std::signbit(cuda::std::numeric_limits::max()) == false); + assert(cuda::std::signbit(1) == false); + assert(cuda::std::signbit(-1) == true); + assert(cuda::std::signbit(cuda::std::numeric_limits::max()) == false); + assert(cuda::std::signbit(cuda::std::numeric_limits::min()) == true); +} + +template +__host__ __device__ void test_isfinite(float val) +{ + static_assert((cuda::std::is_same::value), ""); + assert(cuda::std::isfinite(T(val)) == true); + assert(cuda::std::isfinite(T(-1.0f)) == true); + assert(cuda::std::isfinite(T(1.0f)) == true); + assert(cuda::std::isfinite(T(NAN)) == false); + assert(cuda::std::isfinite(T(INFINITY)) == false); + assert(cuda::std::isfinite(-T(INFINITY)) == false); +} + +__host__ __device__ void test_isfinite(float val) +{ + test_isfinite(val); + test_isfinite(val); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + test_isfinite(); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + test_isfinite<__half>(val); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + test_isfinite<__nv_bfloat16>(val); +#endif // _LIBCUDACXX_HAS_NVBF16 + + assert(cuda::std::isfinite(0) == true); + assert(cuda::std::isfinite(1) == true); + assert(cuda::std::isfinite(-1) == true); + assert(cuda::std::isfinite(cuda::std::numeric_limits::max()) == true); + assert(cuda::std::isfinite(cuda::std::numeric_limits::min()) == true); +} + +__host__ __device__ _CCCL_CONSTEXPR_ISFINITE bool test_constexpr_isfinite(float val) +{ + return cuda::std::isfinite(val); +} + +template +__host__ __device__ void test_isnormal(float val) +{ + static_assert((cuda::std::is_same::value), ""); + assert(cuda::std::isnormal(T(val)) == false); + assert(cuda::std::isnormal(T(-1.0f)) == true); + assert(cuda::std::isnormal(T(1.0f)) == true); + assert(cuda::std::isnormal(T(0.0f)) == false); + assert(cuda::std::isnormal(T(NAN)) == false); + assert(cuda::std::isnormal(T(INFINITY)) == false); + assert(cuda::std::isnormal(-T(INFINITY)) == false); +} + +__host__ __device__ void test_isnormal(float val) +{ + test_isnormal(val); + test_isnormal(val); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + test_isnormal(); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + test_isnormal<__half>(val); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + test_isnormal<__nv_bfloat16>(val); +#endif // _LIBCUDACXX_HAS_NVBF16 + + assert(cuda::std::isnormal(0) == false); + assert(cuda::std::isnormal(1) == true); + assert(cuda::std::isnormal(-1) == true); + assert(cuda::std::isnormal(cuda::std::numeric_limits::max()) == true); + assert(cuda::std::isnormal(cuda::std::numeric_limits::min()) == true); +} + +__host__ __device__ void test_isgreater(float val) +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), + ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), + ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::isgreater(-1.0, 0.F) == false); +} + +__host__ __device__ void test_isgreaterequal(float val) +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), + ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), + ""); + static_assert( + (cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert( + (cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), + ""); + static_assert((cuda::std::is_same::value), + ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::isgreaterequal(-1.0, 0.F) == false); +} + +__host__ __device__ void test_isinf(float val) +{ + static_assert((cuda::std::is_same::value), ""); + + typedef decltype(cuda::std::isinf((double) 0)) DoubleRetType; + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::isinf(-1.0) == false); + assert(cuda::std::isinf(0) == false); + assert(cuda::std::isinf(1) == false); + assert(cuda::std::isinf(-1) == false); + assert(cuda::std::isinf(cuda::std::numeric_limits::max()) == false); + assert(cuda::std::isinf(cuda::std::numeric_limits::min()) == false); +} + +__host__ __device__ _CCCL_CONSTEXPR_ISINF bool test_constexpr_isinf(float val) +{ + return cuda::std::isinf(val); +} + +__host__ __device__ void test_isless(float val) +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), + ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::isless(-1.0, 0.F) == true); +} + +__host__ __device__ void test_islessequal(float val) +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), + ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert( + (cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::islessequal(-1.0, 0.F) == true); +} + +__host__ __device__ void test_islessgreater(float val) +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), + ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert( + (cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), + ""); + static_assert((cuda::std::is_same::value), + ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::islessgreater(-1.0, 0.F) == true); +} + +__host__ __device__ void test_isnan(float val) +{ + static_assert((cuda::std::is_same::value), ""); + + typedef decltype(cuda::std::isnan((double) 0)) DoubleRetType; + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::isnan(-1.0) == false); + assert(cuda::std::isnan(0) == false); + assert(cuda::std::isnan(1) == false); + assert(cuda::std::isnan(-1) == false); + assert(cuda::std::isnan(cuda::std::numeric_limits::max()) == false); + assert(cuda::std::isnan(cuda::std::numeric_limits::min()) == false); +} + +__host__ __device__ _CCCL_CONSTEXPR_ISNAN bool test_constexpr_isnan(float val) +{ + return cuda::std::isnan(val); +} + +__host__ __device__ void test_isunordered(float val) +{ + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#if !defined(_LIBCUDACXX_HAS_NO_LONG_DOUBLE) + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), + ""); +#endif // !_LIBCUDACXX_HAS_NO_LONG_DOUBLE +#ifdef _LIBCUDACXX_HAS_NVFP16 + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVFP16 +#ifdef _LIBCUDACXX_HAS_NVBF16 + static_assert( + (cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); + static_assert((cuda::std::is_same::value), ""); +#endif // _LIBCUDACXX_HAS_NVBF16 + assert(cuda::std::isunordered(-1.0, 0.F) == false); +} + +__host__ __device__ void test(float val) +{ + test_signbit(val); + test_isfinite(val); + test_isnormal(val); + test_isgreater(val); + test_isgreaterequal(val); + test_isinf(val); + test_isless(val); + test_islessequal(val); + test_islessgreater(val); + test_isnan(val); + test_isunordered(val); +} + +__global__ void test_global_kernel(float* val) +{ + test(*val); +} + +int main(int, char**) +{ + volatile float val = 0.0f; + test(val); + +#if defined(_CCCL_BUILTIN_ISNAN) + static_assert(!test_constexpr_isnan(0.0f), ""); +#endif // _CCCL_BUILTIN_ISNAN + +#if defined(_CCCL_BUILTIN_ISINF) + static_assert(!test_constexpr_isinf(0.0f), ""); +#endif // _CCCL_BUILTIN_ISINF + +#if defined(_CCCL_BUILTIN_ISFINITE) || (defined(_CCCL_BUILTIN_ISINF) && defined(_CCCL_BUILTIN_ISNAN)) + static_assert(test_constexpr_isfinite(0.0f), ""); +#endif // _CCCL_BUILTIN_ISFINITE|| (_CCCL_BUILTIN_ISINF && _CCCL_BUILTIN_ISNAN) + + return 0; +}