[CUB] Refactor DeviceHistogram::MultiHistogramEven to always take an environment

[miscco]

Description

closes

Checklist

• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[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.

[miscco]

/ok to test e74bf49

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 6e0f57da-964b-47f9-a867-004625039106

📥 Commits

Reviewing files that changed from the base of the PR and between e74bf49 and d93e20c.

📒 Files selected for processing (2)

• cub/cub/device/device_histogram.cuh
• cub/test/catch2_test_device_histogram_env.cu

🚧 Files skipped from review as they are similar to previous changes (2)

• cub/test/catch2_test_device_histogram_env.cu
• cub/cub/device/device_histogram.cuh

---

Note: CodeRabbit is enabled on this repository as a convenience for maintainers
and contributors. Use your best judgment when considering its review comments and
suggestions — a suggested change may be inadequate, unnecessary, or safe to ignore.
Contributors are not expected to address every comment. Human reviews are what
ultimately matter for merging.

Summary

• Refactored cub::DeviceHistogram::MultiHistogramEven to thread an execution environment (EnvT) through the implementation instead of relying on a cudaStream_t parameter.
• Updated public overloads to take const EnvT& env = {} and route dispatch through detail::dispatch_with_env_and_tuning(...).
• Adjusted Doxygen/template docs to reflect the environment-based API.
• Added tests covering MultiHistogramEven with user-provided temporary storage and multiple environment variants, including streams, stream refs, default env, and GPU execution policies.

Testing

• Added catch2 coverage for temp-storage sizing and histogram correctness across the new environment-based call paths.

Walkthrough

DeviceHistogram::MultiHistogramEven now takes const EnvT& env = {} across its 1D and 2D overloads, with the 2D path routed through detail::dispatch_with_env_and_tuning. The test suite adds coverage for caller-provided temporary storage and several env variants.

Changes

MultiHistogramEven env dispatch migration and tests

Layer / File(s)	Summary
1D MultiHistogramEven signature and forwarding cub/cub/device/device_histogram.cuh	Adds EnvT to the 1D MultiHistogramEven API docs and signatures, replaces cudaStream_t with const EnvT& env = {}, and forwards env through the internal and deprecated wrapper paths.
2D MultiHistogramEven signature and env dispatch cub/cub/device/device_histogram.cuh	Adds EnvT to the 2D MultiHistogramEven API docs and signatures, replaces cudaStream_t with const EnvT& env = {}, and routes the implementation through detail::dispatch_with_env_and_tuning.
MultiHistogramEven env test coverage cub/test/catch2_test_device_histogram_env.cu	Adds includes and a TEST_CASE that checks temp-storage sizing and histogram outputs across multiple env/stream/policy variants.

Suggested reviewers

• srinivasyadav18
• gonidelis
• pauleonix

---

_{Comment @coderabbitai help to get the list of available commands.}

[coderabbitai]

Actionable comments posted: 1

🧹 Nitpick comments (2)

cub/test/catch2_test_device_histogram_env.cu (2)

1086-1088: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick win

important: cudaDeviceSynchronize() masks stream propagation failures. If this overload ignores env and launches on the default stream, the stream sections can still pass. Use stream-specific synchronization or a stream-ordered event/sentinel for the stream env cases. As per path instructions, focus on stream behavior in cub/**/*.

Source: Path instructions

---

1125-1135: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

important: add a section with an actual custom histogram tuning/policy selector. cuda::execution::gpu only proves the policy object is accepted; it does not prove the user-provided-memory overload reads tuning from env, which is the stated PR objective.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 4fcc6884-b84f-4c5b-9be5-3b67d855da77

📥 Commits

Reviewing files that changed from the base of the PR and between edc3cc8 and e74bf49.

📒 Files selected for processing (2)

• cub/cub/device/device_histogram.cuh
• cub/test/catch2_test_device_histogram_env.cu

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 1h 43m: Pass: 100%/287 | Total: 2d 19h | Max: 56m 42s | Hits: 89%/223793

See results here.