Skip to content

Rename warp_threads in tuning policies#9485

Merged
bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:rename_warp_threads
Jun 16, 2026
Merged

Rename warp_threads in tuning policies#9485
bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:rename_warp_threads

Conversation

@bernhardmgruber

Copy link
Copy Markdown
Contributor

Prefer threads_per_warp for consistency with other tuning policy members.

@bernhardmgruber bernhardmgruber requested a review from a team as a code owner June 16, 2026 08:39
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Jun 16, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jun 16, 2026
@coderabbitai

coderabbitai Bot commented Jun 16, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

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: 25c584d9-421e-420f-a64f-173e9750e053

📥 Commits

Reviewing files that changed from the base of the PR and between 910bea4 and a0fc5e0.

📒 Files selected for processing (4)
  • cub/cub/device/dispatch/kernels/kernel_segmented_reduce.cuh
  • cub/cub/device/dispatch/kernels/kernel_segmented_sort.cuh
  • cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh
  • cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh

Note: CodeRabbit is enabled on this repository as a convenience for maintainers
and contributors. Use your best judgment when considering its review comments and
suggestions — a suggested change may be inadequate, unnecessary, or safe to ignore.
Contributors are not expected to address every comment. Human reviews are what
ultimately matter for merging.

Overview

This pull request establishes naming consistency across tuning policies by renaming the warp_threads member to threads_per_warp. The change propagates through both tuning policy definitions and their kernel-level consumers.

Changes

Tuning Policies

tuning_segmented_reduce.cuh and tuning_segmented_sort.cuh

Core member renames in the tuning policy structs:

  • cub::detail::segmented_reduce::warp_reduce_policy::warp_threadsthreads_per_warp
  • cub::detail::segmented_sort::sub_warp_merge_sort_policy::warp_threadsthreads_per_warp

Dependent logic updated to use the renamed member:

  • items_per_tile() and segments_per_block() methods now reference threads_per_warp
  • Equality operators (operator==) updated to compare the renamed member
  • Host stream output (operator<<) updated to emit .threads_per_warp

Kernel Implementations

kernel_segmented_reduce.cuh and kernel_segmented_sort.cuh

Kernels updated to pass threads_per_warp instead of warp_threads when instantiating policy templates:

  • AgentWarpReducePolicy in both DeviceSegmentedReduceKernel and DeviceFixedSizeSegmentedReduceKernel
  • AgentSubWarpMergeSortPolicy in medium-segment and small-kernel variants

Impact

  • Public API Changes: Struct member renamed in warp_reduce_policy and sub_warp_merge_sort_policy
  • Code Size: 21 lines modified across 4 files
  • Functional Impact: No changes to control flow or algorithm behavior—purely a naming standardization

Walkthrough

Renames the warp_threads struct member to threads_per_warp in warp_reduce_policy (segmented reduce tuning) and sub_warp_merge_sort_policy (segmented sort tuning), updating accessor methods, equality operators, and stream output. Propagates the rename to four CUDA kernel files that consume these policy fields when instantiating agent templates.

Changes

warp_threads → threads_per_warp rename

Layer / File(s) Summary
Policy struct member rename
cub/cub/device/dispatch/tuning/tuning_segmented_reduce.cuh, cub/cub/device/dispatch/tuning/tuning_segmented_sort.cuh
warp_reduce_policy::warp_threads renamed to threads_per_warp; sub_warp_merge_sort_policy::warp_threads renamed to threads_per_warp. Both structs update segments_per_block(), items_per_tile(), operator==, and host operator<< to use the new name.
Kernel consumer updates
cub/cub/device/dispatch/kernels/kernel_segmented_reduce.cuh, cub/cub/device/dispatch/kernels/kernel_segmented_sort.cuh
All four kernels (DeviceSegmentedReduceKernel, DeviceFixedSizeSegmentedReduceKernel, DeviceSegmentedSortFallbackKernel, DeviceSegmentedSortKernelSmall) updated to read .threads_per_warp instead of .warp_threads when instantiating AgentWarpReducePolicy / AgentSubWarpMergeSortPolicy and when computing small_threads_per_warp / medium_threads_per_warp constexprs.

Suggested reviewers

  • shwina
  • NaderAlAwar

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

@bernhardmgruber bernhardmgruber enabled auto-merge (squash) June 16, 2026 10:43
@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 2h 23m: Pass: 100%/287 | Total: 11d 03h | Max: 2h 23m | Hits: 19%/964027

See results here.

@bernhardmgruber bernhardmgruber merged commit 90ddbc6 into NVIDIA:main Jun 16, 2026
309 checks passed
@bernhardmgruber bernhardmgruber deleted the rename_warp_threads branch June 16, 2026 11:07
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.

2 participants