[CUB] Track warp reduction optimizations #11999

The issue consolidates implemented and potential optimizations for cub::WarpReduce.

The goal is to keep track of optimizations and provide a common place to track follow-up work.

Context:

• [FEA]: Add missing cub::WarpReduce overloadings and optimizations #3853
• Optimize Warp Reduce #4312

Implemented

Hardware reduction (redux)

• Centralize hardware reduction implementations and capability traits in warp_redux.cuh, simplify dispatch, and introduce native floating-point min/max support.
  • [CUB] Fully support warp redux  #9441
• Extend native floating-point min/max reduction support to SM107.
  • Add Warp redux min/max for Rubin #11100
• Improve dispatch between hardware reduction and shuffle reduction. Use hardware reduction for full 32-thread logical warps when the operator/type is supported; support partial valid counts through identity padding where available. Smaller logical warps retain shuffle reduction.
  • [CUB] A more efficient WarpReduce dispatch #11164]
• Refine integer eligibility traits for hardware reduction.
  • warp_redux uses is_integer instead of is_integral #11951
• Prefer integer-mapped floating-point min/max over native FP32 redux: compare both approaches and disable the slower native path. The native implementation remains in the source, but its capability trait is disabled.
  • Optimize WarpReduce for floating-point types (Min/Max) #11974

Operator and type specializations

• Wide integer bitwise reductions: reduce independent 32-bit chunks with hardware instructions. Reported improvements: approximately 4× for 64-bit and 25× for 128-bit integers.
  • Optimize WarpReduce for bitwise operators #11279
• Extend the bitwise specialization to signed integer types.
  • Fix [QA Agent] Cooperative reduction - Warp #11645
• Wide integer sum: use carry-aware decomposition with three hardware reductions for 64-bit integers and five for 128-bit integers. Reported improvements: 2.4× and 5.8×, respectively.
  • Optimize WarpReduce for large integers (Sum) #11623
• Wide integer min/max: reduce words from most significant to least significant, retaining candidate lanes. Use two hardware reductions for 64-bit integers and four for 128-bit integers. Reported improvements: 4× and 20×, respectively.
  • Optimize WarpReduce for large integers (Min/Max) #11827
• Floating-point min/max: map values to order-preserving integer representations and reuse integer hardware reductions. Cover __half, __nv_bfloat16, float, double, and __float128, subject to target/type availability.
  • Optimize WarpReduce for floating-point types (Min/Max) #11974
• Boolean reductions: use warp vote instructions for logical_and and logical_or, including smaller power-of-two logical warps and partial valid counts. Reported results on RTX A6000: 9× lower latency and 6× higher throughput.
  • Use vote instructions in WarpReduce #11878

Related work

• Add Min/Max and multiple-items-per-thread overloads.
  • Add new WarpReduce overloadings #3884
• Add internal integer splitting/merging and floating-point comparison mappings.
  • CUB - Add internal integer utils and tests (Split WarpReduce PR) #5314
• Extend internal thread operators, SIMD support, and logical warp utilities.
  • CUB - Add internal thread and warp utils (Split WarpReduce PR) #5317
• Clean up the WarpReduce implementation, centralize existing optimizations, and add hardware specializations for bitwise AND/OR.
  • Improve our WarpReduce implementation #6814
• Add device-side WarpReduce throughput benchmarks covering integer, floating-point, and complex types, including extended floating-point types; validate generated instructions.
  • Add WarpReduce Device-Side Benchmarks #6431
• Add BlockReduce benchmarks to compare its algorithms and evaluate downstream performance.
  • Add BlockReduce benchmark #11755
• Improve generic cuda::device::warp_shuffle code generation and register usage.
  • Optimize cuda::device::warp_shuffle #10035
• Reenable HW warp redux min/max f32
  • Reenable HW warp redux min/max f32 #12004

Potential follow-up work

• Segmented hardware reductions: evaluate hardware reduction and vote paths for head-/tail-segmented reductions, accounting for segment-mask construction and segment size. The current shuffle implementation still performs reduction steps.
  • Optimize Warp Reduce #4312
• Non-power-of-two logical warps: revisit the shared-memory fallback and evaluate shuffle-based alternatives, preserving participation and ordering requirements.
  • Optimize Warp Reduce #4312
• Shuffle fallback: revisit predication, compile-time masks, reduction-tree structure, and recursive template instantiation for unsupported hardware combinations.
• Packed/SIMD reductions: evaluate specialized sum and min/max paths for __half2, __nv_bfloat162, float2, short2, and ushort2, including DPX instructions where applicable.
  • Optimize Warp Reduce #4312
• Complex sum: evaluate mapping complex values to packed/vector representations and reusing optimized component-wise reductions.
  • Optimize Warp Reduce #4312
• Absolute min/max operators: consider introducing suitable operators and enabling the corresponding hardware instructions.
  • TODO in warp_redux.cuh
• Result broadcast: evaluate an explicit API for returning the aggregate to every participating lane.
  • Optimize Warp Reduce #4312