Repository navigation
Add env-based cub::DeviceSelect algorithms - #5335
Conversation
|
Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually. Contributors can view more details about this message here. |
This comment has been minimized.
This comment has been minimized.
f709708 to
861f0c0
Compare
| ::cuda::execution::determinism::run_to_run_t>; | ||
|
|
||
| // Static assert to reject gpu_to_gpu determinism since it makes no sense for DeviceSelect | ||
| static_assert(!::cuda::std::is_same_v<requested_determinism_t, ::cuda::execution::determinism::gpu_to_gpu_t>, |
There was a problem hiding this comment.
I am defaulting deterministic requirements for select-if to only run-to-run. Not sure if gpu-to-gpu would make sense. Please ack @NaderAlAwar @fbusato
There was a problem hiding this comment.
I think you can default it to gpu_to_gpu, as select_op is a unary predicate and doesn't interact with neighbouring elements (will not be dealing with Associativity).
We any how promise to maintain relative ordering previously (mentioned in documentation above https://github.com/gonidelis/cccl/blob/select_env/cub/cub/device/device_select.cuh#L403). But in any case, both new API and old might have side effects if the select_op is used in bad way, such as depending on global state or have some randomness such that results are not-deterministic from the functor itself. But, its totally on user. IMO, In short if select_op is deterministic, API will be deterministic.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
| cudaStream_t custom_stream; | ||
| REQUIRE(cudaSuccess == cudaStreamCreate(&custom_stream)); |
There was a problem hiding this comment.
Suggestion: can we just use cuda::stream instead of manual stream management?
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
| ::cuda::execution::determinism::run_to_run_t>; | ||
|
|
||
| // Static assert to reject gpu_to_gpu determinism since it makes no sense for DeviceSelect | ||
| static_assert(!::cuda::std::is_same_v<requested_determinism_t, ::cuda::execution::determinism::gpu_to_gpu_t>, |
There was a problem hiding this comment.
I think you can default it to gpu_to_gpu, as select_op is a unary predicate and doesn't interact with neighbouring elements (will not be dealing with Associativity).
We any how promise to maintain relative ordering previously (mentioned in documentation above https://github.com/gonidelis/cccl/blob/select_env/cub/cub/device/device_select.cuh#L403). But in any case, both new API and old might have side effects if the select_op is used in bad way, such as depending on global state or have some randomness such that results are not-deterministic from the functor itself. But, its totally on user. IMO, In short if select_op is deterministic, API will be deterministic.
cafc921 to
7a0dbf2
Compare
🥳 CI Workflow Results🟩 Finished in 3h 51m: Pass: 100%/249 | Total: 9d 20h | Max: 3h 50m | Hits: 54%/153668See results here. |
partially fixes #7552
Adds
Select::IfandSelect::Flaggedcub::DeviceSelect::FlaggedIf()cub::DeviceSelect::Unique()andcub::DeviceSelect::UniqueByKey()to be added in a subsequent PR.