Add sm107 tunings for DevicePartition::Flagged by gonidelis · Pull Request #11197 · NVIDIA/cccl

gonidelis · 2026-09-04T02:14:55Z

closes https://github.com/NVIDIA-dev/cccl_private/issues/721?issue=NVIDIA-dev%7Ccccl_private%7C831

perf verification results

copy-pr-bot · 2026-09-04T02:14:58Z

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-09-04T21:25:37Z

📝 Summary

Summary by CodeRabbit

Performance
- Added SM107-specific tuning for flagged partition operations using 4-byte offsets.
- Improved handling for supported 1-, 4-, 8-, and 16-byte data inputs through specialized execution strategies.
- Added dispatch support for these optimizations when flags and rejects are enabled.
- 2-byte inputs continue using existing behavior.

Walkthrough

Changes

The SM107 dispatch path now supports benchmark-selected tuning for flagged partition operations with 4-byte offsets. Supported inputs include primitive and 128-bit types of 1, 4, 8, and 16 bytes. Two-byte inputs remain untuned.

SM107 flagged partition tuning

Layer / File(s)	Summary
Flagged partition policy and dispatch integration `cub/cub/device/dispatch/tuning/tuning_select_if.cuh`	`get_sm107_flagged_partition_tuning` validates supported partition and input types, then returns specialized loading, cache, and lookback-delay policies. SM107 selection uses this tuning when flags and rejects are enabled and falls back otherwise.

Suggested reviewers: bernhardmgruber, miscco

Merge Risk: 🔵 Low · up to 143fe

This change adds SM107-specific flagged-partition tuning selection for supported inputs. The remaining risk is limited to aligning the new helper's exception specification with adjacent code before merge.

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

coderabbitai

Actionable comments posted: 1

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 972a6288-5d7c-4cc7-8a17-7b8fcdabf96d

📥 Commits

Reviewing files that changed from the base of the PR and between 486de1c and 143fef2.

📒 Files selected for processing (1)

cub/cub/device/dispatch/tuning/tuning_select_if.cuh

Included review availability: Your plan provides up to 12 included reviews per hour; 8 remain after this review.

coderabbitai · 2026-09-04T21:25:40Z


+  // tunings from cub/benchmarks/bench/partition/flagged.cu, which benchmarks with 4-byte offsets. These are raw
+  // measured values and must not be passed through nominal_4B_items_to_items.
+  [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto get_sm107_flagged_partition_tuning() const


📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win

important: Add noexcept to get_sm107_flagged_partition_tuning. The function only reads fields and constructs value types. The adjacent get_sm107_partition_tuning has this contract. As per coding guidelines, “Functions that do not throw exceptions must be marked noexcept.”

Source: Coding guidelines

@gonidelis, please apply this change

github-actions · 2026-09-04T21:47:52Z

🔬 CUB benchmark SASS comparison

⚠️ The SASS changed for 5 of 84 CUB benchmark target(s). A benchmark run may be necessary

How to request a benchmark run

Request a CUB benchmark run for this PR:

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

benchmarks:
  filters:
    cub:
      - '^cub\.bench\.segmented_topk\.fixed\.keys\.base$'
      - '^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$'
  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	`486de1c44daf7d1a343cbbf0f5477e9bea8ad600`
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.fixed.keys.base`	`sm_75`, `sm_120`, `sm_100`, `sm_110`, `sm_80`, `sm_90`
`cub.bench.segmented_topk.variable.indexed.base`	`sm_75`, `sm_120`, `sm_100`, `sm_110`, `sm_80`, `sm_90`
`cub.bench.segmented_topk.variable.indexed.cluster.base`	`sm_75`, `sm_120`, `sm_100`, `sm_110`, `sm_80`, `sm_90`
`cub.bench.segmented_topk.variable.keys.base`	`sm_75`, `sm_120`, `sm_100`, `sm_110`, `sm_80`, `sm_90`
`cub.bench.segmented_topk.variable.keys.cluster.base`	`sm_75`, `sm_120`, `sm_100`, `sm_110`, `sm_80`, `sm_90`

‼️ Summary of Differences ‼️

Showing 5/5 summaries.

cub.bench.segmented_topk.fixed.keys.base - sm_75

Showing 40/16284 diff lines, 15595 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.fixed.keys.base.sm_75
+++ test/cub.bench.segmented_topk.fixed.keys.base.sm_75
@@ -39345,573 +39345,302 @@
 BRA <+0x0>;
 Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::device_batched_topk_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::policy_selector_from_types<float, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType, (long)128, (long)1024, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0, cuda::args::__4::constant<(int)1024, int>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cuda::args::__4::constant<(int)128, int>, const cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cuda::args::__4::constant<(int)1024, int>, cuda::args::__4::constant<(int)128, int>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0>(T2, T3, T4, T5, T6, T7, T8, T9, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::baseline_kernel_args<cuda::args::__4::__traits<T9>::element_type, T10>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::cluster_kernel_args)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
-S2R R13, SR_CTAID.X ;
-ISETP.GE.U32.AND P0, PT, R13, c[0x0][0x198], PT ;
-SHF.R.S32.HI R0, RZ, 0x1f, R13 ;
+S2R R15, SR_CTAID.X ;
+ISETP.GE.U32.AND P0, PT, R15, c[0x0][0x198], PT ;
+SHF.R.S32.HI R0, RZ, 0x1f, R15 ;
 ISETP.GE.AND.EX P0, PT, R0, c[0x0][0x19c], PT, P0 ;
 @P0 EXIT ;
-S2R R12, SR_TID.X ;
+S2R R17, SR_TID.X ;
 IMAD R4, R0, c[0x0][0x168], RZ ;
-IMAD R7, R13, c[0x0][0x16c], R4 ;
-IMAD.SHL.U32 R15, R12, 0x4, RZ ;
-LOP3.LUT R15, R15, 0xffffff80, RZ, 0xc0, !PT ;
-LOP3.LUT R2, R15, 0x1f, R12, 0xf8, !PT ;
+IMAD R5, R15, c[0x0][0x16c], R4 ;
+IMAD.SHL.U32 R14, R17, 0x4, RZ ;
+LOP3.LUT R14, R14, 0xffffff80, RZ, 0xc0, !PT ;
+LOP3.LUT R2, R14, 0x1f, R17, 0xf8, !PT ;
 SHF.R.S32.HI R3, RZ, 0x1f, R2 ;
-IMAD.WIDE.U32 R4, R13, c[0x0][0x168], R2 ;
-IMAD.IADD R5, R5, 0x1, R7 ;
-LEA R8, P0, R4, c[0x0][0x160], 0x2 ;
-LEA.HI.X R9, R4, c[0x0][0x164], R5, 0x2, P0 ;
-LDG.E.SYS R10, [R8] ;
-LDG.E.SYS R11, [R8+0x80] ;
-LDG.E.SYS R18, [R8+0x100] ;
-LDG.E.SYS R19, [R8+0x180] ;
-SHF.R.U32.HI R17, RZ, 0x5, R12 ;
-IMAD.MOV.U32 R20, RZ, RZ, -0x80000000 ;
-ISETP.NE.AND P5, PT, R12.reuse, RZ, PT ;
-S2R R14, SR_LANEID ;
-ISETP.GT.AND P4, PT, R12, 0xff, PT ;
-IMAD.MOV.U32 R16, RZ, RZ, 0x1 ;

cub.bench.segmented_topk.variable.indexed.base - sm_75

Showing 40/40838 diff lines, 40233 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.indexed.base.sm_75
+++ test/cub.bench.segmented_topk.variable.indexed.base.sm_75
@@ -39048,709 +39048,380 @@
 BRA <+0x0>;
 Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::device_batched_topk_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::policy_selector_from_types<float, int, (long)1024, (long)1024, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0, cuda::args::__4::deferred_sequence<const int *, cuda::args::__4::static_bounds<(int)1, (int)1024>>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::constant_iterator<cuda::__4::counting_iterator<int, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<int *, long, (int)0, (int)0, (int)0>, long>, cuda::args::__4::constant<(int)1024, int>, const cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::constant_iterator<cuda::__4::counting_iterator<int, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<int *, long, (int)0, (int)0, (int)0>, long>, cuda::args::__4::deferred_sequence<const int *, cuda::args::__4::static_bounds<(int)1, (int)1024>>, cuda::args::__4::constant<(int)1024, int>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0>(T2, T3, T4, T5, T6, T7, T8, T9, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::baseline_kernel_args<cuda::args::__4::__traits<T9>::element_type, T10>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::cluster_kernel_args)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
-S2R R21, SR_CTAID.X ;
-ISETP.GE.U32.AND P0, PT, R21, c[0x0][0x1b8], PT ;
-SHF.R.S32.HI R0, RZ, 0x1f, R21 ;
-ISETP.GE.AND.EX P0, PT, R0, c[0x0][0x1bc], PT, P0 ;
+S2R R13, SR_CTAID.X ;
+ISETP.GE.U32.AND P0, PT, R13, c[0x0][0x1b8], PT ;
+SHF.R.S32.HI R12, RZ, 0x1f, R13 ;
+ISETP.GE.AND.EX P0, PT, R12, c[0x0][0x1bc], PT, P0 ;
 @P0 EXIT ;
-LEA R2, P0, R21, c[0x0][0x1a0], 0x2 ;
-LEA.HI.X R3, R21, c[0x0][0x1a4], R0, 0x2, P0 ;
-LDG.E.SYS R12, [R2] ;
-IMNMX.U32 R20, R12, 0x400, PT ;
-ISETP.NE.AND P0, PT, R20, RZ, PT ;
+LEA R2, P0, R13, c[0x0][0x1a0], 0x2 ;
+LEA.HI.X R3, R13, c[0x0][0x1a4], R12, 0x2, P0 ;
+LDG.E.SYS R23, [R2] ;
+IMNMX.U32 R18, R23, 0x400, PT ;
+ISETP.NE.AND P0, PT, R18, RZ, PT ;
 @!P0 EXIT ;
-S2R R22, SR_TID.X ;
-IMAD R4, R0, c[0x0][0x168], RZ ;
-IMAD R5, R21, c[0x0][0x16c], R4 ;
-IMAD.SHL.U32 R23, R22, 0x4, RZ ;
-LOP3.LUT R25, R23, 0xffffff80, RZ, 0xc0, !PT ;
-LOP3.LUT R2, R25, 0x1f, R22, 0xf8, !PT ;
+S2R R19, SR_TID.X ;
+IMAD R0, R12, c[0x0][0x168], RZ ;
+IMAD R5, R13, c[0x0][0x16c], R0 ;
+IMAD.SHL.U32 R21, R19, 0x4, RZ ;
+LOP3.LUT R20, R21, 0xffffff80, RZ, 0xc0, !PT ;
+LOP3.LUT R2, R20, 0x1f, R19, 0xf8, !PT ;
 SHF.R.S32.HI R3, RZ, 0x1f, R2.reuse ;
-IADD3 R24, R2.reuse, 0x20, RZ ;

cub.bench.segmented_topk.variable.indexed.cluster.base - sm_75

Showing 40/137 diff lines, 30 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.indexed.cluster.base.sm_75
+++ test/cub.bench.segmented_topk.variable.indexed.cluster.base.sm_75
@@ -104114,7 +104114,7 @@
 @P0 BRA <-0xe0> ;
 EXIT ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 MOV R1, c[0x0][0x28] ;
 S2R R0, SR_CTAID.X ;
 LDC.S8 R2, c[0x0][0x16c] ;
@@ -114692,7 +114692,7 @@
 MOV R5, R8 ;
 RET.REL.NODEC R6 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)2048>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)2048>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 MOV R1, c[0x0][0x28] ;
 S2R R0, SR_CTAID.X ;
 LDC.S8 R2, c[0x0][0x16c] ;
@@ -125270,7 +125270,7 @@
 MOV R5, R8 ;
 RET.REL.NODEC R6 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)1024, (int)512>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)1024, (int)512>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 MOV R1, c[0x0][0x28] ;
 S2R R0, SR_CTAID.X ;
 LDC.S8 R2, c[0x0][0x16c] ;
@@ -135848,7 +135848,7 @@
 MOV R5, R8 ;
 RET.REL.NODEC R6 0x0 ;
 BRA <+0x0>;
-Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> gen_data<(int)2048, (int)1024>(int, pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
+Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>, thrust::_V_300600_SM_750_800_900_1000_1100_1200::detail::normal_iterator<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_ptr<float>>>, long, cuda::__4::always_true, cuda::__4::__callable_permitting_copied_arguments<thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_vector<float, thrust::_V_300600_SM_750_800_900_1000_1100_1200::device_allocator<float>> <unnamed>::gen_data<(int)2048, (int)1024>(int, <unnamed>::pattern_kind, const int *)::[lambda(unsigned long) (instance 1)]>, float *, cuda::__4::counting_iterator<long, __int128, (int)0, (int)0, (int)0>>(T2, int, bool, T3, T4, T5, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 MOV R1, c[0x0][0x28] ;
 S2R R0, SR_CTAID.X ;
 LDC.S8 R2, c[0x0][0x16c] ;
@@ -146426,7 +146426,7 @@
 MOV R5, R8 ;

cub.bench.segmented_topk.variable.keys.base - sm_75

Showing 40/35312 diff lines, 34623 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.keys.base.sm_75
+++ test/cub.bench.segmented_topk.variable.keys.base.sm_75
@@ -39048,23312 +39048,12443 @@
 BRA <+0x0>;
 Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::device_batched_topk_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::policy_selector_from_types<float, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType, (long)1024, (long)1024, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0, cuda::args::__4::deferred_sequence<const int *, cuda::args::__4::static_bounds<(int)1, (int)1024>>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cuda::args::__4::constant<(int)1024, int>, const cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cuda::args::__4::deferred_sequence<const int *, cuda::args::__4::static_bounds<(int)1, (int)1024>>, cuda::args::__4::constant<(int)1024, int>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0>(T2, T3, T4, T5, T6, T7, T8, T9, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::baseline_kernel_args<cuda::args::__4::__traits<T9>::element_type, T10>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::cluster_kernel_args)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
-S2R R17, SR_CTAID.X ;
-ISETP.GE.U32.AND P0, PT, R17, c[0x0][0x1a8], PT ;
-SHF.R.S32.HI R16, RZ, 0x1f, R17 ;
-ISETP.GE.AND.EX P0, PT, R16, c[0x0][0x1ac], PT, P0 ;
+S2R R11, SR_CTAID.X ;
+ISETP.GE.U32.AND P0, PT, R11, c[0x0][0x1a8], PT ;
+SHF.R.S32.HI R14, RZ, 0x1f, R11 ;
+ISETP.GE.AND.EX P0, PT, R14, c[0x0][0x1ac], PT, P0 ;
 @P0 EXIT ;
-LEA R2, P0, R17, c[0x0][0x190], 0x2 ;
-LEA.HI.X R3, R17, c[0x0][0x194], R16, 0x2, P0 ;
-LDG.E.SYS R23, [R2] ;
-IMNMX.U32 R20, R23, 0x400, PT ;
-ISETP.NE.AND P0, PT, R20, RZ, PT ;
+LEA R2, P0, R11, c[0x0][0x190], 0x2 ;
+LEA.HI.X R3, R11, c[0x0][0x194], R14, 0x2, P0 ;
+LDG.E.SYS R17, [R2] ;
+IMNMX.U32 R13, R17, 0x400, PT ;
+ISETP.NE.AND P0, PT, R13, RZ, PT ;
 @!P0 EXIT ;
-S2R R21, SR_TID.X ;
-IMAD R4, R16, c[0x0][0x168], RZ ;
-IMAD R7, R17, c[0x0][0x16c], R4 ;
-IMAD.SHL.U32 R19, R21, 0x4, RZ ;
-LOP3.LUT R18, R19, 0xffffff80, RZ, 0xc0, !PT ;
-LOP3.LUT R2, R18, 0x1f, R21, 0xf8, !PT ;
+S2R R15, SR_TID.X ;
+IMAD R0, R14, c[0x0][0x168], RZ ;
+IMAD R7, R11, c[0x0][0x16c], R0 ;
+IMAD.SHL.U32 R16, R15, 0x4, RZ ;
+LOP3.LUT R18, R16, 0xffffff80, RZ, 0xc0, !PT ;
+LOP3.LUT R2, R18, 0x1f, R15, 0xf8, !PT ;
 SHF.R.S32.HI R3, RZ, 0x1f, R2.reuse ;
-IADD3 R0, R2.reuse, 0x20, RZ ;

cub.bench.segmented_topk.variable.keys.cluster.base - sm_75

Showing 40/35312 diff lines, 34623 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_topk.variable.keys.cluster.base.sm_75
+++ test/cub.bench.segmented_topk.variable.keys.cluster.base.sm_75
@@ -39048,23312 +39048,12443 @@
 BRA <+0x0>;
 Function : void cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::device_batched_topk_kernel<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::policy_selector_from_types<float, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType, (long)1024, (long)1024, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0, cuda::args::__4::deferred_sequence<const int *, cuda::args::__4::static_bounds<(int)1, (int)1024>>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cuda::args::__4::constant<(int)1024, int>, const cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cuda::__4::strided_iterator<cuda::__4::counting_iterator<float *, long, (int)0, (int)0, (int)0>, long>, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cub::_V_300600_SM_750_800_900_1000_1100_1200::NullType **, cuda::args::__4::deferred_sequence<const int *, cuda::args::__4::static_bounds<(int)1, (int)1024>>, cuda::args::__4::constant<(int)1024, int>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::params::static_discrete_param<cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select, (cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::topk::select)1>, cuda::args::__4::immediate<long, cuda::args::__4::no_bounds>, long, (cuda::execution::__4::determinism::__determinism_t)0, (cuda::execution::__4::tie_break::__tie_break_t)0>(T2, T3, T4, T5, T6, T7, T8, T9, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::baseline_kernel_args<cuda::args::__4::__traits<T9>::element_type, T10>, cub::_V_300600_SM_750_800_900_1000_1100_1200::detail::batched_topk::cluster_kernel_args)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ;
-S2R R17, SR_CTAID.X ;
-ISETP.GE.U32.AND P0, PT, R17, c[0x0][0x1a8], PT ;
-SHF.R.S32.HI R16, RZ, 0x1f, R17 ;
-ISETP.GE.AND.EX P0, PT, R16, c[0x0][0x1ac], PT, P0 ;
+S2R R11, SR_CTAID.X ;
+ISETP.GE.U32.AND P0, PT, R11, c[0x0][0x1a8], PT ;
+SHF.R.S32.HI R14, RZ, 0x1f, R11 ;
+ISETP.GE.AND.EX P0, PT, R14, c[0x0][0x1ac], PT, P0 ;
 @P0 EXIT ;
-LEA R2, P0, R17, c[0x0][0x190], 0x2 ;
-LEA.HI.X R3, R17, c[0x0][0x194], R16, 0x2, P0 ;
-LDG.E.SYS R23, [R2] ;
-IMNMX.U32 R20, R23, 0x400, PT ;
-ISETP.NE.AND P0, PT, R20, RZ, PT ;
+LEA R2, P0, R11, c[0x0][0x190], 0x2 ;
+LEA.HI.X R3, R11, c[0x0][0x194], R14, 0x2, P0 ;
+LDG.E.SYS R17, [R2] ;
+IMNMX.U32 R13, R17, 0x400, PT ;
+ISETP.NE.AND P0, PT, R13, RZ, PT ;
 @!P0 EXIT ;
-S2R R21, SR_TID.X ;
-IMAD R4, R16, c[0x0][0x168], RZ ;
-IMAD R7, R17, c[0x0][0x16c], R4 ;
-IMAD.SHL.U32 R19, R21, 0x4, RZ ;
-LOP3.LUT R18, R19, 0xffffff80, RZ, 0xc0, !PT ;
-LOP3.LUT R2, R18, 0x1f, R21, 0xf8, !PT ;
+S2R R15, SR_TID.X ;
+IMAD R0, R14, c[0x0][0x168], RZ ;
+IMAD R7, R11, c[0x0][0x16c], R0 ;
+IMAD.SHL.U32 R16, R15, 0x4, RZ ;
+LOP3.LUT R18, R16, 0xffffff80, RZ, 0xc0, !PT ;
+LOP3.LUT R2, R18, 0x1f, R15, 0xf8, !PT ;
 SHF.R.S32.HI R3, RZ, 0x1f, R2.reuse ;
-IADD3 R0, R2.reuse, 0x20, RZ ;

github-actions · 2026-09-05T00:12:11Z

🥳 CI Workflow Results

🟩 Finished in 2h 46m: Pass: 100%/176 | Total: 8d 11h | Max: 2h 46m | Hits: 28%/638149

See results here.

github-project-automation Bot added this to CCCL Sep 4, 2026

github-project-automation Bot moved this to Todo in CCCL Sep 4, 2026

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

gonidelis changed the title ~~Add sm107 tunings for DevicePartition::Flagged~~ Add sm107 tunings for DevicePartition::Flagged Sep 4, 2026

cub: add sm107 tunings for DevicePartition::Flagged

143fef2

gonidelis force-pushed the partition_flagged_sm107 branch from 0e515f9 to 143fef2 Compare September 4, 2026 19:37

gonidelis marked this pull request as ready for review September 4, 2026 21:22

gonidelis requested a review from a team as a code owner September 4, 2026 21:22

gonidelis requested a review from srinivasyadav18 September 4, 2026 21:22

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

coderabbitai Bot reviewed Sep 4, 2026

View reviewed changes

bernhardmgruber approved these changes Sep 5, 2026

View reviewed changes

Navigation Menu

Uh oh!

Add sm107 tunings for DevicePartition::Flagged - #11197

Add sm107 tunings for DevicePartition::Flagged#11197
gonidelis wants to merge 1 commit into
NVIDIA:mainfrom
gonidelis:partition_flagged_sm107

gonidelis commented Sep 4, 2026 •

edited

Loading

Uh oh!

copy-pr-bot Bot commented Sep 4, 2026

Uh oh!

coderabbitai Bot commented Sep 4, 2026 •

edited

Loading

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Uh oh!

coderabbitai Bot Sep 4, 2026

Uh oh!

bernhardmgruber Sep 5, 2026

Uh oh!

github-actions Bot commented Sep 4, 2026

Uh oh!

github-actions Bot commented Sep 5, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

gonidelis commented Sep 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

copy-pr-bot Bot commented Sep 4, 2026

Uh oh!

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

Uh oh!

Summary by CodeRabbit

Walkthrough

Changes

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot Sep 4, 2026

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Sep 5, 2026

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Sep 4, 2026

🔬 CUB benchmark SASS comparison

‼️ Summary of Differences ‼️

Uh oh!

github-actions Bot commented Sep 5, 2026

🥳 CI Workflow Results

🟩 Finished in 2h 46m: Pass: 100%/176 | Total: 8d 11h | Max: 2h 46m | Hits: 28%/638149

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

gonidelis commented Sep 4, 2026 •

edited

Loading

coderabbitai Bot commented Sep 4, 2026 •

edited

Loading