Skip to content

[libcu++] Support complex type mixing in cuda::isclose #9891

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
davebayer wants to merge 1 commit into
base: main
Choose a base branch
from
+35 −42
Open

[libcu++] Support complex type mixing in cuda::isclose #9891

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
17 changes: 8 additions & 9 deletions docs/libcudacxx/extended_api/numeric/isclose.rst
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,20 +21,20 @@ Defined in ``<cuda/numeric>`` header.
[[nodiscard]] __host__ __device__
bool isclose(T lhs, T rhs, float relative_tol, T absolute_tol) noexcept;

template <class Complex>
template <class T>
[[nodiscard]] __host__ __device__
bool isclose(const Complex& lhs, const Complex& rhs) noexcept;
bool isclose(const /*complex-type*/<T>& lhs, const /*complex-type*/<T>& rhs) noexcept;

@Jacobfaib Jacobfaib Jul 16, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this functionally any different from the previous form? From your change it seems like what we actually support is

template <class LhsComplex, class RhsComplex>
isclose(const LhsComplex&, const RhsComplex&);


template <class Complex>
template <class T>
[[nodiscard]] __host__ __device__
bool isclose(const Complex& lhs, const Complex& rhs, float relative_tol) noexcept;
bool isclose(const /*complex-type*/<T>& lhs, const /*complex-type*/<T>& rhs, float relative_tol) noexcept;

template <class Complex, class AbsTol>
template <class T>
[[nodiscard]] __host__ __device__
bool isclose(const Complex& lhs,
const Complex& rhs,
bool isclose(const /*complex-type*/<T>& lhs,
const /*complex-type*/<T>& rhs,
float relative_tol,
AbsTol absolute_tol) noexcept;
T absolute_tol) noexcept;
Comment on lines +24 to +37

@coderabbitai coderabbitai Bot Jul 16, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

important: Document cuda::complex<T> and independently selectable lhs/rhs complex types. The signatures repeat one /*complex-type*/ placeholder and Line 70 lists only cuda::std::complex<T> and std::complex<T>, but the new tested API supports cuda::std::complex<T> mixed with cuda::complex<T>. State that both operands may use different supported complex implementations with the same scalar T. As per path instructions, documentation changes must be technically accurate and API-consistent.

Also applies to: 70-70

Source: Path instructions


} // namespace cuda

Expand Down Expand Up @@ -68,7 +68,6 @@ Defined in ``<cuda/numeric>`` header.

- Scalar overloads require ``lhs``, ``rhs``, ``absolute_tol`` to have the same arithmetic type (integer or floating point).
- Complex overloads accept ``cuda::std::complex<T>`` and ``std::complex<T>`` operands.
- ``AbsTol`` must be the same type as the complex value type.

**Special values**

Expand Down
60 changes: 27 additions & 33 deletions libcudacxx/include/cuda/__numeric/isclose.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -89,27 +89,27 @@ __isclose_fp_impl(const _Tp __lhs, const _Tp __rhs, const float __rel_tol, const
return __diff <= ::cuda::std::fmax(__abs_tol, __rel_value);
}

template <typename _ComplexType, typename _AbsTol>
template <class _Tp, template <class> class _ComplexType>
[[nodiscard]] _CCCL_HOST_DEVICE_API bool __isclose_complex_impl(
const _ComplexType& __lhs, const _ComplexType& __rhs, const float __rel_tol, const _AbsTol __abs_tol) noexcept
const _ComplexType<_Tp>& __lhs, const _ComplexType<_Tp>& __rhs, const float __rel_tol, const _Tp __abs_tol) noexcept
{
using __scalar_t _CCCL_NODEBUG_ALIAS = typename _ComplexType::value_type;
using __compare_t _CCCL_NODEBUG_ALIAS = __isclose_compare_t<__scalar_t>;
static_assert(::cuda::is_floating_point_v<__scalar_t>, "cuda::isclose: __scalar_t must be a floating point type");
static_assert(::cuda::is_floating_point_v<_Tp>, "cuda::isclose: _Tp must be a floating point type");
#if _CCCL_HAS_FLOAT128()
// __float128 is not supported because cuda::std::hypot is not implemented for this type
static_assert(!::cuda::std::is_same_v<__scalar_t, __float128>, "cuda::isclose: __float128 is not supported");
static_assert(!::cuda::std::is_same_v<_Tp, __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},
_CCCL_ASSERT(::cuda::std::isfinite(__abs_tol) && __abs_tol >= _Tp{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);
using __compare_t _CCCL_NODEBUG_ALIAS = __isclose_compare_t<_Tp>;
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)
{
Expand Down Expand Up @@ -279,24 +279,18 @@ _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> || ::cuda::is_floating_poin

// Complex overloads

template <typename _Tp, typename _AbsTol, bool = __is_any_complex_v<_Tp>>
inline constexpr bool __isclose_complex_comparison_v = false;

template <typename _Tp, typename _AbsTol>
inline constexpr bool __isclose_complex_comparison_v<_Tp, _AbsTol, true> =
::cuda::std::is_same_v<typename _Tp::value_type, _AbsTol>;

//! @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
_CCCL_TEMPLATE(class _Tp, template <class> class _ComplexType)
_CCCL_REQUIRES(__is_any_complex_v<_ComplexType<_Tp>> _CCCL_AND(
::cuda::std::__cccl_is_integer_v<_Tp> || ::cuda::is_floating_point_v<_Tp>))
[[nodiscard]] _CCCL_HOST_DEVICE_API bool isclose(
const _ComplexType<_Tp>& __lhs, const _ComplexType<_Tp>& __rhs, const float __rel_tol, const _Tp __abs_tol) noexcept
{
return ::cuda::__isclose_complex_impl(__lhs, __rhs, __rel_tol, __abs_tol);
}
Expand All @@ -307,26 +301,26 @@ isclose(const _ComplexType& __lhs, const _ComplexType& __rhs, const float __rel_
//! @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>)
_CCCL_TEMPLATE(class _Tp, template <class> class _ComplexType)
_CCCL_REQUIRES(__is_any_complex_v<_ComplexType<_Tp>> _CCCL_AND(
::cuda::std::__cccl_is_integer_v<_Tp> || ::cuda::is_floating_point_v<_Tp>))
[[nodiscard]] _CCCL_HOST_DEVICE_API bool
isclose(const _ComplexType& __lhs, const _ComplexType& __rhs, const float __rel_tol) noexcept
isclose(const _ComplexType<_Tp>& __lhs, const _ComplexType<_Tp>& __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});
return ::cuda::isclose(__lhs, __rhs, __rel_tol, _Tp{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
_CCCL_TEMPLATE(class _Tp, template <class> class _ComplexType)
_CCCL_REQUIRES(__is_any_complex_v<_ComplexType<_Tp>> _CCCL_AND(
::cuda::std::__cccl_is_integer_v<_Tp> || ::cuda::is_floating_point_v<_Tp>))
[[nodiscard]] _CCCL_HOST_DEVICE_API bool isclose(const _ComplexType<_Tp>& __lhs, const _ComplexType<_Tp>& __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});
return ::cuda::isclose(__lhs, __rhs, ::cuda::__isclose_default_relative_tolerance<_Tp>(), _Tp{0});
}

_CCCL_END_NAMESPACE_CUDA
Expand Down
Loading