[FEA]: Add runtime-width group merge-sort primitive for segmented sort - #11030
[FEA]: Add runtime-width group merge-sort primitive for segmented sort#11030roro2006 wants to merge 1 commit into
Conversation
1c36eeb to
fdf5478
Compare
|
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:
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangessuggestion: The PR adds a runtime-width, stable segmented merge-sort primitive for keys and key-value pairs. It supports partial tiles, sentinels, custom comparators, shared storage, and runtime dispatch. Group merge sort
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 6
🧹 Nitpick comments (3)
cub/cub/detail/group_merge_sort.cuh (3)
107-114: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: rename the
_Unrolltemplate parameter.An identifier that starts with an underscore followed by an uppercase letter is reserved for the implementation in every scope. The project style also requires PascalCase for template parameter names. Rename to
Unrolland update the uses at lines 257, 316, 517, 530, 612.The private helpers
Sync,StoreKeys,ExchangeItems,SortPartialTile, andMergeRoundsare internal tocub::detail, so snake_case applies to them; only the publicSort/StableSortsurface keeps PascalCase.As per coding guidelines: "Use PascalCase for template parameter names" and "Use snake_case for all other symbols, except that the CUB public API uses PascalCase."
Source: Coding guidelines
229-237: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: mark the accessors
[[nodiscard]]andnoexcept.Both functions return a value and have no side effects.
As per coding guidelines: "Most functions with a non-void return type should use
[[nodiscard]]" and "Functions that do not throw exceptions must be markednoexcept."Proposed change
- _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int get_member_tid() const + [[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int get_member_tid() const noexcept { return member_tid; } - _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int get_group_id() const + [[nodiscard]] _CCCL_DEVICE _CCCL_FORCEINLINE unsigned int get_group_id() const noexcept { return group_id; }Source: Coding guidelines
16-19: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: remove
cub/util_ptx.cuhandcuda/std/__utility/swap.h;group_merge_sort.cuhuses neither header directly.MergePathis already declared by the includedcub/block/block_merge_sort.cuh.Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: f06af7a5-c990-49a1-af2f-adcff877c701
📒 Files selected for processing (2)
cub/cub/detail/group_merge_sort.cuhcub/test/catch2_test_group_merge_sort.cu
Included review availability: Your plan provides up to 12 included reviews per hour; 11 remain after this review.
fdf5478 to
794ac94
Compare
There was a problem hiding this comment.
Actionable comments posted: 2
♻️ Duplicate comments (1)
cub/cub/detail/group_merge_sort.cuh (1)
451-460: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick winimportant: barrier ids still collide when a CTA hosts 16 groups.
1u + (group_id % 15u)mapsgroup_id0 andgroup_id15 to barrier 1. This path runs only forthreads_per_group > 32, and 1024 CTA threads withthreads_per_group == 64produce exactly 16 groups. Both groups then executebar.sync 1, 64, so 128 threads arrive at one barrier that expects 64 arrivals. The barrier releases mixed groups, which breaks the shared-memory ordering inmerge_roundsand can also hang the CTA.Either derive the id from an assert-checked group count, or document and enforce a maximum of 15 multi-warp groups per CTA.
Proposed change
else { // Named barrier 0 is reserved for __syncthreads(); use ids 1..15. + // Distinct groups must map to distinct ids, so a CTA may host at most 15 multi-warp groups. + _CCCL_ASSERT(group_id < 15u, "GroupMergeSort supports at most 15 multi-warp groups per CTA"); const unsigned int barrier_id = 1u + (group_id % 15u);
🧹 Nitpick comments (2)
cub/test/catch2_test_group_merge_sort.cu (2)
130-223: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Four pair delegates are never used, so the key-value paths they cover stay untested.
Only
group_sort_pairs_partial_tile_tis launched.group_sort_pairs_full_tile_t,group_sort_pairs_partial_tile_no_sentinel_t,group_sort_pairs_descending_t, andgroup_sort_call_helper_pairs_tare dead code. The PR states that key-value runtime dispatch is covered; no test exercises the key-valuecall_group_merge_runtimeoverload, the full-tile pairSort, the no-sentinel pairSort, or descending pair ordering.Add pair tests that reuse
group_merge_sort_pairs_kernelwith these delegates, or remove the unused delegates.Also applies to: 647-695
736-764: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: The TempStorage test does not check the tail beyond
valid_items.The test truncates both vectors to
valid_itemsbefore comparison, so it does not assert anything about positions>= valid_items. The no-sentinelSortoverload only guarantees the firstvalid_itemspositions, so the truncation is correct. State that guarantee in a comment so a later reader does not restore a full-range comparison.Also consider parameterizing
GROUP_THREADSandvalid_items(including 0 andTILE_ITEMS) instead of the single hardcoded 180, to cover the boundary cases of the partial-tile path.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: b02482d5-fb29-4250-ace8-1c91403b0853
📒 Files selected for processing (2)
cub/cub/detail/group_merge_sort.cuhcub/test/catch2_test_group_merge_sort.cu
Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review.
| template <typename CompareOp> | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void Sort(KeyT (&keys)[ITEMS_PER_THREAD], CompareOp compare_op) | ||
| { | ||
| Sort(keys, compare_op, default_group_threads); | ||
| } | ||
|
|
||
| template <typename CompareOp> | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void Sort(KeyT (&keys)[ITEMS_PER_THREAD], CompareOp compare_op, int valid_items) | ||
| { | ||
| Sort(keys, compare_op, default_group_threads, valid_items); | ||
| } | ||
|
|
||
| template <typename CompareOp> | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void | ||
| Sort(KeyT (&keys)[ITEMS_PER_THREAD], CompareOp compare_op, int valid_items, KeyT oob_default) | ||
| { | ||
| Sort(keys, compare_op, default_group_threads, valid_items, oob_default); | ||
| } | ||
|
|
||
| template <typename CompareOp> | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void | ||
| Sort(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&items)[ITEMS_PER_THREAD], CompareOp compare_op) | ||
| { | ||
| Sort(keys, items, compare_op, default_group_threads); | ||
| } | ||
|
|
||
| template <typename CompareOp> | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void | ||
| Sort(KeyT (&keys)[ITEMS_PER_THREAD], ValueT (&items)[ITEMS_PER_THREAD], CompareOp compare_op, int valid_items) | ||
| { | ||
| Sort(keys, items, compare_op, default_group_threads, valid_items); | ||
| } | ||
|
|
||
| template <typename CompareOp> | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void | ||
| Sort(KeyT (&keys)[ITEMS_PER_THREAD], | ||
| ValueT (&items)[ITEMS_PER_THREAD], | ||
| CompareOp compare_op, | ||
| int valid_items, | ||
| KeyT oob_default) | ||
| { | ||
| Sort(keys, items, compare_op, default_group_threads, valid_items, oob_default); | ||
| } |
There was a problem hiding this comment.
🎯 Functional Correctness | 🔴 Critical | ⚡ Quick win
🔎 Supported by static analysis
🏁 Script executed:
#!/bin/bash
# List all Sort declarations in the header to confirm the colliding signatures.
fd -t f 'group_merge_sort.cuh' cub | xargs -r rg -n -A4 'void[[:space:]]*$|Sort\(' -g '*.cuh'Repository: NVIDIA/cccl
Length of output: 7582
critical: Rename or remove the default-width Sort overloads because they redeclare existing members. The overload at line 384 matches line 259, and the key-value overload at line 405 matches line 322. When KeyT is int, the sentinel overloads also match lines 281 and 341. C++ cannot overload on parameter names, so the class template is ill-formed. Rename all conflicting entry points; changing only one overload is insufficient.
794ac94 to
95b100c
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
♻️ Duplicate comments (1)
cub/cub/detail/group_merge_sort.cuh (1)
370-412: 🎯 Functional Correctness | 🔴 Critical | ⚡ Quick wincritical: the convenience overloads still redeclare existing members, so the class template does not compile.
- Line 377
Sort(KeyT (&keys)[ITEMS_PER_THREAD], CompareOp, int valid_items)has the same signature as line 248Sort(KeyT (&keys)[ITEMS_PER_THREAD], CompareOp, int threads_per_group).- Line 398 has the same signature as line 313.
- Line 384 has the same signature as line 272 when
KeyTisint, and line 405 matches line 333 in the same case.Parameter names do not participate in overloading. Rename these entry points (for example
SortDefaultWidth) or remove them.
🧹 Nitpick comments (2)
cub/cub/detail/group_merge_sort.cuh (2)
640-644: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: align the template parameter order of the dispatch helpers.
The keys-only helpers use
<KeyT, ITEMS_PER_THREAD, MAX_GROUP_THREADS, ValueT, CompareOp>. The key-value helpers use<KeyT, ValueT, ITEMS_PER_THREAD, MAX_GROUP_THREADS, CompareOp>. A caller that setsMAX_GROUP_THREADSexplicitly must remember two different orders. Use one order for all overloads.Also applies to: 765-769
152-159: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valuesuggestion: apply the required function annotations.
The guidelines require
_CCCL_DEVICE_API(or another_CCCL_*_APImacro) instead of bare_CCCL_DEVICE _CCCL_FORCEINLINE, and requirenoexcepton functions that do not throw. The members here, includingsyncat line 428 and theSortoverloads, use_CCCL_DEVICE _CCCL_FORCEINLINEand omitnoexcept; only the two getters at lines 225 and 230 are markednoexcept.As per coding guidelines: "Functions must be marked with
_CCCL_HOST_API,_CCCL_DEVICE_API,_CCCL_HOST_DEVICE_API,_CCCL_TILE_API, or_CCCL_API" and "Functions that do not throw exceptions must be markednoexcept".Also applies to: 428-428
Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 6bb18625-9429-47c2-ac84-7360a8974acd
📒 Files selected for processing (2)
cub/cub/detail/group_merge_sort.cuhcub/test/catch2_test_group_merge_sort.cu
Included review availability: Your plan provides up to 12 included reviews per hour; 9 remain after this review.
95b100c to
63ec31d
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 170ecd76-6a54-4d37-b4e4-e8d58043716a
📒 Files selected for processing (1)
cub/cub/detail/group_merge_sort.cuh
Included review availability: Your plan provides up to 12 included reviews per hour; 8 remain after this review.
63ec31d to
7993b3d
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (2)
cub/cub/detail/group_merge_sort.cuh (2)
271-282: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winsuggestion: document the barrier requirement before temporary storage is reused.
The last shared-memory read of a
Sortcall happens insideserial_merge, and no group synchronization follows it. A thread can return fromSortwhile other group members still readkeys_shared. A caller that reusestemp_storage(or the raw shared slice) directly afterSortreturns then races.Add the requirement to the class doc: call the same group synchronization used by the group width before reusing the storage.
444-455: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winsuggestion: Do not use
__syncthreads()as the multi-warp fallback inGroupMergeSort::sync.
merge_rounds<true>stops whensize >= valid_items, so groups with differentvalid_itemsvalues can callsynca different number of times. The CTA-wide fallback can then deadlock when one group reaches a barrier that another group no longer reaches.Use the named barrier on supported targets. Reject multi-warp groups on targets that cannot provide it.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 5e1d28da-a5c1-4120-b3c0-0d494a14323f
📒 Files selected for processing (1)
cub/cub/detail/group_merge_sort.cuh
Included review availability: Your plan provides up to 12 included reviews per hour; 7 remain after this review.
| /** | ||
| * @brief Sorts the first valid_items keys across threads_per_group threads using an out-of-bounds default sentinel. | ||
| * | ||
| * @tparam CompareOp Functor type having member `bool operator()(KeyT lhs, KeyT rhs)` (Strict Weak Ordering). | ||
| * @param[in,out] keys Keys array held by this thread. | ||
| * @param[in] compare_op Comparison function object returning true if lhs < rhs. | ||
| * @param[in] threads_per_group Number of cooperating threads in this group (must be a power of two). | ||
| * @param[in] valid_items Total number of valid keys in this group segment. | ||
| * @param[in] oob_default Sentinel value ordered after any valid key in the segment. | ||
| */ | ||
| template <typename CompareOp> | ||
| _CCCL_DEVICE_API _CCCL_FORCEINLINE void | ||
| Sort(KeyT (&keys)[ITEMS_PER_THREAD], | ||
| CompareOp compare_op, | ||
| int threads_per_group, | ||
| int valid_items, | ||
| KeyT oob_default) noexcept | ||
| { | ||
| static_assert(KEYS_ONLY, "Keys-only Sort requires ValueT == NullType; use the key-value overload instead"); | ||
| _CCCL_ASSERT(threads_per_group > 0, "threads_per_group must be greater than 0"); | ||
| _CCCL_ASSERT(::cuda::is_power_of_two(threads_per_group), "threads_per_group must be a power of two"); | ||
| _CCCL_ASSERT(valid_items >= 0, "valid_items must be non-negative"); | ||
| ValueT items[ITEMS_PER_THREAD]; | ||
| sort_partial_tile<true>(keys, items, compare_op, threads_per_group, valid_items, oob_default); | ||
| } |
There was a problem hiding this comment.
🗄️ Data Integrity & Integration | 🟡 Minor | ⚡ Quick win
suggestion: document that tile positions at or after valid_items are unspecified.
sort_partial_tile applies the oob_default padding only inside if (ITEMS_PER_THREAD * member_tid < valid_items). A thread whose whole item range is out of bounds keeps its caller-supplied keys. merge_rounds<true> clamps every run to valid_items, so those keys never reach a valid output position, but they also are not replaced by oob_default. In addition, threads that over-consume a clamped run write duplicated keys into output slots at or after valid_items.
A caller that migrates from BlockMergeSort::Sort(keys, compare_op, valid_items, oob_default) can expect the tail of the tile to hold oob_default. State the actual guarantee in the doc block of this overload and of the key-value sentinel overload at lines 387-412.
NVIDIA#10944) Address CodeRabbit review suggestions: - Enforce strict <= 15 multi-warp groups per CTA with bar.sync (1 + group_id) and assertions. - Eliminate unsafe __syncthreads() fallback in sync() on unsupported targets to avoid CTA deadlock. - Validate threads_per_group > 0, <= MAX_GROUP_THREADS, and power-of-two in constructors and dispatch helpers. - Document that tile output positions at or after valid_items are unspecified. - Document storage lifetime and synchronization requirements prior to TempStorage reuse. - Remove redundant convenience Sort overloads to eliminate signature redeclaration conflicts. - Align template parameter ordering across keys-only and key-value dispatch helpers. - Apply _CCCL_DEVICE_API and noexcept annotations to all methods according to guidelines. - Add test coverage for all key-value pair delegates and boundary conditions. - Apply clang-format and verify all pre-commit formatting checks pass.
7993b3d to
f54953b
Compare
1 participant
Adds a reusable runtime-width group merge-sort primitive (
cub::detail::GroupMergeSortandcub::detail::call_group_merge_runtime)supporting:
- Sub-warp (
< 32), warp (== 32), and multi-warp (> 32) group widths selected at runtime.- Non-interfering boundary-safe synchronization using sub-warp masks and hardware named barriers (
bar.sync).- Runtime
valid_itemsclamping with early termination of redundant merge passes.- Full support for keys-only and key-value pairs across arbitrary Strict Weak Ordering comparators.
- Complete Catch2 unit test coverage in
cub/test/catch2_test_group_merge_sort.cu.*I'm not the most experienced, just saw the issue and thought it would be interesting, so any advice would be greatly appreciated