Skip to content

[libcu++] Make constant_wrapper work with nvrtc and without warnings suppressions - #10709

Open
davebayer wants to merge 1 commit into
NVIDIA:mainfrom
davebayer:constant_wrapper_nvrtc
Open

[libcu++] Make constant_wrapper work with nvrtc and without warnings suppressions#10709
davebayer wants to merge 1 commit into
NVIDIA:mainfrom
davebayer:constant_wrapper_nvrtc

Conversation

@davebayer

@davebayer davebayer commented Aug 7, 2026

Copy link
Copy Markdown
Contributor

This PR enables constant_wrapper for nvrtc and removes warnings suppression from the implementation. There are some limitations though. For some reason nvrtc doesn't like the static constexpr value member being instantiated by any non-trivial types without -default-device flag. So basically it has the same limitations as C++17 mode.

Removing warnings suppressions regarding accessing __host__ variable from __device__/__global__ functions have very similar effect. I found a workaround with a __get() static member function that returns the reference to the template parameter instead.

If we merge this PR, we can make constant_wrapper public. There will be just some compiler limitations, which are not fixable from our side.

@davebayer
davebayer requested a review from a team as a code owner August 7, 2026 10:26
@davebayer
davebayer requested a review from fbusato August 7, 2026 10:26
@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
@coderabbitai

coderabbitai Bot commented Aug 7, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Summary by CodeRabbit

  • Bug Fixes

    • Improved constant_wrapper behavior and compatibility when compiling with NVRTC.
    • Added consistent access to wrapped values through a public accessor.
    • Improved support for callable objects, indexing, conversions, comparisons, and other operators.
    • Added support for both single-argument and multi-argument subscript operations where available.
  • Tests

    • Updated coverage for constant-wrapper operations and documented current NVRTC limitations and required workarounds.

Walkthrough

Changes

The constant-wrapper implementation adds __get() and routes callable and subscript operations through it. Tests replace direct internal-value access and add targeted NVRTC exclusions for non-trivial static data member limitations.

Constant wrapper implementation

Layer / File(s) Summary
Wrapper accessor and operator implementation
libcudacxx/include/cuda/std/__utility/constant_wrapper.h
__constant_wrapper adds __get(). Conversion, call, and subscript operations use the accessor. NVRTC-specific implementation guards are removed.

Constant wrapper tests

Layer / File(s) Summary
Test migration to __get()
libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/{assign,binary_ops,convert,pseudo_mutators,types.compile,unary_ops}.pass.cpp
Operation and conversion assertions access wrapped values through __get() or temporary wrapper objects.
NVRTC test compatibility guards
libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/*.pass.cpp
NVRTC diagnostics now describe static-data-member limitations. Affected C++20 test sections are excluded, while host-launched kernel requirements remain documented.

Possibly related PRs

  • NVIDIA/cccl#9702: Both changes modify __constant_wrapper and its tests. This PR adds __get() and NVRTC handling, while #9702 changes the underlying representation and removes string support.

Suggested reviewers: ericniebler, miscco, jacobfaib


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.

🧹 Nitpick comments (2)
libcudacxx/include/cuda/std/__utility/constant_wrapper.h (2)

368-371: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Add [[nodiscard]] to __get(). It returns a reference and has no side effect. As per coding guidelines, “Most functions with a non-void return type should use [[nodiscard]], except for functions with known side effects.”

Sources: Coding guidelines, Path instructions


323-345: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Remove negated conditions from annotated #else and #endif comments. Repeat the original #if condition exactly.

  • libcudacxx/include/cuda/std/__utility/constant_wrapper.h#L323-L345: Update the comments on Lines 329, 334, 342, and 345 to use only _CCCL_HAS_MULTIARG_OPERATOR_BRACKETS().
  • libcudacxx/include/cuda/std/__utility/constant_wrapper.h#L373-L457: Update the comments on Lines 378, 380, 391, 393, 407, 421, 428, 436, 450, and 457 to use only their corresponding original condition.

Based on learnings, annotated #else and #endif comments must repeat the exact condition text and must not use an inverted condition.

Source: Learnings


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 5ba5c1c1-020d-43cd-b7cc-112693135d75

📥 Commits

Reviewing files that changed from the base of the PR and between 10d2278 and 3cb16a7.

📒 Files selected for processing (17)
  • libcudacxx/include/cuda/std/__utility/constant_wrapper.h
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/adl.compile.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/assign.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/binary_ops.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/call.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/comma.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/comp.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/convert.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/general.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/kernel_parameter.array.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/kernel_parameter.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/mem_ptr.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/pseudo_mutators.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/subscript.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/types.compile.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/unary_ops.pass.cpp
💤 Files with no reviewable changes (1)
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/general.pass.cpp

@miscco miscco 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.

I love me my code crimes

@github-actions

github-actions Bot commented Aug 7, 2026

Copy link
Copy Markdown
Contributor

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

Result: 2 regression row(s), 4 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

TU total compilation

-f total-compilation inclusive --sort total

🔴 TU total compilation — Regressions
Rank Regression impact Selected Δ Baseline Current Event Matched traces
1 5.703105 5.703105 14.191645 19.894750 Total Compilation Time: thrust/mr/tls_pool.h 1
🟢 TU total compilation — Improvements
Rank Improvement impact Selected Δ Baseline Current Event Matched traces
1 3.843993 -3.843993 11.875494 8.031501 Total Compilation Time: cub/device/dispatch/dispatch_three_way_partition.cuh 1

Direct file processing

-f file-processing exclusive --sort total

🔴 Direct file processing — Regressions
Rank Regression impact Selected Δ Baseline Current Event Matched traces
1 0.753608 0.753608 22.557706 23.311314 Processing Header File: libcudacxx/include/cuda/__device/physical_device.h 94
🟢 Direct file processing — Improvements
Rank Improvement impact Selected Δ Baseline Current Event Matched traces
1 1.190352 -1.190352 10.221910 9.031558 Processing Header File: libcudacxx/include/cuda/std/__cccl/prologue.h 550
2 0.514571 -0.514571 3.742839 3.228268 Processing Header File: libcudacxx/include/cuda/std/__cccl/epilogue.h 550
3 0.238853 -0.238853 6.905754 6.666901 Processing Header File: libcudacxx/include/cuda/std/__tuple_dir/vector_types.h 238

@github-actions

github-actions Bot commented Aug 7, 2026

Copy link
Copy Markdown
Contributor

😬 CI Workflow Results

🟥 Finished in 3h 21m: Pass: 91%/115 | Total: 4d 23h | Max: 3h 20m | Hits: 51%/1107006

See results here.

AI failure analysis

1. `constant_wrapper` conversion no longer constant-evaluates under MSVC · 5 jobs

Explanation: Five nvcc/MSVC configurations reject `const.wrap.class` static assertions after the wrapper conversion and accessor return types changed from `decltype(value)` to `decltype((_Xp))`. This breaks compile-time use of otherwise correctly typed arithmetic results and conversions.

Evidence:

libcu++ nvcc MSVC / [CTK12.0 MSVC14.39 C++20] Build(amd64), step 4

2026-08-07T12:09:58.8343756Z ******************** TEST 'libcu++ :: std/utilities/const.wrap.class/binary_ops.pass.cpp' FAILED ********************
2026-08-07T12:09:58.8505876Z C:\cccl\libcudacxx\test\libcudacxx\std\utilities\const.wrap.class\binary_ops.pass.cpp(444): error C2131: expression did not evaluate to a constant

libcu++ nvcc MSVC / [CTK13.3 MSVC14.44 C++20] Build(amd64), step 4

2026-08-07T13:02:43.1344949Z ******************** TEST 'libcu++ :: std/utilities/const.wrap.class/binary_ops.pass.cpp' FAILED ********************

Root cause: The PR changed both the conversion operator and new `__get()` accessor to return `decltype((_Xp))` directly. nvcc-generated MSVC host code does not preserve the prior constant-evaluation behavior for these expressions, while the original `decltype(value)` return type encoded the wrapper's intended value-versus-reference semantics. Sources: libcudacxx/include/cuda/std/__utility/constant_wrapper.h:363, libcudacxx/include/cuda/std/__utility/constant_wrapper.h:368, libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/binary_ops.pass.cpp:444, libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/convert.pass.cpp:99.

Suggested next steps: Restore `decltype(value)` as the return type of the conversion operator and `__get()`, while continuing to return `(_Xp)` so device code does not read the static member. Run the focused `const.wrap.class` tests with one nvcc/MSVC configuration before rechecking the remaining Windows matrix entries.

Copy this prompt into a coding agent
Repository: https://github.com/NVIDIA/cccl
Workflow run: https://github.com/NVIDIA/cccl/actions/runs/31170052987
Failure group: `constant_wrapper` conversion no longer constant-evaluates under MSVC
Affected jobs:
- libcu++ nvcc MSVC / [CTK12.0 MSVC14.39 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751184
- libcu++ nvcc MSVC / [CTK13.3 MSVC14.44 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751191
- libcu++ nvcc MSVC / [CTK13.0 MSVC14.44 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751207
- libcu++ nvcc MSVC / [CTK12.9 MSVC14.44 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751230
- libcu++ nvcc MSVC / [CTK13.3 MSVC14.50 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751404

Investigate the nvcc/MSVC constant-evaluation regression in `libcudacxx/include/cuda/std/__utility/constant_wrapper.h`. Reproduce narrowly with `std/utilities/const.wrap.class/binary_ops.pass.cpp` or `convert.pass.cpp`, then change the conversion operator and `__get()` return types from `decltype((_Xp))` to `decltype(value)` while retaining `return (_Xp);`; verify that scalar wrappers remain values, class NTTP wrappers remain references, and the NVRTC-enablement goal is preserved. Run focused libcu++ validation under an available nvcc/MSVC configuration and adjust tests only if required by the intended API semantics.

Jobs:

2. Address-of test takes the wrapper temporary instead of its value · 4 jobs

Explanation: Four nvcc/GCC configurations fail the same newly modified identity test because `&cuda::std::__constant_wrapper<V>{}` applies built-in address-of to a temporary wrapper rather than obtaining the address of the wrapped template-parameter object. Compilation stops before the static assertion can be evaluated.

Evidence:

libcu++ nvcc GCC / [CTK13.0 GCC11 C++20] Build(amd64), step 4

2026-08-07T11:36:41.1202650Z ******************** TEST 'libcu++ :: std/utilities/const.wrap.class/types.compile.pass.cpp' FAILED ********************
2026-08-07T11:36:41.1229535Z /home/coder/cccl/libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/types.compile.pass.cpp:61:17: error: taking address of rvalue [-fpermissive]

libcu++ nvcc GCC / [CTK13.3 GCC11 C++20] Build(amd64), step 4

2026-08-07T11:46:00.6755663Z ******************** TEST 'libcu++ :: std/utilities/const.wrap.class/types.compile.pass.cpp' FAILED ********************

Root cause: The PR replaced `&cuda::std::__constant_wrapper<V>::value` with `&cuda::std::__constant_wrapper<V>{}`. GCC resolves this as the address of the temporary wrapper, producing both an rvalue-address error and a pointer-type mismatch instead of testing the wrapped object's identity. Sources: libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/types.compile.pass.cpp:61, libcudacxx/include/cuda/std/__utility/constant_wrapper.h:368.

Suggested next steps: Change the comparison to take the address of the accessor result, for example `&V == &cuda::std::__constant_wrapper<V>::__get()`, or restore the original `.value` expression if static-member access is acceptable. Compile only `types.compile.pass.cpp` with nvcc/GCC first, then run the focused `const.wrap.class` suite.

Copy this prompt into a coding agent
Repository: https://github.com/NVIDIA/cccl
Workflow run: https://github.com/NVIDIA/cccl/actions/runs/31170052987
Failure group: Address-of test takes the wrapper temporary instead of its value
Affected jobs:
- libcu++ nvcc GCC / [CTK13.0 GCC11 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751581
- libcu++ nvcc GCC / [CTK13.3 GCC11 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751618
- libcu++ nvcc GCC / [CTK12.0 GCC12 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751634
- libcu++ nvcc GCC / [CTK13.3 GCC12 C++20] Build(amd64): https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751703

Fix the invalid address comparison in `libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/types.compile.pass.cpp`. Reproduce the GCC diagnostic at line 61, replace the address of the temporary wrapper with the address of its wrapped object—prefer `return &V == &cuda::std::__constant_wrapper<V>::__get();` to exercise the new accessor without reading the static member directly—and verify the static assertion under focused nvcc/GCC testing. Also confirm the accessor returns a stable `const S&` to the class template-parameter object and coordinate the change with any return-type fix in `constant_wrapper.h`.

Jobs:

3. Windows CUB job could not acquire AWS cache credentials · 1 job

Explanation: The CUB job failed before fetching its container or starting the build because the pinned AWS credentials action ended unsuccessfully. The saved log contains no AWS diagnostic, so the exact OIDC, IAM, or action failure cannot be determined from this run.

Evidence:

CUB nvcc MSVC / [CTK13.3 MSVC14.50 C++20] BuildNoLaunch(amd64): sm{86}, step 4

2026-08-07T10:28:22.4399995Z ##[start-action display=Get AWS credentials for sccache bucket;id=__self.__aws-actions_configure-aws-credentials]
2026-08-07T10:28:22.4539455Z ##[group]Run aws-actions/configure-aws-credentials@254c19bd240aabef8777f48595e9d2d7b972184b
2026-08-07T10:28:23.0246254Z ##[end-action id=__self.__aws-actions_configure-aws-credentials;outcome=failure;conclusion=failure;duration_ms=584]

Root cause: Credential setup failed inside `aws-actions/configure-aws-credentials`, but the action emitted no saved error message. Because other jobs using the same role proceeded, this is most consistent with a transient runner/OIDC action failure, though missing debug output prevents confirmation. Sources: .github/actions/workflow-run-job-windows/action.yml:89.

Suggested next steps: Rerun only this CUB matrix job. If it fails again, enable action and runner debug logging and verify the workflow's `id-token: write` permission, OIDC token availability, and IAM trust conditions before changing build code.

Copy this prompt into a coding agent
Repository: https://github.com/NVIDIA/cccl
Workflow run: https://github.com/NVIDIA/cccl/actions/runs/31170052987
Failure group: Windows CUB job could not acquire AWS cache credentials
Affected jobs:
- CUB nvcc MSVC / [CTK13.3 MSVC14.50 C++20] BuildNoLaunch(amd64): sm{86}: https://github.com/NVIDIA/cccl/actions/runs/31170052987/job/92839751069

Investigate the pre-build AWS credential failure in the Windows CUB job. First rerun or reproduce only the affected matrix entry and capture debug output from `aws-actions/configure-aws-credentials` in `.github/actions/workflow-run-job-windows/action.yml`; verify `id-token: write`, OIDC token issuance, the configured role trust policy, and whether the pinned action behaves correctly on this Windows runner. Do not modify CUB sources; implement a workflow fix only if the credential failure is reproducible, then validate credential acquisition and the subsequent container-fetch step.

Jobs:

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: In Review

Development

Successfully merging this pull request may close these issues.

2 participants