Skip to content

Optimize FP64 fabs on GPUs with low FP64 performance - #12294

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

davebayer merged 8 commits into
NVIDIA:mainfrom
lennartvoelz:libcudacxx/feat/fp64_fabs_int_path

Conversation

@lennartvoelz

@lennartvoelz lennartvoelz commented Oct 27, 2026 •

Copy link
Copy Markdown
Contributor

Description

closes
Part of #11894
Follow up to #12231

  • fabs for double compiles to DADD which is slow on GPUs with low FP64 throughput -> use integer instructions instead
  • __fp64_fabs_int uses integer instructions to clear the sign bit, while conforming to IEEE754 NaN handling

Note: I first tried a mask approach which worked perfectly fine on my machine (sm_87) but the compiler recognized every form of clearing the sign bit as fabs on targets SM100+ -> DADD
Clearing the high 32-bit half with __double2hiint / __hiloint2double avoids that

The issue points out that NaNs need careful treatment for fabs. The integer path conforms to IEEE754 abs, so only the sign bit is cleared, NaN payloads are preserved and sNaNs are not quieted. This matches the host and comp-time eval and the existing test, which asserts !signbit(fabs(-NaN)).
On main, fabs is not consistent. At runtime, fabs is compiled to DADD, which keeps the sign of the NaN and quiets sNaNs. The existing test only passed because its input was folded.
I'd need guidance on which behaviour is preferred here! @fbusato

Benchmarks:

Micro benchmarks:

main this PR speedup
724.97us 83.73us 8.66x

Down stream benchmarks in CUB:

The changes in this PR alone produce results within noise of main.
More interesting is the combination of this PR with #12231
To reproduce the results, use this branch

Elements deterministic reduction F64 deferred_deterministic reduction F64
2^16 -3.61% -2.10%
2^20 -5.60% -5.67%
2^24 -7.91% -7.90%
2^28 -8.13% -8.11%

The table above shows the results for the int32 offsets, int64 is similar.

Checklist

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

@copy-pr-bot

copy-pr-bot Bot commented Oct 27, 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 27, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Oct 27, 2026
- sign and payload are implementation defined on device
@fbusato

fbusato commented Oct 27, 2026

Copy link
Copy Markdown
Contributor

Regarding

On main, fabs is not consistent. At runtime, fabs is compiled to DADD, which keeps the sign of the NaN and quiets sNaNs. The existing test only passed because its input was folded.

This is what we report in the CUDA prog guide https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/mathematical-functions.html#cuda-and-c-c-compliance

Floating-point operations may alter the bit patterns of input NaN payloads. Operations such as absolute value and negation may also not comply with the IEEE 754 requirement, which could result in the sign of a NaN being updated in an implementation-defined manner.

This gives you the freedom to explore the fastest implementation without being 100% IEEE754 compliant

lennartvoelz and others added 2 commits October 27, 2026 23:40
- update comments
- check sign of fabs(-0.0) in new test case
- add fabs benchmark (adopted from isnan benchmark)
@lennartvoelz

lennartvoelz commented Oct 27, 2026 •

Copy link
Copy Markdown
Contributor Author

@fbusato Thanks, I relaxed the new NaN tests accordingly

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

coderabbitai Bot commented Oct 27, 2026 •

Copy link
Copy Markdown
Contributor

Review in Change Stack →

Important

Review skipped

Review was skipped as selected files did not have any reviewable changes.

Run configuration
  • Source mode: Auto
  • Repository: NVIDIA/cccl
  • Branch: main
  • File: .coderabbit.yaml
  • Commit: 8004db6
  • Layers: Repository YAML
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: 1af59178-27e6-4446-aae4-eefbab8bc1ec
Commits

Reviewing files that changed from the base of the PR and between 0133ad4 and 56ecf2d.

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

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: 06b483c
  • Layers: Repository YAML
  • Review profile: CHILL
  • Plan: Enterprise
  • Run ID: 8ce39ae2-338c-47b4-9353-fe0c5b8341ac
Commits

Reviewing files that changed from the base of the PR and between 43c1c39 and 0133ad4.

Files selected for processing: 1
  • libcudacxx/include/cuda/std/__cmath/abs.h
Files skipped from review as they are similar to previous changes: 1
  • libcudacxx/include/cuda/std/__cmath/abs.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 cuda::std::fabs for double-precision values on certain CUDA targets. The operation preserves NaN status and clears the sign of negative zero, with a more efficient implementation on targets with slower FP64 support.
  • Tests
    • Added runtime checks for NaN values and negative zero.
  • Benchmarks
    • Added a benchmark for double-precision fabs.

The change adds an integer-based runtime implementation of cuda::std::fabs for double on slow-FP64 targets. It also adds runtime checks for NaNs and negative zero, and a CUDA benchmark.

Changes: Double fabs implementation and validation

Layer / File(s) Summary
Integer-based fabs dispatch
libcudacxx/include/cuda/std/__cmath/abs.h
The helper clears the sign bit of a double using integer operations. Runtime calls use this path on slow-FP64 targets; constant evaluation and other targets retain the existing paths.
Runtime correctness checks
libcudacxx/test/libcudacxx/std/numerics/c.math/fp_abs.pass.cpp
Runtime checks cover quiet and signaling NaNs with varied signs and payloads, and verify that fabs(-0.0) has no sign bit.
CUDA fabs benchmark
libcudacxx/benchmarks/bench/fabs/basic.cu
The benchmark applies cuda::std::fabs to a double, combines the result storage with rotated input storage, and configures its launch from occupancy and device SM count.

Possibly related issues

  • NVIDIA/cccl issue 11894 requests integer-path optimizations for double-precision <cmath> functions, including fabs. This change adds an integer path for double fabs.

Suggested reviewers: miscco, davebayer

Change: Refactor

Merge Risk: Minimal · previous review 0133a

No actionable issue remains; this change is mergeable after normal checks.

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

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

fbusato commented Nov 3, 2026

Copy link
Copy Markdown
Contributor

/ok to test 56ecf2d

@github-actions

github-actions Bot commented Nov 3, 2026

Copy link
Copy Markdown
Contributor

🔬 Benchmark SASS comparison

⚠️ The SASS changed for 9 of 100 CUB and 2 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\.merge_sort\.keys\.base$'
      - '^cub\.bench\.merge_sort\.pairs\.base$'
      - '^cub\.bench\.reduce\.custom\.base$'
      - '^cub\.bench\.reduce\.deferred_deterministic\.base$'
      - '^cub\.bench\.reduce\.deterministic\.base$'
      - '^cub\.bench\.scan\.exclusive\.custom\.base$'
      - '^cub\.bench\.segmented_reduce\.custom\.base$'
      - '^cub\.bench\.segmented_scan\.custom\.base$'
      - '^cub\.bench\.transform\.complex_cmp\.base$'
      - '^thrust\.bench\.scan\.exclusive\.max\.base$'
      - '^thrust\.bench\.scan\.inclusive\.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.merge_sort.keys.base sm_110, sm_120, sm_75
cub.bench.merge_sort.pairs.base sm_110, sm_120, sm_75
cub.bench.reduce.custom.base sm_110, sm_120, sm_75
cub.bench.reduce.deferred_deterministic.base sm_110, sm_120, sm_75
cub.bench.reduce.deterministic.base sm_110, sm_120, sm_75
cub.bench.scan.exclusive.custom.base sm_110, sm_120, sm_75
cub.bench.segmented_reduce.custom.base sm_110, sm_120, sm_75
cub.bench.segmented_scan.custom.base sm_110, sm_120, sm_75
cub.bench.transform.complex_cmp.base sm_110, sm_120, sm_75
thrust.bench.scan.exclusive.max.base sm_120, sm_110, sm_75
thrust.bench.scan.inclusive.max.base sm_120, sm_110, sm_75

‼️ Summary of Differences ‼️

Showing 10/11 summaries.

cub.bench.merge_sort.keys.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/50085 diff lines, 35200 changes. - ⬇️ Full diff

--- base/cub.bench.merge_sort.keys.base.sm_110
+++ test/cub.bench.merge_sort.keys.base.sm_110
@@ -1935,7 +1935,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, unsigned int, less_t, cuda::std::__4::complex<float>, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 UIADD3 UR4, UPT, UPT, UR4, -0x1, URZ ;
 UISETP.GE.U32.AND UP0, UPT, UR5, UR4, UPT ;
 USHF.L.U32 UR10, UR5, 0xb, URZ ;
-BRA.U UP0, <+0x25c40> ;
+BRA.U UP0, <+0x25e90> ;
 S2R R51, SR_TID.X ;
 LDCU.64 UR4, c[0x0][0x30] ;
 ACQBULK ;
@@ -1943,10 +1943,10 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, unsigned int, less_t, cuda::std::__4::complex<float>, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 LOP3.LUT R0, R0, 0x1f00, RZ, 0xc0, !PT ;
 LOP3.LUT R2, R0, 0x1f, R51, 0xf8, !PT ;
 IADD3 R2, P0, PT, R2, UR10, RZ ;
-SHF.L.U32 R34, R2, 0x3, RZ ;
 IMAD.X R33, RZ, RZ, RZ, P0 ;
+IMAD.SHL.U32 R34, R2, 0x8, RZ ;
+SHF.L.U64.HI R33, R2, 0x3, R33 ;
 IADD3 R18, P0, PT, R34, UR4, RZ ;
-SHF.L.U64.HI R33, R2, 0x3, R33 ;
 IADD3.X R19, PT, PT, R33, UR5, RZ, P0, !PT ;
 LDG.E.64 R16, desc[UR8][R18.64] ;
 LDG.E.64 R14, desc[UR8][R18.64+0x100] ;
@@ -1960,9 +1960,9 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, unsigned int, less_t, cuda::std::__4::complex<float>, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 UMOV UR4, 0x400 ;
 S2R R21, SR_LANEID ;
 ULEA UR4, UR5, UR4, 0x18 ;
-IMAD.IADD R0, R0, 0x1, R21 ;
+IADD3 R0, PT, PT, R0, R21, RZ ;
+LEA.HI.SX32 R50, R0.reuse, R0, 0x1b ;
 IMAD R20, R21, 0x7, R0 ;
-LEA.HI.SX32 R50, R0.reuse, R0, 0x1b ;
 IADD3 R21, PT, PT, R0.reuse, 0x20, RZ ;
 IADD3 R23, PT, PT, R0.reuse, 0x40, RZ ;
 IADD3 R25, PT, PT, R0.reuse, 0x60, RZ ;
@@ -2044,7 +2044,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType *, unsigned int, less_t, cuda::std::__4::complex<float>, cub::_V_300700_SM_750_800_900_1000_1100_1200::NullType>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 @!P0 BRA <+0x50> ;
 IMAD.MOV.U32 R0, RZ, RZ, R3 ;
 MOV R2, 0x760 ;
cub.bench.merge_sort.pairs.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/119748 diff lines, 57352 changes. - ⬇️ Full diff

--- base/cub.bench.merge_sort.pairs.base.sm_110
+++ test/cub.bench.merge_sort.pairs.base.sm_110
@@ -8784,7 +8784,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, int *, cuda::std::__4::complex<float> *, int *, unsigned int, less_t, cuda::std::__4::complex<float>, int>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 UIADD3 UR4, UPT, UPT, UR4, -0x1, URZ ;
 UISETP.GE.U32.AND UP0, UPT, UR5, UR4, UPT ;
 USHF.L.U32 UR6, UR5, 0xb, URZ ;
-BRA.U UP0, <+0x27140> ;
+BRA.U UP0, <+0x27390> ;
 S2R R47, SR_TID.X ;
 LDCU.64 UR4, c[0x0][0x38] ;
 ACQBULK ;
@@ -8939,7 +8939,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, int *, cuda::std::__4::complex<float> *, int *, unsigned int, less_t, cuda::std::__4::complex<float>, int>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 ISETP.GT.U32.AND P0, PT, R2, 0x727fffff, PT ;
 @!P0 BRA <+0x40> ;
 MOV R2, 0xa40 ;
-CALL.REL.NOINC <+0x54350> ;
+CALL.REL.NOINC <+0x54850> ;
 BRA <+0x50> ;
 FMUL.FTZ R9, R0, R3 ;
 FMUL.FTZ R3, R3, 0.5 ;
@@ -8970,7 +8970,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, int *, cuda::std::__4::complex<float> *, int *, unsigned int, less_t, cuda::std::__4::complex<float>, int>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 @!P0 BRA <+0x50> ;
 IMAD.MOV.U32 R0, RZ, RZ, R9 ;
 MOV R2, 0xc30 ;
-CALL.REL.NOINC <+0x54160> ;
+CALL.REL.NOINC <+0x54660> ;
 BRA <+0x50> ;
 FMUL.FTZ R0, R9, R10 ;
 FMUL.FTZ R2, R10, 0.5 ;
@@ -8979,7 +8979,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, int *, cuda::std::__4::complex<float> *, int *, unsigned int, less_t, cuda::std::__4::complex<float>, int>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 NOP ;
 BSYNC.RECONVERGENT B0 ;
 SHF.L.U32 R2, R14, 0x1, RZ ;
-BSSY.RECONVERGENT B2, <+0xcc0> ;
+BSSY.RECONVERGENT B2, <+0xcd0> ;
 FSETP.NEU.AND P1, PT, R8, RZ, PT ;
 IMAD.MOV.U32 R28, RZ, RZ, R31 ;
 ISETP.GT.U32.AND P0, PT, R2, RZ, PT ;
@@ -8991,7 +8991,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::DeviceMergeSortBlockSortKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::merge_sort::policy_selector_from_types<cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float> *, int *, cuda::std::__4::complex<float> *, int *, unsigned int, less_t, cuda::std::__4::complex<float>, int>(bool, T2, T3, T4, T5, T6, T8 *, T9 *, T7, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::vsmem_t)
 @P1 FMUL R13, R3, R0 ;
cub.bench.reduce.custom.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/237090 diff lines, 222140 changes. - ⬇️ Full diff

--- base/cub.bench.reduce.custom.base.sm_110
+++ test/cub.bench.reduce.custom.base.sm_110
@@ -2742,18 +2742,18 @@ 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<cuda::std::__4::complex<float>, unsigned int, max_t, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, unsigned int, unsigned int, max_t, cuda::std::__4::complex<float>, cuda::std::__4::complex<float>, 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 R3, SR_CTAID.X ;
 LDCU UR5, c[0x0][0x28] ;
-BSSY.RECONVERGENT B0, <+0x66a80> ;
+BSSY.RECONVERGENT B0, <+0x67120> ;
 S2R R2, SR_TID.X ;
 LDCU UR4, c[0x0][0x54] ;
 LDCU.64 UR6, c[0x0][0x0] ;
 ULOP3.LUT UP0, URZ, UR5, 0x7, URZ, 0xc0, !UPT ;
 USHF.L.U32 UR4, UR4, 0xb, URZ ;
 IMAD.SHL.U32 R14, R3, 0x800, RZ ;
-BRA.U UP0, <+0x334c0> ;
+BRA.U UP0, <+0x33810> ;
 LDCU UR8, c[0x0][0x50] ;
 IADD3 R17, PT, PT, -R14, UR8, RZ ;
 ISETP.GT.U32.AND P0, PT, R17, 0x7ff, PT ;
-@P0 BRA <+0x18440> ;
+@P0 BRA <+0x185d0> ;
 ISETP.GE.U32.AND P0, PT, R2, R17, PT ;
 BSSY.RECONVERGENT B1, <+0xb0> ;
 IMAD.MOV.U32 R18, RZ, RZ, R2 ;
@@ -2767,8 +2767,8 @@ 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<cuda::std::__4::complex<float>, unsigned int, max_t, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, unsigned int, unsigned int, max_t, cuda::std::__4::complex<float>, cuda::std::__4::complex<float>, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 NOP ;
 BSYNC.RECONVERGENT B1 ;
 ISETP.GE.U32.AND P0, PT, R18, R17, PT ;
-BSSY.RECONVERGENT B3, <+0xff0> ;
-@P0 BRA <+0xfc0> ;
+BSSY.RECONVERGENT B3, <+0x1000> ;
+@P0 BRA <+0xfd0> ;
 LDC.64 R20, c[0x0][0x28] ;
 IMAD.IADD R14, R14, 0x1, R18 ;
 IMAD.WIDE.U32 R22, R14, 0x8, R20 ;
@@ -2789,7 +2789,7 @@ 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<cuda::std::__4::complex<float>, unsigned int, max_t, (cuda::execution::__4::determinism::__determinism_t)1>, (bool)1, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, unsigned int, unsigned int, max_t, cuda::std::__4::complex<float>, cuda::std::__4::complex<float>, cuda::std::__4::identity>(T3, T4, T6, cub::_V_300700_SM_750_800_900_1000_1100_1200::GridEvenShare<T5>, T7, T9, T10)
 ISETP.GT.U32.AND P0, PT, R0, 0x727fffff, PT ;
 @!P0 BRA <+0x40> ;
 MOV R0, 0x330 ;
-CALL.REL.NOINC <+0x667f0> ;
cub.bench.reduce.deferred_deterministic.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/269 diff lines, 168 changes. - ⬇️ Full diff

--- base/cub.bench.reduce.deferred_deterministic.base.sm_110
+++ test/cub.bench.reduce.deferred_deterministic.base.sm_110
@@ -13158,7 +13158,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, R31, R2, PT ;
-BSSY.RECONVERGENT B1, <+0x2460> ;
+BSSY.RECONVERGENT B1, <+0x2480> ;
 CS2R R20, SRZ ;
 STL.64 [R1+0x10], RZ ;
 CS2R R18, SRZ ;
@@ -13168,21 +13168,21 @@ 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)
 CS2R R4, SRZ ;
 CS2R R14, SRZ ;
 BAR.SYNC.DEFER_BLOCKING 0x0 ;
-@P0 BRA <+0x23a0> ;
-BSSY.RECONVERGENT B0, <+0x21d0> ;
+@P0 BRA <+0x23c0> ;
+BSSY.RECONVERGENT B0, <+0x21f0> ;
 IMAD.MOV.U32 R8, RZ, RZ, RZ ;
 CS2R R16, SRZ ;
 CS2R R14, SRZ ;
+ISETP.GE.U32.AND P0, PT, R31, R2.reuse, PT ;
 IMAD R7, R10.reuse, 0x80, R31 ;
-ISETP.GE.U32.AND P0, PT, R31, R2.reuse, PT ;
 CS2R R26, SRZ ;
 CS2R R18, SRZ ;
 IMAD R35, R10, 0x80, R7 ;
 ISETP.GE.U32.AND P1, PT, R7, R2, PT ;
 IMAD R9, R10, 0x80, R35 ;
 ISETP.GE.U32.AND P2, PT, R35, R2, PT ;
+@!P0 LDC.64 R28, c[0x0][0x88] ;
 ISETP.GE.U32.AND P3, PT, R9, R2, PT ;
-@!P0 LDC.64 R28, c[0x0][0x88] ;
 @!P1 LDC.64 R32, c[0x0][0x88] ;
 @!P2 LDC.64 R4, c[0x0][0x88] ;
 @!P0 IMAD.WIDE.U32 R28, R31, 0x8, R28 ;
@@ -13192,50 +13192,52 @@ 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)
 ISETP.GE.U32.AND P4, PT, R31, R2, PT ;
 @!P3 LDC.64 R6, c[0x0][0x88] ;
cub.bench.reduce.deterministic.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/292 diff lines, 206 changes. - ⬇️ Full diff

--- base/cub.bench.reduce.deterministic.base.sm_110
+++ test/cub.bench.reduce.deterministic.base.sm_110
@@ -143,7 +143,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>, 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)
 BSYNC.RECONVERGENT B0 ;
 LDCU UR4, c[0x0][0x98] ;
 STL.64 [R1], RZ ;
-BSSY.RECONVERGENT B1, <+0x25d0> ;
+BSSY.RECONVERGENT B1, <+0x2600> ;
 CS2R R12, SRZ ;
 LDCU.64 UR6, c[0x0][0x60] ;
 STL.64 [R1+0x10], RZ ;
@@ -156,8 +156,8 @@ 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)
 CS2R R14, SRZ ;
 BAR.SYNC.DEFER_BLOCKING 0x0 ;
 ISETP.GE.U32.AND P0, PT, R18, UR4, PT ;
-@P0 BRA <+0x24e0> ;
-BSSY.RECONVERGENT B0, <+0x2300> ;
+@P0 BRA <+0x2510> ;
+BSSY.RECONVERGENT B0, <+0x2330> ;
 IMAD.MOV.U32 R6, RZ, RZ, RZ ;
 CS2R R4, SRZ ;
 CS2R R14, SRZ ;
@@ -165,66 +165,69 @@ 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)
 LDCU UR4, c[0x0][0x98] ;
 CS2R R26, SRZ ;
 ISETP.GE.U32.AND P0, PT, R18, UR4, PT ;
-IMAD R11, R7, 0x80, R18 ;
-IMAD R20, R7, 0x80, R11 ;
-ISETP.GE.U32.AND P1, PT, R11, UR4, PT ;
+IMAD R10, R7, 0x80, R18 ;
+IMAD R11, R7, 0x80, R10 ;
+ISETP.GE.U32.AND P1, PT, R10, UR4, PT ;
 @!P0 LDC.64 R12, c[0x0][0x88] ;
-ISETP.GE.U32.AND P2, PT, R20, UR4, PT ;
-IMAD R22, R7, 0x80, R20 ;
-IMAD R10, R7, 0x80, R22 ;
+ISETP.GE.U32.AND P2, PT, R11, UR4, PT ;
+IMAD R22, R7, 0x80, R11 ;
 ISETP.GE.U32.AND P3, PT, R22, UR4, PT ;
 @!P1 LDC.64 R16, c[0x0][0x88] ;
cub.bench.scan.exclusive.custom.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/98018 diff lines, 37691 changes. - ⬇️ Full diff

--- base/cub.bench.scan.exclusive.custom.base.sm_110
+++ test/cub.bench.scan.exclusive.custom.base.sm_110
@@ -22310,11 +22310,11 @@ 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 cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 BAR.SYNC.DEFER_BLOCKING 0x0 ;
 S2R R4, SR_LANEID ;
 UISETP.GT.U32.AND UP0, UPT, UR5, 0x3, UPT ;
-BRA.U UP0, <+0x1715a0> ;
+BRA.U UP0, <+0x172d00> ;
 UISETP.GT.U32.AND UP0, UPT, UR5, 0x1, UPT ;
 LDCU UR18, c[0x0][0x50] ;
 ULOP3.LUT UR12, UR17, 0xfffffff8, URZ, 0xc0, !UPT ;
-BRA.U UP0, <+0xb4f20> ;
+BRA.U UP0, <+0xb5ac0> ;
 UPRMT UR4, UR18, 0x7770, URZ ;
 ISETP.NE.U32.AND P0, PT, RZ, UR4, PT ;
 @P0 LDC R3, c[0x0][0x5c] ;
@@ -22373,7 +22373,7 @@ 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 cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 LOP3.LUT R198, R5, 0x800000, RZ, 0xfc, !PT ;
 @!P2 BRA <+0x40> ;
 MOV R0, 0x1040 ;
-CALL.REL.NOINC <+0x19c280> ;
+CALL.REL.NOINC <+0x19dc80> ;
 BRA <+0x50> ;
 FMUL.FTZ R13, R3, R6 ;
 FMUL.FTZ R2, R6, 0.5 ;
@@ -22394,7 +22394,7 @@ 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 cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 IMAD.MOV.U32 R3, RZ, RZ, R9 ;
 MOV R0, R204 ;
 MOV R6, 0x1190 ;
-CALL.REL.NOINC <+0x19c290> ;
+CALL.REL.NOINC <+0x19dc90> ;
 IMAD.MOV.U32 R215, RZ, RZ, R0 ;
 NOP ;
 BSYNC.RECONVERGENT B0 ;
@@ -22406,7 +22406,7 @@ 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 cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 BSSY.RECONVERGENT B1, <+0x60> ;
 IMAD.MOV.U32 R3, RZ, RZ, R10 ;
 MOV R0, 0x1250 ;
-CALL.REL.NOINC <+0x19c070> ;
+CALL.REL.NOINC <+0x19da70> ;
cub.bench.segmented_reduce.custom.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/152822 diff lines, 135342 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_reduce.custom.base.sm_110
+++ test/cub.bench.segmented_reduce.custom.base.sm_110
@@ -7748,7 +7748,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_reduce::DeviceFixedSizeSegmentedReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_reduce::policy_selector_from_types<cuda::std::__4::complex<float>, int, max_t>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, int, max_t, cuda::std::__4::complex<float>, cuda::std::__4::complex<float>>(T2, T3, T4, int, T5, T6, T7 *, int, int)
 S2R R14, SR_TID.X ;
 IMAD.U32 R3, RZ, RZ, UR4 ;
 ISETP.GT.AND P0, PT, R3, 0x8, PT ;
-@P0 BRA <+0x10850> ;
+@P0 BRA <+0x10960> ;
 LDCU UR4, c[0x0][0x3c] ;
 IMAD R14, R15, 0x100, R14 ;
 ISETP.GE.AND P0, PT, R14, UR4, PT ;
@@ -7764,7 +7764,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_reduce::DeviceFixedSizeSegmentedReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_reduce::policy_selector_from_types<cuda::std::__4::complex<float>, int, max_t>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, int, max_t, cuda::std::__4::complex<float>, cuda::std::__4::complex<float>>(T2, T3, T4, int, T5, T6, T7 *, int, int)
 LDCU.64 UR4, c[0x0][0x28] ;
 IMAD.WIDE.U32 R6, R14, R3, RZ ;
 SHF.R.S32.HI R5, RZ, 0x1f, R3 ;
-BSSY.RECONVERGENT B0, <+0xf7d0> ;
+BSSY.RECONVERGENT B0, <+0xf8d0> ;
 SHF.L.U32 R8, R6, 0x3, RZ ;
 IMAD R5, R5, R14, RZ ;
 IMAD.IADD R5, R7, 0x1, R5 ;
@@ -7772,16 +7772,16 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_reduce::DeviceFixedSizeSegmentedReduceKernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_reduce::policy_selector_from_types<cuda::std::__4::complex<float>, int, max_t>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, int, max_t, cuda::std::__4::complex<float>, cuda::std::__4::complex<float>>(T2, T3, T4, int, T5, T6, T7 *, int, int)
 IADD3 R4, P1, PT, R8, UR4, RZ ;
 LOP3.LUT P0, RZ, R4, 0x7, RZ, 0xc0, !PT ;
 IADD3.X R5, PT, PT, R9, UR5, RZ, P1, !PT ;
-@P0 BRA <+0x7b90> ;
+@P0 BRA <+0x7c10> ;
 ISETP.NE.U32.AND P0, PT, R3, 0x8, PT ;
-@!P0 BRA <+0x1050> ;
+@!P0 BRA <+0x1060> ;
 ISETP.GE.AND P0, PT, R3, 0x1, PT ;
 CS2R R10, SRZ ;
 IMAD.MOV.U32 R0, RZ, RZ, RZ ;
 @P0 MOV R0, 0x1 ;
 @P0 LDG.E.64.CONSTANT R10, desc[UR6][R4.64] ;
 ISETP.GE.AND P0, PT, R0, R3, PT ;
-@P0 BRA <+0xf6a0> ;
+@P0 BRA <+0xf7a0> ;
 IMAD.WIDE.U32 R4, R0, 0x8, R8 ;
 IMAD.IADD R15, R0, 0x1, -R3 ;
 IADD3 R12, P0, PT, R4, UR4, RZ ;
cub.bench.segmented_scan.custom.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/25477 diff lines, 12328 changes. - ⬇️ Full diff

--- base/cub.bench.segmented_scan.custom.base.sm_110
+++ test/cub.bench.segmented_scan.custom.base.sm_110
@@ -9542,7 +9542,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::device_segmented_scan_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::policy_selector_from_types<cuda::std::__4::complex<float>>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, const int *, const int *, const int *, int, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float>, (bool)0, cuda::std::__4::complex<float>>(T2, T3, T4, T5, T6, T7, T8, T9, int)
 UIMAD.WIDE.U32 UR6, UR4, 0x4, UR12 ;
 UIMAD.WIDE.U32 UR14, UR4, 0x4, UR14 ;
 UIMAD.WIDE.U32 UR10, UR4, 0x4, UR10 ;
-BRA.U UP0, <+0x19880> ;
+BRA.U UP0, <+0x19a20> ;
 IMAD.U32 R4, RZ, RZ, UR6 ;
 IMAD.U32 R2, RZ, RZ, UR14 ;
 IMAD.U32 R5, RZ, RZ, UR7 ;
@@ -9558,7 +9558,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::device_segmented_scan_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::policy_selector_from_types<cuda::std::__4::complex<float>>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, const int *, const int *, const int *, int, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float>, (bool)0, cuda::std::__4::complex<float>>(T2, T3, T4, T5, T6, T7, T8, T9, int)
 ISETP.NE.U32.AND P0, PT, RZ, UR5, PT ;
 @!P0 EXIT ;
 IMAD.U32 R2, RZ, RZ, UR10 ;
-MOV R3, UR11 ;
+IMAD.U32 R3, RZ, RZ, UR11 ;
 LDG.E R45, desc[UR8][R2.64] ;
 FMUL R0, R8.reuse, 0.5 ;
 FMUL R4, R9.reuse, 0.5 ;
@@ -9567,7 +9567,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::device_segmented_scan_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::policy_selector_from_types<cuda::std::__4::complex<float>>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, const int *, const int *, const int *, int, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float>, (bool)0, cuda::std::__4::complex<float>>(T2, T3, T4, T5, T6, T7, T8, T9, int)
 FADD R0, |R0|, -RZ ;
 FADD R4, |R4|, -RZ ;
 S2R R38, SR_TID.X ;
-IMAD.MOV.U32 R32, RZ, RZ, 0x3f6ee581 ;
+HFMA2 R32, -RZ, RZ, 1.857421875, -1409 ;
 IMNMX R35, R36, R5, !PT ;
 IMAD.MOV.U32 R30, RZ, RZ, 0x4016cbe4 ;
 R2UR UR6, R0 ;
@@ -9621,7 +9621,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::device_segmented_scan_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::policy_selector_from_types<cuda::std::__4::complex<float>>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, const int *, const int *, const int *, int, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float>, (bool)0, cuda::std::__4::complex<float>>(T2, T3, T4, T5, T6, T7, T8, T9, int)
 BSSY.RECONVERGENT B0, <+0x60> ;
 IMAD.MOV.U32 R3, RZ, RZ, R6 ;
 MOV R0, 0x6c0 ;
-CALL.REL.NOINC <+0x5c570> ;
+CALL.REL.NOINC <+0x5cad0> ;
 NOP ;
 BSYNC.RECONVERGENT B0 ;
 BRA <+0x50> ;
@@ -9644,8 +9644,8 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::device_segmented_scan_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::segmented_scan::policy_selector_from_types<cuda::std::__4::complex<float>>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, const int *, const int *, const int *, int, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, cuda::std::__4::complex<float>, (bool)0, cuda::std::__4::complex<float>>(T2, T3, T4, T5, T6, T7, T8, T9, int)
 UIADD3 UR10, UPT, UPT, -UR20, UR10, URZ ;
cub.bench.transform.complex_cmp.base - sm_110 - ⚠️ Significant

Floating-point absolute values use explicit sign masks The diffs add LOP3.LUT instructions that clear sign bits and remove absolute-value modifiers from floating-point comparisons, selections, or conversions. This signature occurs in all 27 changed pairs across nine targets on sm_75, sm_110, and sm_120.

Showing 40/2150 diff lines, 733 changes. - ⬇️ Full diff

--- base/cub.bench.transform.complex_cmp.base.sm_110
+++ test/cub.bench.transform.complex_cmp.base.sm_110
@@ -156,14 +156,14 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<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<cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>>, 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<bool>>>, long, cuda::__4::always_true, less_t, bool *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *>(T2, int, bool, T3, T4, T5, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 SYNCS.PHASECHK.TRANS64.TRYWAIT P0, [R9+URZ], R0 ;
 @!UP0 ULOP3.LUT UP3, URZ, UR5, 0x7f, URZ, 0xc0, !UPT ;
 @!UP0 UPLOP3.LUT UP1, UPT, UPT, UPT, UP3, 0x40, 0x4 ;
-@!P0 BRA <+0x12eb0> ;
-BRA.U !UP1, 0x11130 ;
+@!P0 BRA <+0x12fe0> ;
+BRA.U !UP1, 0x11240 ;
 USHF.R.S32.HI UR5, URZ, 0x1f, UR17 ;
-BSSY.RECONVERGENT B0, <+0xfa00> ;
+BSSY.RECONVERGENT B0, <+0xfb00> ;
 ULEA.HI UR5, UR5, UR17, URZ, 0x4 ;
 USHF.R.S32.HI UR5, URZ, 0x4, UR5 ;
 ISETP.GE.AND P0, PT, R5, UR5, PT ;
-@P0 BRA <+0xf9a0> ;
+@P0 BRA <+0xfaa0> ;
 LDCU UR6, c[0x0][0x58] ;
 MOV R6, UR20 ;
 IMAD.U32 R3, RZ, RZ, UR4 ;
@@ -211,7 +211,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<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<cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>>, 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<bool>>>, long, cuda::__4::always_true, less_t, bool *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *>(T2, int, bool, T3, T4, T5, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 @!P0 BRA <+0x50> ;
 MOV R0, R7 ;
 MOV R72, 0x860 ;
-CALL.REL.NOINC <+0x12b70> ;
+CALL.REL.NOINC <+0x12ca0> ;
 BRA <+0x50> ;
 FMUL.FTZ R76, R7, R72 ;
 FMUL.FTZ R0, R72, 0.5 ;
@@ -242,7 +242,7 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<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<cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>>, 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<bool>>>, long, cuda::__4::always_true, less_t, bool *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *>(T2, int, bool, T3, T4, T5, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
 @!P0 BRA <+0x50> ;
 IMAD.MOV.U32 R0, RZ, RZ, R73 ;
 MOV R72, 0xa50 ;
-CALL.REL.NOINC <+0x12980> ;
+CALL.REL.NOINC <+0x12ab0> ;
 BRA <+0x50> ;
 FMUL.FTZ R76, R73, R72 ;
 FMUL.FTZ R0, R72, 0.5 ;
@@ -251,12 +251,12 @@ void cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::transform_kernel<cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::policy_selector_from_types<(bool)0, (bool)1, cuda::std::__4::tuple<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<cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>>, 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<bool>>>, long, cuda::__4::always_true, less_t, bool *, cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *>(T2, int, bool, T3, T4, T5, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::transform::kernel_arg<T6>...)
thrust.bench.scan.exclusive.max.base - sm_120 - ⚠️ Significant

Sign-bit masks added to max scan kernels The DeviceScanKernel code adds hundreds of LOP3.LUT sign-bit masks and uses unmodified operands where absolute-value modifiers were used before. This instruction-count and instruction-mix change repeats in exclusive and inclusive max scans for sm_75, sm_110, and sm_120.

Showing 40/41874 diff lines, 17057 changes. - ⬇️ Full diff

--- base/thrust.bench.scan.exclusive.max.base.sm_120
+++ test/thrust.bench.scan.exclusive.max.base.sm_120
@@ -4438,10 +4438,10 @@ 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<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 cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 UISETP.GT.U32.AND UP0, UPT, UR17, 0x7, UPT ;
 LDCU UR4, c[0x0][0x60] ;
 UMOV UR5, URZ ;
-BRA.U UP0, <+0x9a3a0> ;
+BRA.U UP0, <+0x9ad50> ;
 UISETP.GT.U32.AND UP0, UPT, UR17, 0x3, UPT ;
 LDCU UR7, c[0x0][0x50] ;
-BRA.U UP0, <+0x4ac50> ;
+BRA.U UP0, <+0x4b100> ;
 UPRMT UR6, UR7, 0x7770, URZ ;
 ISETP.NE.U32.AND P0, PT, RZ, UR6, PT ;
 @P0 LDC R5, c[0x0][0x5c] ;
@@ -4503,7 +4503,7 @@ 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<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 cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 @!P2 BRA <+0x50> ;
 MOV R4, R10 ;
 MOV R0, 0xf20 ;
-CALL.REL.NOINC <+0xb2ca0> ;
+CALL.REL.NOINC <+0xb37d0> ;
 BRA <+0x50> ;
 FMUL.FTZ R9, R10, R5 ;
 FMUL.FTZ R4, R5, 0.5 ;
@@ -4523,7 +4523,7 @@ 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<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 cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 @!P0 BRA <+0x50> ;
 MOV R6, R87 ;
 MOV R4, 0x1060 ;
-CALL.REL.NOINC <+0xb2cc0> ;
+CALL.REL.NOINC <+0xb37f0> ;
 MOV R97, R0 ;
 NOP ;
 BSYNC.RECONVERGENT B0 ;
@@ -4535,7 +4535,7 @@ 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<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 cuda::std::__4::complex<float>>>, 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<cuda::std::__4::complex<float>>>, cuda::std::__4::complex<float>, unsigned int, max_t, (bool)0, (bool)0>, const cuda::std::__4::complex<float> *, cuda::std::__4::complex<float> *, cub::_V_300700_SM_750_800_900_1000_1100_1200::ScanTileState<cuda::std::__4::complex<float>, (bool)1>, max_t, cub::_V_300700_SM_750_800_900_1000_1100_1200::detail::InputValue<cuda::std::__4::complex<float>, cuda::std::__4::complex<float> *>, unsigned int, unsigned int, cuda::std::__4::complex<float>, (bool)0, (bool)0, (bool)0, cuda::std::__4::complex<float>>(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)
 BSSY.RECONVERGENT B1, <+0x60> ;
 MOV R4, R16 ;
 MOV R0, 0x1120 ;
-CALL.REL.NOINC <+0xb2aa0> ;
+CALL.REL.NOINC <+0xb35d0> ;
 NOP ;

@github-actions

This comment has been minimized.

@fbusato

fbusato commented Nov 4, 2026

Copy link
Copy Markdown
Contributor

/sass diff ok 56ecf2d

@github-actions

github-actions Bot commented Nov 4, 2026

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 3h 05m: Pass: 100%/230 | Total: 6d 05h | Max: 3h 04m | Hits: 41%/2276764

See results here.

@davebayer
davebayer merged commit 2edc582 into NVIDIA:main Nov 6, 2026
495 of 497 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.

3 participants