Skip to content

[Tile] Actually rewrite complex for extended floating point types once more - #10717

Merged
miscco merged 1 commit into
NVIDIA:mainfrom
miscco:tile_complex_once_more
Aug 12, 2026
Merged

[Tile] Actually rewrite complex for extended floating point types once more#10717
miscco merged 1 commit into
NVIDIA:mainfrom
miscco:tile_complex_once_more

Conversation

@miscco

@miscco miscco commented Aug 7, 2026

Copy link
Copy Markdown
Contributor

This ensures that it works fine, although it has a minor perf regression

@miscco
miscco requested a review from a team as a code owner August 7, 2026 15:37
@miscco
miscco requested a review from pciolkosz August 7, 2026 15:37
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Aug 7, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Aug 7, 2026
@miscco
miscco force-pushed the tile_complex_once_more branch from 0f8a5b3 to 5079a40 Compare August 7, 2026 15:38
@coderabbitai

coderabbitai Bot commented Aug 7, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 6da3f4b6-b68c-4d73-871c-0a381932fc2e

📥 Commits

Reviewing files that changed from the base of the PR and between 8548977 and b8b8c7c.

📒 Files selected for processing (80)
  • libcudacxx/include/cuda/std/__cccl/extended_data_types.h
  • libcudacxx/include/cuda/std/__complex/complex.h
  • libcudacxx/include/cuda/std/__complex/nvbf16.h
  • libcudacxx/include/cuda/std/__complex/nvfp16.h
  • libcudacxx/test/libcudacxx/std/numerics/bit/bit.cast/bit_cast.trivially_copyable.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cases.h
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/conj.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/real.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/assignment_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/assignment_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/minus_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/minus_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/plus_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/plus_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/times_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/times_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.members/construct.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.members/real_imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_equals_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_minus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_minus_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_not_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_not_equals_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_plus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_plus_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_times_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_times_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_minus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_not_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_plus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_times_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/stream_input.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/stream_output.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/unary_minus.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/unary_plus.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/get.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/tuple_element_compiles.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/tuple_size_compiles.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/conj.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/real.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/abi_latest.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/traits.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/types.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/layout.pass.cpp
🚧 Files skipped from review as they are similar to previous changes (75)
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_times_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_minus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/unary_minus.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_not_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/real.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/minus_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/layout.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_minus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/assignment_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_minus_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/traits.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/minus_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_equals_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.members/construct.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_plus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.members/real_imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/tuple_element_compiles.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/times_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cases.h
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/times_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_plus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_not_equals_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/plus_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_times_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_not_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/plus_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpp
  • libcudacxx/include/cuda/std/__cccl/extended_data_types.h
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/get.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/unary_plus.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/abi_latest.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/tuple_size_compiles.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/types.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_plus_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/real.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/conj.pass.cpp
  • libcudacxx/include/cuda/std/__complex/complex.h
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_times_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/conj.pass.cpp
  • libcudacxx/include/cuda/std/__complex/nvfp16.h
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpp

📝 Walkthrough

Summary by CodeRabbit

  • New Features
    • Improved complex-number support for half-precision and bfloat16 values during tile compilation.
    • Added broader cross-type construction and assignment support for extended floating-point complex values.
    • Improved tile-compatible arithmetic and component-wise operations for half and bfloat16 complex numbers.
  • Bug Fixes
    • Corrected scalar complex division assignment availability across compilation targets.
  • Tests
    • Expanded complex arithmetic, transcendental, conversion, tuple, layout, and bit-casting coverage for tile compilation and supported extended floating-point types.

Walkthrough

Extended floating-point detection and cuda::std::complex half and bfloat16 support now include tile compilation. Tile-specific representations and component-wise operations support these builds. Tests for bit casts, types, operators, values, and transcendental functions now run when the feature macros are enabled.

Changes

Extended complex tile support

Layer / File(s) Summary
Complex representations and APIs
libcudacxx/include/cuda/std/__cccl/extended_data_types.h, libcudacxx/include/cuda/std/__complex/complex.h, libcudacxx/include/cuda/std/__complex/nvbf16.h, libcudacxx/include/cuda/std/__complex/nvfp16.h
Float128 detection no longer directly excludes tile compilation. Half and bfloat16 complex types use tile-specific representations and component-wise operations. Cross-type constraints and API annotations were updated.
Type and bit-cast coverage
libcudacxx/test/libcudacxx/std/numerics/bit/bit.cast/bit_cast.trivially_copyable.pass.cpp, libcudacxx/test/libcudacxx/std/numerics/complex.number/cases.h, libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/*.pass.cpp, libcudacxx/test/libcudacxx/std/numerics/complex.number/layout.pass.cpp
Bit-cast, comparison, layout, trait, ABI, and type tests no longer exclude tile compilation.
Complex operators and accessors
libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/*, libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/*, libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.members/*, libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/*, libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/*
Half and bfloat16 tests for construction, assignment, arithmetic, comparisons, streams, accessors, and tuple operations now run during tile compilation.
Complex math coverage
libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/*, libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/*
Half and bfloat16 value and transcendental tests, including edge cases and expected-value calculations, no longer exclude tile compilation.

Suggested labels: libcu++

Suggested reviewers: pciolkosz, gevtushenko


Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

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.

Actionable comments posted: 2

🧹 Nitpick comments (1)
libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/plus_equal_scalar.pass.cpp (1)

45-50: 🎯 Functional Correctness | 🔵 Trivial

suggestion: Run the affected tile tests. Use the relevant CCCL CMake preset and targeted test targets to compile and execute the new __half and __nv_bfloat16 instantiations. The supplied context does not include build or test results. As per coding guidelines, validate changes with targeted builds and tests using the provided CMake presets. As per path instructions, use .agent/skills/cccl-test/SKILL.md and its routed references for libcudacxx complex-number tests.

Sources: Coding guidelines, Path instructions


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: cf1d5208-a708-426c-864b-2362370af637

📥 Commits

Reviewing files that changed from the base of the PR and between 0c7c3ec and 5079a40.

📒 Files selected for processing (80)
  • libcudacxx/include/cuda/std/__cccl/extended_data_types.h
  • libcudacxx/include/cuda/std/__complex/complex.h
  • libcudacxx/include/cuda/std/__complex/nvbf16.h
  • libcudacxx/include/cuda/std/__complex/nvfp16.h
  • libcudacxx/test/libcudacxx/std/numerics/bit/bit.cast/bit_cast.trivially_copyable.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cases.h
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/conj.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/real.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/assignment_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/assignment_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/minus_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/minus_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/plus_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/plus_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/times_equal_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/times_equal_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.members/construct.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.members/real_imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_equals_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_minus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_minus_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_not_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_not_equals_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_plus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_plus_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_times_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_times_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_minus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_not_equals_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_plus_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_times_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/stream_input.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/stream_output.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/unary_minus.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/unary_plus.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/get.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/tuple_element_compiles.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.tuple/tuple_size_compiles.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/conj.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/imag.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/real.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/abi_latest.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/traits.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/complex/types.pass.cpp
  • libcudacxx/test/libcudacxx/std/numerics/complex.number/layout.pass.cpp

Comment thread libcudacxx/include/cuda/std/__complex/nvbf16.h
Comment thread libcudacxx/include/cuda/std/__complex/nvbf16.h
@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@griwes griwes left a comment

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.

The question below, and the direct code suggestion, should be resolved before merging. The other suggestion is up to you.

Tests look good, except apparently you've not updated complex.value.opts/proj.pass.cpp; divide_equal_complex.pass.cpp also has an issue, but that one got cleanly caught by CI.

Comment thread libcudacxx/include/cuda/std/__complex/complex.h
__nv_bfloat16 x;
__nv_bfloat16 y;
};
using __complex_nv_bfloat_repr_t = __complex_fake_nv_bfloat162;

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.

Suggestion: I started five separate comments recommending to un-special-case this and all of them failed. I hate it here. Best I can think of is having a separate template<typename T> struct __type_to_storage_vector { using __type = typename __type_to_vector<T>::__type; };, in these tile paths specializing for half and bf16, and then using __type_to_storage_vector throughout instead of the special cases here and in half? But honestly I am not sure if that is at all better. Sigh. Though it does mean we're not typedefing the same name to different types in the different branches.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

I mean there are non different branches, there is either tile mode or not

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.

What I mean is that there's two separate definitions of __complex_nv_bfloat_repr_t, and we could potentially fold it, but this is very much not a blocker.

Comment thread libcudacxx/include/cuda/std/__complex/nvfp16.h Outdated
…e more

This ensures that it works fine, although it has a minor perf regression
@miscco
miscco force-pushed the tile_complex_once_more branch from 5079a40 to b8b8c7c Compare August 11, 2026 09:15
@coderabbitai

coderabbitai Bot commented Aug 11, 2026

Copy link
Copy Markdown
Contributor

Note

GitHub couldn't provide a complete incremental comparison for this pull request, so CodeRabbit is performing a full review instead. This review may take a little longer.

@github-actions

Copy link
Copy Markdown
Contributor

⏱️ CCCL compile-time benchmark comparison: Public headers compile-time bench

Result: 0 regression row(s), 6 improvement row(s) above threshold.

Run Value
Config public-headers-gcc13
Baseline origin/main
Preset all-dev
Targets cub.headers.base, thrust.cpp.cuda.headers.base, libcudacxx.test.public_headers
GPU / launch args rtx2080 / --cuda 13.3 --host gcc13

Artifacts: reports and traces

Direct file processing

-f file-processing exclusive --sort total

🟢 Direct file processing — Improvements
Rank Improvement impact Selected Δ Baseline Current Event Matched traces
1 0.772741 -0.772741 5.267038 4.494297 Processing Header File: libcudacxx/include/cuda/std/__cccl/prologue.h 550
2 0.406431 -0.406431 6.648210 6.241779 Processing Header File: libcudacxx/include/cuda/std/__type_traits/type_list.h 477
3 0.354418 -0.354418 6.285254 5.930836 Processing Header File: libcudacxx/include/cuda/std/__iterator/concepts.h 437
4 0.305156 -0.305156 4.231467 3.926311 Processing Header File: cub/cub/util_type.cuh 279
5 0.302361 -0.302361 1.872687 1.570326 Processing Header File: libcudacxx/include/cuda/std/__cccl/epilogue.h 550
6 0.255213 -0.255213 10.535898 10.280685 Processing Header File: libcudacxx/include/cuda/__device/physical_device.h 94

@github-actions

github-actions Bot commented Aug 11, 2026

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 4h 39m: Pass: 100%/115 | Total: 4d 20h | Max: 4h 38m | Hits: 45%/1280044

See results here.

AI failure analysis

1. RAPIDS environment merges incompatible rapids-logger pins · 2 jobs

Explanation: Both RAPIDS matrices fail before compilation because the generated Conda environment simultaneously requires incompatible `rapids-logger` 0.2 and 0.3 series. The cugraph/wholegraph matrix additionally exposes CUDA 13.3 compatibility constraints while resolving cuML, but its viable candidates still require `rapids-logger` 0.2 and are blocked by the same 0.3 pin.

Evidence:

Build RAPIDS (optional) / rmm ucxx raft cuvs cugraph wholegraph, step 6

2026-08-11T09:19:05.7100753Z error    libmamba Could not solve for environment specs
2026-08-11T09:19:05.7107689Z     ├─ cuda-version =13.3 * is requested and can be installed;
2026-08-11T09:19:05.7126993Z     │  └─ cuml [26.10.00a52|26.10.00a54|26.10.00a55] would require

Root cause: The selected RAPIDS 26.10 development branches contribute dependency files pinned to different major-minor `rapids-logger` series, and the build utility merges them into one unsatisfiable environment. The logs do not attribute each pin to its originating cloned repository, so the exact upstream dependency file must be identified during reproduction; the PR diff only changes libcudacxx code and does not modify this CI configuration. Sources: .github/workflows/build-rapids.yml:67, .github/workflows/build-rapids.yml:68, ci/rapids/post-create-command.sh:38.

Suggested next steps: Regenerate each failing matrix environment while adding one RAPIDS repository at a time to identify which dependency files request `rapids-logger==0.2.*` and `==0.3.*`, then align those repositories on one compatible series or temporarily select mutually compatible branch revisions. Verify both sets with `RAPIDS_LIBS='rmm ucxx raft cuvs nvforest cuml' .devcontainer/launch.sh -d -c 13.3 -H rapids-conda -- ./ci/rapids/rapids-entrypoint.sh` and the equivalent command using `RAPIDS_LIBS='rmm ucxx raft cuvs cugraph wholegraph'`.

Copy this prompt into a coding agent
Repository: https://github.com/NVIDIA/cccl
Workflow run: https://github.com/NVIDIA/cccl/actions/runs/31476880502
Failure group: RAPIDS environment merges incompatible rapids-logger pins
Affected jobs:
- Build RAPIDS (optional) / rmm ucxx raft cuvs cugraph wholegraph: https://github.com/NVIDIA/cccl/actions/runs/31476880502/job/93732688509
- Build RAPIDS (optional) / rmm ucxx raft cuvs nvforest cuml: https://github.com/NVIDIA/cccl/actions/runs/31476880502/job/93732688511

Investigate the RAPIDS Conda environment failure for the CUDA 13.3 matrices `rmm ucxx raft cuvs nvforest cuml` and `rmm ucxx raft cuvs cugraph wholegraph`. Reproduce only environment generation/solving, inspect the cloned repositories' dependency YAML files to identify the sources of `rapids-logger==0.2.*,>=0.0.0a0` and `rapids-logger==0.3.*`, and confirm that a common logger series also leaves a CUDA 13.3-compatible cuML candidate. Implement the smallest fix by aligning dependency pins or selecting compatible RAPIDS branch revisions in the CI configuration; do not change the libcudacxx PR code to mask this dependency failure. Run focused environment-generation validation for both affected `RAPIDS_LIBS` sets and report the files changed and solver result.

Jobs:

@miscco
miscco merged commit ca2e70a into NVIDIA:main Aug 12, 2026
141 of 143 checks passed
@miscco
miscco deleted the tile_complex_once_more branch August 12, 2026 07:10
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Aug 12, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

2 participants