Exposes cuda::execution::guarantee#10022
Conversation
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
🚧 Files skipped from review as they are similar to previous changes (2)
📝 WalkthroughSummary by CodeRabbit
WalkthroughAdds ChangesExecution guarantees
Assessment against linked issues
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 3
🧹 Nitpick comments (3)
libcudacxx/include/cuda/__execution/guarantee.h (1)
67-71: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Complete the Doxygen contract for
guarantee. Add@param[in]documentation for__guaranteesand@returndocumentation for the non-void result, including the by-value storage and noexcept behavior.Source: Coding guidelines
libcudacxx/test/libcudacxx/cuda/execution/guarantee.pass.cpp (2)
45-71: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion: Assert the advertised
constexprandnoexceptcontract directly.static_assert(test())covers constexpr evaluation, but add astatic_assert(noexcept(exec::guarantee(...)))check and a mutation check proving bundled guarantees are stored by value.
31-42: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Use
TEST_FUNCfor these host/device test helpers instead of raw_CCCL_HOST_DEVICE. This keeps the test aligned with the harness’ portability and compiler-mode handling.Sources: Coding guidelines, Path instructions
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 909cf52e-48dd-43b4-93d0-c357293aeb0c
📒 Files selected for processing (5)
libcudacxx/include/cuda/__execution/guarantee.hlibcudacxx/include/cuda/executionlibcudacxx/include/cuda/execution.guarantee.hlibcudacxx/test/libcudacxx/cuda/execution/guarantee.fail.cpplibcudacxx/test/libcudacxx/cuda/execution/guarantee.pass.cpp
|
|
||
| ::cuda::std::execution::env<_Guarantees...> __env{__guarantees...}; | ||
|
|
||
| return ::cuda::std::execution::prop{__get_guarantees_t{}, __env}; |
There was a problem hiding this comment.
| return ::cuda::std::execution::prop{__get_guarantees_t{}, __env}; | |
| return ::cuda::std::execution::prop{__get_guarantees, __env}; |
May as well reuse the global here?
There was a problem hiding this comment.
I was curious why we might follow the same __get_requirements_t{} for the requirements and consulted Claude, and the justification sounds reasonable - and something I wasn't aware of.
In device passes
_CCCL_GLOBAL_CONSTANTis__device__ constexprwithoutinline. Copying it inside this inline host-device function would odr-use a per-TU object — the static-linkage device-memory emission #7682 is trying to get rid of, plus formally divergent inline-function definitions across TUs. Since prop only uses the key's type, the temporary gets the identical result without naming the object; same reason require.h/tune.h spell it this way.
I had also Claude verify that mere declaration does not, while odr-use will actually cause materialization:
Mere declaration costs nothing; only odr-use materializes the object. Here's the probe (clang CUDA, -fgpu-rdc, device pass):
struct CPO { __device__ void operator()() const noexcept {} };
__device__ constexpr CPO unused_cpo{}; // declared, never touched
__device__ constexpr CPO used_cpo{}; // called + address taken in a kernel
Result in the generated PTX:
- unused_cpo — absent entirely. No symbol, no storage, at any optimization level. A namespace-scope constexpr variable that's never odr-used doesn't have to exist as an object, and the compiler doesn't create it. So the declaration line in guarantee.h is free.
- used_cpo — .const .align 1 .b8 used_cpo[1]; — one byte of actual device storage, materialized because the code needed the object itself.
I'm leaving a comment in the code.
There was a problem hiding this comment.
Why do we need _CCCL_GLOBAL_CONSTANT and cannot just use constexpr inline?
There was a problem hiding this comment.
Iirc, inline constexpr variables aren't accessible from device code.
| [[nodiscard]] _CCCL_NODEBUG_API constexpr auto guarantee(_Guarantees... __guarantees) noexcept | ||
| { | ||
| static_assert((::cuda::std::is_base_of_v<__guarantee, _Guarantees> && ...), | ||
| "Only guarantees can be passed to guarantee"); |
There was a problem hiding this comment.
Even though the is_base_of hints at it quite strongly, we should be explicit here that being a "guarantee" means publicly subclassing from cuda::__guarantee (and consider dropping the __ then, it seems a bit weird to ask users to subclass a dunder class). Perhaps cuda::guarantee_base to mimic std::views::view_base.
There was a problem hiding this comment.
This base class is just a means to improve detecting wrong usage of the API, right? I would not need it public then, but it also does not hurt.
One argument could be that users could define their own guarantees if they had access to the base class.
There was a problem hiding this comment.
I think if we want to allow users to define their own guarantees this makes sense. I would limit API surface for now and only open up for deriving from a guarantee_base in follow-up work, when there's demand or we deem it useful? What do you think?
This comment has been minimized.
This comment has been minimized.
| [[nodiscard]] _CCCL_NODEBUG_API constexpr auto guarantee(_Guarantees... __guarantees) noexcept | ||
| { | ||
| static_assert((::cuda::std::is_base_of_v<__guarantee, _Guarantees> && ...), | ||
| "Only guarantees can be passed to guarantee"); |
There was a problem hiding this comment.
This base class is just a means to improve detecting wrong usage of the API, right? I would not need it public then, but it also does not hurt.
One argument could be that users could define their own guarantees if they had access to the base class.
This comment has been minimized.
This comment has been minimized.
🥳 CI Workflow Results🟩 Finished in 4h 56m: Pass: 100%/120 | Total: 4d 07h | Max: 2h 39m | Hits: 46%/1320699See results here. |
3 participants
Closes #10017
This PR pulls the guarantees mechanism out of #9278 so that upcoming work like
runs_on(see #9814) can build on it, while the first concrete guarantee (with the future and shape ofmax_total_num_itemsstill being discussed). The idea is thatguaranteelets a caller promise properties about the problem that an algorithm may exploit.The mechanism mirrors
require, with two differences: guarantees may carry runtime state (e.g., a runtime bound), soguarantee(...)stores its arguments by value instead of discarding them, and, following review feedback on #9278,guarantee(...)isconstexprandnoexcept.