[Draft] Document CUB environments for maintainers#10013
[Draft] Document CUB environments for maintainers#10013bernhardmgruber wants to merge 2 commits into
Conversation
|
I think it's really hard to explain how environments are handled in CUB without the many internal facilities we have for them. That makes me wonder whether some of them should become public. For example:
We don't need those in the public to implement CUB, but if a user wants to inspect environments themselves, they are going to have a hard time. |
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesCUB environments and requirements
Suggested reviewers: Comment |
There was a problem hiding this comment.
Actionable comments posted: 5
🧹 Nitpick comments (2)
docs/thrust/system.rst (1)
1-11: 🧹 Nitpick | 🔵 Trivialsuggestion: Run the repository’s targeted documentation build and
pre-commitchecks before merging, and verify that bothsystem/systemsandsystem/diagnosticsresolve in the generated navigation. These checks were not run in this file-only review.Source: Coding guidelines
docs/thrust/algorithms/reductions/transformed.rst (1)
10-10: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Validate the revised API glob paths with the repository’s targeted documentation build and
pre-commit. The generated API tree and required build instructions were not supplied, so path resolution cannot be confirmed from this diff alone. As per coding guidelines, reviewCONTRIBUTING.mdand the CI overview before running validation.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 93891432-8bbd-4e1c-9c78-45dadcbf6a13
📒 Files selected for processing (65)
docs/cub/api.rstdocs/cub/benchmarking.rstdocs/cub/block_wide.rstdocs/cub/determinism.rstdocs/cub/developer_overview.rstdocs/cub/device_topk_requirements.rstdocs/cub/device_wide.rstdocs/cub/environments.rstdocs/cub/index.rstdocs/cub/thread_level.rstdocs/cub/tuning.rstdocs/cub/tuning_infra.rstdocs/cub/warp_wide.rstdocs/thrust/algorithms.rstdocs/thrust/algorithms/copying.rstdocs/thrust/algorithms/copying/gather.rstdocs/thrust/algorithms/copying/scatter.rstdocs/thrust/algorithms/merging.rstdocs/thrust/algorithms/prefix_sum.rstdocs/thrust/algorithms/prefix_sums/segmented.rstdocs/thrust/algorithms/prefix_sums/transformed.rstdocs/thrust/algorithms/reductions.rstdocs/thrust/algorithms/reductions/comparisons.rstdocs/thrust/algorithms/reductions/counting.rstdocs/thrust/algorithms/reductions/extrema.rstdocs/thrust/algorithms/reductions/logical.rstdocs/thrust/algorithms/reductions/predicates.rstdocs/thrust/algorithms/reductions/transformed.rstdocs/thrust/algorithms/reordering.rstdocs/thrust/algorithms/reordering/partitioning.rstdocs/thrust/algorithms/reordering/shuffling.rstdocs/thrust/algorithms/reordering/stream_compaction.rstdocs/thrust/algorithms/searching.rstdocs/thrust/algorithms/searching/binary_search.rstdocs/thrust/algorithms/set_operations.rstdocs/thrust/algorithms/sorting.rstdocs/thrust/algorithms/transformations.rstdocs/thrust/algorithms/transformations/filling.rstdocs/thrust/algorithms/transformations/modifying.rstdocs/thrust/algorithms/transformations/replacing.rstdocs/thrust/api.rstdocs/thrust/containers.rstdocs/thrust/function_objects.rstdocs/thrust/function_objects/adaptors.rstdocs/thrust/function_objects/placeholder.rstdocs/thrust/function_objects/predefined.rstdocs/thrust/index.rstdocs/thrust/iterators.rstdocs/thrust/memory_management.rstdocs/thrust/memory_management/allocators.rstdocs/thrust/memory_management/memory_resources.rstdocs/thrust/numerics.rstdocs/thrust/parallel_execution_policies.rstdocs/thrust/random.rstdocs/thrust/random/distributions.rstdocs/thrust/random/engine.rstdocs/thrust/random/engine_adaptors.rstdocs/thrust/random/predefined.rstdocs/thrust/system.rstdocs/thrust/system/diagnostics.rstdocs/thrust/system/systems.rstdocs/thrust/utility.rstdocs/thrust/utility/pair.rstdocs/thrust/utility/tuple.rstdocs/thrust/utility/type_traits.rst
💤 Files with no reviewable changes (2)
- docs/thrust/api.rst
- docs/cub/api.rst
There was a problem hiding this comment.
Caution
Inline review comments failed to post. This is likely due to GitHub's internal server error or limits when posting large numbers of comments. If you are seeing this consistently it is likely a permissions issue. Please check "Moderation" -> "Code review limits" under your organization settings.
Actionable comments posted: 5
🧹 Nitpick comments (2)
docs/thrust/system.rst (1)
1-11: 🧹 Nitpick | 🔵 Trivialsuggestion: Run the repository’s targeted documentation build and
pre-commitchecks before merging, and verify that bothsystem/systemsandsystem/diagnosticsresolve in the generated navigation. These checks were not run in this file-only review.Source: Coding guidelines
docs/thrust/algorithms/reductions/transformed.rst (1)
10-10: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Validate the revised API glob paths with the repository’s targeted documentation build and
pre-commit. The generated API tree and required build instructions were not supplied, so path resolution cannot be confirmed from this diff alone. As per coding guidelines, reviewCONTRIBUTING.mdand the CI overview before running validation.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 93891432-8bbd-4e1c-9c78-45dadcbf6a13
📒 Files selected for processing (65)
docs/cub/api.rstdocs/cub/benchmarking.rstdocs/cub/block_wide.rstdocs/cub/determinism.rstdocs/cub/developer_overview.rstdocs/cub/device_topk_requirements.rstdocs/cub/device_wide.rstdocs/cub/environments.rstdocs/cub/index.rstdocs/cub/thread_level.rstdocs/cub/tuning.rstdocs/cub/tuning_infra.rstdocs/cub/warp_wide.rstdocs/thrust/algorithms.rstdocs/thrust/algorithms/copying.rstdocs/thrust/algorithms/copying/gather.rstdocs/thrust/algorithms/copying/scatter.rstdocs/thrust/algorithms/merging.rstdocs/thrust/algorithms/prefix_sum.rstdocs/thrust/algorithms/prefix_sums/segmented.rstdocs/thrust/algorithms/prefix_sums/transformed.rstdocs/thrust/algorithms/reductions.rstdocs/thrust/algorithms/reductions/comparisons.rstdocs/thrust/algorithms/reductions/counting.rstdocs/thrust/algorithms/reductions/extrema.rstdocs/thrust/algorithms/reductions/logical.rstdocs/thrust/algorithms/reductions/predicates.rstdocs/thrust/algorithms/reductions/transformed.rstdocs/thrust/algorithms/reordering.rstdocs/thrust/algorithms/reordering/partitioning.rstdocs/thrust/algorithms/reordering/shuffling.rstdocs/thrust/algorithms/reordering/stream_compaction.rstdocs/thrust/algorithms/searching.rstdocs/thrust/algorithms/searching/binary_search.rstdocs/thrust/algorithms/set_operations.rstdocs/thrust/algorithms/sorting.rstdocs/thrust/algorithms/transformations.rstdocs/thrust/algorithms/transformations/filling.rstdocs/thrust/algorithms/transformations/modifying.rstdocs/thrust/algorithms/transformations/replacing.rstdocs/thrust/api.rstdocs/thrust/containers.rstdocs/thrust/function_objects.rstdocs/thrust/function_objects/adaptors.rstdocs/thrust/function_objects/placeholder.rstdocs/thrust/function_objects/predefined.rstdocs/thrust/index.rstdocs/thrust/iterators.rstdocs/thrust/memory_management.rstdocs/thrust/memory_management/allocators.rstdocs/thrust/memory_management/memory_resources.rstdocs/thrust/numerics.rstdocs/thrust/parallel_execution_policies.rstdocs/thrust/random.rstdocs/thrust/random/distributions.rstdocs/thrust/random/engine.rstdocs/thrust/random/engine_adaptors.rstdocs/thrust/random/predefined.rstdocs/thrust/system.rstdocs/thrust/system/diagnostics.rstdocs/thrust/system/systems.rstdocs/thrust/utility.rstdocs/thrust/utility/pair.rstdocs/thrust/utility/tuple.rstdocs/thrust/utility/type_traits.rst
💤 Files with no reviewable changes (2)
- docs/thrust/api.rst
- docs/cub/api.rst
🛑 Comments failed to post (5)
docs/cub/device_topk_requirements.rst (1)
45-75: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
rg -n "Device(Batched)?TopK|static_assert|determinism::|tie_break::|output_ordering::" cub docs/cub/device_topk_requirements.rstRepository: NVIDIA/cccl
Length of output: 50368
🏁 Script executed:
sed -n '35,85p;190,330p' docs/cub/device_topk_requirements.rst | cat -nRepository: NVIDIA/cccl
Length of output: 9749
🏁 Script executed:
sed -n '186,330p' docs/cub/device_topk_requirements.rst | cat -nRepository: NVIDIA/cccl
Length of output: 6984
important: Split current support from the long-term contract
docs/cub/device_topk_requirements.rst:45-75,202-203,304-323still presents the deterministic default and suggested pinned/sorted requirements as available today, but the current-support note says onlynot_guaranteed/unspecified/unsortedcompiles now and thatcub::DeviceTopKstill ignorestie_break. Mark the unsupported rows as future behavior or trim the table to combinations that compile today.Sources: Coding guidelines, Path instructions
docs/cub/environments.rst (3)
48-49: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
pre-commit run --files docs/cub/environments.rstRepository: NVIDIA/cccl
Length of output: 197
🏁 Script executed:
#!/bin/bash set -euo pipefail printf '\n== repo files ==\n' git ls-files | rg '(^|/)(CONTRIBUTING\.md|CI|ci|\.github/workflows/|docs/cub/environments\.rst)$' printf '\n== CONTRIBUTING.md (if present) ==\n' if [ -f CONTRIBUTING.md ]; then cat -n CONTRIBUTING.md | sed -n '1,220p' fi printf '\n== CI overview candidates ==\n' for f in $(git ls-files | rg '(^|/)(CI|ci|continuous|workflow|workflows).*(md|rst|txt)$|(^|/)(ci|CI)\.md$'); do echo "--- $f ---" wc -l "$f" done printf '\n== docs/cub/environments.rst (around the cited lines) ==\n' cat -n docs/cub/environments.rst | sed -n '35,70p' printf '\n== search for mr_env snippet and related constructs ==\n' rg -n "mr_env|get_memory_resource|resource_ref|prop\{cuda::mr::get_memory_resource" docs/cub/environments.rstRepository: NVIDIA/cccl
Length of output: 18255
🏁 Script executed:
#!/bin/bash set -euo pipefail printf '\n== surrounding context in docs/cub/environments.rst ==\n' cat -n docs/cub/environments.rst | sed -n '1,90p' printf '\n== occurrences of alloc in docs/cub/environments.rst ==\n' rg -n '\balloc\b' docs/cub/environments.rstRepository: NVIDIA/cccl
Length of output: 4990
important: Add the missing semicolon and use
mrinstead ofallocinmr_env; this example doesn't compile as written.Sources: Coding guidelines, Path instructions
134-137: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
rg -n "__get_requirements_t|static_assert|determinism::run_to_run" docs/cub/environments.rstRepository: NVIDIA/cccl
Length of output: 649
🏁 Script executed:
sed -n '120,155p' docs/cub/environments.rst | cat -nRepository: NVIDIA/cccl
Length of output: 2202
important: Fix this example before publishing. It should call
cuda::execution::__get_requirements_tconsistently, and the assertion should compare the queried value directly, e.g.static_assert(determinism == cuda::execution::determinism::run_to_run);Sources: Coding guidelines, Path instructions
159-159: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash set -euo pipefail printf '\n== repo root ==\n' pwd printf '\n== top-level files ==\n' git ls-files | sed -n '1,120p' printf '\n== CONTRIBUTING.md ==\n' if git ls-files CONTRIBUTING.md >/dev/null 2>&1; then cat -n CONTRIBUTING.md | sed -n '1,220p' else echo 'CONTRIBUTING.md not found' fi printf '\n== CI overview candidates ==\n' git ls-files | rg 'CI|ci|workflow|github' || true printf '\n== docs/cub/environments.rst excerpt ==\n' cat -n docs/cub/environments.rst | sed -n '145,170p' printf '\n== docs/cub/index.rst excerpt ==\n' cat -n docs/cub/index.rst | sed -n '470,495p'Repository: NVIDIA/cccl
Length of output: 50369
🏁 Script executed:
#!/bin/bash set -euo pipefail echo "== CONTRIBUTING exists? ==" git ls-files CONTRIBUTING.md echo "== CI guidance files ==" git ls-files | rg '(^|/)(CI|ci|workflow|github).*(md|rst)$|(^|/)CI\.md$|(^|/)ci\.md$' || true echo "== docs/cub/environments.rst lines 155-162 ==" sed -n '155,162p' docs/cub/environments.rst | cat -n echo "== docs/cub/index.rst lines 482-489 ==" sed -n '482,489p' docs/cub/index.rst | cat -nRepository: NVIDIA/cccl
Length of output: 1499
important: Fix the malformed reST link markup in both docs files.
docs/cub/environments.rst: the second link is missing its opening backtick (tuning <cub-policy-selectors>_).docs/cub/index.rst: add the opening backtick before the external URL so the link renders correctly.📍 Affects 2 files
docs/cub/environments.rst#L159-L159(this comment)docs/cub/index.rst#L486-L486Sources: Coding guidelines, Path instructions
docs/thrust/function_objects.rst (1)
53-56: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: this
thrust::transformcall passes only one input range, but the lambda takes two parameters, so the example does not compile. Addb.begin()as the second input range, or make the lambda unary.
63e8f5e to
480b21a
Compare
There was a problem hiding this comment.
Actionable comments posted: 2
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: be009a6d-24e4-488f-8a25-273de87e0f64
📒 Files selected for processing (2)
docs/cub/environments.rstdocs/cub/index.rst
🚧 Files skipped from review as they are similar to previous changes (1)
- docs/cub/environments.rst
| warp_wide | ||
| block_wide | ||
| device_wide | ||
| environments |
There was a problem hiding this comment.
🗄️ Data Integrity & Integration | 🟠 Major | 🏗️ Heavy lift
important: Do not expose the environments page as user-facing guidance until its internal APIs have an explicit support contract. The page directs readers to __query_or, __get_requirements_t, and __get_tuning_t; either document these as unstable implementation details or provide public replacements before adding this navigation entry. As per path instructions, documentation changes must address technical accuracy and API/version consistency.
Source: Path instructions
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
important: Fix the non-compiling C++ example in docs/cub/environments.rst before exposing it. The prop{...} declaration for mr_env is missing a terminating semicolon, violating the buildable-examples requirement.
Source: Path instructions
|
Reflecting a bit after a discussion with @gonidelis and reading #7712 again, I think we actually need two documentation pages. One that shows the user how to build an environment to be passed to a CUB algorithm. We don't explain CPOs or queries or what else you can do with an environment. Then a second page (maybe in the developer section) what else we can do with an environment. How to query it, what operations can be performed etc. @gonidelis can you make a new PR and extract the above described subset addressing #7712? Thx! |
|
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. |
480b21a to
93b4396
Compare
|
|
||
| Almost all CUB device-wide algorithms accept an execution environment as their last argument. | ||
| Environments are objects responding to queries which return values. | ||
| A "query" in this case is usually a (tag) type naming the query, for example `cuda::execution::__get_tuning_t`. |
There was a problem hiding this comment.
Should probably list the most commonly used ones
| A "query" in this case is usually a (tag) type naming the query, for example `cuda::execution::__get_tuning_t`. | |
| A "query" in this case is usually a (tag) type naming the query, for example `cuda::execution::__get_tuning_t`, `cuda::get_stream`, or `cuda::mr::get_memory_resource`. |
| The purpose of an environment is to provide properties that govern the execution of an algorithm, | ||
| like the used CUDA stream, memory resource, required determinism, requested tuning, etc. |
There was a problem hiding this comment.
Personally the explanation that really illuminated environments was @gevtushenko likening them to **kwargs arguments in Python. i.e. some kind of optional, extensible hook that users can attach arbitrary information to which the algorithm interprets as needed.
Perhaps we can snake this in here somewhere?
| See also the documentation on `determinism <cub-determinism>`_ and tuning <cub-policy-selectors>`_. | ||
|
|
||
|
|
||
| Properties of environments |
There was a problem hiding this comment.
IMO this section should go first. The most important thing for maintainers is how to pass envs around and what they should not do with them, so knowing the following upfront is the most useful information:
- Environment queries should all be immutable so we should always pass by
const &. - Algorithms do not need to concern themselves with env lifetimes.
Also the following questions come to mind here:
- Are environments required to be idempotent? Is it OK to query an env multiple times for the same information or should we cache results?
- What is the general policy for handling environments that contain properties you do not handle? Should you
static_assert()? Let it pass?
| If no such value is found, it will return the fallback value, `fallback_stream`, instead. | ||
|
|
||
| .. | ||
| TODO(bgruber) Does cuda::std::execution::__query_or(cudaStream_t{}, cuda::get_stream, fallback_stream); work? |
There was a problem hiding this comment.
It does not.
auto r = cuda::std::execution::__query_or(cudaStream_t{}, cuda::get_stream, 1.0f);
static_assert(same_as<decltype(r), float>);But cuda::get_stream(cudaStream_t{}) works. cuda::__call_or() also works, returning the cudaStream_t instead of the float.
| Querying environments | ||
| -------------------------------------------------------------------------------- | ||
|
|
||
| There are two ways to get a value out of an environment again. |
There was a problem hiding this comment.
*four, also cuda::__call_or and cuda::__lazy_call_or. Actually this is a good point to bring this up, we should really unify all of these, I myself don't actually know which is better.
| cudaStream_t stream = cuda::get_stream(env); | ||
|
|
||
| If `env` would not contain any value responding to the `cuda::get_stream` query, the above would not compile. | ||
| Using a CPO to query an environment is preferred over using `__query_or` in general. |
There was a problem hiding this comment.
This statement seems counterintuitive. The environments are usually optional, so it's not known (unless you check) whether an env can satisfy a query, so __query_or seems like the safest general method no?
| static_assert(is_same_v<decltype(determinism), cuda::execution::determinism::run_to_run>): | ||
|
|
||
| .. | ||
| TODO(bgruber): Should we mention `__queryable_with`? |
There was a problem hiding this comment.
Yes, I think so. It would round out the discussion by documenting how to check whether you can query without doing it.
This PR tries to extend the CUB documentation on what environments are and how we use them.
Fixes: #7712
Fixes: #7276