Skip to content

Use signed offset type for DevicePartition#8971

Merged
davebayer merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:partition_offset_t_fix
May 13, 2026
Merged

Use signed offset type for DevicePartition#8971
davebayer merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:partition_offset_t_fix

Conversation

@bernhardmgruber
Copy link
Copy Markdown
Contributor

The env-overloads in cub::DevicePartition used unsigned offset types, while the non-env-overloads used signed ones. This was presumably an oversight in #7459.

Fixes: #8968

The env-overloads in cub::DevicePartition used unsigned offset types, while the non-env-overloads used signed ones. This was presumably an oversight in NVIDIA#7459.

Fixes: NVIDIA#8968
@bernhardmgruber bernhardmgruber requested a review from a team as a code owner May 13, 2026 18:33
@bernhardmgruber bernhardmgruber requested a review from elstehle May 13, 2026 18:33
@github-project-automation github-project-automation Bot moved this to Todo in CCCL May 13, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL May 13, 2026
@coderabbitai
Copy link
Copy Markdown

coderabbitai Bot commented May 13, 2026

Review Change Stack

📝 Walkthrough

Summary by CodeRabbit

  • Bug Fixes
    • Improved offset computation in DevicePartition's Flagged and If APIs for better handling of large datasets.

Walkthrough

cub::DevicePartition environment-based Flagged and If overloads are updated to use signed offset type selection (detail::choose_signed_offset<NumItemsT>) instead of unsigned, matching the behavior of non-environment overloads and adding the ChooseOffsetT::is_exceeding_offset_type bounds check.

Changes

Offset type alignment

Layer / File(s) Summary
Offset type switch in Flagged and If overloads
cub/cub/device/device_partition.cuh
Flagged (lines 324–331) and If (lines 575–582) environment-based overloads switch from unsigned detail::choose_offset_t<NumItemsT> to signed detail::choose_signed_offset<NumItemsT>. Both now perform the exceeding-offset-type bounds check using the signed offset type, aligning with non-environment overloads.

Assessment against linked issues

Objective Addressed Explanation
Unify env-overloads to use signed offset types (#8968)

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

Copy link
Copy Markdown

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

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

🧹 Nitpick comments (1)
cub/cub/device/device_partition.cuh (1)

324-331: ⚡ Quick win

suggestion: Add regression tests for env overload offset-bound behavior.
The signed offset switch and is_exceeding_offset_type guard look aligned with the non-env path; add targeted tests that cover boundary (max signed offset) and overflow (max + 1 for unsigned NumItemsT) inputs for both env Flagged and env If overloads to prevent regressions.

As per coding guidelines, "cub/**/*: Focus on algorithm correctness, temporary-storage protocol, dispatch/policy selection, stream behavior, CUDA error handling, synchronization, memory access safety, performance regressions, and test coverage."

Also applies to: 575-582


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 2d52e5ac-8795-4344-b77f-a99a02b6f277

📥 Commits

Reviewing files that changed from the base of the PR and between e49bdfa and 98092ec.

📒 Files selected for processing (1)
  • cub/cub/device/device_partition.cuh

@github-actions
Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 3h 28m: Pass: 100%/283 | Total: 6d 16h | Max: 1h 53m | Hits: 53%/430269

See results here.

@davebayer davebayer merged commit e15d87a into NVIDIA:main May 13, 2026
602 of 607 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL May 13, 2026
@github-actions
Copy link
Copy Markdown
Contributor

Successfully created backport PR for branch/3.4.x:

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

cub::DevicePartition env-overloads using unsigned offset type while non-env overloads use signed one

2 participants