[Tile] Fix complex interop with extended floating point types - #10550
Merged
Conversation
Contributor
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesThe library now provides host/device-qualified half and bfloat16 math and complex operations, applies execution-check suppression to selected complex operations, and separates extended-floating-point complex conversions. Tests add host/device annotations, update force-tile expectations, and avoid unsupported half/bfloat16 tile instantiations. Complex and math library interfaces
Host/device test coverage
Possibly related PRs
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 2
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 3710ac2e-2c40-49a9-9cc2-021f5f7441ce
📒 Files selected for processing (53)
libcudacxx/include/cuda/__complex/get_real_imag.hlibcudacxx/include/cuda/std/__cmath/logarithms.hlibcudacxx/include/cuda/std/__cmath/modulo.hlibcudacxx/include/cuda/std/__cmath/rounding_functions.hlibcudacxx/include/cuda/std/__complex/complex.hlibcudacxx/include/cuda/std/__complex/nvbf16.hlibcudacxx/include/cuda/std/__complex/nvfp16.hlibcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.assign.pass.cpplibcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.bad_macros.pass.cpplibcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.cons.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/error_functions.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/exponential_functions.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/fdim.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/fma.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/gamma.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/hyperbolic_functions.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/hypot.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/inverse_hyperbolic_functions.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/inverse_trigonometric_functions.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/modulo.pass.cpplibcudacxx/test/libcudacxx/std/numerics/c.math/rounding.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpplibcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpplibcudacxx/test/support/test_macros.h
Comment on lines
+98
to
+116
| 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); |
Contributor
There was a problem hiding this comment.
📐 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/stdRepository: 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/__complexRepository: 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 -nRepository: 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
Comment on lines
+292
to
+306
| [[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; | ||
| } |
Contributor
There was a problem hiding this comment.
📐 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: markconjandprojnoexcept; replacecomplex<__nv_bfloat16>(...)with brace initialization.libcudacxx/include/cuda/std/__complex/nvfp16.h#L292-L305: markconjandprojnoexcept; replacecomplex<__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
Contributor
🥳 CI Workflow Results🟩 Finished in 4h 07m: Pass: 100%/115 | Total: 6d 05h | Max: 4h 03m | Hits: 44%/2435910See results here. |
miscco
enabled auto-merge (squash)
July 30, 2026 16:05
Jacobfaib
approved these changes
Jul 30, 2026
davebayer
pushed a commit
to davebayer/cccl
that referenced
this pull request
Aug 4, 2026
…#10550) * [Tile] Fix complex interop with extended floating point types * [Tile] Mark more math functions as host device
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
While we are at it also mark some of the cmath functions that clipped through as host device