Skip to content

[libcu++] Implement std::constant_wrapper#9046

Open
davebayer wants to merge 13 commits into
NVIDIA:mainfrom
davebayer:libcuxx_constant_wrapper
Open

[libcu++] Implement std::constant_wrapper#9046
davebayer wants to merge 13 commits into
NVIDIA:mainfrom
davebayer:libcuxx_constant_wrapper

Conversation

@davebayer
Copy link
Copy Markdown
Contributor

@davebayer davebayer commented May 17, 2026

This PR implements std::constant_wrapper in cuda::std:: namespace backported to C++20. The implementation is not that complicated but it almost completely breaks nvcc and older compilers.

Fixes #9109.

@davebayer davebayer requested a review from a team as a code owner May 17, 2026 13:32
@davebayer davebayer requested a review from gevtushenko May 17, 2026 13:32
@github-project-automation github-project-automation Bot moved this to Todo in CCCL May 17, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL May 17, 2026
@coderabbitai
Copy link
Copy Markdown
Contributor

coderabbitai Bot commented May 17, 2026

Review Change Stack

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review
📝 Walkthrough

Walkthrough

important: Adds cuda::std::constant_wrapper (implementation, feature gates, public exposure) and a comprehensive test suite exercising construction, call, subscript, operators, pseudo-mutators, and related feature-detection macros.

Changes

important:

Constant Wrapper Feature

Layer / File(s) Summary
Feature detection macros
libcudacxx/include/cuda/std/__cccl/dialect.h
Lowered multi-arg operator[] threshold and added _CCCL_HAS_STATIC_SUBSCRIPT_OPERATOR() and _CCCL_HAS_STATIC_CALL_OPERATOR() with NVCC/NVRTC and CUDA-toolchain guards.
Core wrapper implementation
libcudacxx/include/cuda/std/__utility/constant_wrapper.h
Adds __cw_fixed_value (scalar/array) and deduction guide, __constexpr_param and cw variable template, __cw_operators (consteval unary/binary/comparisons), constant_wrapper template with conversion and consteval assignment, call operator(s), feature-controlled operator[] paths, pseudo-mutators and compound-assignment operators, plus diagnostics and include-guard epilogue.
Public API integration
libcudacxx/include/cuda/std/utility, libcudacxx/include/cuda/std/version
Includes constant_wrapper.h in public utility header and registers __cccl_lib_constant_wrapper feature-test macro.
Test infrastructure and behavior checks
libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/*
Adds helpers and comprehensive tests (ADL, assignment, CTAD, cw instantiation, call/subscript/member-pointer/unary/binary/comparison/pseudo-mutators/comma, fixed-array construction, kernel-parameter, and compile-time/runtime validation).
  • Possibly related PRs:

    • NVIDIA/cccl#8988: Updates multi-arg operator[] feature detection and mdspan tests gated on that macro.
  • Suggested reviewers:

    • gevtushenko
    • griwes
    • fbusato
    • bernhardmgruber

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

Copy link
Copy Markdown
Contributor

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

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

Actionable comments posted: 5


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 651d9cca-a982-4876-9fe3-8e0bed904d0a

📥 Commits

Reviewing files that changed from the base of the PR and between fecc3a3 and 6f8f1e9.

📒 Files selected for processing (22)
  • libcudacxx/include/cuda/std/__cccl/dialect.h
  • libcudacxx/include/cuda/std/__utility/constant_wrapper.h
  • libcudacxx/include/cuda/std/utility
  • libcudacxx/include/cuda/std/version
  • 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/ctad.compile.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw_fixed.array.ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw_fixed.ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/general.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/helpers.h
  • 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

Comment thread libcudacxx/include/cuda/std/__utility/constant_wrapper.h
Comment on lines +25 to +32
# include <cuda/std/__cstddef/types.h>
# include <cuda/std/__functional/invoke.h>
# include <cuda/std/__type_traits/is_constructible.h>
# include <cuda/std/__type_traits/remove_cvref.h>
# include <cuda/std/__utility/declval.h>
# include <cuda/std/__utility/forward.h>
# include <cuda/std/__utility/integer_sequence.h>

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.

🛠️ Refactor suggestion | 🟠 Major | ⚡ Quick win

suggestion: Add a direct include for is_invocable_v usage.
is_invocable_v is used but its defining header is not included directly; please add the precise type-traits header to avoid transitive-include fragility.

As per coding guidelines: “include only directly-required headers (no transitive includes) with the most precise internal header paths.”

Also applies to: 296-306

Comment thread libcudacxx/include/cuda/std/version
@davebayer davebayer force-pushed the libcuxx_constant_wrapper branch from 6f8f1e9 to 4518f37 Compare May 17, 2026 13:41
Copy link
Copy Markdown
Contributor

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

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

Actionable comments posted: 1


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 7e0e3205-d62f-456c-a899-1eca2f817c7d

📥 Commits

Reviewing files that changed from the base of the PR and between 6f8f1e9 and 4518f37.

📒 Files selected for processing (22)
  • libcudacxx/include/cuda/std/__cccl/dialect.h
  • libcudacxx/include/cuda/std/__utility/constant_wrapper.h
  • libcudacxx/include/cuda/std/utility
  • libcudacxx/include/cuda/std/version
  • 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/ctad.compile.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw_fixed.array.ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw_fixed.ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/general.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/helpers.h
  • 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 skipped from review due to trivial changes (1)
  • libcudacxx/include/cuda/std/utility
🚧 Files skipped from review as they are similar to previous changes (3)
  • libcudacxx/include/cuda/std/version
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/helpers.h
  • libcudacxx/include/cuda/std/__cccl/dialect.h

Comment thread libcudacxx/include/cuda/std/__utility/constant_wrapper.h
@davebayer davebayer force-pushed the libcuxx_constant_wrapper branch from 4518f37 to 2346681 Compare May 17, 2026 14:46
Copy link
Copy Markdown
Contributor

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

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

Actionable comments posted: 2

♻️ Duplicate comments (1)
libcudacxx/include/cuda/std/__utility/constant_wrapper.h (1)

81-82: 🛠️ Refactor suggestion | 🟠 Major | ⚡ Quick win

important: make cw an inline constexpr variable template to satisfy header ODR/guideline expectations.

As per coding guidelines: “all namespace/global constexpr variables must be inline.”

🧹 Nitpick comments (1)
libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/types.compile.pass.cpp (1)

46-50: ⚡ Quick win

suggestion: Line 47 disables the only direct string-literal value-content check, so this file currently validates shape but not payload. Add a temporary compile-time fallback (for example, per-index static_assert checks on value[0..4]) until ranges::equal is usable.


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 1f32351b-47cf-48c0-9a61-2e03bb1fb3f4

📥 Commits

Reviewing files that changed from the base of the PR and between 4518f37 and 2346681.

📒 Files selected for processing (22)
  • libcudacxx/include/cuda/std/__cccl/dialect.h
  • libcudacxx/include/cuda/std/__utility/constant_wrapper.h
  • libcudacxx/include/cuda/std/utility
  • libcudacxx/include/cuda/std/version
  • 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/ctad.compile.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw_fixed.array.ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/cw_fixed.ctor.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/general.pass.cpp
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/helpers.h
  • 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 skipped from review due to trivial changes (2)
  • libcudacxx/include/cuda/std/utility
  • libcudacxx/include/cuda/std/version
🚧 Files skipped from review as they are similar to previous changes (2)
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/helpers.h
  • libcudacxx/include/cuda/std/__cccl/dialect.h

Comment thread libcudacxx/include/cuda/std/__utility/constant_wrapper.h Outdated
Comment on lines +45 to +49
TEST_FUNC void final_phase(auto gathered, auto available)
{
if constexpr (gathered == available)
{
profit();
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.

⚠️ Potential issue | 🔴 Critical | ⚡ Quick win

🧩 Analysis chain

🌐 Web query:

In C++20, is if constexpr (x == y)valid insidevoid f(auto x, auto y)whenxandy are ordinary function parameters?

💡 Result:

In C++20, the function void f(auto x, auto y) is a function template (using the abbreviated function template syntax) [1]. Inside this function, if constexpr (x == y) is valid only if the expression x == y is a constant expression [2]. Because x and y are function parameters, they are not constant expressions at compile time, even though the function is a template [3][2]. if constexpr requires its condition to be a constant expression that can be evaluated at compile time [2]. Since the values of x and y are only known at runtime, x == y cannot be evaluated at compile time, and therefore if constexpr (x == y) will result in a compilation error [3][2]. If you need to perform conditional logic based on the values of function parameters, you must use a standard if statement [2]. If you need to perform conditional logic based on the types of x and y at compile time, you can use if constexpr with type traits, such as if constexpr (std::is_same_v<decltype(x), decltype(y)>) [1].

Citations:


🏁 Script executed:

cat -n libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/general.pass.cpp | sed -n '40,55p'

Repository: NVIDIA/cccl

Length of output: 491


critical: if constexpr (gathered == available) is ill-formed because function parameters are not constant expressions. The condition in if constexpr must be evaluable at compile time. Use a regular if statement, or if you need compile-time branching based on types, use if constexpr with type traits like std::is_same_v<decltype(gathered), decltype(available)>.

Copy link
Copy Markdown
Contributor

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

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

Actionable comments posted: 2


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 7ed90de6-f814-4373-ac6a-905fe28dc678

📥 Commits

Reviewing files that changed from the base of the PR and between 3a52eca and 0e08305.

📒 Files selected for processing (1)
  • libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/kernel_parameter.pass.cpp

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

Copy link
Copy Markdown
Contributor

@Jacobfaib Jacobfaib left a comment

Choose a reason for hiding this comment

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

Why not C++17?

// a host member cannot be directly read in a __device__/__global__ function
_CCCL_BEGIN_NV_DIAG_SUPPRESS(342, 20094)

// todo: Remove this once #9019 is merged.
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.

Suggested change
// todo: Remove this once #9019 is merged.
// TODO(dbayer): Remove this once #9019 is merged.

Note the name is not an assignment of the TODO, it is so that future devs know who to ask about this for additional context when fixing in the future

Comment thread libcudacxx/include/cuda/std/__utility/constant_wrapper.h
Comment thread libcudacxx/include/cuda/std/__utility/constant_wrapper.h Outdated
@davebayer
Copy link
Copy Markdown
Contributor Author

Why not C++17?

Because the constant_wrapper's NTTP is of __cw_fixed_value type, which is not supported in C++17. We might be able to backport it with some more limitations to C++17, but I would start with the simpler case for now.

This is definitely not the first PR regarding this feature

@davebayer davebayer force-pushed the libcuxx_constant_wrapper branch 2 times, most recently from 7ae24b1 to 7fb0b88 Compare May 20, 2026 19:12
@davebayer davebayer force-pushed the libcuxx_constant_wrapper branch from 7fb0b88 to e04c6f5 Compare May 20, 2026 20:20
@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

Comment thread libcudacxx/test/libcudacxx/std/utilities/const.wrap.class/call.pass.cpp Outdated
@github-actions
Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 1h 11m: Pass: 100%/116 | Total: 1d 06h | Max: 49m 33s | Hits: 99%/333961

See results here.

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.

[FEA]: Initial cuda::std::constant_wrapper implementation

2 participants