Skip to content

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

Merged
bernhardmgruber merged 9 commits into
NVIDIA:mainfrom
bernhardmgruber:use_tuning_segmented_radix
May 18, 2026
Merged

bernhardmgruber merged 9 commits into
NVIDIA:mainfrom
bernhardmgruber:use_tuning_segmented_radix

Conversation

@bernhardmgruber

@bernhardmgruber bernhardmgruber commented May 12, 2026 •

Copy link
Copy Markdown
Contributor

Fixes: #8926

@bernhardmgruber
bernhardmgruber requested review from a team as code owners May 12, 2026 15:58
@bernhardmgruber
bernhardmgruber requested a review from gonidelis May 12, 2026 15:58
@bernhardmgruber
bernhardmgruber requested a review from pauleonix May 12, 2026 15:58
@github-project-automation github-project-automation Bot moved this to Todo in CCCL May 12, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL May 12, 2026
@bernhardmgruber
bernhardmgruber marked this pull request as draft May 12, 2026 15:59
@copy-pr-bot

copy-pr-bot Bot commented May 12, 2026

Copy link
Copy Markdown
Contributor

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.

@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Review to In Progress in CCCL May 12, 2026
@github-actions

This comment has been minimized.

@bernhardmgruber
bernhardmgruber force-pushed the use_tuning_segmented_radix branch from ea0103c to 41f97ab Compare May 15, 2026 19:56
@bernhardmgruber
bernhardmgruber marked this pull request as ready for review May 15, 2026 21:55
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL May 15, 2026
@coderabbitai

coderabbitai Bot commented May 15, 2026 •

Copy link
Copy Markdown
Contributor

Review Change Stack

Note

Reviews paused

It 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 reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review
📝 Walkthrough

Walkthrough

important: Migrates segmented-radix-sort dispatch to derive policy from a tuning environment; public overloads forward the environment; benchmark updated to call DeviceSegmentedRadixSort::SortKeys; adds tuning tests that assert selected block sizes.

Changes

Segmented Radix Sort Tuning API

Layer / File(s) Summary
Dispatch API refactor to environment-driven policy selection
cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh
Internal dispatch template accepts TuningEnvT (default ::cuda::std::execution::env<>), derives policy_selector_t from it, validates via static_assert, constructs kernel source/launcher internally, and computes active_policy by invoking the derived selector.
Public DeviceSegmentedRadixSort overloads forward tuning environment
cub/cub/device/device_segmented_radix_sort.cuh
All environment-based overloads (SortPairs/SortPairsDescending and SortKeys/SortKeysDescending; pointer and DoubleBuffer variants) use dispatch_with_env and forward tuning_env plus a decomposer {} into segmented_radix_sort::dispatch.
Benchmark refactor to use public SortKeys API
cub/benchmarks/bench/segmented_radix_sort/keys.cu
seg_radix_sort drops the named type-list parameter, removes unused segment/value aliases and value buffer wiring, simplifies key-pointer setup, and switches temp-storage sizing and execution to call DeviceSegmentedRadixSort::SortKeys directly.
Comprehensive tuning tests for block-size selection
cub/test/catch2_test_device_segmented_radix_sort_env.cu
Adds #if TEST_LAUNCH == 0 tuning tests with a tuning-policy functor that forces threads_per_block; parameterized tests (64,128) verify the tuned blockDim.x for SortPairs/SortPairsDescending/SortKeys/SortKeysDescending and DoubleBuffer variants by recording d_block_size[0].

Assessment against linked issues

Objective Addressed Explanation
Use the new tuning API internally for detail::segmented_radix_sort::dispatch [#8926] ✅

Suggested reviewers

  • pauleonix
  • oleksandr-pavlyk

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

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 2


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: ba363737-cffc-471b-83a3-af69cfc27682

📥 Commits

Reviewing files that changed from the base of the PR and between db90838 and 3c3a5f9.

📒 Files selected for processing (4)
  • cub/benchmarks/bench/segmented_radix_sort/keys.cu
  • cub/cub/device/device_segmented_radix_sort.cuh
  • cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh
  • cub/test/catch2_test_device_segmented_radix_sort_env.cu

@bernhardmgruber
bernhardmgruber force-pushed the use_tuning_segmented_radix branch 2 times, most recently from 33dd692 to 0466d9a Compare May 15, 2026 22:39
using offset_t = OffsetT;
using begin_offset_it_t = const offset_t*;
using end_offset_it_t = const offset_t*;
using segment_size_t = offset_t;

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This causes the SASS diffs and I think the benchmark definition is wrong. The public API only allows a segment size type of int32.

@bernhardmgruber
bernhardmgruber force-pushed the use_tuning_segmented_radix branch from 0466d9a to 2e2fdeb Compare May 16, 2026 00:01
@github-actions

This comment has been minimized.

@bernhardmgruber
bernhardmgruber force-pushed the use_tuning_segmented_radix branch from 2e2fdeb to 4817007 Compare May 18, 2026 07:19
@bernhardmgruber
bernhardmgruber enabled auto-merge (squash) May 18, 2026 07:47
@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 2h 31m: Pass: 100%/285 | Total: 3d 13h | Max: 59m 30s | Hits: 65%/329558

See results here.

@bernhardmgruber
bernhardmgruber merged commit ac84829 into NVIDIA:main May 18, 2026
308 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL May 18, 2026
@bernhardmgruber
bernhardmgruber deleted the use_tuning_segmented_radix branch May 18, 2026 09:53
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

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

3 participants