Fix BlockScan accumulator type handling #6443
Open
+81
−7
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Aminsed
force-pushed
the
fix-blockscan-accum
branch
from
bc27d15 to
c122fad
Compare
fbusato
requested changes
Nov 3, 2025
cub/cub/block/block_reduce.cuh
Outdated
| // Reduce partials | ||
| T partial = cub::ThreadReduce(inputs, reduction_op); | ||
| T partial = | ||
| cub::ThreadReduce<::cuda::std::remove_reference_t<decltype(inputs)>, ReductionOp, T, T>(inputs, reduction_op); |
There was a problem hiding this comment.
this looks a partial solution. It could regress for small integer types. For example, reduction/scan over int8_t. It is better to perform the computation with 32-bit and cast back at the end
There was a problem hiding this comment.
Good catch. I dropped the explicit T and taught ThreadReduce to keep its __accumulator_t promotion, so int8_t still widens to 32-bit.
cub/cub/block/block_reduce.cuh
Outdated
| { | ||
| // Reduce partials | ||
| T partial = cub::ThreadReduce(inputs, ::cuda::std::plus<>{}); | ||
| T partial = cub::ThreadReduce<::cuda::std::remove_reference_t<decltype(inputs)>, ::cuda::std::plus<>, T, T>( |
There was a problem hiding this comment.
nit. please isolate the first template parameter with using to improve readability
There was a problem hiding this comment.
@fbusato Thanks! I reverted that spot to plain ThreadReduce(inputs, …), so there’s nothing left to alias. Let me know if you’d still like a using helper there.
Aminsed
force-pushed
the
fix-blockscan-accum
branch
from
85c2484 to
0bcd084
Compare
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Summary
ThreadReduceaccumulator types pinned to the block value type acrossBlockScanandBlockReduceBlockScanwith a functor returning a wider typeMotivation
#5668shows thatBlockScanwidens the accumulator when the scan functor returns a wider type than the block value. That implicit widening breaks user code that relies on the original type and can even hit deleted overloads.Explanation
ThreadReducewas deducing its accumulator type from the functor instead of the block valueT. The patch explicitly instantiatesThreadReducewithAccumT = TeverywhereBlockScanandBlockReducedispatch through it, including the raking specialization. The new unit test exercises an operator that returnslong longforintinputs and verifies the accumulator remainsint.Rationale
ThreadReducecall sites; public APIs and template parameters stay the same.Testing
pre-commit run --files cub/cub/block/block_scan.cuh cub/cub/block/block_reduce.cuh cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh cub/test/catch2_test_block_scan.cu