Skip to content

[Tile] Fix complex interop with extended floating point types #10550

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

Merged
miscco merged 2 commits into from
Jul 30, 2026
Merged

[Tile] Fix complex interop with extended floating point types #10550

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0250658
[Tile] Fix complex interop with extended floating point types
miscco Jul 30, 2026
7d54466
[Tile] Mark more math functions as host device
miscco Jul 30, 2026
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
4 changes: 4 additions & 0 deletions libcudacxx/include/cuda/__complex/get_real_imag.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,24 +28,28 @@

_CCCL_BEGIN_NAMESPACE_CUDA

_CCCL_EXEC_CHECK_DISABLE
template <class _Tp>
[[nodiscard]] _CCCL_API constexpr _Tp __get_real(const complex<_Tp>& __c) noexcept
{
return __c.real();
}

_CCCL_EXEC_CHECK_DISABLE
template <class _Tp>
[[nodiscard]] _CCCL_API constexpr _Tp __get_imag(const complex<_Tp>& __c) noexcept
{
return __c.imag();
}

_CCCL_EXEC_CHECK_DISABLE
template <class _Tp>
[[nodiscard]] _CCCL_API constexpr _Tp __get_real(const ::cuda::std::complex<_Tp>& __c) noexcept
{
return __c.real();
}

_CCCL_EXEC_CHECK_DISABLE
template <class _Tp>
[[nodiscard]] _CCCL_API constexpr _Tp __get_imag(const ::cuda::std::complex<_Tp>& __c) noexcept
{
Expand Down
4 changes: 2 additions & 2 deletions libcudacxx/include/cuda/std/__cmath/logarithms.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -201,15 +201,15 @@ template <class _Integer, enable_if_t<is_integral_v<_Integer>, int> = 0>
#endif // _CCCL_HAS_LONG_DOUBLE()

#if _LIBCUDACXX_HAS_NVFP16()
[[nodiscard]] _CCCL_API inline __half log10(__half __x) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half log10(__half __x) noexcept
{
NV_IF_ELSE_TARGET(
NV_PROVIDES_SM_53, (return ::hlog10(__x);), (return __float2half(::cuda::std::log10f(__half2float(__x)));))
}
#endif // _LIBCUDACXX_HAS_NVFP16()

#if _LIBCUDACXX_HAS_NVBF16()
[[nodiscard]] _CCCL_API inline __nv_bfloat16 log10(__nv_bfloat16 __x) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 log10(__nv_bfloat16 __x) noexcept
{
NV_IF_ELSE_TARGET(
NV_IS_DEVICE, (return ::hlog10(__x);), (return __float2bfloat16(::cuda::std::log10f(__bfloat162float(__x)));))
Expand Down
4 changes: 2 additions & 2 deletions libcudacxx/include/cuda/std/__cmath/modulo.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -180,7 +180,7 @@ template <class _A1, class _A2, enable_if_t<is_arithmetic_v<_A1> && is_arithmeti
#endif // _CCCL_HAS_LONG_DOUBLE()

#if _LIBCUDACXX_HAS_NVFP16()
[[nodiscard]] _CCCL_API inline __half modf(__half __x, __half* __y) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half modf(__half __x, __half* __y) noexcept
{
const __half __integral_part = ::cuda::std::trunc(__x);
*__y = __integral_part;
Expand All @@ -191,7 +191,7 @@ template <class _A1, class _A2, enable_if_t<is_arithmetic_v<_A1> && is_arithmeti
#endif // _LIBCUDACXX_HAS_NVFP16()

#if _LIBCUDACXX_HAS_NVBF16()
[[nodiscard]] _CCCL_API inline __nv_bfloat16 modf(__nv_bfloat16 __x, __nv_bfloat16* __y) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 modf(__nv_bfloat16 __x, __nv_bfloat16* __y) noexcept
{
const __nv_bfloat16 __integral_part = ::cuda::std::trunc(__x);
*__y = __integral_part;
Expand Down
8 changes: 4 additions & 4 deletions libcudacxx/include/cuda/std/__cmath/rounding_functions.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -802,14 +802,14 @@ _CCCL_API inline double nexttoward(_Integer __x, long double __y) noexcept
#endif // _CCCL_HAS_LONG_DOUBLE()

#if _LIBCUDACXX_HAS_NVFP16()
[[nodiscard]] _CCCL_API inline __half rint(__half __x) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half rint(__half __x) noexcept
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hrint(__x);), (return __float2half(::cuda::std::rint(__half2float(__x)));))
}
#endif // _LIBCUDACXX_HAS_NVFP16()

#if _LIBCUDACXX_HAS_NVBF16()
[[nodiscard]] _CCCL_API inline __nv_bfloat16 rint(__nv_bfloat16 __x) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 rint(__nv_bfloat16 __x) noexcept
{
NV_IF_ELSE_TARGET(
NV_IS_DEVICE, (return ::hrint(__x);), (return __float2bfloat16(::cuda::std::rint(__bfloat162float(__x)));))
Expand Down Expand Up @@ -953,14 +953,14 @@ template <class _Integer, enable_if_t<is_integral_v<_Integer>, int> = 0>
#endif // _CCCL_HAS_LONG_DOUBLE()

#if _LIBCUDACXX_HAS_NVFP16()
[[nodiscard]] _CCCL_API inline __half trunc(__half __x) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half trunc(__half __x) noexcept
{
NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::htrunc(__x);), (return __float2half(::cuda::std::trunc(__half2float(__x)));))
}
#endif // _LIBCUDACXX_HAS_NVFP16()

#if _LIBCUDACXX_HAS_NVBF16()
[[nodiscard]] _CCCL_API inline __nv_bfloat16 trunc(__nv_bfloat16 __x) noexcept
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 trunc(__nv_bfloat16 __x) noexcept
{
NV_IF_ELSE_TARGET(
NV_IS_DEVICE, (return ::htrunc(__x);), (return __float2bfloat16(::cuda::std::trunc(__bfloat162float(__x)));))
Expand Down
13 changes: 11 additions & 2 deletions libcudacxx/include/cuda/std/__complex/complex.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,35 +95,44 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _LIBCUDACXX_COMPLEX_ALIGNAS complex
_CCCL_HIDE_FROM_ABI constexpr complex& operator=(const complex&) noexcept(is_nothrow_copy_assignable_v<_Tp>) = default;
_CCCL_HIDE_FROM_ABI constexpr complex& operator=(complex&&) noexcept(is_nothrow_move_assignable_v<_Tp>) = default;

template <class _Up, enable_if_t<__cccl_internal::__is_non_narrowing_convertible<_Tp, _Up>::value, int> = 0>
template <class _Up,
enable_if_t<!__is_extended_floating_point_v<_Up>, int> = 0,
enable_if_t<__cccl_internal::__is_non_narrowing_convertible<_Tp, _Up>::value, int> = 0>
_CCCL_API constexpr complex(const complex<_Up>& __c)
: __re_(static_cast<_Tp>(__c.real()))
, __im_(static_cast<_Tp>(__c.imag()))
{}

template <class _Up,
enable_if_t<!__is_extended_floating_point_v<_Up>, int> = 0,
enable_if_t<!__cccl_internal::__is_non_narrowing_convertible<_Tp, _Up>::value, int> = 0,
enable_if_t<is_constructible_v<_Tp, _Up>, int> = 0>
_CCCL_API explicit constexpr complex(const complex<_Up>& __c)
: __re_(static_cast<_Tp>(__c.real()))
, __im_(static_cast<_Tp>(__c.imag()))
{}

template <class _Up, enable_if_t<__is_extended_floating_point_v<_Up>, int> = 0>
_CCCL_HOST_DEVICE_API complex(const complex<_Up>& __c);
Comment on lines +98 to +116

@coderabbitai coderabbitai Bot Jul 30, 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.

📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win

🧩 Analysis chain

🏁 Script executed:

#!/usr/bin/env bash
set -euo pipefail

rg -n -C2 '_CCCL_TEMPLATE|_CCCL_REQUIRES' libcudacxx/include/cuda/std/__complex
rg -n -C2 '__is_extended_floating_point_v' libcudacxx/include/cuda/std

Repository: NVIDIA/cccl

Length of output: 16671

🏁 Script executed:

#!/usr/bin/env bash
set -euo pipefail

sed -n '1,180p' libcudacxx/include/cuda/std/__complex/complex.h
printf '\n---- nearby complex overload patterns ----\n'
rg -n -C2 '_CCCL_TEMPLATE|_CCCL_REQUIRES|enable_if_t<.*__is_extended_floating_point_v' libcudacxx/include/cuda/std/__complex

Repository: NVIDIA/cccl

Length of output: 13902

🏁 Script executed:

#!/usr/bin/env bash
set -euo pipefail

rg -n -C3 '_CCCL_TEMPLATE\(class _[A-Za-z0-9_]+\).*constructor|_CCCL_REQUIRES.*complex<|_CCCL_HOST_DEVICE_API.*complex<|_CCCL_TEMPLATE\(class _Up\)' libcudacxx/include/cuda/std | sed -n '1,220p'

Repository: NVIDIA/cccl

Length of output: 20892

🏁 Script executed:

#!/usr/bin/env bash
set -euo pipefail

sed -n '1,170p' libcudacxx/include/cuda/std/__complex/complex.h | cat -n

Repository: NVIDIA/cccl

Length of output: 7610

important: Replace these overload constraints with _CCCL_TEMPLATE/_CCCL_REQUIRES to match the rest of libcudacxx/include/cuda/std/__complex/complex.h; keep the implicit/explicit split and apply the same change to the assignment overloads below.

Source: Coding guidelines


_CCCL_API constexpr complex& operator=(const value_type& __re)
{
__re_ = __re;
__im_ = value_type();
return *this;
}

template <class _Up>
template <class _Up, enable_if_t<!__is_extended_floating_point_v<_Up>, int> = 0>
_CCCL_API constexpr complex& operator=(const complex<_Up>& __c)
{
__re_ = __c.real();
__im_ = __c.imag();
return *this;
}

template <class _Up, enable_if_t<__is_extended_floating_point_v<_Up>, int> = 0>
_CCCL_HOST_DEVICE_API constexpr complex& operator=(const complex<_Up>& __c);

#if _CCCL_HOSTED()
template <class _Up>
_CCCL_API inline complex(const ::std::complex<_Up>& __other)
Expand Down
33 changes: 31 additions & 2 deletions libcudacxx/include/cuda/std/__complex/nvbf16.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,11 +120,15 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__nv_bfloat162)) compl
{}

# if !_CCCL_COMPILER(GCC, <, 10) // Old GCC considers those as deleted
_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex(const complex&) noexcept = default;
_CCCL_HIDE_FROM_ABI complex(complex&&) noexcept = default;
_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex(complex&&) noexcept = default;

_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex& operator=(const complex&) noexcept = default;
_CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept = default;
_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept = default;
# endif // !_CCCL_COMPILER(GCC, <, 10)

template <class _Up, enable_if_t<__cccl_internal::__is_non_narrowing_convertible<value_type, _Up>::value, int> = 0>
Expand Down Expand Up @@ -276,6 +280,31 @@ _CCCL_HOST_DEVICE_API inline complex<double>& complex<double>::operator=(const c
return *this;
}

[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 real(const complex<__nv_bfloat16>& __c) noexcept
{
return __c.real();
}
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 imag(const complex<__nv_bfloat16>& __c) noexcept
{
return __c.imag();
}

[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__nv_bfloat16> conj(const complex<__nv_bfloat16>& __c)
{
return complex<__nv_bfloat16>(__c.real(), ::__hneg(__c.imag()));
}

[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__nv_bfloat16> proj(const complex<__nv_bfloat16>& __c)
{
complex<__nv_bfloat16> __r = __c;
if (::cuda::std::isinf(__c.real()) || ::cuda::std::isinf(__c.imag()))
{
__r = complex<__nv_bfloat16>(
numeric_limits<__nv_bfloat16>::infinity(), ::cuda::std::copysign(::__float2bfloat16(0.0f), __c.imag()));
}
return __r;
}
Comment on lines +292 to +306

@coderabbitai coderabbitai Bot Jul 30, 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.

📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win

important: Make these non-throwing overloads noexcept and use brace initialization for constructed complex values.

  • libcudacxx/include/cuda/std/__complex/nvbf16.h#L292-L306: mark conj and proj noexcept; replace complex<__nv_bfloat16>(...) with brace initialization.
  • libcudacxx/include/cuda/std/__complex/nvfp16.h#L292-L305: mark conj and proj noexcept; replace complex<__half>(...) with brace initialization.

As per coding guidelines, “Functions that do not throw exceptions must be marked noexcept” and “Use uniform initialization for class constructors.”

📍 Affects 2 files
  • libcudacxx/include/cuda/std/__complex/nvbf16.h#L292-L306 (this comment)
  • libcudacxx/include/cuda/std/__complex/nvfp16.h#L292-L305

Source: Coding guidelines


template <>
struct __get_complex_impl<__nv_bfloat16>
{
Expand Down
32 changes: 30 additions & 2 deletions libcudacxx/include/cuda/std/__complex/nvfp16.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -120,11 +120,15 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__half2)) complex<__ha
{}

# if !_CCCL_COMPILER(GCC, <, 10) // Old GCC considers those as deleted
_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex(const complex&) noexcept = default;
_CCCL_HIDE_FROM_ABI complex(complex&&) noexcept = default;
_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex(complex&&) noexcept = default;

_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex& operator=(const complex&) noexcept = default;
_CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept = default;
_CCCL_EXEC_CHECK_DISABLE
_CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept = default;
# endif // !_CCCL_COMPILER(GCC, <, 10)

template <class _Up, enable_if_t<__cccl_internal::__is_non_narrowing_convertible<value_type, _Up>::value, int> = 0>
Expand Down Expand Up @@ -276,6 +280,30 @@ _CCCL_HOST_DEVICE_API inline complex<double>& complex<double>::operator=(const c
return *this;
}

[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half real(const complex<__half>& __c) noexcept
{
return __c.real();
}
[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half imag(const complex<__half>& __c) noexcept
{
return __c.imag();
}

[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__half> conj(const complex<__half>& __c)
{
return complex<__half>(__c.real(), ::__hneg(__c.imag()));
}

[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__half> proj(const complex<__half>& __c)
{
complex<__half> __r = __c;
if (::cuda::std::isinf(__c.real()) || ::cuda::std::isinf(__c.imag()))
{
__r = complex<__half>(numeric_limits<__half>::infinity(), ::cuda::std::copysign(::__float2half(0.0f), __c.imag()));
}
return __r;
}

template <>
struct __get_complex_impl<__half>
{
Expand Down
4 changes: 2 additions & 2 deletions libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.assign.pass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@
#include "test_macros.h"

template <class T, class U>
TEST_FUNC void test_assignment()
TEST_HOST_DEVICE_FUNC void test_assignment()
{
cuda::std::complex<T> from_only_real{static_cast<T>(-1.0), static_cast<T>(1.0)};
cuda::std::complex<T> from_only_imag{static_cast<T>(-1.0), static_cast<T>(1.0)};
Expand All @@ -38,7 +38,7 @@ TEST_FUNC void test_assignment()
assert(from_real_imag.imag() == static_cast<T>(112.0));
}

TEST_FUNC void test()
TEST_HOST_DEVICE_FUNC void test()
{
#if _LIBCUDACXX_HAS_NVFP16()
test_assignment<__half, float>();
Expand Down
4 changes: 2 additions & 2 deletions libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.bad_macros.pass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,15 +20,15 @@
#include "test_macros.h"

template <class T, class U>
TEST_FUNC void test_assignment(cuda::std::complex<U> v = {})
TEST_HOST_DEVICE_FUNC void test_assignment(cuda::std::complex<U> v = {})
{
cuda::std::complex<T> converting(v);

cuda::std::complex<T> assigning{};
assigning = v;
}

TEST_FUNC void test()
TEST_HOST_DEVICE_FUNC void test()
{
#if _LIBCUDACXX_HAS_NVFP16()
test_assignment<__half, float>();
Expand Down
4 changes: 2 additions & 2 deletions libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.cons.pass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@
#include "test_macros.h"

template <class T, class U>
TEST_FUNC void test_construction()
TEST_HOST_DEVICE_FUNC void test_construction()
{
const cuda::std::complex<U> only_real{static_cast<U>(42.0), static_cast<U>(0.0)};
const cuda::std::complex<U> only_imag{static_cast<U>(0.0), static_cast<U>(42.0)};
Expand All @@ -34,7 +34,7 @@ TEST_FUNC void test_construction()
assert(from_real_imag.imag() == static_cast<T>(112.0));
}

TEST_FUNC void test()
TEST_HOST_DEVICE_FUNC void test()
{
#if _LIBCUDACXX_HAS_NVFP16()
test_construction<__half, float>();
Expand Down
12 changes: 6 additions & 6 deletions libcudacxx/test/libcudacxx/std/numerics/c.math/error_functions.pass.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,8 +7,8 @@
//
//===----------------------------------------------------------------------===//

// XFAIL: enable-tile
// nvbug6077402: error: "call to non-tile function not supported!"
// UNSUPPORTED: force-tile
// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed

// <cmath>

Expand All @@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4244) // conversion from 'const double' to 'int', possib
TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned

template <typename T>
TEST_FUNC void test_erf(T val)
TEST_HOST_DEVICE_FUNC void test_erf(T val)
{
using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
static_assert(cuda::std::is_same_v<decltype(cuda::std::erf(T{})), ret>);
Expand Down Expand Up @@ -75,7 +75,7 @@ TEST_FUNC void test_erf(T val)
}

template <typename T>
TEST_FUNC void test_erfc(T val)
TEST_HOST_DEVICE_FUNC void test_erfc(T val)
{
using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
static_assert(cuda::std::is_same_v<decltype(cuda::std::erfc(T{})), ret>);
Expand Down Expand Up @@ -130,13 +130,13 @@ TEST_FUNC void test_erfc(T val)
}

template <typename T>
TEST_FUNC void test(const T val)
TEST_HOST_DEVICE_FUNC void test(const T val)
{
test_erf<T>(val);
test_erfc<T>(val);
}

TEST_FUNC void test(const float val)
TEST_HOST_DEVICE_FUNC void test(const float val)
{
test<float>(val);
test<double>(val);
Expand Down
Loading
Loading