Skip to content

[CUB] Vectorize DeviceReduce loads for any contiguous iterator - #12326

Merged
bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:reduce-contiguous-vectorization
Oct 31, 2026
Merged

bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:reduce-contiguous-vectorization

Conversation

@bernhardmgruber

Copy link
Copy Markdown
Contributor

AgentReduce only vectorized loads for raw pointers. It now unwraps any contiguous iterator (e.g. thrust::device_ptr, device_vector iterators) and applies the cache modifier on top of the resulting pointer, resolving the TODO in agent_reduce.cuh.

🤖 Generated with Claude Code

@bernhardmgruber
bernhardmgruber requested a review from a team as a code owner October 29, 2026 08:56
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Oct 29, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Oct 29, 2026
@coderabbitai

coderabbitai Bot commented Oct 29, 2026 •

Copy link
Copy Markdown
Contributor

Review in Change Stack →

Important

Review skipped

The saved review base belongs to an older reviewed commit. This saved history cannot establish the base for an incremental review. Comment @coderabbitai full review to establish a new review baseline. No full review was started, and the last reviewed checkpoint was preserved.

You can disable this status message by setting the reviews.review_status to false in the CodeRabbit configuration file.

Use the checkbox below for a quick retry:

  • 🔍 Trigger review
📝 Summary

Summary by CodeRabbit

  • Performance
    • Expanded vectorized loading eligibility to include contiguous iterators while retaining existing compatibility checks.
    • This broadens the cases that can use vectorized loading without changing the existing requirements for policy, item count, data types, or alignment. Scalar and partial-tile loading continue to support the adapted input.

AgentReduce now adapts contiguous input iterators for cache-modified access and uses the adapted iterator for vectorization checks, alignment checks, and tile loads. The benchmark configuration also changes selected Thrust targets and enabled GPUs.

Changes: AgentReduce iterator adaptation

Layer / File(s) Summary
Adapt input iterators and use them for loads
cub/cub/agent/agent_reduce.cuh
AgentReduceImpl stores a cache-modified iterator adapted from contiguous inputs. Vectorization eligibility accepts contiguous iterators while retaining existing checks. Alignment checks and vectorized, scalar, and partial-tile loads use the adapted iterator.

Benchmark configuration

Layer / File(s) Summary
Configure benchmark targets
ci/bench.yaml
Three Thrust benchmark target filters are active. The GPU list enables gb300, t4, rtxa6000, l4, h100, and rtxpro6000; rtx2080 and rtx4090 remain commented out.

Possibly related PRs

  • NVIDIA/cccl#12049 — Adds contiguous-iterator vectorization and rebasing support in AgentReduce.
  • NVIDIA/cccl#12215 — Changes AgentReduce vectorization eligibility and load paths.
  • NVIDIA/cccl#12043 — Adds striped vectorized-load machinery and updates AgentReduce’s vectorized full-tile path.

Suggested reviewers: griwes

Change: Feature

Merge Risk: Minimal · previous review 619c8

No confirmed merge-blocking issue remains. Verify GB300 runner availability during normal CI checks.

- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

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

@github-actions

This comment has been minimized.

@miscco miscco 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.

This looks good to me, it also fixed the dropping of the cache modifier in the fallback path which is nice :)

Comment on lines +221 to +223
&& THRUST_NS_QUALIFIER::is_contiguous_iterator_v<InputIteratorT>
// TODO(bgruber): remove the check for is_primitive<ValueT> in CCCL 4.0
&& (is_primitive<InputT>::value || ::cuda::is_trivially_copyable_v<InputT>)

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.

Forgot to send: This is still allowing vectorization for volatile base types.

Should we check for that?

@bernhardmgruber bernhardmgruber Oct 29, 2026 •

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.

If volatile is emitted via the cache modifier, then this should be retained by the rewarp of the vectorized pointer. Otherwise, we drop it I believe, which is what the old code did as well.

Edit: Claude confirms

@github-actions

This comment has been minimized.

@bernhardmgruber
bernhardmgruber enabled auto-merge (squash) October 29, 2026 20:57
@github-actions

This comment has been minimized.

@bernhardmgruber

Copy link
Copy Markdown
Contributor Author

SASS diff analysis

I downloaded the SASS diff artifacts and went through them.

Segmented top-k (4 targets): there are no SASS changes in any top-k kernel. The only changed kernels in those targets are DeviceReduceKernel and DeviceReduceSingleTileKernel for const int over a thrust::device_ptr. They come from the thrust::reduce(d_segment_sizes.begin(), d_segment_sizes.end()) call in the benchmark setup (keys_common.cuh / indexed_common.cuh), which sums the segment sizes outside the timed region. These diffs can be ignored.

Reductions (thrust.bench.reduce.basic, thrust.bench.transform_reduce.{sum,unary}): the SASS change is confined to DeviceReduceKernel and DeviceReduceSingleTileKernel instantiated with device_ptr, on all archs (sm_75 to sm_120). In short:

  • Loads now get the cache modifier: LDG.E* becomes LDG.E*.CONSTANT. For __int128 this is the only change.
  • Wider loads appear: LDG.E.128.CONSTANT for narrower types, and 64-bit loads for short/char.
  • The instruction stream gets bigger, about 2-3x more lines for int. There are many more address computations, branches and remainder paths around the vectorized loop.

Register counts, spills and occupancy can't be read from the diff artifacts, so a benchmark run is needed to judge the performance impact.

Analysis done by Claude.

@bernhardmgruber
bernhardmgruber requested a review from a team as a code owner October 29, 2026 22:30

@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.

🔇 Additional comments: 1
ci/bench.yaml-48-48: 1

48-48: Major · Stability & Availability · Quick win

⚠️ Unverified finding
Verification ran but could not confirm this finding. It is shown for review, not as a verified issue.

important: Do not enable gb300 with the current benchmark runner mapping.

The workflow hard-codes linux-amd64-gpu-${{ inputs.gpu }}-latest-1, so this entry selects linux-amd64-gpu-gb300-latest-1. The repository's GB300 lanes use arm64, and the runner-label template includes the CPU architecture. This benchmark job will target the wrong runner instead of the GB300 pool. Carry the CPU architecture through benchmark dispatch, or remove gb300 until the workflow supports the arm64 label. (github.com)

As per path instructions, fix the matrix-to-runner mapping before enabling this GPU.

Source: Path instructions


Review info
Run configuration
  • Source mode: Auto
  • Repository: NVIDIA/cccl
  • Branch: main
  • File: .coderabbit.yaml
  • Commit: 5939983
  • Layers: Repository YAML
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: 8ba92aef-6f3c-40d9-b35d-0c3624823a43
Commits

Reviewing files that changed from the base of the PR and between 43fc27c and 619c8a4.

Files selected for processing: 1
  • ci/bench.yaml

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.

@github-actions

This comment has been minimized.

@bernhardmgruber

bernhardmgruber commented Oct 30, 2026 •

Copy link
Copy Markdown
Contributor Author

NVBench robust-intervals outcomes from bench run 40966440359 (base: origin/main).

GPU Benchmark FAST SAME SLOW AMBG UNKN Total
h100 reduce.basic 3 4 0 21 0 28
h100 transform_reduce.sum 3 1 0 24 0 28
h100 transform_reduce.unary 3 1 0 24 0 28
l4 reduce.basic 5 4 0 19 0 28
l4 transform_reduce.sum 1 2 0 25 0 28
l4 transform_reduce.unary 1 3 0 24 0 28
t4 reduce.basic 8 0 0 20 0 28
t4 transform_reduce.sum 2 0 0 26 0 28
t4 transform_reduce.unary 3 0 0 25 0 28
rtxa6000 reduce.basic 8 2 0 18 0 28
rtxa6000 transform_reduce.sum 9 2 0 17 0 28
rtxa6000 transform_reduce.unary 7 3 0 18 0 28
rtxpro6000 reduce.basic 0 6 0 22 0 28
rtxpro6000 transform_reduce.sum 0 1 0 27 0 28
rtxpro6000 transform_reduce.unary 0 7 0 21 0 28
All All 53 36 0 331 0 420

No regressions (SLOW) on any GPU.

🤖 Generated with Claude Code

@bernhardmgruber

bernhardmgruber commented Oct 30, 2026 •

Copy link
Copy Markdown
Contributor Author

Archived robust compare-intervals output of thrust.bench.reduce.basic.base for every GPU from bench run 40966440359 (base: origin/main), kept here in case the CI artifacts expire.

h100
# base

## [0] NVIDIA H100 NVL

|  T{ct}  |  Elements  |                         Ref |                          Cmp |    Change |  Status  |
|---------|------------|-----------------------------|------------------------------|-----------|----------|
|   I8    |    2^16    |  40.992 [-1.184, +0.384] us |  40.512  [-1.216, +0.320] us |           |   AMBG   |
|   I8    |    2^20    |  40.960 [-0.928, +0.416] us |  40.480  [-2.400, +0.576] us |           |   AMBG   |
|   I8    |    2^24    |  49.792 [-1.504, +0.480] us |  49.344  [-7.328, +0.416] us |           |   AMBG   |
|   I8    |    2^28    | 164.576 [-3.520, +4.064] us | 120.160  [-1.632, +6.208] us | <= -27.4% |   FAST   |
|   I16   |    2^16    |  40.992 [-0.864, +0.384] us |  41.088  [-2.880, +0.576] us |           |   AMBG   |
|   I16   |    2^20    |  41.248 [-1.184, +0.416] us |  41.088  [-1.152, +0.448] us |           |   AMBG   |
|   I16   |    2^24    |  58.976 [-5.120, +0.576] us |  49.888  [-2.656, +0.608] us |  <= -6.7% |   FAST   |
|   I16   |    2^28    | 217.344 [-6.720, +1.632] us | 190.720  [-2.048, +3.200] us |  <= -8.6% |   FAST   |
|   I32   |    2^16    |  41.120 [-1.216, +0.960] us |  40.736  [-0.704, +0.512] us |           |   AMBG   |
|   I32   |    2^20    |  41.536 [-3.072, +0.448] us |  40.960  [-1.024, +0.480] us |           |   AMBG   |
|   I32   |    2^24    |  67.616 [-8.864, +0.544] us |  67.584 [-14.176, +2.368] us |           |   AMBG   |
|   I32   |    2^28    | 340.224 [-7.712, +1.088] us | 332.448  [-3.232, +5.280] us |           |   AMBG   |
|   I64   |    2^16    |  40.992 [-0.736, +0.512] us |  40.608  [-0.864, +0.416] us |           |   AMBG   |
|   I64   |    2^20    |  41.184 [-0.896, +0.416] us |  40.864  [-3.392, +0.416] us |           |   AMBG   |
|   I64   |    2^24    |  85.344 [-2.592, +0.768] us |  85.088  [-2.048, +0.704] us |           |   SAME   |
|   I64   |    2^28    | 621.440 [-7.104, +1.856] us | 621.152  [-8.192, +1.632] us |           |   SAME   |
|  I128   |    2^16    |  41.120 [-0.704, +0.448] us |  40.896  [-1.280, +0.576] us |           |   AMBG   |
|  I128   |    2^20    |  50.464 [-8.192, +0.448] us |  49.792  [-8.160, +0.480] us |           |   AMBG   |
|  I128   |    2^24    | 121.024 [-2.048, +4.512] us | 120.320  [-1.472, +2.432] us |           |   AMBG   |
|  I128   |    2^28    |        1.194 ms [-9, +4] us |        1.194 ms [-10, +4] us |           |   SAME   |
|   F32   |    2^16    |  41.120 [-0.992, +0.512] us |  40.672  [-1.760, +0.576] us |           |   AMBG   |
|   F32   |    2^20    |  41.440 [-1.088, +0.416] us |  40.896  [-1.504, +0.480] us |           |   AMBG   |
|   F32   |    2^24    |  67.712 [-7.808, +0.544] us |  67.168  [-6.176, +0.512] us |           |   AMBG   |
|   F32   |    2^28    | 340.320 [-8.192, +0.992] us | 333.120  [-3.136, +4.832] us |           |   AMBG   |
|   F64   |    2^16    |  41.152 [-0.896, +0.416] us |  40.992  [-0.800, +0.384] us |           |   SAME   |
|   F64   |    2^20    |  41.376 [-2.784, +0.512] us |  40.800  [-1.152, +0.512] us |           |   AMBG   |
|   F64   |    2^24    |  85.632 [-2.624, +0.672] us |  84.928  [-2.592, +0.928] us |           |   AMBG   |
|   F64   |    2^28    | 621.984 [-7.552, +2.144] us | 621.184  [-6.432, +1.568] us |           |   AMBG   |

# Summary

- Total Matches: 28
  - Unchanged   (classified as SAME): 4
  - Improvement (clear timing gap, %Diff < 0): 3
  - Regression  (clear timing gap, %Diff > 0): 0
  - Ambiguous (comparison requires more evidence): 21
    - Reasons:
      - centers_not_close: 17 (timing centers are not close enough to declare same)
      - bulk_time_support_mismatch: 3 (sample: min(ref=97.7%, cmp=89.2%) >= 97.0%; support: min(ref=91.5%, cmp=78.5%) >= 80.0%)
      - bulk_cycle_support_mismatch: 1 (sample: min(ref=94.5%, cmp=98.6%) >= 97.0%; support: min(ref=94.5%, cmp=98.6%) >= 80.0%)
  - Unknown     (input data unavailable or unusable): 0
l4
# base

## [0] NVIDIA L4

|  T{ct}  |  Elements  |                          Ref |                          Cmp |    Change |  Status  |
|---------|------------|------------------------------|------------------------------|-----------|----------|
|   I8    |    2^16    |  41.792  [-8.320, +0.512] us |  42.624  [-2.112, +8.480] us |           |   AMBG   |
|   I8    |    2^20    |  42.720  [-1.696, +1.120] us |  43.040  [-3.328, +8.640] us |           |   AMBG   |
|   I8    |    2^24    | 138.720  [-2.944, +1.056] us | 130.304  [-3.360, +1.280] us |           |   AMBG   |
|   I8    |    2^28    |        1.560 ms [-15, +5] us |        1.235 ms [-16, +3] us | <= -24.8% |   FAST   |
|   I16   |    2^16    |  42.592  [-6.752, +0.384] us |  43.648  [-2.656, +8.352] us |           |   AMBG   |
|   I16   |    2^20    |  51.136 [-10.144, +0.544] us |  44.160  [-3.488, +8.064] us |           |   AMBG   |
|   I16   |    2^24    | 235.360  [-8.416, +0.896] us | 218.816  [-2.848, +4.768] us |           |   AMBG   |
|   I16   |    2^28    |        2.602 ms [-26, +5] us |        2.281 ms [-11, +3] us | <= -12.8% |   FAST   |
|   I32   |    2^16    |  42.880  [-1.728, +0.480] us |  43.104  [-6.496, +8.864] us |           |   AMBG   |
|   I32   |    2^20    |  60.192  [-1.888, +1.600] us |  60.352  [-4.000, +7.072] us |           |   AMBG   |
|   I32   |    2^24    | 411.264  [-8.736, +1.600] us | 404.288  [-7.552, +4.704] us |           |   AMBG   |
|   I32   |    2^28    |        4.705 ms [-11, +4] us |        4.386 ms [-12, +3] us |  <= -7.0% |   FAST   |
|   I64   |    2^16    |  42.944  [-1.856, +0.544] us |  43.328 [-2.048, +11.040] us |           |   AMBG   |
|   I64   |    2^20    |  86.368  [-4.736, +1.088] us |  86.400  [-5.376, +0.864] us |           |   SAME   |
|   I64   |    2^24    | 709.312  [-8.064, +1.632] us | 700.064  [-8.704, +2.784] us |           |   AMBG   |
|   I64   |    2^28    |        8.597 ms [-38, +6] us |        8.338 ms [-14, +8] us |  <= -2.6% |   FAST   |
|  I128   |    2^16    |  44.096  [-2.688, +5.280] us |  44.864 [-3.392, +12.576] us |           |   AMBG   |
|  I128   |    2^20    | 131.040  [-4.160, +3.808] us | 131.488  [-3.904, +1.216] us |           |   SAME   |
|  I128   |    2^24    |        1.228 ms [-11, +4] us |         1.226 ms [-8, +3] us |           |   SAME   |
|  I128   |    2^28    |      16.975 ms [-65, +15] us |      16.974 ms [-86, +13] us |           |   SAME   |
|   F32   |    2^16    |  42.848  [-3.840, +0.864] us |  42.976  [-2.272, +2.048] us |           |   AMBG   |
|   F32   |    2^20    |  60.288 [-2.144, +10.848] us |  60.448  [-2.208, +2.144] us |           |   AMBG   |
|   F32   |    2^24    | 411.168  [-8.128, +1.440] us | 403.712  [-6.720, +4.480] us |           |   AMBG   |
|   F32   |    2^28    |        4.704 ms [-11, +3] us |        4.387 ms [-21, +3] us |  <= -6.9% |   FAST   |
|   F64   |    2^16    |  42.464  [-1.728, +0.832] us |  43.072  [-1.888, +9.024] us |           |   AMBG   |
|   F64   |    2^20    |  86.336  [-3.520, +0.928] us |  86.624  [-3.200, +0.800] us |           |   AMBG   |
|   F64   |    2^24    | 711.040  [-7.488, +1.728] us | 701.152  [-7.584, +1.856] us |           |   AMBG   |
|   F64   |    2^28    |        8.606 ms [-31, +8] us |       8.343 ms [-17, +10] us |           |   AMBG   |

# Summary

- Total Matches: 28
  - Unchanged   (classified as SAME): 4
  - Improvement (clear timing gap, %Diff < 0): 5
  - Regression  (clear timing gap, %Diff > 0): 0
  - Ambiguous (comparison requires more evidence): 19
    - Reasons:
      - centers_not_close: 12 (timing centers are not close enough to declare same)
      - bulk_cycle_gap_not_confirmed: 3 (clear timing gap was not confirmed by bulk cycle intervals)
      - bulk_time_support_mismatch: 3 (sample: min(ref=94.9%, cmp=96.2%) >= 97.0%; support: min(ref=83.2%, cmp=89.8%) >= 80.0%)
      - bulk_cycle_support_mismatch: 1 (sample: min(ref=98.3%, cmp=96.8%) >= 97.0%; support: min(ref=98.3%, cmp=96.8%) >= 80.0%)
  - Unknown     (input data unavailable or unusable): 0
t4
# base

## [0] Tesla T4

|  T{ct}  |  Elements  |                          Ref |                          Cmp |    Change |  Status  |
|---------|------------|------------------------------|------------------------------|-----------|----------|
|   I8    |    2^16    |  24.320  [-0.832, +0.864] us |  20.096  [-1.024, +0.448] us |           |   AMBG   |
|   I8    |    2^20    |  30.496  [-1.280, +0.416] us |  25.600  [-0.960, +0.448] us |           |   AMBG   |
|   I8    |    2^24    | 104.032  [-1.472, +0.544] us |  93.216  [-1.664, +0.640] us |           |   AMBG   |
|   I8    |    2^28    |         1.205 ms [-8, +3] us |         1.047 ms [-5, +2] us | <= -14.1% |   FAST   |
|   I16   |    2^16    |  19.744  [-0.608, +0.320] us |  18.816  [-2.080, +0.832] us |           |   AMBG   |
|   I16   |    2^20    |  28.192  [-1.056, +0.576] us |  28.128  [-1.216, +0.320] us |           |   AMBG   |
|   I16   |    2^24    | 165.600  [-1.536, +0.448] us | 153.952  [-2.336, +0.608] us |  <= -6.1% |   FAST   |
|   I16   |    2^28    |        2.558 ms [-17, +4] us |         2.049 ms [-6, +2] us | <= -23.9% |   FAST   |
|   I32   |    2^16    |  19.840  [-1.216, +0.672] us |  22.016  [-0.768, +0.352] us |           |   AMBG   |
|   I32   |    2^20    |  40.576  [-1.184, +0.320] us |  41.248  [-1.312, +0.480] us |           |   AMBG   |
|   I32   |    2^24    | 289.792  [-2.048, +0.736] us | 278.752  [-2.624, +0.672] us |  <= -3.0% |   FAST   |
|   I32   |    2^28    |        4.662 ms [-23, +5] us |        4.017 ms [-19, +6] us | <= -15.3% |   FAST   |
|   I64   |    2^16    |  20.352  [-1.536, +0.512] us |  22.016  [-2.720, +0.736] us |           |   AMBG   |
|   I64   |    2^20    |  61.184  [-1.216, +0.480] us |  61.024  [-1.344, +0.480] us |           |   AMBG   |
|   I64   |    2^24    | 528.000  [-2.592, +0.736] us | 527.648  [-4.512, +1.408] us |           |   AMBG   |
|   I64   |    2^28    |        8.138 ms [-20, +7] us |        7.917 ms [-28, +8] us |  <= -2.4% |   FAST   |
|  I128   |    2^16    |  24.032  [-1.024, +0.416] us |  24.544  [-2.912, +0.768] us |           |   AMBG   |
|  I128   |    2^20    |  93.312 [-2.400, +10.720] us |  92.800  [-1.792, +4.832] us |           |   AMBG   |
|  I128   |    2^24    |         1.027 ms [-4, +1] us |         1.027 ms [-4, +1] us |           |   AMBG   |
|  I128   |    2^28    |      15.918 ms [-42, +12] us |      15.955 ms [-39, +10] us |           |   AMBG   |
|   F32   |    2^16    |  18.464  [-1.824, +0.704] us |  18.656  [-1.952, +0.928] us |           |   AMBG   |
|   F32   |    2^20    |  42.688  [-1.792, +0.480] us |  41.248  [-1.920, +0.384] us |           |   AMBG   |
|   F32   |    2^24    | 290.272  [-1.952, +4.128] us | 278.208  [-2.240, +0.768] us |  <= -3.3% |   FAST   |
|   F32   |    2^28    |        4.664 ms [-22, +6] us |        4.018 ms [-16, +5] us | <= -15.4% |   FAST   |
|   F64   |    2^16    |  22.464  [-1.760, +0.736] us |  23.168  [-2.432, +3.104] us |           |   AMBG   |
|   F64   |    2^20    |  64.544  [-1.536, +0.704] us |  64.096  [-1.984, +0.640] us |           |   AMBG   |
|   F64   |    2^24    | 543.680 [-10.272, +2.176] us | 539.648 [-10.848, +1.632] us |           |   AMBG   |
|   F64   |    2^28    |       8.504 ms [-60, +11] us |        7.873 ms [-17, +5] us |           |   AMBG   |

# Summary

- Total Matches: 28
  - Unchanged   (classified as SAME): 0
  - Improvement (clear timing gap, %Diff < 0): 8
  - Regression  (clear timing gap, %Diff > 0): 0
  - Ambiguous (comparison requires more evidence): 20
    - Reasons:
      - centers_not_close: 10 (timing centers are not close enough to declare same)
      - bulk_cycle_gap_not_confirmed: 5 (clear timing gap was not confirmed by bulk cycle intervals)
      - bulk_cycle_support_mismatch: 4 (sample: min(ref=51.9%, cmp=3.4%) >= 97.0%; support: min(ref=51.9%, cmp=3.4%) >= 80.0%)
      - weak_interval_overlap: 1 (timing intervals do not overlap strongly enough to declare same)
  - Unknown     (input data unavailable or unusable): 0
rtxa6000
# base

## [0] NVIDIA RTX A6000

|  T{ct}  |  Elements  |                         Ref |                         Cmp |    Change |  Status  |
|---------|------------|-----------------------------|-----------------------------|-----------|----------|
|   I8    |    2^16    |  14.848 [-0.512, +0.128] us |  14.272 [-0.736, +0.128] us |           |   AMBG   |
|   I8    |    2^20    |  17.088 [-0.640, +0.128] us |  16.544 [-0.544, +0.128] us |           |   AMBG   |
|   I8    |    2^24    |  51.520 [-0.704, +0.224] us |  47.616 [-0.768, +0.160] us |  <= -6.4% |   FAST   |
|   I8    |    2^28    | 558.048 [-2.688, +1.056] us | 452.128 [-1.504, +1.312] us | <= -22.5% |   FAST   |
|   I16   |    2^16    |  14.752 [-0.864, +0.160] us |  14.592 [-0.544, +0.096] us |           |   AMBG   |
|   I16   |    2^20    |  18.816 [-0.576, +0.160] us |  18.624 [-0.576, +0.128] us |           |   AMBG   |
|   I16   |    2^24    |  79.904 [-0.640, +0.224] us |  75.296 [-0.512, +0.192] us |  <= -5.0% |   FAST   |
|   I16   |    2^28    | 999.904 [-4.032, +1.088] us | 879.328 [-3.552, +1.056] us | <= -13.1% |   FAST   |
|   I32   |    2^16    |  14.528 [-0.448, +0.224] us |  14.496 [-0.576, +0.160] us |           |   AMBG   |
|   I32   |    2^20    |  23.072 [-0.608, +0.128] us |  23.072 [-0.544, +0.160] us |           |   SAME   |
|   I32   |    2^24    | 130.048 [-0.640, +0.288] us | 127.872 [-0.640, +0.288] us |           |   AMBG   |
|   I32   |    2^28    |        1.851 ms [-4, +1] us |        1.744 ms [-7, +1] us |  <= -5.8% |   FAST   |
|   I64   |    2^16    |  15.392 [-0.576, +0.160] us |  15.392 [-0.640, +0.192] us |           |   AMBG   |
|   I64   |    2^20    |  31.904 [-0.576, +0.192] us |  31.648 [-0.640, +0.160] us |           |   AMBG   |
|   I64   |    2^24    | 234.848 [-0.960, +0.352] us | 232.768 [-1.024, +0.256] us |           |   AMBG   |
|   I64   |    2^28    |       3.480 ms [-10, +2] us |        3.383 ms [-2, +1] us |  <= -2.6% |   FAST   |
|  I128   |    2^16    |  16.736 [-0.544, +0.224] us |  16.864 [-0.608, +0.192] us |           |   AMBG   |
|  I128   |    2^20    |  48.096 [-0.704, +0.544] us |  48.160 [-0.768, +0.544] us |           |   AMBG   |
|  I128   |    2^24    | 447.808 [-1.312, +0.352] us | 448.160 [-1.408, +0.352] us |           |   AMBG   |
|  I128   |    2^28    |       6.904 ms [-30, +6] us |       6.904 ms [-31, +5] us |           |   AMBG   |
|   F32   |    2^16    |  14.880 [-0.512, +0.192] us |  14.816 [-0.672, +0.192] us |           |   SAME   |
|   F32   |    2^20    |  23.232 [-0.608, +0.160] us |  23.712 [-0.992, +0.256] us |           |   AMBG   |
|   F32   |    2^24    | 130.560 [-0.992, +7.168] us | 128.320 [-0.768, +3.040] us |           |   AMBG   |
|   F32   |    2^28    |        1.852 ms [-5, +1] us |        1.745 ms [-5, +1] us |  <= -5.8% |   FAST   |
|   F64   |    2^16    |  16.576 [-0.576, +0.224] us |  16.480 [-0.576, +0.160] us |           |   AMBG   |
|   F64   |    2^20    |  33.408 [-0.800, +0.192] us |  33.120 [-0.640, +0.160] us |           |   AMBG   |
|   F64   |    2^24    | 238.464 [-1.568, +0.512] us | 235.904 [-1.536, +0.608] us |           |   AMBG   |
|   F64   |    2^28    |       3.486 ms [-12, +3] us |        3.386 ms [-4, +2] us |  <= -2.5% |   FAST   |

# Summary

- Total Matches: 28
  - Unchanged   (classified as SAME): 2
  - Improvement (clear timing gap, %Diff < 0): 8
  - Regression  (clear timing gap, %Diff > 0): 0
  - Ambiguous (comparison requires more evidence): 18
    - Reasons:
      - centers_not_close: 12 (timing centers are not close enough to declare same)
      - bulk_cycle_support_mismatch: 3 (sample: min(ref=96.9%, cmp=19.1%) >= 97.0%; support: min(ref=96.9%, cmp=19.2%) >= 80.0%)
      - bulk_time_support_mismatch: 2 (sample: min(ref=99.3%, cmp=93.4%) >= 97.0%; support: min(ref=95.0%, cmp=72.2%) >= 80.0%)
      - bulk_cycle_gap_not_confirmed: 1 (clear timing gap was not confirmed by bulk cycle intervals)
  - Unknown     (input data unavailable or unusable): 0
rtxpro6000
# base

## [0] NVIDIA RTX PRO 6000 Blackwell Server Edition

|  T{ct}  |  Elements  |                          Ref |                          Cmp |   Change |  Status  |
|---------|------------|------------------------------|------------------------------|----------|----------|
|   I8    |    2^16    |  47.328  [-3.776, +0.672] us |  47.296  [-4.576, +0.672] us |          |   AMBG   |
|   I8    |    2^20    |  47.392  [-5.600, +0.832] us |  47.296  [-4.576, +0.672] us |          |   AMBG   |
|   I8    |    2^24    |  69.696  [-4.000, +0.864] us |  69.504 [-10.528, +0.928] us |          |   AMBG   |
|   I8    |    2^28    | 294.656 [-10.560, +2.272] us | 281.248  [-5.984, +1.504] us |          |   AMBG   |
|   I16   |    2^16    |  47.520  [-5.088, +0.608] us |  47.616  [-5.440, +0.736] us |          |   AMBG   |
|   I16   |    2^20    |  47.424  [-5.120, +0.608] us |  47.456  [-3.840, +0.704] us |          |   AMBG   |
|   I16   |    2^24    |  84.416  [-7.584, +0.736] us |  84.480  [-8.672, +6.592] us |          |   AMBG   |
|   I16   |    2^28    | 460.192  [-8.320, +3.008] us | 456.288  [-6.560, +1.920] us |          |   AMBG   |
|   I32   |    2^16    |  46.976  [-4.544, +1.184] us |  47.104  [-7.904, +1.184] us |          |   AMBG   |
|   I32   |    2^20    |  54.272 [-11.296, +1.376] us |  54.176 [-12.416, +1.312] us |          |   AMBG   |
|   I32   |    2^24    | 113.024 [-10.880, +1.216] us | 113.056  [-9.248, +1.248] us |          |   AMBG   |
|   I32   |    2^28    | 807.936  [-8.480, +2.624] us | 806.528  [-8.032, +2.208] us |          |   SAME   |
|   I64   |    2^16    |  47.744  [-5.152, +0.416] us |  47.936  [-5.504, +0.416] us |          |   AMBG   |
|   I64   |    2^20    |  55.264  [-6.080, +0.384] us |  55.200  [-6.048, +0.384] us |          |   AMBG   |
|   I64   |    2^24    | 172.576  [-6.272, +0.896] us | 172.160  [-5.536, +1.408] us |          |   SAME   |
|   I64   |    2^28    |        1.503 ms [-10, +3] us |         1.503 ms [-8, +2] us |          |   SAME   |
|  I128   |    2^16    |  47.008  [-4.416, +1.184] us |  47.200  [-4.992, +1.120] us |          |   AMBG   |
|  I128   |    2^20    |  68.960 [-10.176, +1.184] us |  69.024  [-8.032, +1.312] us |          |   AMBG   |
|  I128   |    2^24    | 282.208  [-6.240, +1.920] us | 280.736  [-6.816, +1.440] us |          |   AMBG   |
|  I128   |    2^28    |         2.899 ms [-8, +2] us |         2.899 ms [-9, +2] us |          |   SAME   |
|   F32   |    2^16    |  49.152  [-4.064, +0.512] us |  47.424  [-4.896, +0.608] us |          |   AMBG   |
|   F32   |    2^20    |  56.224 [-11.040, +0.928] us |  54.656 [-11.072, +0.576] us |          |   AMBG   |
|   F32   |    2^24    | 111.296 [-11.456, +1.120] us | 113.088  [-8.576, +0.864] us |          |   AMBG   |
|   F32   |    2^28    | 807.680  [-7.264, +3.232] us | 805.312  [-7.424, +3.040] us |          |   SAME   |
|   F64   |    2^16    |  49.216  [-4.000, +0.608] us |  53.600 [-12.032, +0.864] us |          |   AMBG   |
|   F64   |    2^20    |  57.088  [-3.808, +0.576] us |  54.944  [-4.704, +6.112] us |          |   AMBG   |
|   F64   |    2^24    | 180.000 [-12.128, +1.120] us | 177.376  [-9.952, +1.056] us |          |   AMBG   |
|   F64   |    2^28    |         1.510 ms [-9, +3] us |        1.509 ms [-11, +2] us |          |   SAME   |

# Summary

- Total Matches: 28
  - Unchanged   (classified as SAME): 6
  - Improvement (clear timing gap, %Diff < 0): 0
  - Regression  (clear timing gap, %Diff > 0): 0
  - Ambiguous (comparison requires more evidence): 22
    - Reasons:
      - bulk_time_support_mismatch: 12 (sample: min(ref=98.4%, cmp=92.8%) >= 97.0%; support: min(ref=94.7%, cmp=79.7%) >= 80.0%)
      - centers_not_close: 9 (timing centers are not close enough to declare same)
      - bulk_cycle_support_mismatch: 1 (sample: min(ref=95.6%, cmp=95.8%) >= 97.0%; support: min(ref=95.6%, cmp=95.8%) >= 80.0%)
  - Unknown     (input data unavailable or unusable): 0

🤖 Generated with Claude Code

@bernhardmgruber
bernhardmgruber force-pushed the reduce-contiguous-vectorization branch from 619c8a4 to 43fc27c Compare October 30, 2026 20:59
Unwrap contiguous iterators in AgentReduce instead of requiring raw pointers, and apply the cache modifier on top of the unwrapped pointer.

Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
@bernhardmgruber
bernhardmgruber force-pushed the reduce-contiguous-vectorization branch from 43fc27c to 4de9b4d Compare October 30, 2026 21:00
@Jacobfaib

Copy link
Copy Markdown
Contributor

/sass diff ok 4de9b4d

@github-actions

Copy link
Copy Markdown
Contributor

Error: @Jacobfaib, the SASS approval job is not ready. Wait for the "Require SASS approval" job to finish, then post /sass diff ok 4de9b4d797b8772af32df75ab686913055d35eb8 again if it fails.

@github-actions

github-actions Bot commented Oct 30, 2026 •

Copy link
Copy Markdown
Contributor

🔬 Benchmark SASS comparison

⚠️ The SASS changed for 4 of 100 CUB and 3 of 75 Thrust benchmark target(s). A benchmark run may be necessary

How to request a benchmark run
Request a benchmark run for this PR:

1. Replace the `benchmarks:` block of ci/bench.yaml with exactly this:

benchmarks:
  filters:
    cpp:
      - '^cub\.bench\.segmented_topk\.variable\.indexed\.base$'
      - '^cub\.bench\.segmented_topk\.variable\.indexed\.cluster\.base$'
      - '^cub\.bench\.segmented_topk\.variable\.keys\.base$'
      - '^cub\.bench\.segmented_topk\.variable\.keys\.cluster\.base$'
      - '^thrust\.bench\.reduce\.basic\.base$'
      - '^thrust\.bench\.transform_reduce\.sum\.base$'
      - '^thrust\.bench\.transform_reduce\.unary\.base$'
  gpus:
    - "h100"   # pick the GPUs that this change can affect

2. Commit with `[bench-only]` at the end of the commit summary, so that
   the unrelated CI jobs are skipped. Then push.

ci/bench.yaml must match ci/bench.template.yaml before the PR can merge.
Reset it once the measurement is done.
Run Value
Baseline c7428b87e51a60cf83515ff22c17578e1e49a6e1
Tested HEAD
Architectures 75-real;80-real;90-real;100-real;110-real;120-real;120-virtual
Targets with a SASS change
Target Architectures with a SASS change
cub.bench.segmented_topk.variable.indexed.base sm_120, sm_100, sm_75, sm_110, sm_80, sm_90
cub.bench.segmented_topk.variable.indexed.cluster.base sm_120, sm_100, sm_75, sm_110, sm_80, sm_90
cub.bench.segmented_topk.variable.keys.base sm_120, sm_100, sm_75, sm_110, sm_80, sm_90
cub.bench.segmented_topk.variable.keys.cluster.base sm_120, sm_100, sm_75, sm_110, sm_80, sm_90
thrust.bench.reduce.basic.base sm_80, sm_75, sm_100, sm_110, sm_120, sm_90
thrust.bench.transform_reduce.sum.base sm_80, sm_75, sm_100, sm_110, sm_120, sm_90
thrust.bench.transform_reduce.unary.base sm_80, sm_75, sm_100, sm_110, sm_120, sm_90

‼️ Summary of Differences ‼️

Showing 7/7 summaries.

cub.bench.segmented_topk.variable.indexed.base - sm_120 - ⚠️ Significant

DeviceReduce kernels add instructions, memory operations, and control flow The diffs replace large regions of DeviceReduceKernel and DeviceReduceSingleTileKernel with more instructions, loads, stores, synchronization operations, and control instructions. This signature occurs in all four segmented_topk variable targets on all six architectures.

Showing 40/4439 diff lines, 4241 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.indexed.base.sm_120
+++ test/cub.bench.segmented_topk.variable.indexed.base.sm_120
@@ -31550,698 +31550,83 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 Function : void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 LDC R1, c[0x0][0x24] ;
 S2R R0, SR_CTAID.X ;
-LDCU.64 UR4, c[0x0][0x50] ;
-BSSY.RECONVERGENT B0, <+0x2160> ;
-S2R R14, SR_TID.X ;
-LDCU.64 UR6, c[0x0][0x0] ;
-MOV R7, UR4 ;
+LDCU UR4, c[0x0][0x28] ;
+BSSY.RECONVERGENT B0, <+0x2430> ;
+S2R R2, SR_TID.X ;
+LDCU UR5, c[0x0][0x54] ;
+LDCU.64 UR8, c[0x0][0x0] ;
+ULOP3.LUT UP0, URZ, UR4, 0xf, URZ, 0xc0, !UPT ;
 USHF.L.U32 UR5, UR5, 0xc, URZ ;
-IMAD.SHL.U32 R4, R0, 0x1000, RZ ;
-IADD R3, -R4, R7 ;
-ISETP.GT.U32.AND P0, PT, R3, 0xfff, PT ;
-@P0 BRA <+0xf60> ;
-ISETP.GE.U32.AND P0, PT, R14, R3, PT ;
+IMAD.SHL.U32 R6, R0, 0x1000, RZ ;
+BRA.U UP0, <+0x10c0> ;
+LDCU UR6, c[0x0][0x50] ;
+IADD R25, -R6, UR6 ;
+ISETP.GT.U32.AND P0, PT, R25, 0xfff, PT ;
+@P0 BRA <+0x860> ;
+ISETP.GE.U32.AND P0, PT, R2, R25, PT ;
 BSSY.RECONVERGENT B1, <+0xb0> ;
-HFMA2 R10, -RZ, RZ, 0, 0 ;
-MOV R9, R14 ;
+HFMA2 R8, -RZ, RZ, 0, 0 ;
+MOV R11, R2 ;
 @P0 BRA <+0x60> ;
-LDC.64 R10, c[0x0][0x28] ;
-IADD R9, R4, R14 ;
-IMAD.WIDE.U32 R10, R9, 0x4, R10 ;
-LDG.E R10, desc[UR6][R10.64] ;
cub.bench.segmented_topk.variable.indexed.cluster.base - sm_120 - ⚠️ Significant

DeviceReduce kernels add instructions, memory operations, and control flow The diffs replace large regions of DeviceReduceKernel and DeviceReduceSingleTileKernel with more instructions, loads, stores, synchronization operations, and control instructions. This signature occurs in all four segmented_topk variable targets on all six architectures.

Showing 40/4435 diff lines, 4233 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.indexed.cluster.base.sm_120
+++ test/cub.bench.segmented_topk.variable.indexed.cluster.base.sm_120
@@ -117474,698 +117474,83 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 Function : void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 LDC R1, c[0x0][0x24] ;
 S2R R0, SR_CTAID.X ;
-LDCU.64 UR4, c[0x0][0x50] ;
-BSSY.RECONVERGENT B0, <+0x2160> ;
-S2R R14, SR_TID.X ;
-LDCU.64 UR6, c[0x0][0x0] ;
-MOV R7, UR4 ;
+LDCU UR4, c[0x0][0x28] ;
+BSSY.RECONVERGENT B0, <+0x2430> ;
+S2R R2, SR_TID.X ;
+LDCU UR5, c[0x0][0x54] ;
+LDCU.64 UR8, c[0x0][0x0] ;
+ULOP3.LUT UP0, URZ, UR4, 0xf, URZ, 0xc0, !UPT ;
 USHF.L.U32 UR5, UR5, 0xc, URZ ;
-IMAD.SHL.U32 R4, R0, 0x1000, RZ ;
-IADD R3, -R4, R7 ;
-ISETP.GT.U32.AND P0, PT, R3, 0xfff, PT ;
-@P0 BRA <+0xf60> ;
-ISETP.GE.U32.AND P0, PT, R14, R3, PT ;
+IMAD.SHL.U32 R6, R0, 0x1000, RZ ;
+BRA.U UP0, <+0x10c0> ;
+LDCU UR6, c[0x0][0x50] ;
+IADD R25, -R6, UR6 ;
+ISETP.GT.U32.AND P0, PT, R25, 0xfff, PT ;
+@P0 BRA <+0x860> ;
+ISETP.GE.U32.AND P0, PT, R2, R25, PT ;
 BSSY.RECONVERGENT B1, <+0xb0> ;
-HFMA2 R10, -RZ, RZ, 0, 0 ;
-MOV R9, R14 ;
+HFMA2 R8, -RZ, RZ, 0, 0 ;
+MOV R11, R2 ;
 @P0 BRA <+0x60> ;
-LDC.64 R10, c[0x0][0x28] ;
-IADD R9, R4, R14 ;
-IMAD.WIDE.U32 R10, R9, 0x4, R10 ;
-LDG.E R10, desc[UR6][R10.64] ;
cub.bench.segmented_topk.variable.keys.base - sm_120 - ⚠️ Significant

DeviceReduce kernels add instructions, memory operations, and control flow The diffs replace large regions of DeviceReduceKernel and DeviceReduceSingleTileKernel with more instructions, loads, stores, synchronization operations, and control instructions. This signature occurs in all four segmented_topk variable targets on all six architectures.

Showing 40/4439 diff lines, 4241 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.keys.base.sm_120
+++ test/cub.bench.segmented_topk.variable.keys.base.sm_120
@@ -29033,698 +29033,83 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 Function : void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 LDC R1, c[0x0][0x24] ;
 S2R R0, SR_CTAID.X ;
-LDCU.64 UR4, c[0x0][0x50] ;
-BSSY.RECONVERGENT B0, <+0x2160> ;
-S2R R14, SR_TID.X ;
-LDCU.64 UR6, c[0x0][0x0] ;
-MOV R7, UR4 ;
+LDCU UR4, c[0x0][0x28] ;
+BSSY.RECONVERGENT B0, <+0x2430> ;
+S2R R2, SR_TID.X ;
+LDCU UR5, c[0x0][0x54] ;
+LDCU.64 UR8, c[0x0][0x0] ;
+ULOP3.LUT UP0, URZ, UR4, 0xf, URZ, 0xc0, !UPT ;
 USHF.L.U32 UR5, UR5, 0xc, URZ ;
-IMAD.SHL.U32 R4, R0, 0x1000, RZ ;
-IADD R3, -R4, R7 ;
-ISETP.GT.U32.AND P0, PT, R3, 0xfff, PT ;
-@P0 BRA <+0xf60> ;
-ISETP.GE.U32.AND P0, PT, R14, R3, PT ;
+IMAD.SHL.U32 R6, R0, 0x1000, RZ ;
+BRA.U UP0, <+0x10c0> ;
+LDCU UR6, c[0x0][0x50] ;
+IADD R25, -R6, UR6 ;
+ISETP.GT.U32.AND P0, PT, R25, 0xfff, PT ;
+@P0 BRA <+0x860> ;
+ISETP.GE.U32.AND P0, PT, R2, R25, PT ;
 BSSY.RECONVERGENT B1, <+0xb0> ;
-HFMA2 R10, -RZ, RZ, 0, 0 ;
-MOV R9, R14 ;
+HFMA2 R8, -RZ, RZ, 0, 0 ;
+MOV R11, R2 ;
 @P0 BRA <+0x60> ;
-LDC.64 R10, c[0x0][0x28] ;
-IADD R9, R4, R14 ;
-IMAD.WIDE.U32 R10, R9, 0x4, R10 ;
-LDG.E R10, desc[UR6][R10.64] ;
cub.bench.segmented_topk.variable.keys.cluster.base - sm_120 - ⚠️ Significant

DeviceReduce kernels add instructions, memory operations, and control flow The diffs replace large regions of DeviceReduceKernel and DeviceReduceSingleTileKernel with more instructions, loads, stores, synchronization operations, and control instructions. This signature occurs in all four segmented_topk variable targets on all six architectures.

Showing 40/4439 diff lines, 4241 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.keys.cluster.base.sm_120
+++ test/cub.bench.segmented_topk.variable.keys.cluster.base.sm_120
@@ -29033,698 +29033,83 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 Function : void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<int, unsigned int, cuda::std::__4::plus<int>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<const int>>, int *, unsigned int, unsigned int, cuda::std::__4::plus<int>, int, int, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 LDC R1, c[0x0][0x24] ;
 S2R R0, SR_CTAID.X ;
-LDCU.64 UR4, c[0x0][0x50] ;
-BSSY.RECONVERGENT B0, <+0x2160> ;
-S2R R14, SR_TID.X ;
-LDCU.64 UR6, c[0x0][0x0] ;
-MOV R7, UR4 ;
+LDCU UR4, c[0x0][0x28] ;
+BSSY.RECONVERGENT B0, <+0x2430> ;
+S2R R2, SR_TID.X ;
+LDCU UR5, c[0x0][0x54] ;
+LDCU.64 UR8, c[0x0][0x0] ;
+ULOP3.LUT UP0, URZ, UR4, 0xf, URZ, 0xc0, !UPT ;
 USHF.L.U32 UR5, UR5, 0xc, URZ ;
-IMAD.SHL.U32 R4, R0, 0x1000, RZ ;
-IADD R3, -R4, R7 ;
-ISETP.GT.U32.AND P0, PT, R3, 0xfff, PT ;
-@P0 BRA <+0xf60> ;
-ISETP.GE.U32.AND P0, PT, R14, R3, PT ;
+IMAD.SHL.U32 R6, R0, 0x1000, RZ ;
+BRA.U UP0, <+0x10c0> ;
+LDCU UR6, c[0x0][0x50] ;
+IADD R25, -R6, UR6 ;
+ISETP.GT.U32.AND P0, PT, R25, 0xfff, PT ;
+@P0 BRA <+0x860> ;
+ISETP.GE.U32.AND P0, PT, R2, R25, PT ;
 BSSY.RECONVERGENT B1, <+0xb0> ;
-HFMA2 R10, -RZ, RZ, 0, 0 ;
-MOV R9, R14 ;
+HFMA2 R8, -RZ, RZ, 0, 0 ;
+MOV R11, R2 ;
 @P0 BRA <+0x60> ;
-LDC.64 R10, c[0x0][0x28] ;
-IADD R9, R4, R14 ;
-IMAD.WIDE.U32 R10, R9, 0x4, R10 ;
-LDG.E R10, desc[UR6][R10.64] ;
thrust.bench.reduce.basic.base - sm_80 - ⚠️ Significant

DeviceReduceKernel instruction, memory, and control-flow expansion Across all three changed reduction targets and all six architectures, DeviceReduceKernel specializations replace large instruction regions and increase the instruction count. The diffs add address arithmetic, global loads, predicates, branches, and synchronization operations, so the instruction mix, memory operations, and control-flow structure differ.

Showing 40/37674 diff lines, 34802 changes. - ⬇️ Full diff

--- base/thrust.bench.reduce.basic.base.sm_80
+++ test/thrust.bench.reduce.basic.base.sm_80
@@ -242,222 +242,249 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<__int128, unsigned int, cuda::std::__4::plus<__int128>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<__int128>>, __int128 *, unsigned int, unsigned int, cuda::std::__4::plus<__int128>, __int128, __int128, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x0] ;
 S2R R0, SR_CTAID.X ;
 ULDC.64 UR6, c[0x0][0xf0] ;
-BSSY B0, <+0x3230> ;
+BSSY B0, <+0x33e0> ;
 S2R R3, SR_TID.X ;
 IMAD.SHL.U32 R28, R0, 0x400, RZ ;
 IADD3 R2, -R28, c[0x0][0x160], RZ ;
 ISETP.GE.U32.AND P0, PT, R2, 0x400, PT ;
-@!P0 BRA <+0x1770> ;
-SHF.R.S32.HI R20, RZ, 0x1f, R3 ;
-IADD3 R4, P0, R28, R3, RZ ;
-IMAD.X R5, RZ, RZ, R20, P0 ;
-LEA R24, P0, R4, c[0x0][0x138], 0x4 ;
-LEA.HI.X R25, R4, c[0x0][0x13c], R5, 0x4, P0 ;
-LDG.E.128 R4, [R24.64] ;
-LDG.E.128 R8, [R24.64+0x1000] ;
-LDG.E.128 R12, [R24.64+0x2000] ;
-LDG.E.128 R16, [R24.64+0x3000] ;
-IMAD.MOV.U32 R23, RZ, RZ, c[0x0][0x164] ;
-IMAD.SHL.U32 R23, R23, 0x400, RZ ;
+@!P0 BRA <+0x1920> ;
+IADD3 R35, R3.reuse, 0x100, RZ ;
+SHF.R.S32.HI R24, RZ, 0x1f, R3 ;
+IADD3 R5, P0, R28.reuse, R3, RZ ;
+SHF.R.S32.HI R23, RZ, 0x1f, R35 ;
+IADD3 R9, P1, R28, R35, RZ ;
+IMAD.X R16, RZ, RZ, R24, P0 ;
+IADD3 R33, R3, 0x200, RZ ;
+IADD3 R29, R3, 0x300, RZ ;
+IMAD.X R10, RZ, RZ, R23, P1 ;
+SHF.R.S32.HI R22, RZ, 0x1f, R33 ;
+IADD3 R6, P2, R28, R33, RZ ;
+LEA R4, P0, R5, c[0x0][0x138], 0x4 ;
+LEA R8, P1, R9, c[0x0][0x138], 0x4 ;
+IMAD.X R7, RZ, RZ, R22, P2 ;
+SHF.R.S32.HI R20, RZ, 0x1f, R29 ;
thrust.bench.transform_reduce.sum.base - sm_80 - ⚠️ Significant

DeviceReduceKernel instruction, memory, and control-flow expansion Across all three changed reduction targets and all six architectures, DeviceReduceKernel specializations replace large instruction regions and increase the instruction count. The diffs add address arithmetic, global loads, predicates, branches, and synchronization operations, so the instruction mix, memory operations, and control-flow structure differ.

Showing 40/48168 diff lines, 45277 changes. - ⬇️ Full diff

--- base/thrust.bench.transform_reduce.sum.base.sm_80
+++ test/thrust.bench.transform_reduce.sum.base.sm_80
@@ -242,331 +242,357 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<__int128, unsigned int, cuda::std::__4::plus<__int128>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<__int128>>, __int128 *, unsigned int, unsigned int, cuda::std::__4::plus<__int128>, __int128, __int128, square_t<__int128>>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x0] ;
 S2R R0, SR_CTAID.X ;
 ULDC.64 UR6, c[0x0][0xf0] ;
-BSSY B1, <+0x2d50> ;
+BSSY B1, <+0x2f00> ;
 S2R R5, SR_TID.X ;
-IMAD.SHL.U32 R2, R0, 0x400, RZ ;
-IADD3 R24, -R2, c[0x0][0x160], RZ ;
+IMAD.SHL.U32 R22, R0, 0x400, RZ ;
+IADD3 R24, -R22, c[0x0][0x160], RZ ;
 ISETP.GE.U32.AND P0, PT, R24, 0x400, PT ;
-@!P0 BRA <+0x17e0> ;
-SHF.R.S32.HI R3, RZ, 0x1f, R5 ;
-IADD3 R4, P0, R2, R5, RZ ;
+@!P0 BRA <+0x1980> ;
+SHF.R.S32.HI R2, RZ, 0x1f, R5 ;
+IADD3 R3, P0, R22, R5, RZ ;
+IMAD.X R4, RZ, RZ, R2, P0 ;
+LEA R8, P0, R3, c[0x0][0x138], 0x4 ;
+LEA.HI.X R9, R3, c[0x0][0x13c], R4, 0x4, P0 ;
+IADD3 R4, R5, 0x100, RZ ;
+LDG.E.128.CONSTANT R8, [R8.64] ;
+SHF.R.S32.HI R3, RZ, 0x1f, R4 ;
+IADD3 R6, P0, R22, R4, RZ ;
 IMAD.X R7, RZ, RZ, R3, P0 ;
-LEA R28, P0, R4, c[0x0][0x138], 0x4 ;
-LEA.HI.X R29, R4, c[0x0][0x13c], R7, 0x4, P0 ;
-LDG.E.128 R12, [R28.64] ;
-LDG.E.128 R8, [R28.64+0x1000] ;
-IMAD.WIDE.U32 R6, R13, R12, RZ ;
+LEA R12, P0, R6, c[0x0][0x138], 0x4 ;
+LEA.HI.X R13, R6, c[0x0][0x13c], R7, 0x4, P0 ;
+LDG.E.128.CONSTANT R12, [R12.64] ;
+IMAD.WIDE.U32 R6, R9, R8, RZ ;
+IMAD.WIDE.U32 R28, R8, R8, RZ ;
+IMAD.WIDE.U32 R6, P0, R8, R9, R6 ;
+IMAD.X R17, RZ, RZ, RZ, P0 ;
thrust.bench.transform_reduce.unary.base - sm_80 - ⚠️ Significant

DeviceReduceKernel instruction, memory, and control-flow expansion Across all three changed reduction targets and all six architectures, DeviceReduceKernel specializations replace large instruction regions and increase the instruction count. The diffs add address arithmetic, global loads, predicates, branches, and synchronization operations, so the instruction mix, memory operations, and control-flow structure differ.

Showing 40/39029 diff lines, 35818 changes. - ⬇️ Full diff

--- base/thrust.bench.transform_reduce.unary.base.sm_80
+++ test/thrust.bench.transform_reduce.unary.base.sm_80
@@ -242,23 +242,41 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<__int128, unsigned int, cuda::std::__4::plus<__int128>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, thrust::_V_300700_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300700_SM_750_800_900_1000_1100_1200::device_ptr<__int128>>, __int128 *, unsigned int, unsigned int, cuda::std::__4::plus<__int128>, __int128, int, plus_one<__int128>>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x0] ;
 S2R R0, SR_CTAID.X ;
 ULDC.64 UR6, c[0x0][0xf0] ;
-BSSY B0, <+0x35f0> ;
+BSSY B0, <+0x37a0> ;
 S2R R3, SR_TID.X ;
 IMAD.SHL.U32 R28, R0, 0x400, RZ ;
 IADD3 R2, -R28, c[0x0][0x160], RZ ;
 ISETP.GE.U32.AND P0, PT, R2, 0x400, PT ;
-@!P0 BRA <+0x1950> ;
-SHF.R.S32.HI R20, RZ, 0x1f, R3 ;
-IADD3 R4, P0, R28, R3, RZ ;
-IMAD.X R5, RZ, RZ, R20, P0 ;
-LEA R22, P0, R4, c[0x0][0x138], 0x4 ;
-LEA.HI.X R23, R4, c[0x0][0x13c], R5, 0x4, P0 ;
-LDG.E.128 R4, [R22.64] ;
-LDG.E.128 R8, [R22.64+0x1000] ;
-LDG.E.128 R12, [R22.64+0x2000] ;
-LDG.E.128 R16, [R22.64+0x3000] ;
-IMAD.MOV.U32 R29, RZ, RZ, c[0x0][0x164] ;
-IMAD.SHL.U32 R29, R29, 0x400, RZ ;
+@!P0 BRA <+0x1b00> ;
+IADD3 R33, R3.reuse, 0x100, RZ ;
+SHF.R.S32.HI R26, RZ, 0x1f, R3 ;
+IADD3 R5, P0, R28.reuse, R3, RZ ;
+SHF.R.S32.HI R24, RZ, 0x1f, R33 ;
+IADD3 R9, P1, R28, R33, RZ ;
+IMAD.X R16, RZ, RZ, R26, P0 ;
+IADD3 R31, R3, 0x200, RZ ;
+IADD3 R29, R3, 0x300, RZ ;
+IMAD.X R10, RZ, RZ, R24, P1 ;
+SHF.R.S32.HI R22, RZ, 0x1f, R31 ;
+IADD3 R6, P2, R28, R31, RZ ;
+LEA R4, P0, R5, c[0x0][0x138], 0x4 ;
+LEA R8, P1, R9, c[0x0][0x138], 0x4 ;
+IMAD.X R7, RZ, RZ, R22, P2 ;
+SHF.R.S32.HI R20, RZ, 0x1f, R29 ;

@Jacobfaib

Copy link
Copy Markdown
Contributor

/sass diff ok 4de9b4d

@github-actions

Copy link
Copy Markdown
Contributor

Error: @Jacobfaib, no visible matching SASS report precedes your approval command. Review the current SASS report, make sure it is visible, then post /sass diff ok 4de9b4d797b8772af32df75ab686913055d35eb8 again.

@Jacobfaib

Copy link
Copy Markdown
Contributor

/sass diff ok 4de9b4d

@github-actions

Copy link
Copy Markdown
Contributor

@Jacobfaib, your SASS approval is recorded, but GitHub rejects CI rerun requests while it is running. After the current CI run finishes, rerun its failed jobs to apply the approval.

@github-actions

This comment has been minimized.

@bernhardmgruber
bernhardmgruber merged commit 66a6b2f into NVIDIA:main Oct 31, 2026
464 of 466 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Oct 31, 2026
@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 3h 21m: Pass: 100%/214 | Total: 9d 18h | Max: 3h 21m | Hits: 34%/820730

See results here.

@bernhardmgruber
bernhardmgruber deleted the reduce-contiguous-vectorization branch October 31, 2026 21:46
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

3 participants