Tuning policy cleanup 2/2#9712
Conversation
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (27)
✅ Files skipped from review due to trivial changes (20)
🚧 Files skipped from review as they are similar to previous changes (6)
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesThis change documents CUB policy and dispatch deprecations since version 3.5, updates guidance to the tuning APIs, moves scan-by-key policy implementation into Deprecation and policy migration
Possibly related PRs
Suggested reviewers: Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (4)
cub/test/catch2_test_device_adjacent_difference_env.cu (1)
277-281: 📐 Maintainability & Code Quality | 🔵 Trivialsuggestion: This
to_stringlambda is duplicated verbatim across ~19 catch2 policy-property test files in this cohort. Consider hoisting it into a shared test header (e.g. acatch2_test_helper.hor similar) to avoid copy-paste drift.
[optional_and_nitpick]cub/test/catch2_test_device_memcpy_env.cu (1)
244-251: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: the
to_stringlambda plusREQUIRE(!to_string(...).empty())boilerplate is duplicated verbatim across nearly every one of the ~20 catch2 policy test files touched in this cohort. Consider extracting a shared helper (e.g.to_string_via_ostreamin a common test utility header) to avoid this repeated pattern.cub/test/catch2_test_device_radix_sort_env.cu (1)
1569-1579: 🎯 Functional Correctness | 🔵 Trivial | ⚡ Quick winsuggestion:
p1_scan/p2_scan(cub::ScanPolicy) is built and used in the aggregate comparison viap1 == p2, but unlike every other sub-policy here it gets no dedicatedSTATIC_REQUIRE(p1_scan == p2_scan)and noREQUIRE(!to_string(p1_scan).empty()).ScanPolicyhas its ownoperator<</operator==overloads (seetuning_radix_sort.cuh), so this leaves that sub-policy's streaming/equality contract untested in isolation, unlike histogram/exclusive_sum/onesweep/downsweep/alt_downsweep/upsweep/alt_upsweep/single_tile.Add missing scan sub-policy checks
STATIC_REQUIRE(p1_onesweep == p2_onesweep); STATIC_REQUIRE_FALSE(p1_onesweep != p2_onesweep); + STATIC_REQUIRE(p1_scan == p2_scan); + STATIC_REQUIRE_FALSE(p1_scan != p2_scan); + STATIC_REQUIRE(p1_downsweep == p2_downsweep);REQUIRE(!to_string(p1_onesweep).empty()); + REQUIRE(!to_string(p1_scan).empty()); REQUIRE(!to_string(p1_downsweep).empty());Also applies to: 1613-1622, 1675-1717
cub/test/catch2_test_device_segmented_sort_keys_env.cu (1)
669-674: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚖️ Poor tradeoffsuggestion: the
to_stringlambda (stream tostd::ostringstream, return.str()) is duplicated verbatim across ~20 catch2 policy test files in this cohort. Consider hoisting it into a shared test utility header to avoid the copy-paste drift risk.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: c935cb35-4f69-4b9d-8aae-c6bc67ce8749
📒 Files selected for processing (46)
cub/cub/agent/agent_adjacent_difference.cuhcub/cub/agent/agent_radix_sort_downsweep.cuhcub/cub/agent/agent_radix_sort_histogram.cuhcub/cub/agent/agent_radix_sort_onesweep.cuhcub/cub/agent/agent_radix_sort_upsweep.cuhcub/cub/agent/agent_scan.cuhcub/cub/agent/agent_scan_by_key.cuhcub/cub/agent/agent_select_if.cuhcub/cub/device/dispatch/dispatch_adjacent_difference.cuhcub/cub/device/dispatch/dispatch_batch_memcpy.cuhcub/cub/device/dispatch/dispatch_histogram.cuhcub/cub/device/dispatch/dispatch_merge_sort.cuhcub/cub/device/dispatch/dispatch_radix_sort.cuhcub/cub/device/dispatch/dispatch_reduce.cuhcub/cub/device/dispatch/dispatch_reduce_by_key.cuhcub/cub/device/dispatch/dispatch_scan.cuhcub/cub/device/dispatch/dispatch_scan_by_key.cuhcub/cub/device/dispatch/dispatch_segmented_radix_sort.cuhcub/cub/device/dispatch/dispatch_segmented_reduce.cuhcub/cub/device/dispatch/dispatch_segmented_sort.cuhcub/cub/device/dispatch/dispatch_select_if.cuhcub/cub/device/dispatch/dispatch_three_way_partition.cuhcub/cub/device/dispatch/dispatch_unique_by_key.cuhcub/cub/device/dispatch/tuning/tuning_merge_sort.cuhcub/cub/device/dispatch/tuning/tuning_scan_by_key.cuhcub/test/catch2_test_device_adjacent_difference_env.cucub/test/catch2_test_device_find_bound_sorted_values_env.cucub/test/catch2_test_device_find_env.cucub/test/catch2_test_device_for_env.cucub/test/catch2_test_device_histogram_env.cucub/test/catch2_test_device_memcpy_env.cucub/test/catch2_test_device_merge_env.cucub/test/catch2_test_device_merge_sort_env.cucub/test/catch2_test_device_partition_env.cucub/test/catch2_test_device_radix_sort_env.cucub/test/catch2_test_device_reduce_env.cucub/test/catch2_test_device_run_length_encode_env.cucub/test/catch2_test_device_scan_by_key_custom_policy_hub.cucub/test/catch2_test_device_scan_by_key_env.cucub/test/catch2_test_device_scan_env.cucub/test/catch2_test_device_segmented_radix_sort_env.cucub/test/catch2_test_device_segmented_reduce_env.cucub/test/catch2_test_device_segmented_scan_env.cucub/test/catch2_test_device_segmented_sort_keys_env.cucub/test/catch2_test_device_select_env.cucub/test/catch2_test_device_transform_env.cu
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
policy_cleanup_2
branch
from
4058816 to
18cc85c
Compare
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
policy_cleanup_2
branch
from
18cc85c to
01871df
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (1)
cub/cub/device/dispatch/dispatch_scan_by_key.cuh (1)
686-855: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚖️ Poor tradeoffFree
dispatch()body duplicates__invoke/temp-storage logic verbatim.Line-by-line matches the class method's temp-storage sizing/aliasing and kernel launch sequence (385-502), so no drift today, but two independently-maintained copies of the decoupled-lookback temp-storage protocol is a risk if one is patched without the other. Given this mirrors the same transitional pattern used elsewhere in the ongoing tuning-policy migration (class retained only for the deprecated alias), leaving as-is for now, but worth collapsing once
DispatchScanByKeyis removed in CCCL 4.0.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: c401991c-6d50-4c2b-9479-9f253c7d4a46
📒 Files selected for processing (28)
cub/cub/agent/agent_adjacent_difference.cuhcub/cub/agent/agent_radix_sort_downsweep.cuhcub/cub/agent/agent_radix_sort_histogram.cuhcub/cub/agent/agent_radix_sort_onesweep.cuhcub/cub/agent/agent_radix_sort_upsweep.cuhcub/cub/agent/agent_scan.cuhcub/cub/agent/agent_scan_by_key.cuhcub/cub/agent/agent_select_if.cuhcub/cub/device/dispatch/dispatch_adjacent_difference.cuhcub/cub/device/dispatch/dispatch_batch_memcpy.cuhcub/cub/device/dispatch/dispatch_histogram.cuhcub/cub/device/dispatch/dispatch_merge_sort.cuhcub/cub/device/dispatch/dispatch_radix_sort.cuhcub/cub/device/dispatch/dispatch_reduce.cuhcub/cub/device/dispatch/dispatch_reduce_by_key.cuhcub/cub/device/dispatch/dispatch_scan.cuhcub/cub/device/dispatch/dispatch_scan_by_key.cuhcub/cub/device/dispatch/dispatch_segmented_radix_sort.cuhcub/cub/device/dispatch/dispatch_segmented_reduce.cuhcub/cub/device/dispatch/dispatch_segmented_sort.cuhcub/cub/device/dispatch/dispatch_select_if.cuhcub/cub/device/dispatch/dispatch_three_way_partition.cuhcub/cub/device/dispatch/dispatch_unique_by_key.cuhcub/cub/device/dispatch/tuning/tuning_merge_sort.cuhcub/cub/device/dispatch/tuning/tuning_scan_by_key.cuhcub/test/catch2_test_device_partition_env.cucub/test/catch2_test_device_reduce_env.cucub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu
✅ Files skipped from review due to trivial changes (18)
- cub/cub/agent/agent_radix_sort_upsweep.cuh
- cub/cub/device/dispatch/dispatch_adjacent_difference.cuh
- cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
- cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
- cub/cub/agent/agent_select_if.cuh
- cub/cub/agent/agent_radix_sort_onesweep.cuh
- cub/cub/agent/agent_radix_sort_histogram.cuh
- cub/cub/device/dispatch/dispatch_segmented_sort.cuh
- cub/cub/agent/agent_scan.cuh
- cub/cub/agent/agent_adjacent_difference.cuh
- cub/cub/device/dispatch/dispatch_merge_sort.cuh
- cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh
- cub/cub/device/dispatch/dispatch_histogram.cuh
- cub/cub/device/dispatch/dispatch_unique_by_key.cuh
- cub/cub/agent/agent_radix_sort_downsweep.cuh
- cub/cub/device/dispatch/dispatch_segmented_reduce.cuh
- cub/cub/device/dispatch/dispatch_reduce.cuh
- cub/cub/device/dispatch/dispatch_radix_sort.cuh
🚧 Files skipped from review as they are similar to previous changes (7)
- cub/cub/device/dispatch/dispatch_batch_memcpy.cuh
- cub/cub/agent/agent_scan_by_key.cuh
- cub/cub/device/dispatch/dispatch_three_way_partition.cuh
- cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu
- cub/cub/device/dispatch/dispatch_scan.cuh
- cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
- cub/cub/device/dispatch/dispatch_select_if.cuh
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: 1
🧹 Nitpick comments (1)
cub/cub/device/dispatch/dispatch_scan_by_key.cuh (1)
686-855: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚖️ Poor tradeoffFree
dispatch()body duplicates__invoke/temp-storage logic verbatim.Line-by-line matches the class method's temp-storage sizing/aliasing and kernel launch sequence (385-502), so no drift today, but two independently-maintained copies of the decoupled-lookback temp-storage protocol is a risk if one is patched without the other. Given this mirrors the same transitional pattern used elsewhere in the ongoing tuning-policy migration (class retained only for the deprecated alias), leaving as-is for now, but worth collapsing once
DispatchScanByKeyis removed in CCCL 4.0.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: c401991c-6d50-4c2b-9479-9f253c7d4a46
📒 Files selected for processing (28)
cub/cub/agent/agent_adjacent_difference.cuhcub/cub/agent/agent_radix_sort_downsweep.cuhcub/cub/agent/agent_radix_sort_histogram.cuhcub/cub/agent/agent_radix_sort_onesweep.cuhcub/cub/agent/agent_radix_sort_upsweep.cuhcub/cub/agent/agent_scan.cuhcub/cub/agent/agent_scan_by_key.cuhcub/cub/agent/agent_select_if.cuhcub/cub/device/dispatch/dispatch_adjacent_difference.cuhcub/cub/device/dispatch/dispatch_batch_memcpy.cuhcub/cub/device/dispatch/dispatch_histogram.cuhcub/cub/device/dispatch/dispatch_merge_sort.cuhcub/cub/device/dispatch/dispatch_radix_sort.cuhcub/cub/device/dispatch/dispatch_reduce.cuhcub/cub/device/dispatch/dispatch_reduce_by_key.cuhcub/cub/device/dispatch/dispatch_scan.cuhcub/cub/device/dispatch/dispatch_scan_by_key.cuhcub/cub/device/dispatch/dispatch_segmented_radix_sort.cuhcub/cub/device/dispatch/dispatch_segmented_reduce.cuhcub/cub/device/dispatch/dispatch_segmented_sort.cuhcub/cub/device/dispatch/dispatch_select_if.cuhcub/cub/device/dispatch/dispatch_three_way_partition.cuhcub/cub/device/dispatch/dispatch_unique_by_key.cuhcub/cub/device/dispatch/tuning/tuning_merge_sort.cuhcub/cub/device/dispatch/tuning/tuning_scan_by_key.cuhcub/test/catch2_test_device_partition_env.cucub/test/catch2_test_device_reduce_env.cucub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu
✅ Files skipped from review due to trivial changes (18)
- cub/cub/agent/agent_radix_sort_upsweep.cuh
- cub/cub/device/dispatch/dispatch_adjacent_difference.cuh
- cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh
- cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
- cub/cub/agent/agent_select_if.cuh
- cub/cub/agent/agent_radix_sort_onesweep.cuh
- cub/cub/agent/agent_radix_sort_histogram.cuh
- cub/cub/device/dispatch/dispatch_segmented_sort.cuh
- cub/cub/agent/agent_scan.cuh
- cub/cub/agent/agent_adjacent_difference.cuh
- cub/cub/device/dispatch/dispatch_merge_sort.cuh
- cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh
- cub/cub/device/dispatch/dispatch_histogram.cuh
- cub/cub/device/dispatch/dispatch_unique_by_key.cuh
- cub/cub/agent/agent_radix_sort_downsweep.cuh
- cub/cub/device/dispatch/dispatch_segmented_reduce.cuh
- cub/cub/device/dispatch/dispatch_reduce.cuh
- cub/cub/device/dispatch/dispatch_radix_sort.cuh
🚧 Files skipped from review as they are similar to previous changes (7)
- cub/cub/device/dispatch/dispatch_batch_memcpy.cuh
- cub/cub/agent/agent_scan_by_key.cuh
- cub/cub/device/dispatch/dispatch_three_way_partition.cuh
- cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu
- cub/cub/device/dispatch/dispatch_scan.cuh
- cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh
- cub/cub/device/dispatch/dispatch_select_if.cuh
🛑 Comments failed to post (1)
cub/cub/device/dispatch/dispatch_scan_by_key.cuh (1)
886-899: 📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win
important: wrong API name in deprecation message.
CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceScan")should referenceDeviceScanByKey, notDeviceScan. This is inconsistent with the analogous message inagent_scan_by_key.cuh("Use the tuning API for DeviceScanByKey") and with the correct log strings elsewhere in this same file (lines 612, 718).🩹 Proposed fix
using DispatchScanByKey - CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceScan") = detail::scan_by_key::dispatch_scan_by_key< + CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceScanByKey") = detail::scan_by_key::dispatch_scan_by_key<📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.using DispatchScanByKey CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceScanByKey") = detail::scan_by_key::dispatch_scan_by_key< KeysInputIteratorT, ValuesInputIteratorT, ValuesOutputIteratorT, EqualityOp, ScanOpT, InitValueT, OffsetT, AccumT, PolicyHub, PolicySelector, KernelSource, KernelLauncherFactory>;
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
force-pushed
the
policy_cleanup_2
branch
from
01871df to
b488e45
Compare
This comment has been minimized.
This comment has been minimized.
🥳 CI Workflow Results🟩 Finished in 1h 28m: Pass: 100%/287 | Total: 3d 11h | Max: 1h 09m | Hits: 93%/215664See results here. |
3 participants
Fixes part of #9378