Skip to content

Add block batch reduce algorithm - #5177

Open
pfultz2 wants to merge 10 commits into
developfrom
block-batch-reduce
Open

Add block batch reduce algorithm#5177
pfultz2 wants to merge 10 commits into
developfrom
block-batch-reduce

Conversation

@pfultz2

@pfultz2 pfultz2 commented Aug 24, 2026

Copy link
Copy Markdown
Collaborator

Motivation

Technical Details

Changelog Category

Add a CHANGELOG.md entry for any option other than Not Applicable

    • Added: New functionality.
    • Changed: Changes to existing functionality.
    • Removed: Functionality or support that has been removed. (Compared to a previous release)
    • Optimized: Component performance that has been optimized or improved.
    • Resolved Issues: Known issues from a previous version that have been resolved.
    • Not Applicable: This PR is not to be included in the changelog.

Follow the LLVM AI Tool Use Policy for contributions using AI.

@gh-app-migraphx-bot-pr-write

Copy link
Copy Markdown
Test Batch New Rate (9f9e56) Old Rate (3a503c)* Diff Status
torchvision-resnet50 64 3,320.07 3,264.92 1.69%
torchvision-resnet50_fp16 64 7,873.94 7,548.67 4.31%
torchvision-densenet121 32 2,489.57 2,483.99 0.22%
torchvision-densenet121_fp16 32 5,017.88 5,004.24 0.27%
torchvision-inceptionv3 32 2,070.23 2,058.51 0.57%
torchvision-inceptionv3_fp16 32 4,447.23 4,416.99 0.68%
cadene-inceptionv4 16 817.04 820.61 -0.44%
cadene-resnext64x4 16 782.76 782.78 -0.00%
slim-mobilenet 64 8,374.57 8,386.36 -0.14%
slim-nasnetalarge 64 228.35 228.86 -0.22%
slim-resnet50v2 64 3,235.67 3,180.91 1.72%
bert-mrpc-onnx 8 1,171.69 1,168.84 0.24%
bert-mrpc-tf 1 510.14 498.63 2.31%
pytorch-examples-wlang-gru 1 367.65 473.35 -22.33% 🔴
pytorch-examples-wlang-lstm 1 541.35 384.83 40.67% 🔆
torchvision-resnet50_1 1 1,058.55 1,046.63 1.14%
cadene-dpn92_1 1 452.09 437.32 3.38%
cadene-resnext101_1 1 365.95 365.89 0.02%
onnx-taau-downsample 1 842.22 844.09 -0.22%
dlrm-criteoterabyte 1 32.25 32.42 -0.53%
dlrm-criteoterabyte_fp16 1 51.63 51.80 -0.33%
agentmodel 1 14,910.13 9,209.12 61.91% 🔆
unet_fp16 2 58.35 58.80 -0.76%
resnet50v1_fp16 1 1,461.17 1,366.11 6.96% 🔆
resnet50v1_int8 1 1,801.75 1,883.96 -4.36%
bert_base_cased_fp16 64 1,097.94 1,098.16 -0.02%
bert_large_uncased_fp16 32 345.76 345.59 0.05%
bert_large_fp16 1 207.11 206.59 0.25%
distilgpt2_fp16 16 2,093.52 2,092.89 0.03%
yolov5s 1 561.85 558.33 0.63%
tinyllama 1 45.77 45.83 -0.13%
vicuna-fastchat 1 44.27 44.20 0.15%
whisper-tiny-encoder 1 412.44 411.87 0.14%
whisper-tiny-decoder 1 411.68 408.48 0.78%
llama2_7b 1 20.86 20.84 0.10%
qwen1.5-7b 1 23.66 23.58 0.35%
phi3-3.8b 1 28.36 26.72 6.16% 🔆
llama3-8b 1 22.68 21.80 4.01%
whisper-large-encoder 1 10.16 10.18 -0.14%
whisper-large-decoder 1 107.98 105.30 2.55%
mistral-7b 1 23.80 23.78 0.09%
FLUX.1-schnell 1 786.44 755.22 4.13%

Regressions detected 🔴

* No develop baseline was found for this PR's branch point; compared against the latest available develop run instead.

@gh-app-migraphx-bot-pr-write

Copy link
Copy Markdown
Test Status Result
bert-mrpc-onnx PASSED: MIGraphX meets tolerance
bert-mrpc-tf ERROR - check error output
traceback
Traceback (most recent call last):
File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 377, in
main()
File "/src/AMDMIGraphX/tools/accuracy/accuracy_checker.py", line 313, in main
import tensorflow as tf
File "/usr/local/lib/python3.12/dist-packages/tensorflow/init.py", line 40, in
from tensorflow.python import pywrap_tensorflow as _pywrap_tensorflow # pylint: disable=unused-import
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/usr/local/lib/python3.12/dist-packages/tensorflow/python/pywrap_tensorflow.py", line 37, in
self_check.preload_check()
File "/usr/local/lib/python3.12/dist-packages/tensorflow/python/platform/self_check.py", line 63, in preload_check
from tensorflow.python.platform import _pywrap_cpu_feature_guard
ImportError: libnuma.so.1: cannot open shared object file: No such file or directory
pytorch-examples-wlang-gru 🔴 FAILED: MIGraphX is not within tolerance - check verbose output
pytorch-examples-wlang-lstm 🔴 FAILED: MIGraphX is not within tolerance - check verbose output
dlrm-criteoterabyte PASSED: MIGraphX meets tolerance
agentmodel PASSED: MIGraphX meets tolerance
unet PASSED: MIGraphX meets tolerance
resnet50v1 PASSED: MIGraphX meets tolerance
bert_base_cased_fp16 PASSED: MIGraphX meets tolerance
bert_large_uncased_fp16 🔴 FAILED: MIGraphX is not within tolerance - check verbose output
bert_large PASSED: MIGraphX meets tolerance
yolov5s PASSED: MIGraphX meets tolerance
tinyllama PASSED: MIGraphX meets tolerance
vicuna-fastchat PASSED: MIGraphX meets tolerance
whisper-tiny-encoder PASSED: MIGraphX meets tolerance
whisper-tiny-decoder PASSED: MIGraphX meets tolerance
distilgpt2_fp16 🔴 FAILED: MIGraphX is not within tolerance - check verbose output
llama2_7b PASSED: MIGraphX meets tolerance
qwen1.5-7b PASSED: MIGraphX meets tolerance
phi3-3.8b PASSED: MIGraphX meets tolerance
llama3-8b PASSED: MIGraphX meets tolerance
whisper-large-encoder PASSED: MIGraphX meets tolerance
whisper-large-decoder PASSED: MIGraphX meets tolerance
mistral-7b PASSED: MIGraphX meets tolerance
FLUX.1-schnell PASSED: MIGraphX meets tolerance

Copilot AI 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.

Pull request overview

Adds a batched GPU block-reduction algorithm for tiled fused reductions.

Changes:

  • Adds batched accumulation, slicing, and result assignment.
  • Adds JIT eligibility and tuning logic.
  • Adds a cppcheck-safe math wrapper definition.

Review used a single pass without subagent fan-out. GPU verification on wave32 and wave64 remains necessary.

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated 3 comments.

File Description
reduce.hpp Implements the batched kernel algorithm.
math.hpp Adds cppcheck macro handling.
reduce.cpp Selects and tunes batched reductions.

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment on lines +323 to +327
if(any_of(ins->inputs(), [&](instruction_ref input) {
if(input->name() != "@param")
return false;
return input->get_shape().lens() == rlens or input->get_shape().broadcasted();
}))
Comment on lines +593 to +597
static std::size_t
batch_iterations(std::size_t n_per_block, std::size_t relements, std::size_t block_size)
{
return (n_per_block * relements) / block_size;
}
Comment on lines +275 to +279
/// Like block_reduce but computes N independent reductions in one pass. The
/// iteration space covers N consecutive segments of relements elements and
/// each lane accumulates into the vector slot of the segment it is reading.
template <index_int N, class Op, class T, class Size, class F>
__device__ auto block_batch_reduce(index idx, Op op, T init, Size n, F f)
@pfultz2
pfultz2 marked this pull request as ready for review August 31, 2026 18:20
@pfultz2
pfultz2 requested a review from causten as a code owner August 31, 2026 18:20
@pfultz2

pfultz2 commented Sep 1, 2026

Copy link
Copy Markdown
Collaborator Author

Here is some perf numbers for batch 4 and 2 rows:

Shape N Outputs MB Baseline (ms) Time (ms) GB/s Change
4x2x2x2 2 16 0.0 0.0033 0.0033 0.0 +0.0%
4x2x8x8 8 64 0.0 0.0035 0.0034 0.3 +2.9%
4x2x256x8 8 2048 0.03 0.0033 0.0034 9.7 -2.9%
4x2x1024x4 4 8192 0.07 0.0040 0.0039 16.6 +2.6%
4x2x4096x2 2 32768 0.13 0.0156 0.0154 8.5 +1.3%
4x2x655360x2 2 5242880 20.97 0.0611 0.0608 344.8 +0.5%
4x2x262144x4 4 2097152 16.78 0.0278 0.0280 599.4 -0.7%
4x2x131072x8 8 1048576 16.78 0.0171 0.0170 986.8 +0.6%
4x2x65536x16 16 524288 16.78 0.0172 0.0169 993.4 +1.8%
4x2x32768x32 32 262144 16.78 0.0172 0.0170 986.8 +1.2%
4x2x16384x64 64 131072 16.78 0.0232 0.0186 902.8 +24.7%
4x2x8192x128 128 65536 16.78 0.0140 0.0118 1426.7 +18.6%
4x2x256x256 256 2048 1.05 0.0160 0.0046 229.0 +247.8%
4x2x512x512 512 4096 4.19 0.0056 0.0055 765.9 +1.8%
4x2x1024x1024 1024 8192 16.78 0.0111 0.0077 2169.4 +44.2%
4x2x2048x2048 2048 16384 67.11 0.0256 0.0195 3449.2 +31.3%
4x2x4096x768 768 32768 50.33 0.0187 0.0150 3345.2 +24.7%
4x2x8192x768 768 65536 100.66 0.0314 0.0281 3581.5 +11.7%
4x2x8192x1024 1024 65536 134.22 0.0423 0.0370 3622.7 +14.3%
4x2x16384x256 256 131072 67.11 0.0237 0.0219 3059.2 +8.2%
4x2x32768x128 128 262144 67.11 0.0416 0.0358 1872.9 +16.2%
4x2x2048x4096 4096 16384 134.22 0.0448 0.0371 3617.4 +20.8%
4x2x4096x4096 4096 32768 268.44 0.2148 0.2143 1252.9 +0.2%
4x2x1024x8192 8192 8192 134.22 0.0434 0.0373 3600.2 +16.4%
4x2x8192x512 512 65536 67.11 0.0221 0.0192 3487.6 +15.1%
4x2x16384x1024 1024 131072 268.44 0.2159 0.2150 1248.6 +0.4%
4x2x512x2048 2048 4096 16.78 0.0116 0.0077 2166.2 +50.6%
4x2x256x4096 4096 2048 16.78 0.0103 0.0077 2185.8 +33.8%
4x2x128x8192 8192 1024 16.78 0.0136 0.0077 2187.3 +76.6%
4x2x64x16384 16384 512 16.78 0.0106 0.0080 2105.8 +32.5%
4x2x32x65536 65536 256 33.55 0.0747 0.0745 450.4 +0.3%
4x2x16x131072 131072 128 33.55 0.0434 0.0434 774.0 +0.0%
4x2x8x262144 262144 64 33.55 0.0439 0.0439 764.5 +0.0%
4x2x4x655360 655360 32 41.94 0.0456 0.0473 885.9 -3.6%
4x2x2x655360 655360 16 20.97 0.0280 0.0280 749.3 +0.0%
4x2x2048x32000 32000 16384 1048.58 0.8352 0.8354 1255.2 -0.0%
4x2x1024x50272 50272 8192 823.66 0.6703 0.6704 1228.7 -0.0%
4x2x8192x8192 8192 65536 1073.74 0.8530 0.8504 1262.7 +0.3%
4x2x4096x16384 16384 32768 1073.74 0.8525 0.8504 1262.6 +0.2%

@pfultz2 pfultz2 added the high priority A PR with high priority for review and merging. label Sep 1, 2026

@TedThemistokleous TedThemistokleous left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just one comment but wont block review for this - more style question

auto noutputs = plan.finputs.size() - shapes.size() + 1;
auto tile = find_reduce_tile(
plan.virtual_inputs, noutputs, plan.reduce_output_shape, plan.reduction_shape.lens());
assert(not ins->module_inputs().empty());

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why not put this at the top of the function similar to compile()? No point in doing the work here if module inputs are empty

# Conflicts:
#	src/targets/gpu/jit/reduce.cpp

@bdevorem bdevorem left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

lgtm, but can you address the top two copilot comments? They look convincing

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

high priority A PR with high priority for review and merging.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants