Allow baseline_topk_policy to hold a variable number of worker policies - #11741
Open
bernhardmgruber wants to merge 16 commits into
Open
bernhardmgruber wants to merge 16 commits into
bernhardmgruber wants to merge 16 commits into
Conversation
bernhardmgruber
changed the title
baseline_topk_policy to hold a variable number of worker policies
Contributor
|
Navigate logical layers of code changes, visualize relationships, and explore their blast radius. 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:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Repository: NVIDIA/cccl/.coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (1)
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review. 📝 SummarySummary by CodeRabbit
WalkthroughThe baseline top-k policy now stores worker policies in a fixed-capacity structural vector with room for 10 entries. The default policy and three benchmark initializers retain their six worker policies. ChangesSegmented top-k policy
Priority: ⬇️ Low Change: Feature Merge Risk: ⚪ Minimal · up to The six-entry policy initializers fit the new capacity, and the selector’s default-construction requirement is met. No actionable merge-blocking risk remains in the reviewed changes. Comment |
This comment has been minimized.
This comment has been minimized.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 2
ℹ️ Review info
⚙️ Run configuration
Configuration used: Repository: NVIDIA/cccl/.coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: d846f768-eede-4d04-b245-3008828b27c7
📒 Files selected for processing (2)
cub/cub/device/dispatch/tuning/common.cuhcub/cub/device/dispatch/tuning/tuning_batched_topk.cuh
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 9 remain after this review.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
pauleonix
approved these changes
Sep 30, 2026
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.
miscco
approved these changes
Sep 30, 2026
This comment has been minimized.
This comment has been minimized.
Replace the fixed cuda::std::array<worker_policy, 6> with cuda::std::inplace_vector<worker_policy, 10>, decoupling the number of worker policies a tuning actually specifies from a compile-time capacity ceiling. This lets tunings list fewer (or, up to the new cap, more) worker policies without touching every construction site, while staying constexpr-evaluable down to GCC7 (a heap-allocating container like std::vector cannot be used here: allocation is never permitted in a C++17 constant expression, and even C++20's relaxed rules require any allocation to be freed within the same evaluation, so it could never back state that must persist as a constexpr policy object). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…policies
cub::detail::cc_dispatch instantiates the resolved policy as a non-type
template parameter (policy_constant / cuda::std::integral_constant),
which C++20 permits only for structural types: every base class and
non-static data member must be public. cuda::std::inplace_vector keeps
its storage private, so embedding it in baseline_topk_policy broke that
property for topk_policy and made every CC-dispatched instantiation
fail to compile ("a template parameter of class type must be of
structural class type").
Replace it with a small local worker_policy_list that keeps the same
capacity/size split (a public cuda::std::array plus a public count) and
the same call-site API (size(), operator[], range-for, ==), but with
entirely public state, so it stays a structural type.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Generalize the local worker_policy_list into cub::detail::structural_inplace_vector<T, Capacity> and move it to tuning/common.cuh so other tuning policies needing a dynamically sized, NTTP-safe list can reuse it. Trims the interface to what baseline_topk_policy actually needs: construction from an initializer_list, size(), const operator[], and const begin()/end(). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…n header Address review feedback and round out the type: - Restore a default constructor (dispatch_compute_cap / policy_selector require the resolved policy to satisfy cuda::std::regular). - Assert in the initializer-list constructor that the list doesn't exceed Capacity, instead of silently writing past `elems`. - Add reference/iterator aliases, a mutable operator[], and mutable begin()/end(), matching inplace_vector's interface. - Add at(), throwing std::out_of_range like inplace_vector::at(). - Assert in operator[] on out-of-range access. - Move the type out of the general-purpose tuning/common.cuh into its own tuning/structural_inplace_vector.cuh. All of the above remain usable in constexpr evaluation: verified that an out-of-bounds operator[]/at() access, and an oversized initializer list, each fail to compile under a static_assert, both with and without CCCL_ENABLE_ASSERTIONS. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…or at() elems doesn't need to be a cuda::std::array: operator[], begin()/end(), and == are all hand-rolled already, so a plain C array works and drops the include. Use _CCCL_VERIFY (always-on) instead of a manual throw for at()'s bounds check, matching its "must always check" contract -- unlike operator[]'s _CCCL_ASSERT, which stays gated behind CCCL_ENABLE_ASSERTIONS. This also drops the exception-macro/stdexcept includes. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Add the remaining std::array member types (difference_type, pointer, const_pointer, reverse_iterator, const_reverse_iterator), element access (front(), back(), data()), and capacity queries (empty(), max_size()). front()/back() assert non-empty, matching operator[]'s bounds assertion. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…vector No rbegin()/rend() were added, so the type aliases had no accessors to back them; drop them (and the now-unused reverse_iterator include) rather than carry unused surface area. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
…olsets The MSVC (CTK12.0/12.9/13.0/13.3, MSVC toolset 14.29 and 14.39; 14.44+ all pass) jobs failed with e.g. "block_topk_air<KeyT,0,0,...>" -- a worker_policy with threads_per_block == 0 reached device-agent instantiation, i.e. these older MSVC toolsets failed to correctly constant-fold worker_per_segment_policies somewhere in the CC-dispatch/NTTP chain, silently falling back to a zero-initialized value instead of hard-erroring at the actual failure point. Two simplifications to reduce the constexpr-evaluation burden on these toolsets (both still constexpr-correct, verified locally under -std=c++17/20 with -Wall -Wextra -Werror and with CCCL_ENABLE_ASSERTIONS): - The initializer-list constructor no longer mutates `count` via `count++` inside the copy loop; `count` is now set once via the constructor's member-initializer-list from `ilist.size()`, and the loop only writes `elems[i]`. - Drop the _CCCL_ASSERT from operator[], matching actual std::vector / std::array / cuda::std::inplace_vector semantics (operator[] is unchecked; at() remains bounds-checked). operator[] sits in a template recursion (find_covering_policy_index) evaluated once per candidate policy, so it's the likeliest remaining hot spot for this class of bug. This is a best-effort fix based on matching a known constexpr-evaluation bug window already worked around elsewhere in this codebase for MSVC toolsets below 14.39/14.44 (see cub/util_type.cuh); MSVC isn't available to reproduce locally, so this needs CI to confirm. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Rather than dropping operator[]'s assertion outright, gate all bounds checks (_CCCL_ASSERT/_CCCL_VERIFY call sites) behind _CCCL_COMPILER(MSVC, >=, 19, 44) via header-local _CCCL_SIV_ASSERT/_CCCL_SIV_VERIFY macros. MSVC toolsets below 14.44 (14.29, 14.39 -- see the earlier "Simplify constexpr paths" commit) compile this type with the checks compiled out entirely; every other compiler, and MSVC 14.44+, keeps full bounds checking. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Guarding the bounds-check asserts behind an MSVC-version check had zero effect -- CI still failed at the exact same line with all checks compiled out, proving the asserts were never the cause. The real difference from the known-working cuda::std::array this type replaced: that was a pure aggregate (no user-declared constructors); structural_inplace_vector had a user-declared initializer-list constructor with a loop. Something about that pushes the older MSVC toolsets (14.29, 14.39 -- 14.44+ unaffected) past whatever bug threshold causes the "block_topk_air<KeyT,0,0,...>" failure: a worker_policy silently read back as zero-initialized deep in CC dispatch's NTTP-based policy resolution. Drop both constructors and go back to a pure aggregate, exactly matching the old array-based design's shape. Construction sites (make_baseline_policy() and the three benchmark call sites) now list Capacity elements plus an explicit count, instead of relying on the initializer-list constructor to infer it. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Rather than staying a pure aggregate (requiring callers to redundantly
spell out the element count alongside the element list), try a variadic
constructor instead of the initializer_list<T> one. The pack is spliced
directly into elems's mem-initializer (`elems{static_cast<T>(us)...}`) with
no runtime loop, and initializer_list -- a distinct class type with its own
machinery -- is never involved. This is the next-most-targeted hypothesis
for the MSVC 14.29/14.39 constant-folding failure, now that a pure
aggregate is confirmed to fix it and a loop-free but initializer_list-based
constructor is confirmed not to.
Also add a single-job CI override (ci/matrix.yaml) pinned to the exact
failing configuration (CUB build_nolid, CTK12.0, MSVC 2019 aka 14.29,
C++17) so this can be iterated on without waiting ~40 minutes for the full
matrix each time. Must be reset to an empty override before merging.
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
The targeted CUB build_nolid / CTK12.0 / MSVC2019(14.29) / C++17 job now passes with the parameter-pack constructor. Restore the full PR matrix to confirm the other previously-failing MSVC toolset/CTK combinations (14.39, and CTK12.9/13.0/13.3 paired with 14.29) and the rest of the project. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Now confirmed via CI that the actual culprit was the initializer_list constructor, not these asserts (guarding them earlier had zero effect). Drop the now-unnecessary _CCCL_COMPILER(MSVC, >=, 19, 44) gate and go back to plain _CCCL_ASSERT/_CCCL_VERIFY everywhere, unconditionally. Re-enable the single-job CI override (CUB build_nolid, CTK12.0, MSVC2019, C++17) to verify this revert against the known-problematic configuration without spending a full matrix run; reset to empty again once confirmed. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
The revert of the MSVC-version guard (back to plain _CCCL_ASSERT/_CCCL_VERIFY) passed on the targeted job (CUB build_nolid, CTK12.0, MSVC2019/14.29, C++17). Restore the full PR matrix. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
bernhardmgruber
force-pushed
the
topk-dynamic-worker-policies
branch
from
September 30, 2026 20:38
b195c8c to
b97afb3
Compare
Contributor
😬 CI Workflow Results🟥 Finished in 2h 49m: Pass: 98%/202 | Total: 9d 15h | Max: 2h 48m | Hits: 41%/1888462See results here. AI failure analysis1. Batched TopK policy constructors leave members uninitialized · 1 jobExplanation: The PR replaces the fixed `std::array` member with a default-constructible `structural_inplace_vector`, exposing implicit policy construction that does not initialize `multi_worker_per_segment_policy`, `backend`, or `cluster`. Clang-tidy treats these diagnostics as errors, and the same header diagnostics recur across translation units. Evidence: Copy this prompt into a coding agentJobs: 2. CUDA Yum repository metadata objects return HTTP 404 · 1 jobExplanation: The enabled CUDA repository referenced three repodata objects that returned HTTP 404 across all five retries and multiple CDN addresses. The job failed during generic GCC/ccache installation before either Python wheel build began, indicating a repository or container-metadata failure rather than a source-code failure. Evidence: Copy this prompt into a coding agentJobs: |
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
baseline_topk_policy::worker_per_segment_policieswas a fixedcuda::std::array<worker_policy, 6>, forcing every tuning to specify exactly 6 worker policies. We want a dynamic count instead.std::vectorwon't work since it does not work in C++17, neither does a custom container type relying on dynamic allocation.The actual motivation to do this now is because @anikaj-eng requires dynamic sub policies as well for a new segmented sort implementation, so I am experimenting here whether
inplace_vectoris feasible.Co-Authored-By: Claude Sonnet 5 noreply@anthropic.com