diff --git a/docs/libcudacxx/extended_api/numeric.rst b/docs/libcudacxx/extended_api/numeric.rst index 3a00e877662..7a6957a242c 100644 --- a/docs/libcudacxx/extended_api/numeric.rst +++ b/docs/libcudacxx/extended_api/numeric.rst @@ -11,6 +11,7 @@ Numeric numeric/saturating_add_overflow numeric/div_overflow numeric/saturating_div_overflow + numeric/isclose numeric/mul_overflow numeric/saturating_mul_overflow numeric/narrow @@ -59,6 +60,11 @@ Numeric - CCCL 3.2.0 - CUDA 13.2 + * - :ref:`cuda::isclose ` + - Checks whether two values are approximately equal + - CCCL 3.5.0 + - CUDA 13.5 + * - :ref:`cuda::mul_overflow ` - Performs multiplication with overflow checking - CCCL 3.4.0 diff --git a/docs/libcudacxx/extended_api/numeric/isclose.rst b/docs/libcudacxx/extended_api/numeric/isclose.rst new file mode 100644 index 00000000000..d385533454e --- /dev/null +++ b/docs/libcudacxx/extended_api/numeric/isclose.rst @@ -0,0 +1,102 @@ +.. _libcudacxx-extended-api-numeric-isclose: + +``cuda::isclose`` +================= + +Defined in ```` header. + +.. code:: cpp + + namespace cuda { + + template + [[nodiscard]] __host__ __device__ __tile__ constexpr + bool isclose(T lhs, T rhs) noexcept; + + template + [[nodiscard]] __host__ __device__ __tile__ constexpr + bool isclose(T lhs, T rhs, float relative_tol) noexcept; + + template + [[nodiscard]] __host__ __device__ __tile__ constexpr + bool isclose(T lhs, T rhs, float relative_tol, T absolute_tol) noexcept; + + template + [[nodiscard]] __host__ __device__ constexpr + bool isclose(const Complex& lhs, const Complex& rhs) noexcept; + + template + [[nodiscard]] __host__ __device__ constexpr + bool isclose(const Complex& lhs, const Complex& rhs, float relative_tol) noexcept; + + template + [[nodiscard]] __host__ __device__ constexpr + bool isclose(const Complex& lhs, const Complex& rhs, + float relative_tol, + AbsTol absolute_tol) noexcept; + + } // namespace cuda + +``cuda::isclose`` checks whether two values are approximately equal using the weak symmetric comparison in a similar manner to `PEP 485 `_: + +.. code:: cpp + + abs(lhs - rhs) <= max(absolute_tol, relative_tol * max(abs(lhs), abs(rhs))) + +- The overloads without ``absolute_tol`` use ``absolute_tol == 0``. +- The overloads without ``relative_tol`` use a default relative tolerance based on half of available digits of accuracy. + +**Parameters** + +- ``lhs``: The first value to compare. +- ``rhs``: The second value to compare. +- ``relative_tol``: The relative tolerance. Passing ``0`` performs a purely absolute tolerance check when ``absolute_tol`` is non-zero. +- ``absolute_tol``: The absolute tolerance. This is useful for comparisons near zero. + +**Precision** + +- ``relative_tol``:Must be in the range [0.0, 1.0]. +- ``absolute_tol``: Must be finite and non-negative. + +**Return value** + +- Returns ``true`` if ``lhs`` and ``rhs`` are close to each other, otherwise returns ``false``. + +**Constraints** + +- Scalar overloads require ``lhs``, ``rhs``, ``absolute_tol`` to have the same arithmetic type (integer or floating point). +- Complex overloads accept ``cuda::std::complex``, ``cuda::complex``, ``std::complex`` operands. +- ``AbsTol`` must be the same type as the complex value type. + +**Special values** + +- ``NaN`` is never close to any value, including another ``NaN``. +- Infinity and negative infinity are only close to themselves. + +Example +------- + +.. code:: cuda + + #include + #include + #include + + __global__ void kernel() + { + assert(cuda::isclose( 1.0, 1.0 + 5e-10)); + assert(!cuda::isclose(1.0, 1.0 + 5e-8)); + + assert(!cuda::isclose(0.0, 1e-12)); + assert(cuda::isclose( 0.0, 1e-12, 0.0, 1e-12)); + + cuda::std::complex z1{1.0, 1.0}; + cuda::std::complex z2{2.0, 0.0}; + assert(cuda::isclose(z1, z2, 0.75)); + } + + int main() + { + kernel<<<1, 1>>>(); + cudaDeviceSynchronize(); + } diff --git a/libcudacxx/include/cuda/__numeric/isclose.h b/libcudacxx/include/cuda/__numeric/isclose.h new file mode 100644 index 00000000000..d05f6fc1e6c --- /dev/null +++ b/libcudacxx/include/cuda/__numeric/isclose.h @@ -0,0 +1,280 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// 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) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___NUMERIC_ISCLOSE_H +#define _CUDA___NUMERIC_ISCLOSE_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 +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +template +using __isclose_compare_t _CCCL_NODEBUG_ALIAS = ::cuda::std:: + conditional_t<(::cuda::std::__is_extended_floating_point_v<_Tp> && sizeof(_Tp) <= sizeof(float)), float, _Tp>; + +// compute 10^-(digits10 / 2) +template +[[nodiscard]] _CCCL_API _CCCL_CONSTEVAL float __isclose_default_relative_tolerance() noexcept +{ + constexpr auto __digits = ::cuda::ceil_div(::cuda::std::numeric_limits<_Tp>::max_digits10, 2); + auto __exp = 1.0f; + for (int __i = 0; __i < __digits; ++__i) + { + __exp *= 10.0f; + } + return 1.0f / __exp; +} + +template +[[nodiscard]] _CCCL_API constexpr bool +__isclose_fp_impl(const _Tp __lhs, const _Tp __rhs, const float __rel_tol, const _Tp __abs_tol) noexcept +{ + _CCCL_ASSERT(::cuda::in_range(__rel_tol, 0.0f, 1.0f), + "cuda::isclose: relative tolerance must be in the range [0.0, 1.0]"); + _CCCL_ASSERT(::cuda::std::isfinite(__abs_tol) && __abs_tol >= _Tp{0}, + "cuda::isclose: absolute tolerance must be finite and non-negative"); + if (__lhs == __rhs) + { + return true; + } + if (!::cuda::std::isfinite(__lhs) || !::cuda::std::isfinite(__rhs)) + { + return false; + } + const auto __diff = ::cuda::std::fabs(__lhs - __rhs); + const auto __lhs_abs = ::cuda::std::fabs(__lhs); + const auto __rhs_abs = ::cuda::std::fabs(__rhs); + const auto __rel_value = static_cast<_Tp>(__rel_tol * ::cuda::std::fmax(__lhs_abs, __rhs_abs)); + return __diff <= ::cuda::std::fmax(__abs_tol, __rel_value); +} + +template +[[nodiscard]] _CCCL_HOST_DEVICE_API bool __isclose_complex_impl( + const _ComplexType& __lhs, const _ComplexType& __rhs, const float __rel_tol, const _AbsTol __abs_tol) noexcept +{ + using __scalar_t _CCCL_NODEBUG_ALIAS = typename _ComplexType::value_type; + using __compare_t _CCCL_NODEBUG_ALIAS = __isclose_compare_t<__scalar_t>; +#if _CCCL_HAS_FLOAT128() + static_assert(!::cuda::std::is_same_v<__scalar_t, __float128>, "cuda::isclose: __float128 is not supported"); +#endif // _CCCL_HAS_FLOAT128() + _CCCL_ASSERT(::cuda::in_range(__rel_tol, 0.0f, 1.0f), + "cuda::isclose: relative tolerance must be in the range [0.0, 1.0]"); + _CCCL_ASSERT(::cuda::std::isfinite(__abs_tol) && __abs_tol >= __scalar_t{0}, + "cuda::isclose: absolute tolerance must be finite and non-negative"); + + const auto __lhs_real = static_cast<__compare_t>(::cuda::__get_real(__lhs)); + const auto __lhs_imag = static_cast<__compare_t>(::cuda::__get_imag(__lhs)); + const auto __rhs_real = static_cast<__compare_t>(::cuda::__get_real(__rhs)); + const auto __rhs_imag = static_cast<__compare_t>(::cuda::__get_imag(__rhs)); + const auto __abs = static_cast<__compare_t>(__abs_tol); + + if (__lhs_real == __rhs_real && __lhs_imag == __rhs_imag) + { + return true; + } + if (!::cuda::std::isfinite(__lhs_real) || !::cuda::std::isfinite(__lhs_imag) || !::cuda::std::isfinite(__rhs_real) + || !::cuda::std::isfinite(__rhs_imag)) + { + return false; + } + const auto __diff = ::cuda::std::hypot(__lhs_real - __rhs_real, __lhs_imag - __rhs_imag); + const auto __lhs_abs = ::cuda::std::hypot(__lhs_real, __lhs_imag); + const auto __rhs_abs = ::cuda::std::hypot(__rhs_real, __rhs_imag); + const auto __rel_value = __rel_tol * ::cuda::std::fmax(__lhs_abs, __rhs_abs); + return __diff <= ::cuda::std::fmax(__abs, __rel_value); +} + +template +[[nodiscard]] _CCCL_API constexpr ::cuda::std::make_unsigned_t<_Tp> +__safe_abs_diff(const _Tp __lhs, const _Tp __rhs) noexcept +{ + using __unsigned_t _CCCL_NODEBUG_ALIAS = ::cuda::std::make_unsigned_t<_Tp>; + const auto __lhs_abs = ::cuda::uabs(__lhs); + const auto __rhs_abs = ::cuda::uabs(__rhs); + const auto __is_lhs_negative = ::cuda::std::is_signed_v<_Tp> && __lhs < _Tp{0}; + const auto __is_rhs_negative = ::cuda::std::is_signed_v<_Tp> && __rhs < _Tp{0}; + + if (__is_lhs_negative != __is_rhs_negative) + { + return static_cast<__unsigned_t>(__lhs_abs + __rhs_abs); + } + return (__lhs_abs < __rhs_abs) + ? static_cast<__unsigned_t>(__rhs_abs - __lhs_abs) + : static_cast<__unsigned_t>(__lhs_abs - __rhs_abs); +} + +template +[[nodiscard]] _CCCL_API constexpr bool +__isclose_integer_impl(const _Tp __lhs, const _Tp __rhs, const float __rel_tol, const _Tp __abs_tol) noexcept +{ + _CCCL_ASSERT(::cuda::in_range(__rel_tol, 0.0f, 1.0f), + "cuda::isclose: relative tolerance must be in the range [0.0, 1.0]"); + if constexpr (::cuda::std::__cccl_is_signed_integer_v<_Tp>) + { + _CCCL_ASSERT(__abs_tol >= _Tp{0}, "cuda::isclose: absolute tolerance must be non-negative"); + } + using __unsigned_t _CCCL_NODEBUG_ALIAS = ::cuda::std::make_unsigned_t<_Tp>; + const auto __lhs_abs = ::cuda::uabs(__lhs); + const auto __rhs_abs = ::cuda::uabs(__rhs); + const auto __diff = ::cuda::__safe_abs_diff(__lhs, __rhs); + const auto __abs = static_cast<__unsigned_t>(__abs_tol); + const auto __rel_value = static_cast<__unsigned_t>(__rel_tol * ::cuda::std::max(__lhs_abs, __rhs_abs)); + return __diff <= ::cuda::std::max(__abs, __rel_value); +} + +//---------------------------------------------------------------------------------------------------------------------- +// Public API + +// Scalar overloads + +//! @brief Checks whether two arithmetic values are close to each other using a relative and absolute tolerance. +//! +//! @param __lhs The first value to compare. +//! @param __rhs The second value to compare. +//! @param __rel_tol The relative tolerance. +//! @param __abs_tol The absolute tolerance. +//! @return True if __lhs and __rhs are close to each other, false otherwise. +_CCCL_TEMPLATE(typename _Tp) +_CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> || ::cuda::is_floating_point_v<_Tp>) +[[nodiscard]] _CCCL_API constexpr bool +isclose(const _Tp __lhs, const _Tp __rhs, const float __rel_tol, const _Tp __abs_tol) noexcept +{ + if constexpr (::cuda::std::__cccl_is_integer_v<_Tp>) + { + return ::cuda::__isclose_integer_impl(+__lhs, +__rhs, __rel_tol, +__abs_tol); + } + else + { + using __value_t _CCCL_NODEBUG_ALIAS = __isclose_compare_t<_Tp>; + return ::cuda::__isclose_fp_impl( + static_cast<__value_t>(__lhs), static_cast<__value_t>(__rhs), __rel_tol, static_cast<__value_t>(__abs_tol)); + } +} + +//! @brief Checks whether two arithmetic values are close to each other using a relative tolerance. +//! +//! @param __lhs The first value to compare. +//! @param __rhs The second value to compare. +//! @param __rel_tol The relative tolerance. +//! @return True if __lhs and __rhs are close to each other, false otherwise. +_CCCL_TEMPLATE(typename _Tp) +_CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> || ::cuda::is_floating_point_v<_Tp>) +[[nodiscard]] _CCCL_API constexpr bool isclose(const _Tp __lhs, const _Tp __rhs, const float __rel_tol) noexcept +{ + return ::cuda::isclose(__lhs, __rhs, __rel_tol, _Tp{0}); +} + +//! @brief Checks whether two arithmetic values are close to each other using the default relative tolerance. +//! +//! @param __lhs The first value to compare. +//! @param __rhs The second value to compare. +//! @return True if __lhs and __rhs are close to each other, false otherwise. +_CCCL_TEMPLATE(typename _Tp) +_CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> || ::cuda::is_floating_point_v<_Tp>) +[[nodiscard]] _CCCL_API constexpr bool isclose(const _Tp __lhs, const _Tp __rhs) noexcept +{ + if constexpr (::cuda::std::__cccl_is_integer_v<_Tp>) + { + return __lhs == __rhs; + } + else + { + constexpr auto __rel_tol = ::cuda::__isclose_default_relative_tolerance<_Tp>(); + return ::cuda::isclose(__lhs, __rhs, __rel_tol, _Tp{0}); + } +} + +// Complex overloads + +template > +inline constexpr bool __isclose_complex_comparison_v = false; + +template +inline constexpr bool __isclose_complex_comparison_v<_Tp, _AbsTol, true> = + ::cuda::std::is_same_v; + +//! @brief Checks whether two complex values are close to each other using a relative and absolute tolerance. +//! +//! @param __lhs The first value to compare. +//! @param __rhs The second value to compare. +//! @param __rel_tol The relative tolerance. +//! @param __abs_tol The absolute tolerance. +//! @return True if __lhs and __rhs are close to each other, false otherwise. +_CCCL_TEMPLATE(typename _ComplexType, typename _AbsTol) +_CCCL_REQUIRES(__isclose_complex_comparison_v<_ComplexType, _AbsTol>) +[[nodiscard]] _CCCL_HOST_DEVICE_API bool +isclose(const _ComplexType& __lhs, const _ComplexType& __rhs, const float __rel_tol, const _AbsTol __abs_tol) noexcept +{ + return ::cuda::__isclose_complex_impl(__lhs, __rhs, __rel_tol, __abs_tol); +} + +//! @brief Checks whether two complex values are close to each other using a relative tolerance. +//! +//! @param __lhs The first value to compare. +//! @param __rhs The second value to compare. +//! @param __rel_tol The relative tolerance. +//! @return True if __lhs and __rhs are close to each other, false otherwise. +_CCCL_TEMPLATE(typename _ComplexType) +_CCCL_REQUIRES(__is_any_complex_v<_ComplexType>) +[[nodiscard]] _CCCL_HOST_DEVICE_API bool +isclose(const _ComplexType& __lhs, const _ComplexType& __rhs, const float __rel_tol) noexcept +{ + using __scalar_t _CCCL_NODEBUG_ALIAS = typename _ComplexType::value_type; + return ::cuda::isclose(__lhs, __rhs, __rel_tol, __scalar_t{0}); +} + +//! @brief Checks whether two complex values are close to each other using the default relative tolerance. +//! +//! @param __lhs The first value to compare. +//! @param __rhs The second value to compare. +//! @return True if __lhs and __rhs are close to each other, false otherwise. +_CCCL_TEMPLATE(typename _ComplexType) +_CCCL_REQUIRES(__is_any_complex_v<_ComplexType>) +[[nodiscard]] _CCCL_HOST_DEVICE_API bool isclose(const _ComplexType& __lhs, const _ComplexType& __rhs) noexcept +{ + using __scalar_t _CCCL_NODEBUG_ALIAS = typename _ComplexType::value_type; + return ::cuda::isclose(__lhs, __rhs, ::cuda::__isclose_default_relative_tolerance<__scalar_t>(), __scalar_t{0}); +} + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___NUMERIC_ISCLOSE_H diff --git a/libcudacxx/include/cuda/numeric b/libcudacxx/include/cuda/numeric index f9beaa7fa7f..9cf3feaf42f 100644 --- a/libcudacxx/include/cuda/numeric +++ b/libcudacxx/include/cuda/numeric @@ -22,6 +22,7 @@ #include #include +#include #include #include #include diff --git a/libcudacxx/test/libcudacxx/cuda/numeric/isclose/isclose.pass.cpp b/libcudacxx/test/libcudacxx/cuda/numeric/isclose/isclose.pass.cpp new file mode 100644 index 00000000000..5b329f6e9ba --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/numeric/isclose/isclose.pass.cpp @@ -0,0 +1,218 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// 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) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include +#include +#include + +#if _CCCL_HAS_HOST_STD_LIB() +# include +#endif // _CCCL_HAS_HOST_STD_LIB() + +#include "test_macros.h" + +template +inline constexpr bool has_isclose_abs_tol_v = false; + +template +inline constexpr bool has_isclose_abs_tol_v< + T, + AbsTol, + cuda::std::void_t(), cuda::std::declval(), 0.0f, cuda::std::declval()))>> = true; + +template +TEST_FUNC constexpr bool test_floating_point() +{ + static_assert(cuda::std::is_same_v); + static_assert(cuda::std::is_same_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::isclose(T{}, T{}))); + static_assert(noexcept(cuda::isclose(T{}, T{}, 0.0f))); + static_assert(noexcept(cuda::isclose(T{}, T{}, 0.0f, T{}))); + + constexpr auto tol = cuda::__isclose_default_relative_tolerance(); + assert(cuda::isclose(T{1}, T{1})); + assert(cuda::isclose(T{1}, T{1} + static_cast(tol / 2.0f))); + assert(!cuda::isclose(T{1}, T{1} + static_cast(tol * 2.0f))); + + assert(cuda::isclose(T{10}, T{11}, 0.1f)); + assert(cuda::isclose(T{11}, T{10}, 0.1f)); + assert(!cuda::isclose(T{10}, T{12}, 0.1f)); + + const auto inf = cuda::std::numeric_limits::infinity(); + const auto nan = cuda::std::numeric_limits::quiet_NaN(); + assert(cuda::isclose(inf, inf)); + assert(cuda::isclose(-inf, -inf)); + assert(!cuda::isclose(inf, -inf)); + assert(!cuda::isclose(nan, nan)); + assert(!cuda::isclose(nan, T{})); + return true; +} + +template +TEST_FUNC constexpr bool test_integral() +{ + static_assert(cuda::std::is_same_v); + static_assert(cuda::std::is_same_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::isclose(T{}, T{}))); + static_assert(noexcept(cuda::isclose(T{}, T{}, 0.0f))); + static_assert(noexcept(cuda::isclose(T{}, T{}, 0.0f, T{}))); + + // equal values + assert(cuda::isclose(T{1}, T{1})); + assert(!cuda::isclose(T{1}, T{2})); + + // relative tolerance values + assert(cuda::isclose(T{100}, T{101}, 0.02f)); + assert(cuda::isclose(T{101}, T{100}, 0.02f)); + assert(!cuda::isclose(T{100}, T{103}, 0.02f)); + + // absolute tolerance values, positive/negative values + assert(cuda::isclose(-3, 4, 0.0f, 7)); + assert(!cuda::isclose(-3, 4, 0.0f, 6)); + assert(cuda::isclose(3, -4, 0.0f, 7)); + assert(!cuda::isclose(3, -4, 0.0f, 6)); + + // absolute tolerance values, negative values + assert(cuda::isclose(-3, -5, 0.0f, 2)); + assert(!cuda::isclose(-3, -5, 0.0f, 1)); + assert(cuda::isclose(-5, -3, 0.0f, 2)); + assert(!cuda::isclose(-5, -3, 0.0f, 1)); + + // absolute tolerance values, positive values + assert(cuda::isclose(3u, 5u, 0.0f, 2u)); + assert(!cuda::isclose(3u, 5u, 0.0f, 1u)); + assert(cuda::isclose(5u, 3u, 0.0f, 2u)); + assert(!cuda::isclose(5u, 3u, 0.0f, 1u)); + return true; +} + +template +TEST_FUNC void test_complex() +{ + using T = typename Complex::value_type; + static_assert(cuda::std::is_same_v); + static_assert(cuda::std::is_same_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::isclose(cuda::std::declval(), cuda::std::declval()))); + static_assert(noexcept(cuda::isclose(cuda::std::declval(), cuda::std::declval(), 0.0f))); + static_assert(noexcept( + cuda::isclose(cuda::std::declval(), cuda::std::declval(), 0.0f, cuda::std::declval()))); + + assert(cuda::isclose(Complex{T{1}, T{2}}, Complex{T{1}, T{2}})); + assert(cuda::isclose(Complex{T{3}, T{4}}, Complex{T{3}, T{4.4}}, 0.1f)); + assert(!cuda::isclose(Complex{T{3}, T{4}}, Complex{T{3}, T{5}}, 0.1f)); + + assert(!cuda::isclose(Complex{T{0}, T{0}}, Complex{T{0.3}, T{0.4}})); + auto abs_tol = T{0.5}; +#if _LIBCUDACXX_HAS_NVBF16() + if constexpr (cuda::std::is_same_v) + { + abs_tol = T{0.51}; + } +#endif // _LIBCUDACXX_HAS_NVBF16() + assert(cuda::isclose(Complex{T{0}, T{0}}, Complex{T{0.3}, T{0.4}}, 0.0f, abs_tol)); + assert(!cuda::isclose(Complex{T{0}, T{0}}, Complex{T{0.3}, T{0.4}}, 0.0f, T{0.25})); + + const auto inf = cuda::std::numeric_limits::infinity(); + const auto nan = cuda::std::numeric_limits::quiet_NaN(); + assert(cuda::isclose(Complex{inf, T{1}}, Complex{inf, T{1}})); + assert(!cuda::isclose(Complex{inf, T{1}}, Complex{inf, T{2}}, 1.0f)); + assert(!cuda::isclose(Complex{nan, T{}}, Complex{nan, T{}})); + assert(!cuda::isclose(Complex{nan, T{}}, Complex{})); +} + +TEST_FUNC constexpr void test_invalid_complex_cases() +{ + static_assert(!has_isclose_abs_tol_v, float>); + static_assert(!has_isclose_abs_tol_v, float>); + static_assert(!has_isclose_abs_tol_v, double>); + static_assert(!has_isclose_abs_tol_v, double>); +#if _CCCL_HAS_HOST_STD_LIB() + static_assert(!has_isclose_abs_tol_v, float>); + static_assert(!has_isclose_abs_tol_v, double>); +#endif // _CCCL_HAS_HOST_STD_LIB() +} + +TEST_FUNC constexpr bool test() +{ + test_floating_point(); + test_floating_point(); +#if _CCCL_HAS_LONG_DOUBLE() + test_floating_point(); +#endif // _CCCL_HAS_LONG_DOUBLE() +#if _CCCL_HAS_FLOAT128() + test_floating_point<__float128>(); +#endif // _CCCL_HAS_FLOAT128() + test_integral(); + test_integral(); + test_integral(); + test_integral(); + test_integral(); + test_integral(); + test_integral(); + test_integral(); + test_integral(); + test_integral(); +#if _CCCL_HAS_INT128() + test_integral<__int128_t>(); + test_integral<__uint128_t>(); +#endif // _CCCL_HAS_INT128() + + test_invalid_complex_cases(); + return true; +} + +template