Allow public tuning of the `cub::DeviceReduce::ReduceByKey` and `cub::DeviceRunLengthEncode::Encode` by bernhardmgruber · Pull Request #9329 · NVIDIA/cccl

bernhardmgruber · 2026-06-09T14:30:51Z

New public entities

Entity	Enumerators / Data members
`cub::ReduceByKeyPolicy`	`int threads_per_block` `int items_per_thread` `BlockLoadAlgorithm load_algorithm` `CacheLoadModifier load_modifier` `BlockScanAlgorithm scan_algorithm` `LookbackDelayPolicy lookback_delay`
`cub::RleEncodePolicy`	`int threads_per_block` `int items_per_thread` `BlockLoadAlgorithm load_algorithm` `CacheLoadModifier load_modifier` `BlockScanAlgorithm scan_algorithm` `LookbackDelayPolicy lookback_delay`

copy-pr-bot · 2026-06-09T14:30:54Z

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-06-09T14:54:23Z

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

important:

Walkthrough

Public ReduceByKeyPolicy and RLEEncodePolicy types added; agent policy implementation moved to detail::agent_reduce_by_key_policy with deprecated alias; selectors, policy_hub, dispatch, kernels, DeviceRunLengthEncode, docs, benchmarks, and tests updated to new contracts.

Changes

ReduceByKey tuning API productization

Layer / File(s)	Summary
Core policy type definitions and agent integration `cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh`, `cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh`, `cub/cub/agent/agent_reduce_by_key.cuh`	Add `ReduceByKeyPolicy` and `RLEEncodePolicy` structs; introduce `detail::agent_reduce_by_key_policy` and keep `AgentReduceByKeyPolicy` as a deprecated alias.
Selectors and policy_hub wiring `cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh`, `cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh`	Selector concepts and policy_hub builders updated to construct/return `ReduceByKeyPolicy`/`RLEEncodePolicy`; `lookback_delay` populated from hub delay-constructor types.
RLE encode unification and device adapter `cub/cub/device/device_run_length_encode.cuh`, `cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh`	RLE encode selectors and hub updated to return `RLEEncodePolicy`; DeviceRunLengthEncode maps RLE selectors to `ReduceByKeyPolicy` via an adapter and passes it to reduce-by-key dispatch.
Dispatch layer updates and deprecation `cub/cub/device/dispatch/dispatch_reduce_by_key.cuh`, `cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh`	Kernel and sizing logic use `policy.lookback_delay.*`; `DispatchReduceByKey` annotated deprecated; streaming dispatcher uses `ReduceByKeyPolicy`.
Documentation, benchmarks, and tests `cub/cub/device/device_reduce.cuh`, `cub/benchmarks/...`, `cub/test/*`	Add Doxygen tuning docs; benchmark and test policy-selector return types updated; tests add compile-time checks and new env-based coverage; a test TU suppresses deprecated API warnings.

Assessment against linked issues

Objective	Addressed	Explanation
Expose public tuning policy for `DeviceReduce::ReduceByKey` [`#9326`]	✅
Define policy selector concept and infrastructure [`#9326`]	✅
Update internal dispatch and kernel usage to new policy [`#9326`]	✅
Document tuning API and provide examples [`#9326`]	✅

suggestion:

shwina
gevtushenko
srinivasyadav18

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

coderabbitai

Actionable comments posted: 1

🧹 Nitpick comments (1)

cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh (1)

34-42: ⚡ Quick win

suggestion: Use repository Doxygen form for the new public policy type. This is a new public API surface, but the type comment is plain //! text rather than the repository’s //! @brief`` style, and the field descriptions are only trailing comments. Converting this block to the standard Doxygen form will make the generated API docs consistent.

Source: Coding guidelines

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 7dcce4ca-632f-4559-a542-3f091b2ec8f8

📥 Commits

Reviewing files that changed from the base of the PR and between 3632bcc and 9db6c98.

📒 Files selected for processing (11)

cub/benchmarks/bench/reduce/by_key.cu
cub/benchmarks/bench/run_length_encode/encode.cu
cub/cub/agent/agent_reduce_by_key.cuh
cub/cub/device/device_reduce.cuh
cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh
cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu
cub/test/catch2_test_device_reduce_env.cu
cub/test/catch2_test_device_run_length_encode_env.cu

Fixes: NVIDIA#9326

coderabbitai

Actionable comments posted: 1

🧹 Nitpick comments (2)

cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh (2)

36-37: ⚡ Quick win

suggestion: switch this doc line to Doxygen //! @brief ... format to match the repository documentation rule.

As per coding guidelines, "Use Doxygen-style //! comments with @brief tags for documentation."

Source: Coding guidelines

61-67: ⚡ Quick win

suggestion: annotate the new host stream operator with _CCCL_HOST_API for consistency with function annotation requirements in CUB headers.

As per coding guidelines, "All functions must be marked with _CCCL_HOST_API, _CCCL_DEVICE_API, or _CCCL_API annotation."

Source: Coding guidelines

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 5b3c0ab5-dba5-4e0c-8213-21f089580977

📥 Commits

Reviewing files that changed from the base of the PR and between d00c348 and 9ddc842.

📒 Files selected for processing (13)

cub/benchmarks/bench/reduce/by_key.cu
cub/benchmarks/bench/run_length_encode/encode.cu
cub/cub/agent/agent_reduce_by_key.cuh
cub/cub/device/device_reduce.cuh
cub/cub/device/device_run_length_encode.cuh
cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh
cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh
cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh
cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu
cub/test/catch2_test_device_reduce_by_key_env_api.cu
cub/test/catch2_test_device_reduce_env.cu
cub/test/catch2_test_device_run_length_encode_env.cu

✅ Files skipped from review due to trivial changes (1)

cub/cub/device/device_reduce.cuh

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

cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh
cub/test/catch2_test_device_run_length_encode_env.cu
cub/benchmarks/bench/run_length_encode/encode.cu
cub/test/catch2_test_device_reduce_env.cu
cub/test/catch2_test_device_reduce_by_key_custom_policy_hub.cu
cub/benchmarks/bench/reduce/by_key.cu
cub/cub/agent/agent_reduce_by_key.cuh
cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh

coderabbitai

🧹 Nitpick comments (1)

cub/test/catch2_test_device_run_length_encode_env.cu (1)

306-306: ⚡ Quick win

suggestion: Test category "[reduce][device]" is inconsistent with all other tests in this file, which use "[run_length_encode][device]". Although RLEEncodePolicy is related to reduce-by-key, keeping tags uniform within a test file aids filtering and discoverability. Consider changing to "[run_length_encode][device]" or adding both tags.

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 7fa4eff3-bf1f-4427-80a8-a0c28da32845

📥 Commits

Reviewing files that changed from the base of the PR and between 9ddc842 and 22cfd55.

📒 Files selected for processing (1)

cub/test/catch2_test_device_run_length_encode_env.cu

github-actions · 2026-06-09T23:57:02Z

🥳 CI Workflow Results

🟩 Finished in 2h 26m: Pass: 100%/287 | Total: 10d 09h | Max: 2h 09m | Hits: 31%/676307

See results here.

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

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

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

bernhardmgruber marked this pull request as ready for review June 9, 2026 14:45

bernhardmgruber requested review from a team as code owners June 9, 2026 14:45

bernhardmgruber requested review from gevtushenko and shwina June 9, 2026 14:45

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

coderabbitai Bot reviewed Jun 9, 2026

View reviewed changes

Comment thread cub/test/catch2_test_device_reduce_env.cu Resolved

bernhardmgruber force-pushed the tune_reduce_by_key branch from 9db6c98 to d00c348 Compare June 9, 2026 15:07

NaderAlAwar reviewed Jun 9, 2026

View reviewed changes

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

Comment thread cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh Resolved

This comment has been minimized.

Sign in to view

bernhardmgruber added 5 commits June 9, 2026 20:47

Allow public tuning of the cub::DeviceReduce::ReduceByKey

556b7cb

Fixes: NVIDIA#9326

fixes

fc87f99

Fix test

e467e23

Duplicate RLEEncodePolicy

a19a30a

Add API exmple

9ddc842

bernhardmgruber force-pushed the tune_reduce_by_key branch from d00c348 to 9ddc842 Compare June 9, 2026 18:47

bernhardmgruber commented Jun 9, 2026

View reviewed changes

Comment thread cub/cub/device/dispatch/tuning/tuning_rle_encode.cuh Resolved

RLEEncodePolicy TEST

22cfd55

coderabbitai Bot reviewed Jun 9, 2026

View reviewed changes

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

coderabbitai Bot reviewed Jun 9, 2026

View reviewed changes

bernhardmgruber added 5 commits June 9, 2026 22:40

missing wrapper

270b145

fix

173f231

Rename to RLEEncodePolicy

87422c2

Fix test category

5bdd696

Fix namespace

3f47348

Fix

6015fc8

NaderAlAwar approved these changes Jun 9, 2026

View reviewed changes

bernhardmgruber enabled auto-merge (squash) June 9, 2026 21:29

bernhardmgruber changed the title ~~Allow public tuning of the cub::DeviceReduce::ReduceByKey~~ Allow public tuning of the cub::DeviceReduce::ReduceByKey and cub::DeviceRunLengthEncode::Encode Jun 9, 2026

bernhardmgruber merged commit 087c594 into NVIDIA:main Jun 9, 2026
310 checks passed

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

bernhardmgruber deleted the tune_reduce_by_key branch June 10, 2026 06:52

Navigation Menu

Repository navigation

Allow public tuning of the `cub::DeviceReduce::ReduceByKey` and `cub::DeviceRunLengthEncode::Encode` - #9329

Allow public tuning of the `cub::DeviceReduce::ReduceByKey` and `cub::DeviceRunLengthEncode::Encode`#9329
bernhardmgruber merged 12 commits into
NVIDIA:mainfrom
bernhardmgruber:tune_reduce_by_key

bernhardmgruber commented Jun 9, 2026 •

edited

Loading

Uh oh!

copy-pr-bot Bot commented Jun 9, 2026

Uh oh!

coderabbitai Bot commented Jun 9, 2026 •

edited

Loading

Reviews paused

Uh oh!

coderabbitai Bot left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

coderabbitai Bot left a comment

Uh oh!

Uh oh!

coderabbitai Bot left a comment

Uh oh!

github-actions Bot commented Jun 9, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

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

Uh oh!

New public entities

Uh oh!

copy-pr-bot Bot commented Jun 9, 2026

Uh oh!

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

Uh oh!

Reviews paused

Walkthrough

Changes

Assessment against linked issues

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Jun 9, 2026

🥳 CI Workflow Results

🟩 Finished in 2h 26m: Pass: 100%/287 | Total: 10d 09h | Max: 2h 09m | Hits: 31%/676307

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

bernhardmgruber commented Jun 9, 2026 •

edited

Loading

coderabbitai Bot commented Jun 9, 2026 •

edited

Loading