[libcu++] Make argument namespace and wrappers construction public#9251
Conversation
davebayer
left a comment
davebayer
left a comment
There was a problem hiding this comment.
I'm sorry, but I'll block this PR. We said there would be a discussion about the naming. I'm still strongly against having a separate namespace for these types.
|
Ready to act? Review this PR in Change Stack to turn feedback into patch suggestions you can inspect and refine. Note Reviews pausedIt 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 Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
suggestion: Walkthrough This PR replaces internal cuda::__argument wrapper/bounds/traits names with public cuda::argument / cuda::args equivalents, adds a public forwarding header Changes CUDA argument API public-API migration
Assessment against linked issues
Possibly related PRs
Suggested labels
Suggested reviewers
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: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 4424dd10-8027-4d66-97f4-e3d43a7b6cf4
📒 Files selected for processing (20)
cub/benchmarks/bench/segmented_topk/fixed/keys.cucub/benchmarks/bench/segmented_topk/variable/keys.cucub/cub/agent/agent_batched_topk.cuhcub/cub/detail/segmented_params.cuhcub/cub/device/dispatch/dispatch_batched_topk.cuhcub/cub/device/dispatch/kernels/kernel_batched_topk.cuhcub/test/catch2_test_device_segmented_topk_keys.cucub/test/catch2_test_device_segmented_topk_pairs.culibcudacxx/include/cuda/__argument/argument.hlibcudacxx/include/cuda/__argument/argument_bounds.hlibcudacxx/include/cuda/argumentlibcudacxx/include/cuda/std/__internal/namespaces.hlibcudacxx/test/libcudacxx/cuda/argument/argument_bounds.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/deferred_argument.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpplibcudacxx/test/libcudacxx/cuda/argument/static_bounds_type_mismatch.fail.cpplibcudacxx/test/libcudacxx/cuda/argument/usage_example.pass.cpp
Tagging @gevtushenko |
pciolkosz
force-pushed
the
make_arguments_framework_public
branch
from
6002bc6 to
3210e3d
Compare
This comment has been minimized.
This comment has been minimized.
| auto segment_sizes = ::cuda::__argument::__constant<MaxSegmentSize>{}; | ||
| auto k = ::cuda::__argument::__constant<MaxNumSelected>{}; | ||
| auto select_direction = ::cuda::__argument::__constant<cub::detail::topk::select::max>{}; | ||
| auto segment_sizes = ::cuda::argument::constant<MaxSegmentSize>{}; |
There was a problem hiding this comment.
C++26 has constant_wrapper. It's hard to back-port fully to earlier C++ versions (though there are proposals pending review that might fix this), but for now, it's how C++ spells "compile-time constant." Does this cuda::argument::constant do the same thing? If so, then should we consider just making this an alias to constant_wrapper (or a hypothetical cuda::std::constant_wrapper)?
|
|
||
| using segment_size_val_t = typename ::cuda::__argument::__traits<SegmentSizeParameterT>::element_type; | ||
| using num_segments_val_t = typename ::cuda::__argument::__traits<NumSegmentsParameterT>::element_type; | ||
| using segment_size_val_t = typename ::cuda::argument::__traits<SegmentSizeParameterT>::element_type; |
There was a problem hiding this comment.
Is this library usable without exposing __traits to the public API? If not, then should we consider waiting until __traits is ready?
There was a problem hiding this comment.
The main user facing part is the construction of the wrappers, while our APIs will use the traits to examine them. Long term we want to expose the __traits I think, but we don't have to right now so keeping them private leaves us room to change them easily and make public later
There was a problem hiding this comment.
Thanks for explaining! That makes sense.
| //! @brief Wraps a compile-time constant argument value. | ||
| template <auto _Value> | ||
| struct __constant | ||
| struct constant |
There was a problem hiding this comment.
This looks more or less like constant_wrapper. Should we consider spelling it cuda::std::constant_wrapper?
There was a problem hiding this comment.
We considered that design, cuda::std::constant_wrapper is difficult to back port to older c++ standards and the timeline probably doesn't leave enough room to implement it. I think its fine to just interoperate with constant wrapper later on
| static_assert(__valid_static_bounds_v<__element_type, _StaticBounds>, | ||
| "static argument bounds cannot be represented by the element type"); | ||
|
|
||
| _Arg __arg_; |
There was a problem hiding this comment.
This is a public member. That means it's possible to set the value without validating the bounds. Should we instead consider making the value private, so that every constructed immediate instance has been validated?
There was a problem hiding this comment.
I could make it a private member and add an underscored __access structure that is a friend or I can just underscore the variable and have less code this way. I don't mind either way, but we need a way to access it at the end of the day.
But I think this question does apply to __unwrap, which currently has two version, const and not-const reference. I could see the non-const reference version being not necessary, but it's also not public yet, so we don't have to resolve it right now. I might as well remove it and we can add it when needed
There was a problem hiding this comment.
@pciolkosz Thanks for answering! I don't quite understand the argument why validation is separate from construction.
There was a problem hiding this comment.
Also: Why can't users access the value (by value)? Isn't this just a validated value? Why wouldn't this class just have a value() member?
There was a problem hiding this comment.
I decided to not be stubborn and changed all members of the wrappers to private and added a friend struct to access them. We might end up exposing member functions to access them, but for now there is only __unwrap free function
There was a problem hiding this comment.
Bless you, @pciolkosz ! I can fall asleep with one less worry now : - )
|
|
||
| template <class _ElementType, auto _Lowest, auto _Highest> | ||
| inline constexpr bool __valid_static_bounds_v<_ElementType, __static_bounds<_Lowest, _Highest>> = | ||
| inline constexpr bool __valid_static_bounds_v<_ElementType, static_bounds<_Lowest, _Highest>> = |
There was a problem hiding this comment.
There are runtime_bounds and static_bounds. Do we also need a representation of half-run-time, half-static bounds?
submdspan represents bounds generally as any type for which structured binding into two elements, both convertible to index_type, is well-formed. This would let users pass in bounds however they like. It also would permit a unified representation of any combination of run-time and static bounds.
There was a problem hiding this comment.
Also, the template parameters' types are unconstrained here. What are the actual requirements on these types?
There was a problem hiding this comment.
You can tighten the static bounds with runtime bounds, so in this case you can do static_bounds<0, 1000> and runtime_bounds(100, 1000) which effectively would be a half-run-time, half-static bounds. It avoids overcomplicating the input types at the expense of this case being a bit more verbose, but I think it should be rare anyway.
Also, the template parameters' types are unconstrained here. What are the actual requirements on these types?
Good point, I added type validation
This comment has been minimized.
This comment has been minimized.
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 7cc00549-bb49-459c-96d8-ed4906627049
📒 Files selected for processing (22)
cub/benchmarks/bench/segmented_topk/fixed/keys.cucub/benchmarks/bench/segmented_topk/variable/keys.cucub/cub/agent/agent_batched_topk.cuhcub/cub/detail/segmented_params.cuhcub/cub/device/dispatch/dispatch_batched_topk.cuhcub/cub/device/dispatch/kernels/kernel_batched_topk.cuhcub/test/catch2_test_device_segmented_topk_keys.cucub/test/catch2_test_device_segmented_topk_pairs.culibcudacxx/include/cuda/__argument/argument.hlibcudacxx/include/cuda/__argument/argument_bounds.hlibcudacxx/include/cuda/std/__internal/namespaces.hlibcudacxx/test/libcudacxx/cuda/argument/argument_bounds.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/deferred_argument.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/deferred_sequence_scalar.fail.cpplibcudacxx/test/libcudacxx/cuda/argument/deferred_sequence_scalar_traits.fail.cpplibcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpplibcudacxx/test/libcudacxx/cuda/argument/static_argument_bounds_type.fail.cpplibcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpplibcudacxx/test/libcudacxx/cuda/argument/static_bounds_type_mismatch.fail.cpplibcudacxx/test/libcudacxx/cuda/argument/usage_example.pass.cpp
✅ Files skipped from review due to trivial changes (1)
- libcudacxx/test/libcudacxx/cuda/argument/static_bounds_type_mismatch.fail.cpp
🚧 Files skipped from review as they are similar to previous changes (17)
- libcudacxx/include/cuda/std/__internal/namespaces.h
- libcudacxx/test/libcudacxx/cuda/argument/deferred_sequence_scalar.fail.cpp
- libcudacxx/test/libcudacxx/cuda/argument/deferred_sequence_scalar_traits.fail.cpp
- libcudacxx/test/libcudacxx/cuda/argument/static_bounds_conversion.fail.cpp
- libcudacxx/test/libcudacxx/cuda/argument/static_argument_bounds_type.fail.cpp
- libcudacxx/test/libcudacxx/cuda/argument/argument_traits.pass.cpp
- cub/cub/agent/agent_batched_topk.cuh
- cub/cub/detail/segmented_params.cuh
- cub/cub/device/dispatch/kernels/kernel_batched_topk.cuh
- cub/benchmarks/bench/segmented_topk/fixed/keys.cu
- cub/benchmarks/bench/segmented_topk/variable/keys.cu
- cub/test/catch2_test_device_segmented_topk_pairs.cu
- libcudacxx/test/libcudacxx/cuda/argument/usage_example.pass.cpp
- cub/test/catch2_test_device_segmented_topk_keys.cu
- libcudacxx/include/cuda/__argument/argument_bounds.h
- libcudacxx/test/libcudacxx/cuda/argument/dynamic_argument.pass.cpp
- libcudacxx/test/libcudacxx/cuda/argument/static_argument.pass.cpp
|
Changes from the code review session:
One thing to consider is if we want to expose |
This comment has been minimized.
This comment has been minimized.
This fully depends on whether we want to productize the full framework for users to use in their own libraries. For now I would say the answer to that is no These types are currently only used inside CCCL and for that we should not expose such interfaces |
Yup, I'm with Michael on this: Let's limit API surface to as little as possible for now. Just enough to allow users providing bounds on the arguments. |
pciolkosz
force-pushed
the
make_arguments_framework_public
branch
from
9088870 to
4777595
Compare
This comment has been minimized.
This comment has been minimized.
|
|
||
| _CCCL_BEGIN_NAMESPACE_CUDA_ARGUMENT | ||
|
|
||
| struct __access; |
There was a problem hiding this comment.
We should have cuda/__fwd/argument.h for that
| // ===================================================================== | ||
| // __assert_in_range | ||
| // ===================================================================== | ||
|
|
||
| template <class _To, class _From> | ||
| _CCCL_API constexpr void __assert_in_range([[maybe_unused]] _From __val) noexcept | ||
| { |
pciolkosz
force-pushed
the
make_arguments_framework_public
branch
from
4777595 to
5a0f0fe
Compare
pciolkosz
force-pushed
the
make_arguments_framework_public
branch
from
5a0f0fe to
f4cb111
Compare
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
🥳 CI Workflow Results🟩 Finished in 3h 27m: Pass: 100%/343 | Total: 7d 01h | Max: 2h 29m | Hits: 97%/485146See results here. |
5 participants
Closes #9254
This PR removes underscores from the
cuda::argumentnamespace and the user facing argument wrappers and bound objects construction.Traits and other utility functions remain underscored, those are meant mostly to be used inside CCCL APIs to examine the argument wrappers passed by the user, but they will most likely be public later on, once we are more confident with the design.