Use the new tuning API for `detail::radix_sort::dispatch` by bernhardmgruber · Pull Request #7949 · NVIDIA/cccl

bernhardmgruber · 2026-03-09T21:25:12Z

Merge before: Use tuning policy as tag for tuning environment #7835
No SASS changes for cub.bench.radix_sort.keys.base on SM75;80;90;100
No SASS changes for cub.bench.radix_sort.pairs.base on SM75;80;90;100

Fixes: #7947

copy-pr-bot · 2026-03-09T21:25:15Z

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.

coderabbitai · 2026-05-13T15:45:02Z

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: 87db7a3b-367f-48c4-a7bf-c0ca3f6a4229

📥 Commits

Reviewing files that changed from the base of the PR and between c019054 and fe72f1e.

📒 Files selected for processing (1)

cub/benchmarks/bench/radix_sort/policy_selector.h

📝 Walkthrough

Summary by CodeRabbit

Refactor
- Centralized radix-sort tuning and dispatch for more consistent, maintainable sorting behavior.
- Benchmarks simplified to call the consolidated sort APIs directly, streamlining execution and improving benchmark accounting.
Tests
- Added tests that verify tuning behavior and temporary-storage allocation differences across configurations.

important: ## Walkthrough

DeviceRadixSort dispatch was centralized into select_tuning_and_dispatch and tuning_env is threaded through decomposer and public SortKeys/SortPairs overloads; benchmarks now call DeviceRadixSort directly with env-backed allocator and updated nvbench metrics; tests were added to measure tuning-driven temporary-storage allocations.

important: ## Changes

Radix Sort tuning API integration and dispatch routing

Layer / File(s)	Summary
Core dispatch router and decomposer integration `cub/cub/device/device_radix_sort.cuh`	Add `select_tuning_and_dispatch` helper; update `radix_sort_with_decomposer` to accept/forward `tuning_env` and route decomposer dispatch through the selector.
SortPairs entry point routing `cub/cub/device/device_radix_sort.cuh`	All SortPairs overloads (ascending/descending, pointer/DoubleBuffer, env/non-env, decomposer variants) now dispatch via `select_tuning_and_dispatch`, forwarding `tuning_env` when present.
SortKeys entry point routing `cub/cub/device/device_radix_sort.cuh`	All SortKeys overloads (ascending/descending, pointer/DoubleBuffer, env/non-env, decomposer variants) now dispatch via `select_tuning_and_dispatch`, forwarding `tuning_env` when present.
Policy selector and temp-storage sizing `cub/benchmarks/bench/radix_sort/policy_selector.h`	`max_onesweep_temp_storage_size` derives `active_policy` from `policy_selector`; `max_temp_storage_size` normalizes `offset_t` via `choose_offset_t`.
Benchmark modernization to public APIs `cub/benchmarks/bench/radix_sort/keys.cu`, `cub/benchmarks/bench/radix_sort/pairs.cu`	Benchmarks allocate explicit in/out device buffers, use `caching_allocator_t` with `cub_bench_env`, update nvbench accounting, and call `DeviceRadixSort::SortKeys`/`SortPairs` directly (removed manual temp-storage sizing and internal dispatch).
Tuning behavior validation tests `cub/test/catch2_test_device_radix_sort_env.cu`	Add `tiny_onesweep_policy_selector`, `measure_allocated_bytes`, and TEST_CASEs (guarded by `TEST_LAUNCH == 0`) that assert allocation-size differences when tuning `threads_per_block`.

important: ## Assessment against linked issues

Objective	Addressed	Explanation
Use the new tuning API internally for `detail::radix_sort::dispatch` [`#7947`]	✅

Tip

💬 Introducing Slack Agent: The best way for teams to turn conversations into code.

Slack Agent is built on CodeRabbit's deep understanding of your code, so your team can collaborate across the entire SDLC without losing context.

Generate code and open pull requests
Plan features and break down work
Investigate incidents and troubleshoot customer tickets together
Automate recurring tasks and respond to alerts with triggers
Summarize progress and report instantly

Built for teams:

Shared memory across your entire org—no repeating context
Per-thread sandboxes to safely plan and execute work
Governance built-in—scoped access, auditability, and budget controls

One agent for your entire SDLC. Right inside Slack.

👉 Get started

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

coderabbitai

Actionable comments posted: 2

Caution

Some comments are outside the diff and can’t be posted inline due to platform limitations.

⚠️ Outside diff range comments (1)

cub/cub/device/device_radix_sort.cuh (1)

191-224: ⚠️ Potential issue | 🟠 Major | 🏗️ Heavy lift

Add TuningEnvT parameter to radix_sort_with_decomposer and pass it to select_tuning_and_dispatch. Currently, env-based SortPairs/SortKeys overloads for custom keys extract tuning from the environment but drop it on the floor because radix_sort_with_decomposer has no TuningEnvT parameter. This breaks cuda::execution::tune(...) for decomposer-based env APIs. Either add the parameter to thread tuning through, or switch env-based decomposer overloads to use dispatch_with_env_and_tuning like other algorithms (e.g., adjacent_difference).

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: fa87ecb6-ce9c-4097-9f0f-5f180acf4934

📥 Commits

Reviewing files that changed from the base of the PR and between 72d92b8 and b5cc930.

📒 Files selected for processing (5)

cub/benchmarks/bench/radix_sort/keys.cu
cub/benchmarks/bench/radix_sort/pairs.cu
cub/benchmarks/bench/radix_sort/policy_selector.h
cub/cub/device/device_radix_sort.cuh
cub/test/catch2_test_device_radix_sort_env.cu

coderabbitai · 2026-05-13T15:53:30Z

+  auto l = [&](auto env) {
+    auto data = c2h::device_vector<int>(10'000); // must be larger than the single tile path
+    return cub::DeviceRadixSort::SortPairs(
+      data.data().get(),
+      data.data().get(),
+      data.data().get(),
+      data.data().get(),
+      static_cast<int>(data.size()),
+      0,
+      32,
+      env);


⚠️ Potential issue | 🟠 Major | ⚡ Quick win

important: these tuning tests are invoking unsupported in-place radix-sort variants. The pointer cases alias *_in and *_out, and the DoubleBuffer cases use the same pointer for current and alternate storage, so the allocation-size assertions are currently riding on undefined behavior instead of a valid dispatch path. Use distinct input/output buffers here even if the test only cares about temporary-storage size.

Also applies to: 1168-1172, 1181-1192, 1201-1205, 1214-1218, 1227-1231, 1240-1244, 1253-1257

coderabbitai

♻️ Duplicate comments (1)

cub/test/catch2_test_device_radix_sort_env.cu (1)

1154-1164: ⚠️ Potential issue | 🟠 Major | ⚡ Quick win

important: These tuning tests still exercise unsupported in-place radix-sort paths by aliasing input/output buffers and by using DoubleBuffer with identical current/alternate pointers. Examples include Line 1157-Line 1160, Line 1176, Line 1235, Line 1410-Line 1413, and Line 1466-Line 1469. This can make temp-storage-size assertions depend on undefined behavior instead of a valid dispatch path. Use distinct input/output buffers for pointer overloads, distinct current/alternate buffers for each DoubleBuffer, and separate key/value storage for pairs.

As per coding guidelines: cub/**/*: focus on algorithm correctness and temporary-storage protocol.

Also applies to: 1174-1178, 1187-1198, 1207-1211, 1220-1224, 1233-1237, 1246-1250, 1259-1263, 1274-1278, 1287-1291, 1330-1334, 1343-1347, 1405-1415, 1462-1474, 1484-1494

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: a6307174-f001-4fc0-a1a7-6a2821c0d302

📥 Commits

Reviewing files that changed from the base of the PR and between b5cc930 and 8b996b2.

📒 Files selected for processing (2)

cub/cub/device/device_radix_sort.cuh
cub/test/catch2_test_device_radix_sort_env.cu

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

cub/cub/device/device_radix_sort.cuh

Fixes: NVIDIA#7947

coderabbitai

Actionable comments posted: 1

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 65f2b8ec-e67b-4638-8c7c-b53df36b2f05

📥 Commits

Reviewing files that changed from the base of the PR and between 8b996b2 and c019054.

📒 Files selected for processing (5)

cub/benchmarks/bench/radix_sort/keys.cu
cub/benchmarks/bench/radix_sort/pairs.cu
cub/benchmarks/bench/radix_sort/policy_selector.h
cub/cub/device/device_radix_sort.cuh
cub/test/catch2_test_device_radix_sort_env.cu

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

cub/benchmarks/bench/radix_sort/pairs.cu
cub/benchmarks/bench/radix_sort/keys.cu
cub/test/catch2_test_device_radix_sort_env.cu
cub/cub/device/device_radix_sort.cuh

github-actions · 2026-05-13T21:16:27Z

🥳 CI Workflow Results

🟩 Finished in 1h 40m: Pass: 100%/283 | Total: 4d 00h | Max: 1h 38m | Hits: 98%/199284

See results here.

github-project-automation Bot added this to CCCL Mar 9, 2026

github-project-automation Bot moved this to Todo in CCCL Mar 9, 2026

cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Mar 9, 2026

bernhardmgruber changed the title ~~Use the new tuning API for DeviceRadixSort~~ Use the new tuning API for detail::radix_sort::dispatch Mar 9, 2026

bernhardmgruber force-pushed the use_tuning_radix branch 2 times, most recently from e24e2be to c23f8b8 Compare March 25, 2026 16:57

bernhardmgruber marked this pull request as ready for review March 25, 2026 18:37

bernhardmgruber requested review from a team as code owners March 25, 2026 18:37

bernhardmgruber requested review from davebayer and oleksandr-pavlyk March 25, 2026 18:37

bernhardmgruber commented Mar 25, 2026

View reviewed changes

Comment thread cub/benchmarks/bench/segmented_radix_sort/keys.cu Resolved

cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Mar 25, 2026

bernhardmgruber marked this pull request as draft March 25, 2026 18:39

cccl-authenticator-app Bot moved this from In Review to In Progress in CCCL Mar 25, 2026

This comment has been minimized.

Sign in to view

bernhardmgruber force-pushed the use_tuning_radix branch 2 times, most recently from a4e3081 to dcbf736 Compare March 25, 2026 22:36

bernhardmgruber commented Mar 25, 2026

View reviewed changes

Comment thread cub/benchmarks/bench/radix_sort/keys.cu Resolved

bernhardmgruber commented Mar 25, 2026

View reviewed changes

Comment thread cub/benchmarks/bench/radix_sort/keys.cu Resolved

bernhardmgruber force-pushed the use_tuning_radix branch from d77e959 to 43a6707 Compare March 25, 2026 23:13

miscco approved these changes Mar 30, 2026

View reviewed changes

Comment thread cub/cub/device/device_segmented_radix_sort.cuh Resolved

bernhardmgruber force-pushed the use_tuning_radix branch from 43a6707 to db584e2 Compare March 31, 2026 06:39

bernhardmgruber force-pushed the use_tuning_radix branch from db584e2 to dc279a8 Compare April 7, 2026 08:03

bernhardmgruber force-pushed the use_tuning_radix branch from dc279a8 to ea0103c Compare May 12, 2026 15:41

bernhardmgruber mentioned this pull request May 12, 2026

Use the new tuning API internally for detail::segmented_radix_sort::dispatch #8927

Merged

3 tasks done

bernhardmgruber force-pushed the use_tuning_radix branch 2 times, most recently from 01e69f3 to d959589 Compare May 13, 2026 15:39

bernhardmgruber marked this pull request as ready for review May 13, 2026 15:41

coderabbitai Bot reviewed May 13, 2026

View reviewed changes

This comment has been minimized.

Sign in to view

bernhardmgruber added 16 commits May 13, 2026 20:41

Use the new tuning API for DeviceRadixSort

cb1a5fc

Fixes: NVIDIA#7947

fixes

ca6a185

device layer

64027db

fix arch_id

6f8e250

Fixes

c6d5422

refactor

37e8dca

Fix benchmark

59ac484

fix

71579d7

refactor

0c76b96

revert

efd4668

revert

6dc095e

device layer

7ba0ab8

more tests

51960b2

drop comments

f7b5459

fix unused tuning_env

056327f

Reduce SMEM

c019054

bernhardmgruber force-pushed the use_tuning_radix branch from 8b996b2 to c019054 Compare May 13, 2026 18:42

coderabbitai Bot reviewed May 13, 2026

View reviewed changes

Comment thread cub/benchmarks/bench/radix_sort/policy_selector.h Resolved

use selector over hub

fe72f1e

davebayer approved these changes May 14, 2026

View reviewed changes

bernhardmgruber merged commit f03d350 into NVIDIA:main May 14, 2026
302 of 304 checks passed

bernhardmgruber deleted the use_tuning_radix branch May 14, 2026 21:21

github-project-automation Bot moved this from In Review to Done in CCCL May 14, 2026

coderabbitai Bot mentioned this pull request May 21, 2026

Small segmented radix sort cleanup #9092

Merged

gonidelis mentioned this pull request May 26, 2026

[BUG]: RadixSortKeys benchmark contains an unexpected cudaMalloc in the critical benchmark region #9132

Closed

Navigation Menu

Repository navigation

Use the new tuning API for `detail::radix_sort::dispatch` - #7949

Use the new tuning API for `detail::radix_sort::dispatch`#7949
bernhardmgruber merged 17 commits into
NVIDIA:mainfrom
bernhardmgruber:use_tuning_radix

bernhardmgruber commented Mar 9, 2026 •

edited

Loading

Uh oh!

copy-pr-bot Bot commented Mar 9, 2026

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

Uh oh!

Uh oh!

coderabbitai Bot commented May 13, 2026 •

edited

Loading

Summary by CodeRabbit

Uh oh!

coderabbitai Bot left a comment

Uh oh!

Uh oh!

coderabbitai Bot May 13, 2026

Uh oh!

coderabbitai Bot left a comment

Uh oh!

This comment has been minimized.

coderabbitai Bot left a comment

Uh oh!

Uh oh!

github-actions Bot commented May 13, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

bernhardmgruber commented Mar 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

copy-pr-bot Bot commented Mar 9, 2026

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

Uh oh!

Uh oh!

coderabbitai Bot commented May 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

coderabbitai Bot May 13, 2026

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

This comment has been minimized.

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

github-actions Bot commented May 13, 2026

🥳 CI Workflow Results

🟩 Finished in 1h 40m: Pass: 100%/283 | Total: 4d 00h | Max: 1h 38m | Hits: 98%/199284

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

bernhardmgruber commented Mar 9, 2026 •

edited

Loading

coderabbitai Bot commented May 13, 2026 •

edited

Loading