optimize blockReduceWarp#9101
Conversation
|
Note Reviews pausedIt 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 Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
suggestion: WalkthroughRefactors BlockReduce warp-aggregate application so warp 0 deterministically cooperates to reduce per-warp aggregates (reduce_op_sync fast path, shuffle fallback), removes the atomic non-deterministic path, consolidates Sum/Reduce, and adds NVBench benchmark plus CI bench filter/GPU updates. ChangesWarp-0 Cooperative Reduction
Benchmarks and CI
Assessment against linked issues
Out-of-scope changes
Comment |
|
@charan-003, thanks a lot for the contribution. Even before starting the review, I would like to see if there are actual performance benefits for this approach. |
sure, working on it |
charan-003
force-pushed
the
optimize_blockreduce
branch
from
0b48748 to
cb6dd6c
Compare
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
🧹 Nitpick comments (2)
cub/cub/block/specializations/block_reduce_warp_reductions.cuh (2)
125-125: 💤 Low valuesuggestion: Use
static_cast<int>(warps)instead of C-style castint(warps)for consistency with thestatic_castused in the same expression.- const int num_warps = FullTile ? int(warps) : static_cast<int>(::cuda::ceil_div(num_valid, logical_warp_size)); + const int num_warps = FullTile ? static_cast<int>(warps) : static_cast<int>(::cuda::ceil_div(num_valid, logical_warp_size));
148-148: 💤 Low valuesuggestion: Per coding guidelines, variables not modified should be
const.dummy_storageis only passed to the constructor.- NullType dummy_storage; + const NullType dummy_storage{};
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: ffc63f63-2430-4ac6-8a11-ae99c579665d
📒 Files selected for processing (1)
cub/cub/block/specializations/block_reduce_warp_reductions.cuh
|
@charan-003 thanks for reporting them but I think a better approach would be:
|
Thanks a lot for the guidance! I didn't know the exact process... |
charan-003
force-pushed
the
optimize_blockreduce
branch
from
830c2b4 to
a949e20
Compare
There was a problem hiding this comment.
Actionable comments posted: 3
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 53e3ccfa-d3a5-484a-85fd-7cf18e522152
📒 Files selected for processing (3)
ci/bench.yamlcub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuhcub/benchmarks/bench/reduce/block_reduce_warp_reductions_sum.cu
✅ Files skipped from review due to trivial changes (1)
- cub/benchmarks/bench/reduce/block_reduce_warp_reductions_sum.cu
| template <int BlockThreads> | ||
| struct benchmark_op_t | ||
| { | ||
| template <typename T> |
There was a problem hiding this comment.
🛠️ Refactor suggestion | 🟠 Major | ⚡ Quick win
suggestion: Line 17 and Line 35 use a single-letter template parameter (T). Use _Tp-style naming here to satisfy the template-parameter naming rule for this codebase.
As per coding guidelines "Avoid single letter names for template parameters; use _Tp instead of _T".
Also applies to: 35-35
| __device__ __forceinline__ T operator()(T thread_data) const | ||
| { |
There was a problem hiding this comment.
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
# Read-only verification: confirm missing CCCL API annotations / noexcept in this file
rg -nP '^\s*(template\s*<[^>]+>\s*)?(?!.*_CCCL_(HOST_API|DEVICE_API|API)).*\b(operator\(\)|block_reduce_warp_reductions)\b' cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh
rg -nP '\b(operator\(\)|block_reduce_warp_reductions)\b(?!.*noexcept)' cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuhRepository: NVIDIA/cccl
Length of output: 391
🏁 Script executed:
#!/bin/bash
set -euo pipefail
sed -n '1,120p' cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh | nl -ba | sed -n '10,80p'Repository: NVIDIA/cccl
Length of output: 98
🏁 Script executed:
#!/bin/bash
set -euo pipefail
rg -n "operator\(\)|block_reduce_warp_reductions" cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuhRepository: NVIDIA/cccl
Length of output: 289
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Capture the full function signatures/headers for the two targets (best-effort)
perl -0777 -ne '
my $f=$_;
if($f =~ /(operator\(\)\s*const[^\n]*\n\s*\{)/s){ print "FOUND operator():\n$1\n"; }
if($f =~ /(block_reduce_warp_reductions[^\n]*\n\s*\{)/s){ print "FOUND block_reduce_warp_reductions():\n$1\n"; }
' cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuhRepository: NVIDIA/cccl
Length of output: 170
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Show the relevant top portion with line numbers (no nl dependency)
awk '{printf "%5d:%s\n", NR, $0}' cub/benchmarks/bench/reduce/block_reduce_warp_warp_reductions_base.cuh | sed -n '1,110p'Repository: NVIDIA/cccl
Length of output: 196
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Check for noexcept / CCCL API tokens near the declarations
rg -n "operator\(\)|block_reduce_warp_reductions" -n cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh
rg -n --context=3 "operator\(\)|block_reduce_warp_reductions" cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh
rg -n --context=0 "_CCCL_(HOST_API|DEVICE_API|API)" cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuhRepository: NVIDIA/cccl
Length of output: 1049
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Look for exception-throwing in the relevant functions/bodies
rg -n "\bthrow\b" cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuhRepository: NVIDIA/cccl
Length of output: 37
🏁 Script executed:
#!/bin/bash
set -euo pipefail
FILE="cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh"
# Print a small, line-numbered window around operator() and block_reduce_warp_reductions
sed -n '10,60p' "$FILE" | awk '{printf "%5d:%s\n", NR+10-1, $0}'Repository: NVIDIA/cccl
Length of output: 2048
🏁 Script executed:
#!/bin/bash
set -euo pipefail
FILE="cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh"
# Check for _CCCL_* usage and noexcept in/near the relevant symbols
rg -n "_CCCL_(HOST_API|DEVICE_API|API)" "$FILE"
rg -n "noexcept" "$FILE"
rg -n --context=2 "operator\(\)" "$FILE"
rg -n --context=2 "block_reduce_warp_reductions" "$FILE"Repository: NVIDIA/cccl
Length of output: 37
🏁 Script executed:
#!/bin/bash
set -euo pipefail
FILE="cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh"
# Look for exception throwing inside this file
rg -n "\bthrow\b" "$FILE"Repository: NVIDIA/cccl
Length of output: 37
🏁 Script executed:
#!/bin/bash
set -euo pipefail
FILE="cub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuh"
echo "== _CCCL_* tokens =="
rg -n "_CCCL_(HOST_API|DEVICE_API|API)" "$FILE" || true
echo "== noexcept tokens =="
rg -n "noexcept" "$FILE" || true
echo "== throw tokens =="
rg -n "\bthrow\b" "$FILE" || trueRepository: NVIDIA/cccl
Length of output: 117
important: Add required _CCCL_* API annotations and noexcept to benchmark functions.
benchmark_op_t::operator()(line 18) lacks_CCCL_DEVICE_APIandnoexcept.block_reduce_warp_reductions(line 36) lacks_CCCL_HOST_APIandnoexcept.
Suggested change
- __device__ __forceinline__ T operator()(T thread_data) const
+ _CCCL_DEVICE_API __forceinline__ auto operator()(_Tp thread_data) const noexcept -> _Tp-template <typename T>
-void block_reduce_warp_reductions(nvbench::state& state, nvbench::type_list<T>)
+template <typename _Tp>
+_CCCL_HOST_API void block_reduce_warp_reductions(nvbench::state& state, nvbench::type_list<_Tp>) noexcept
There was a problem hiding this comment.
Actionable comments posted: 1
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 519c38f3-5544-4f5a-9750-fc04a66589e0
📒 Files selected for processing (3)
ci/bench.yamlcub/benchmarks/bench/reduce/block_reduce_warp_reductions_base.cuhcub/benchmarks/bench/reduce/block_reduce_warp_reductions_sum.cu
✅ Files skipped from review due to trivial changes (1)
- ci/bench.yaml
|
/ok to test 80b5ac2 |
This comment has been minimized.
This comment has been minimized.
|
@fbusato the benchmarks shows a slowdown. can we use bit_ceil(warps) for the logical warp size. Does this approach look reasonable, or would you suggest a different direction? |
|
The code looks good and I really like the simplifications. Related to performance, some high-level thoughts:
|
|
/ok to test dae2e4a |
|
@charan-003 sorry for the delay. Please ping me every time you need to trigger the CI |
This comment has been minimized.
This comment has been minimized.
|
@fbusato can you please trigger the CI? |
|
/ok to test 4210958 |
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
| int max_blocks_per_SM = 0; | ||
| NVBENCH_CUDA_CALL_NOEXCEPT( | ||
| cudaOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_SM, kernel, BlockThreads, 0)); | ||
| const int grid_size = max_blocks_per_SM * num_SMs; |
There was a problem hiding this comment.
There are a couple of issues with this benchmarking approach that I tried to address in #8607:
- If another (version of the) compiler decides to use more or less registers in a way that changes occupancy, results become very hard to compare for performance regression testing. My first idea was to just include the grid size in the benchmark output, but that didn't help much as it would still not show up in our current infra's benchmark comparisons. So I went for a bigger hammer and just use a hardcoded workload that is big enough that the tail effect should be negligible on current devices (2^28 items same as the biggest workloads for our device primitive benchmarks).
- Both the current approach and even more so my solution to 1. is measuring throughput. For block and warp primitives we are also interested in their latency in contexts where they do not make up the main workload of a kernel, but are still in the critical path. Therefore I added a second benchmark where only a single warp is launched.
- This one might be less important, but it was a particular concern for
WarpReduceBatched: I also added a benchmarking kernel with full launch bounds to force the compiler to deal with a lot of register pressure without pushing occupancy below a certain value. If a code change or toolchain change should affect the register allocation, this benchmark should show it very clearly since it is much more prone to spilling to local memory.
There was a problem hiding this comment.
@pauleonix thanks for the comments
so from what i understand for this PR
- we also want to replace occupancy based grids with 2^28 workloads.
- Add a single warp latency benchmark alongside the throughput one?
- Add a
WarpReduceBatched benchmarkwith full launch bounds?
is that right? anything else you'd like to see?
There was a problem hiding this comment.
- and 2. we want. For 3. I'm not yet sure if it justifies the additional compilation and benchmarking time.
|
@fbusato @pauleonix, do we want to add anything else? |
|
/ok to test 88b1da0 |
🥳 CI Workflow Results🟩 Finished in 2h 29m: Pass: 100%/287 | Total: 12d 17h | Max: 2h 18m | Hits: 22%/974831See results here. |
|
@fbusato ping |
|
@charan-003 I don't have any additional comment. There are two things to do before merging the PRs:
|
@fbusato |
|
these results don't show a clear benefits. I would add the results for all 3 SMs for DeviceReduce and BlockReduce benchmarks. |
@fbusato |
|
the results show minimal differences. Could you also please add the results for device-level benchmarks? |
added device-level benchmark |
|
sorry @charan-003, maybe I didn't explain it well. |
4 participants
Closes #5167
Replace sequential aggregation with cooperative warp reduction
Benchmark Comparison:
optimize_blockreducevsorigin/mainrtxa6000 — cub.bench.reduce.min.base
[0] NVIDIA RTX A6000
Summary: Total Matches: 56 — Pass: 28 / Failure: 28
rtxa6000 — cub.bench.reduce.sum.base
[0] NVIDIA RTX A6000
Summary: Total Matches: 64 — Pass: 43 / Failure: 21
h100 — cub.bench.reduce.sum.base
[0] NVIDIA H100
Summary: Total Matches: 64 — Pass: 44 / Failure: 20
Details
h100 — cub.bench.reduce.min.base
[0] NVIDIA H100 PCIe
Summary: Total Matches: 56 — Pass: 49 / Failure: 7
rtxpro6000 — cub.bench.reduce.sum.base
[0] NVIDIA RTX Pro 6000
Summary: Total Matches: 64 — Pass: 34 / Unknown: 8 / Failure: 22
rtxpro6000 — cub.bench.reduce.min.base
[0] NVIDIA RTX PRO 6000 Blackwell Server Edition
Summary: Total Matches: 56 — Pass: 39 / Unknown: 8 / Failure: 9