Skip to content

Optimize FP64 comparison functions on GPUs with slow FP64 performance - #12231

Merged
davebayer merged 24 commits into
NVIDIA:mainfrom
lennartvoelz:libcudacxx/feat/fp64_comp_int_fast_path
Nov 6, 2026
Merged

davebayer merged 24 commits into
NVIDIA:mainfrom
lennartvoelz:libcudacxx/feat/fp64_comp_int_fast_path

Conversation

@lennartvoelz

@lennartvoelz lennartvoelz commented Oct 23, 2026 •

Copy link
Copy Markdown
Contributor

Description

Part of #11894

Comparison functions

  • isgreater, isless, isgreaterequal, islessequal and islessgreater compile to DSETP for double, which is slow on GPUs with low FP64 throughput -> idea to convert to int64 representation with the same ordering as double for the comparison
  • __fp64_to_ordered_int maps the bits of double to int64_t with the same ordering:
    I chose a selection form:
    key = x < 0 ? INT64_MIN - x : x -> -0 and +0 map to the same value
    The function is local to traits.h, as discussed in [FEA]: Optimize double-precision <cmath> functionalities with integer path #11894
  • NaN is handled by the existing double isnan fast path
  • the new path only runs on slow FP64 GPUs, guarded by the existing _CCCL_DISPATCH_SLOW_FP64_TARGET macro

Benchmarks:

op main this PR speedup
isless 724.6us 286.3us 2.53x
islessequal 725.0us 286.3us 2.53x
isgreater 725.2us 286.3us 2.53x
isgreaterequal 725.3us 286.5us 2.53x
isgreaterequal 725.2us 268.1us 2.70x

No FP64 instructions are emitted by the compiler on slow FP64 targets.

fmax/fmin

  • __fp64_to_total_order_int maps the bits of a double to int64_t in total order, i.e. -0 < +0. NaNs are placed beyond +-inf
  • __fp64_fmax_fmin_int adds function template that provides branch free min/max paths for doubles with integer instructions. -0 < +0, single NaN ignored and two NaNs return first argument quieted

Benchmarks:

op main this PR speedup
fmax 725.2us 395.8us 1.83x
fmin 725.4us 395.9us 1.83x

fmax / fmin also affect downstream paths in CUB via cuda::std::minimum/maximum:

Benchmark Path to fmin/fmax 2^16 2^20 2^24 2^28
CUB scan/applications/P1/running-min-max cuda::minimum/maximum −21.0% −20.4% −21.1% −21.6%
CUB scan/applications/P1/log-cdf-from-log-pdfs cuda::maximum −5.1% −4.8% −4.9% −4.6%
CUB reduce/min cuda::minimum −1.7% +2.0% −0.4% +0.3%
CUB reduce/deterministic (RFA sum) fmax(fabs(a), fabs(b)) 0.0% +1.4% +2.3% +2.5%

The only downstream benchmark negatively affected is using fabs, which is the reason for the slowdown. I'll address this in a separate PR.

Added tests for comparisons and fmin/fmax that cover edge cases

Checklist

  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

- adds __fp64_to_ordered_int helper that uses a selection form to convert the bit pattern to int64 with the same ordering as fp64
- adds double overloads for isgreater, isless, isgreaterequal, islessequal and islessgreater
- uses _CCCL_DISPATCH_SLOW_FP64_TARGET to dispatch between "slow" and "fast" fp64 arithmetic targets
@copy-pr-bot

copy-pr-bot Bot commented Oct 23, 2026

Copy link
Copy Markdown
Contributor

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@github-project-automation github-project-automation Bot moved this to Todo in CCCL Oct 23, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Oct 23, 2026
lennartvoelz and others added 8 commits October 23, 2026 22:35
quite -> quiet in min_max.h
- adds __fp64_to_total_order_int helper that maps the bit pattern to int64 in IEEE totalOrder (-0 < +0)
- adds branch-free __fp64_fmax_fmin_int helper that matches the hardware result bitwise (NaN handling, quiet bit)
- uses _CCCL_DISPATCH_SLOW_FP64_TARGET in fmax and fmin to dispatch to the integer path for double
- tests for fmin/fmax and the comparison functions isless, islessequal, isgreater, isgreaterequal, islessgreater
…argets the generic fallback (x < y ? y : x) returns either zero for ±0, while the integer path on slow targets orders -0 < +0

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

LGTM but I want @fbusato or @davebayer to weigh in too

@lennartvoelz
lennartvoelz marked this pull request as ready for review October 26, 2026 11:43
@lennartvoelz
lennartvoelz requested review from a team as code owners October 26, 2026 11:43
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Oct 26, 2026
@lennartvoelz

Copy link
Copy Markdown
Contributor Author

Ready for review
One thing that came up with the new test cases:
The integer path reproduces the hardware double min/max behavior: one NaN is ignored and two return NaNs return the first argument with the quiet bit set. For fmax(x, x) and fmin(x, x) with a signaling NaN, the compiler folds this without executing max in hardware, so x is returned unchanged. The hardware path is not self-consistent logically because fmax(x, x) != fmax(x, y) with y == x. I think no integer implementation can match both?
How should this be handled?

@coderabbitai

coderabbitai Bot commented Oct 26, 2026 •

Copy link
Copy Markdown
Contributor

Review in Change Stack →

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review

No actionable comments were generated in the recent review.

Recent review info
Run configuration
  • Source mode: Auto
  • Repository: NVIDIA/cccl
  • Branch: main
  • File: .coderabbit.yaml
  • Commit: 5fbf0f0
  • Layers: Repository YAML
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: 4ee786ea-1736-4093-8443-aac8fff668d1
Commits

Reviewing files that changed from the base of the PR and between 620ccb6 and 9fa5949.

Files selected for processing: 1
  • libcudacxx/include/cuda/std/__cmath/traits.h

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

Summary

Summary by CodeRabbit

  • Bug Fixes
    • Improved double-precision fmin and fmax behavior on supported devices, including handling of signed zero, infinities, denormal values, and NaNs.
    • Improved double-precision comparison behavior for ordering and NaN edge cases, including comparisons involving signed zero and infinities.
    • Added coverage for edge cases involving limits and mixed integer and floating-point arguments.
  • New Features
    • Added benchmarks for double-precision comparison, minimum, and maximum operations.

This PR adds integer-based double comparison and min/max paths for slow-FP64 targets. It also adds edge-case tests and CUDA benchmarks for the affected operations.

Changes: Slow-FP64 math paths

Layer / File(s) Summary
Double comparison predicates
libcudacxx/include/cuda/std/__cmath/traits.h, libcudacxx/test/libcudacxx/std/numerics/c.math/fp_traits.pass.cpp, libcudacxx/benchmarks/bench/comparison/basic.cu
Adds ordered-integer implementations for five double comparison predicates on slow-FP64 targets. Tests cover signed zeros, NaNs, infinities, denormals, limits, and mixed arguments. Adds benchmarks for the predicates.
Double min/max functions
libcudacxx/include/cuda/std/__cmath/min_max.h, libcudacxx/test/libcudacxx/std/numerics/c.math/fp_min_max.pass.cpp, libcudacxx/benchmarks/bench/fmin_fmax/basic.cu
Adds integer-based double selection and slow-FP64 dispatch for fmin and fmax. Tests cover signed zeros, NaNs, infinities, denormals, and mixed arguments. Adds benchmarks for both functions.

Possibly related issues

Suggested reviewers: fbusato

Change: Refactor

Merge Risk: Minimal · previous review 9fa59

No merge-blocking behavior was identified; the change is ready for normal checks.

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

@lennartvoelz

Copy link
Copy Markdown
Contributor Author

All review comments should be addressed now
As expected, with the changes to fabs in #12294 the downstream regression in the CUB deterministic reduction turns into a performance gain (see the description of the #12294)

@lennartvoelz
lennartvoelz requested a review from fbusato November 2, 2026 15:57
@fbusato

fbusato commented Nov 2, 2026

Copy link
Copy Markdown
Contributor

/ok to test e672e96

@github-actions

github-actions Bot commented Nov 2, 2026

Copy link
Copy Markdown
Contributor

🔬 Benchmark SASS comparison

⚠️ The SASS changed for 6 of 100 CUB and 0 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\.reduce\.deferred_deterministic\.base$'
      - '^cub\.bench\.reduce\.deterministic\.base$'
      - '^cub\.bench\.reduce\.min\.base$'
      - '^cub\.bench\.reduce\.warp_reduce_min\.base$'
      - '^cub\.bench\.scan\.applications\.P1\.log\-cdf\-from\-log\-pdfs\.base$'
      - '^cub\.bench\.scan\.applications\.P1\.running\-min\-max\.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 8004db6f23be90b94f7d064cb4c42c22f8b1f259
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.reduce.deferred_deterministic.base sm_120, sm_75, sm_110
cub.bench.reduce.deterministic.base sm_120, sm_75, sm_110
cub.bench.reduce.min.base sm_120, sm_75, sm_110
cub.bench.reduce.warp_reduce_min.base sm_75
cub.bench.scan.applications.P1.log-cdf-from-log-pdfs.base sm_120, sm_75, sm_110
cub.bench.scan.applications.P1.running-min-max.base sm_120, sm_75, sm_110

‼️ Summary of Differences ‼️

Showing 6/6 summaries.

cub.bench.reduce.deferred_deterministic.base - sm_120 - ⚠️ Significant

Deterministic accumulator maximum selection expands The deterministic reduction kernels replace absolute-value DSETP.MAX selection blocks with longer integer classification, comparison, and FSEL sequences. This signature occurs in both deterministic targets on sm_75, sm_110, and sm_120.

Showing 40/3107 diff lines, 2278 changes. - ⬇️ Full diff

--- base/cub.bench.reduce.deferred_deterministic.base.sm_120
+++ test/cub.bench.reduce.deferred_deterministic.base.sm_120
@@ -12475,7 +12475,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeterministicDeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<double, unsigned int, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::deterministic_sum_t<double, (void *)0>, (cuda::execution::__4::determinism::__determinism_t)2>, double *, int *, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::deterministic_sum_t<double, (void *)0>, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::ReproducibleFloatingAccumulator<double, (int)3, (void *)0>, cuda::std::__4::identity>(T2, T5 *, T3, T4, T6, int)
 BSYNC.RECONVERGENT B0 ;
 STL.64 [R1], RZ ;
 ISETP.GE.U32.AND P0, PT, R22, R21, PT ;
-BSSY.RECONVERGENT B1, <+0x22a0> ;
+BSSY.RECONVERGENT B1, <+0x28a0> ;
 MOV.64 R14, RZ ;
 STL.64 [R1+0x10], RZ ;
 MOV.64 R12, RZ ;
@@ -12485,78 +12485,171 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeterministicDeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<double, unsigned int, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::deterministic_sum_t<double, (void *)0>, (cuda::execution::__4::determinism::__determinism_t)2>, double *, int *, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::deterministic_sum_t<double, (void *)0>, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::ReproducibleFloatingAccumulator<double, (int)3, (void *)0>, cuda::std::__4::identity>(T2, T5 *, T3, T4, T6, int)
 MOV.64 R8, RZ ;
 MOV.64 R6, RZ ;
 BAR.SYNC.DEFER_BLOCKING 0x0 ;
-@P0 BRA <+0x21e0> ;
+@P0 BRA <+0x27e0> ;
 MOV.64 R8, RZ ;
 MOV.64 R6, RZ ;
 HFMA2 R0, -RZ, RZ, 0, 0 ;
-BSSY.RECONVERGENT B0, <+0x1fe0> ;
+BSSY.RECONVERGENT B0, <+0x25e0> ;
 IMAD R12, R19.reuse, 0x80, R22 ;
-ISETP.GE.U32.AND P3, PT, R22, R21.reuse, PT ;
+ISETP.GE.U32.AND P0, PT, R22, R21.reuse, PT ;
+MOV.64 R26, RZ ;
 MOV.64 R28, RZ ;
-IMAD R14, R19.reuse, 0x80, R12 ;
-ISETP.GE.U32.AND P2, PT, R12, R21.reuse, PT ;
-MOV.64 R26, RZ ;
-IMAD R34, R19, 0x80, R14 ;
-ISETP.GE.U32.AND P1, PT, R14, R21, PT ;
-ISETP.GE.U32.AND P0, PT, R34, R21, PT ;
-@!P3 LDC.64 R2, c[0x0][0x88] ;
-@!P2 LDC.64 R4, c[0x0][0x88] ;
-@!P1 LDC.64 R10, c[0x0][0x88] ;
-@!P3 IMAD.WIDE.U32 R2, R22, 0x8, R2 ;
+ISETP.GE.U32.AND P1, PT, R12, R21, PT ;
+IMAD R14, R19, 0x80, R12 ;
+ISETP.GE.U32.AND P2, PT, R14, R21, PT ;
cub.bench.reduce.deterministic.base - sm_120 - ⚠️ Significant

Deterministic accumulator maximum selection expands The deterministic reduction kernels replace absolute-value DSETP.MAX selection blocks with longer integer classification, comparison, and FSEL sequences. This signature occurs in both deterministic targets on sm_75, sm_110, and sm_120.

Showing 40/1596 diff lines, 1098 changes. - ⬇️ Full diff

--- base/cub.bench.reduce.deterministic.base.sm_120
+++ test/cub.bench.reduce.deterministic.base.sm_120
@@ -83,58 +83,57 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::DeterministicDeviceReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::reduce::policy_selector_from_types<double, unsigned int, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::deterministic_sum_t<double, (void *)0>, (cuda::execution::__4::determinism::__determinism_t)2>, const double *, int, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::deterministic_sum_t<double, (void *)0>, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::rfa::ReproducibleFloatingAccumulator<double, (int)3, (void *)0>, cuda::std::__4::identity>(T2, T5 *, T3, T4, T6, int)
 HFMA2 R3, -RZ, RZ, 0, 0 ;
 S2R R19, SR_CTAID.X ;
 LDCU UR5, c[0x0][0x4] ;
-BSSY.RECONVERGENT B0, <+0x310> ;
+BSSY.RECONVERGENT B0, <+0x300> ;
 MOV R2, R1 ;
 IADD.64 R16, R2, UR4 ;
 ISETP.GT.U32.AND P0, PT, R18, 0x66, PT ;
-IMAD R20, R19, 0x80, R18 ;
-@P0 BRA <+0x2a0> ;
+IMAD R2, R19, 0x80, R18 ;
+@P0 BRA <+0x290> ;
 ISETP.NE.U32.AND P0, PT, R18, RZ, PT ;
-BSSY.RECONVERGENT B1, <+0x220> ;
-@!P0 MOV.64 R2, 0x5010000000000000 ;
-@!P0 DMUL R2, R2, 0.75 ;
-@!P0 DMUL R2, R2, 5.78960446186580977118e+76 ;
-@!P0 DMUL R2, R2, 5.78960446186580977118e+76 ;
-@!P0 DMUL R2, R2, 5.78960446186580977118e+76 ;
-@!P0 DADD R2, R2, R2 ;
-@!P0 BRA <+0x190> ;
+BSSY.RECONVERGENT B1, <+0x210> ;
+@!P0 MOV.64 R4, 0x5010000000000000 ;
+@!P0 DMUL R4, R4, 0.75 ;
+@!P0 DMUL R4, R4, 5.78960446186580977118e+76 ;
+@!P0 DMUL R4, R4, 5.78960446186580977118e+76 ;
+@!P0 DMUL R4, R4, 5.78960446186580977118e+76 ;
+@!P0 DADD R4, R4, R4 ;
+@!P0 BRA <+0x180> ;
 ISETP.GT.U32.AND P0, PT, R18, 0x33, PT ;
-@P0 BRA <+0x120> ;
-IMAD R2, R18, -0x28, RZ ;
-IADD R3, R2.reuse, 0x4e6 ;
-IADD R0, R2, 0x140a ;
+@P0 BRA <+0x110> ;
+IMAD R4, R18, -0x28, RZ ;
+IADD R3, R4.reuse, 0x4e6 ;
cub.bench.reduce.min.base - sm_120 - ⚠️ Significant

Double extrema selection expands into explicit comparison sequences The scan, minimum-reduction, and log-CDF kernels replace double-precision DSETP.MIN and DSETP.MAX blocks with longer ISETP, LOP3, SEL, and FSEL sequences. The instruction mix and count change across sm_75, sm_110, and sm_120 for all three targets.

Showing 40/29569 diff lines, 29217 changes. - ⬇️ Full diff

--- base/cub.bench.reduce.min.base.sm_120
+++ test/cub.bench.reduce.min.base.sm_120
@@ -5883,137 +5883,3357 @@ 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<double, unsigned int, cuda::__4::minimum<void>, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, double *, double *, unsigned int, unsigned int, cuda::__4::minimum<void>, double, double, 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 UR4, c[0x0][0x28] ;
-BSSY.RECONVERGENT B0, <+0x5090> ;
+LDC R3, c[0x0][0x54] ;
+BSSY.RECONVERGENT B0, <+0xc220> ;
+S2R R10, SR_TID.X ;
+LDCU.64 UR6, c[0x0][0x0] ;
+ULOP3.LUT UP0, URZ, UR4, 0xf, URZ, 0xc0, !UPT ;
+IMAD.SHL.U32 R3, R3, 0x800, RZ ;
+IMAD.SHL.U32 R24, R0, 0x800, RZ ;
+BRA.U UP0, <+0x6040> ;
+LDCU UR5, c[0x0][0x50] ;
+IADD R2, -R24, UR5 ;
+ISETP.GT.U32.AND P0, PT, R2, 0x7ff, PT ;
+@P0 BRA <+0x2b90> ;
+ISETP.GE.U32.AND P0, PT, R10, R2, PT ;
+BSSY.RECONVERGENT B1, <+0xb0> ;
+MOV.64 R6, RZ ;
+MOV R3, R10 ;
+@P0 BRA <+0x60> ;
+LDC.64 R6, c[0x0][0x28] ;
+IADD R3, R24, R10 ;
+IMAD.WIDE.U32 R6, R3, 0x8, R6 ;
+LDG.E.64.CONSTANT R6, desc[UR6][R6.64] ;
+IADD R3, R10, 0x100 ;
+NOP ;
+BSYNC.RECONVERGENT B1 ;
+ISETP.GE.U32.AND P0, PT, R3, R2, PT ;
+BSSY.RECONVERGENT B1, <+0xb60> ;
+@P0 BRA <+0xb30> ;
+LOP3.LUT R4, RZ, R3, RZ, 0x33, !PT ;
+BSSY.RECONVERGENT B2, <+0x330> ;
+IADD R5, R4, UR5 ;
+LEA.HI R4, R5, 0x1, RZ, 0x18 ;
+IADD R5, -R24, R5 ;
+LOP3.LUT P1, R16, R4, 0x3, RZ, 0xc0, !PT ;
cub.bench.reduce.warp_reduce_min.base - sm_75 - ⚠️ Significant

Warp-reduce kernel adds clock-value predicate loop The benchmark kernel adds integer predicates, selects, and a backward branch before the virtual SM identifier read, while two register moves disappear. This signature occurs only in the sm_75 warp-reduce-min pair.

Showing 40/52 diff lines, 38 changes. - ⬇️ Full diff

--- base/cub.bench.reduce.warp_reduce_min.base.sm_75
+++ test/cub.bench.reduce.warp_reduce_min.base.sm_75
@@ -234,13 +234,47 @@ void benchmark_kernel<(int)256, (int)128, benchmark_op_t, double, (bool)0>(T3)
 BRA <+0x0>;
 Function : void benchmark_kernel<(int)256, (int)128, benchmark_op_t, double, (bool)0>(T3)
 IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x0] ;
+IMAD.MOV.U32 R0, RZ, RZ, RZ ;
 S2UR UR4, SR_CLOCKLO ;
 DEPBAR.LE SB0, 0x0 ;
 S2UR UR5, SR_CLOCKLO ;
 DEPBAR.LE SB0, 0x0 ;
+IMAD.U32 R2, RZ, RZ, UR4 ;
+IMAD.U32 R3, RZ, RZ, UR5 ;
+ISETP.GT.U32.AND P1, PT, R2, RZ, PT ;
+LOP3.LUT R4, R3, 0x7fffffff, RZ, 0xc0, !PT ;
+IADD3 R0, R0, 0x1, RZ ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, PT, P1 ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, P0, P1 ;
+SEL R2, R2, RZ, !P0 ;
+SEL R3, R3, 0x7ff80000, !P0 ;
+ISETP.GT.U32.AND P1, PT, R2, RZ, PT ;
+LOP3.LUT R4, R3, 0x7fffffff, RZ, 0xc0, !PT ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, PT, P1 ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, P0, P1 ;
+SEL R2, R2, RZ, !P0 ;
+SEL R3, R3, 0x7ff80000, !P0 ;
+ISETP.GT.U32.AND P1, PT, R2, RZ, PT ;
+LOP3.LUT R4, R3, 0x7fffffff, RZ, 0xc0, !PT ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, PT, P1 ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, P0, P1 ;
+SEL R2, R2, RZ, !P0 ;
+SEL R3, R3, 0x7ff80000, !P0 ;
+ISETP.GT.U32.AND P1, PT, R2, RZ, PT ;
+LOP3.LUT R4, R3, 0x7fffffff, RZ, 0xc0, !PT ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, PT, P1 ;
+ISETP.GT.U32.AND.EX P0, PT, R4, 0x7ff00000, P0, P1 ;
+SEL R2, R2, RZ, !P0 ;
+SEL R3, R3, 0x7ff80000, !P0 ;
+ISETP.GT.U32.AND P1, PT, R2, RZ, PT ;
+LOP3.LUT R4, R3, 0x7fffffff, RZ, 0xc0, !PT ;
cub.bench.scan.applications.P1.log-cdf-from-log-pdfs.base - sm_120 - ⚠️ Significant

Double extrema selection expands into explicit comparison sequences The scan, minimum-reduction, and log-CDF kernels replace double-precision DSETP.MIN and DSETP.MAX blocks with longer ISETP, LOP3, SEL, and FSEL sequences. The instruction mix and count change across sm_75, sm_110, and sm_120 for all three targets.

Showing 40/91053 diff lines, 88995 changes. - ⬇️ Full diff

--- base/cub.bench.scan.applications.P1.log-cdf-from-log-pdfs.base.sm_120
+++ test/cub.bench.scan.applications.P1.log-cdf-from-log-pdfs.base.sm_120
@@ -105,302 +105,345 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::scan::DeviceScanKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::scan::policy_selector_from_types<const double *, double *, double, unsigned int, impl::log_add_plus, (bool)0, (bool)0>, const double *, double *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<double, (bool)1>, impl::log_add_plus, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType, unsigned int, unsigned int, double, (bool)0, (bool)0, (bool)0, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType>(T2, T3, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::scan::tile_state_kernel_arg_t<T4, T9>, int, T5, T6, T8, int)
 LDCU UR5, c[0x0][0x50] ;
 ACQBULK ;
 LDCU.64 UR6, c[0x0][0x0] ;
-LDC R8, c[0x0][0x48] ;
-IADD R8, R8, UR4 ;
-IMAD R44, R8, 0x380, RZ ;
-IADD R0, -R44, UR5 ;
-ISETP.GE.U32.AND P0, PT, R0, 0x381, PT ;
-@!P0 BRA <+0x2a920> ;
-S2R R9, SR_TID.X ;
+LDC R14, c[0x0][0x48] ;
+IADD R14, R14, UR4 ;
+IMAD R62, R14, 0x380, RZ ;
+IADD R11, -R62, UR5 ;
+ISETP.GE.U32.AND P0, PT, R11, 0x381, PT ;
+@!P0 BRA <+0x2f140> ;
+S2R R24, SR_TID.X ;
 LDCU.64 UR4, c[0x0][0x28] ;
-HFMA2 R45, -RZ, RZ, 0, 0 ;
-LOP3.LUT R2, R9.reuse, 0x1f, RZ, 0xc0, !PT ;
-LOP3.LUT R3, R9, 0x3e0, RZ, 0xc0, !PT ;
+HFMA2 R63, -RZ, RZ, 0, 0 ;
+LOP3.LUT R2, R24.reuse, 0x1f, RZ, 0xc0, !PT ;
+LOP3.LUT R3, R24, 0x3e0, RZ, 0xc0, !PT ;
 IMAD R2, R3, 0x7, R2 ;
 MOV R3, RZ ;
-IADD.64 R2, R44, R2 ;
-SHF.L.U64.HI R27, R2.reuse, 0x3, R3 ;
-IMAD.SHL.U32 R26, R2, 0x8, RZ ;
-IADD.64 R4, R26, UR4 ;
+IADD.64 R2, R62, R2 ;
+SHF.L.U64.HI R3, R2.reuse, 0x3, R3 ;
+IMAD.SHL.U32 R2, R2, 0x8, RZ ;
+IADD.64 R4, R2, UR4 ;
 LDG.E.64 R6, desc[UR6][R4.64] ;
-LDG.E.64 R12, desc[UR6][R4.64+0x100] ;
-LDG.E.64 R14, desc[UR6][R4.64+0x200] ;
cub.bench.scan.applications.P1.running-min-max.base - sm_120 - ⚠️ Significant

Double extrema selection expands into explicit comparison sequences The scan, minimum-reduction, and log-CDF kernels replace double-precision DSETP.MIN and DSETP.MAX blocks with longer ISETP, LOP3, SEL, and FSEL sequences. The instruction mix and count change across sm_75, sm_110, and sm_120 for all three targets.

Showing 40/18493 diff lines, 18090 changes. - ⬇️ Full diff

--- base/cub.bench.scan.applications.P1.running-min-max.base.sm_120
+++ test/cub.bench.scan.applications.P1.running-min-max.base.sm_120
@@ -508,3239 +508,6218 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::scan::DeviceScanKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::scan::policy_selector_from_types<cuda::__4::transform_iterator<impl::embed_op<double>, const double *>, impl::min_max_t<double> *, impl::min_max_t<double>, unsigned int, impl::scan_op, (bool)0, (bool)0>, cuda::__4::transform_iterator<impl::embed_op<double>, const double *>, impl::min_max_t<double> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<impl::min_max_t<double>, (bool)0>, impl::scan_op, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType, unsigned int, unsigned int, impl::min_max_t<double>, (bool)0, (bool)0, (bool)0, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType>(T2, T3, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::scan::tile_state_kernel_arg_t<T4, T9>, int, T5, T6, T8, int)
 LDCU UR5, c[0x0][0x58] ;
 ACQBULK ;
 LDCU.64 UR8, c[0x0][0x0] ;
-LDC R40, c[0x0][0x50] ;
-IADD R40, R40, UR4 ;
-IMAD R10, R40, 0x180, RZ ;
-IADD R6, -R10, UR5 ;
-ISETP.GE.U32.AND P0, PT, R6, 0x181, PT ;
-@!P0 BRA <+0x4620> ;
+LDC R7, c[0x0][0x50] ;
+IADD R7, R7, UR4 ;
+IMAD R64, R7, 0x180, RZ ;
+IADD R0, -R64, UR5 ;
+ISETP.GE.U32.AND P0, PT, R0, 0x181, PT ;
+@!P0 BRA <+0xafa0> ;
 S2R R3, SR_TID.X ;
 LDCU.64 UR4, c[0x0][0x28] ;
-HFMA2 R11, -RZ, RZ, 0, 0 ;
-LOP3.LUT R22, R3.reuse, 0x1f, RZ, 0xc0, !PT ;
-LOP3.LUT R23, R3, 0x3e0, RZ, 0xc0, !PT ;
-IMAD R22, R23, 0x3, R22 ;
-MOV R23, RZ ;
-IADD.64 R22, R10, R22 ;
-LEA R4, P0, R22, UR4, 0x3 ;
-LEA.HI.X R5, R22, UR5, R23, 0x3, P0 ;
-LDG.E.64 R24, desc[UR8][R4.64] ;
-LDG.E.64 R28, desc[UR8][R4.64+0x100] ;
-LDG.E.64 R32, desc[UR8][R4.64+0x200] ;
+HFMA2 R65, -RZ, RZ, 0, 0 ;
+LOP3.LUT R4, R3.reuse, 0x1f, RZ, 0xc0, !PT ;
+LOP3.LUT R5, R3, 0x3e0, RZ, 0xc0, !PT ;
+IMAD R4, R5, 0x3, R4 ;
+MOV R5, RZ ;
+IADD.64 R64, R64, R4 ;
+LEA R8, P0, R64, UR4, 0x3 ;
+LEA.HI.X R9, R64, UR5, R65, 0x3, P0 ;
+LDG.E.64 R16, desc[UR8][R8.64] ;

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@fbusato

fbusato commented Nov 4, 2026

Copy link
Copy Markdown
Contributor

/sass diff ok e672e96

@github-actions

github-actions Bot commented Nov 4, 2026

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 1d 00h: Pass: 100%/230 | Total: 6d 12h | Max: 3h 25m | Hits: 54%/1571631

See results here.

@davebayer
davebayer merged commit 758ca4d into NVIDIA:main Nov 6, 2026
738 of 745 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Nov 6, 2026
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.

4 participants