Add env based DeviceFor::* algorithms

[gonidelis]

Adds env based DeviceFor::*

closed #7541

There are not deterministic guarantees imposed for the DeviceFor::* algorithms.

[copy-pr-bot]

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.

[gonidelis]

One thing that might kill backwards compatibility from unifying the env overloads with the old ones is that we tend to make new overloads [[nodiscard]]. Adding this to the "new" env overloads (which just substitute the old) might pop compile erros to users out of the blue. @bernhardmgruber @gevtushenko should I avoid doing the overloads in this pr nodiscard?

[gonidelis]

todo: add docs about memory_resource (?)

[bernhardmgruber]

One thing that might kill backwards compatibility from unifying the env overloads with the old ones is that we tend to make new overloads [[nodiscard]]. Adding this to the "new" env overloads (which just substitute the old) might pop compile erros to users out of the blue. @bernhardmgruber @gevtushenko should I avoid doing the overloads in this pr nodiscard?

Good point. Please omit the [[nodiscard]] in this case.

[gonidelis]

Since we replaced the cudaStream_t stream parameter with EnvT env, we broke backwards compatibility and callers that relied on implicit conversion from nvbench::cuda_stream to cudaStream_t. The extents benchmark showcases that. I see three options:

a) We go back to only adding env overloads (not converting old ones) for DeviceFor
b) We check internally if passed 3rd argument (env) is nvbench::cuda_stream and do the conversion explicitly inside the body
c) We break, and ask for people that were relying on implicit conversion to explicitly .get_stream()

[bernhardmgruber]

Since we replaced the cudaStream_t stream parameter with EnvT env, we broke backwards compatibility and callers that relied on implicit conversion from nvbench::cuda_stream to cudaStream_t. The extents benchmark showcases that.

That's interesting, because we hit the same case in cub::DeviceTransform with RMM's stream type. We solved this in #7266 and added tests in #7278. What we apparently missed are types that are convertible to cudaStream_t and are not copyable, because then you can't pass them to the environment parameter.

Let me ask @miscco

[bernhardmgruber]

Let me ask @miscco

Started an internal Slack thread.

In the meantime, I came up with a fix and pushed a commit.

[bernhardmgruber]

/ok to test 712a5f3

[bernhardmgruber]

pre-commit.ci autofix

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[bernhardmgruber]

/ok to test ceda482

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 3h 49m: Pass: 100%/249 | Total: 7d 02h | Max: 3h 48m | Hits: 84%/155556

See results here.

[gonidelis]

😡😡😡😡😡