Add tiling for reductions by pfultz2 · Pull Request #4954 · ROCm/AMDMIGraphX

pfultz2 · 2026-06-10T21:05:13Z

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.

gh-app-migraphx-bot-pr-write · 2026-06-11T01:16:41Z

Test	Batch	New Rate (567afa)	Old Rate (1a5bb6)*	Diff	Status
torchvision-resnet50	64	3,321.07	3,265.81	1.69%	✅
torchvision-resnet50_fp16	64	3,785.62	7,548.56	-49.85%	🔴
torchvision-densenet121	32	272.37	2,481.85	-89.03%	🔴
torchvision-densenet121_fp16	32	964.66	4,975.40	-80.61%	🔴
torchvision-inceptionv3	32	1,888.25	2,060.05	-8.34%	🔴
torchvision-inceptionv3_fp16	32	4,477.25	4,414.40	1.42%	✅
cadene-inceptionv4	16	665.96	584.00	14.03%	🔆
cadene-resnext64x4	16	503.95	782.84	-35.63%	🔴
slim-mobilenet	64	8,293.45	7,190.62	15.34%	🔆
slim-nasnetalarge	64	174.66	222.84	-21.62%	🔴
slim-resnet50v2	64	887.33	2,017.65	-56.02%	🔴
bert-mrpc-onnx	8	259.77	1,168.85	-77.78%	🔴
bert-mrpc-tf	1	55.44	487.56	-88.63%	🔴
pytorch-examples-wlang-gru	1	141.55	477.33	-70.35%	🔴
pytorch-examples-wlang-lstm	1	151.03	387.83	-61.06%	🔴
torchvision-resnet50_1	1	680.25	1,038.23	-34.48%	🔴
cadene-dpn92_1	1	475.84	449.51	5.86%	🔆
cadene-resnext101_1	1	258.06	365.36	-29.37%	🔴
onnx-taau-downsample	1	711.75	844.73	-15.74%	🔴
dlrm-criteoterabyte	1	0.68	17.44	-96.12%	🔴
dlrm-criteoterabyte_fp16	1	25.24	31.21	-19.13%	🔴
agentmodel	1	1,452.10	5,880.42	-75.31%	🔴
unet_fp16	2	32.80	52.15	-37.11%	🔴
resnet50v1_fp16	1	267.03	1,385.99	-80.73%	🔴
resnet50v1_int8	1	261.76	1,657.40	-84.21%	🔴
bert_base_cased_fp16	64	531.66	1,022.83	-48.02%	🔴
bert_large_uncased_fp16	32	316.92	339.30	-6.60%	🔴
bert_large_fp16	1	204.66	205.64	-0.48%	✅
distilgpt2_fp16	16	1,107.44	2,086.12	-46.91%	🔴
yolov5s	1	126.01	560.85	-77.53%	🔴
tinyllama	1	23.55	45.80	-48.59%	🔴
vicuna-fastchat	1	42.25	44.03	-4.04%	✅
whisper-tiny-encoder	1	412.03	328.75	25.33%	🔆
whisper-tiny-decoder	1	132.17	214.20	-38.30%	🔴
llama2_7b	1	20.82	14.70	41.66%	🔆
qwen1.5-7b	1	23.59	23.49	0.43%	✅
phi3-3.8b	1	26.77	26.78	-0.02%	✅
llama3-8b	1	21.74	21.70	0.20%	✅
whisper-large-encoder	1	7.39	8.77	-15.72%	🔴
whisper-large-decoder	1	56.26	109.89	-48.80%	🔴
mistral-7b	1	23.76	23.65	0.46%	✅
FLUX.1-schnell	1	760.77	172.94	339.91%	🔆

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 · 2026-06-11T01:16:42Z

Test	Status	Result
bert-mrpc-onnx	✅	PASSED: MIGraphX meets tolerance
bert-mrpc-tf	✅	PASSED: MIGraphX meets tolerance
pytorch-examples-wlang-gru	✅	PASSED: MIGraphX meets tolerance
pytorch-examples-wlang-lstm	✅	PASSED: MIGraphX meets tolerance
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	✅	PASSED: MIGraphX meets tolerance
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

codecov · 2026-07-17T18:33:03Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop    #4954      +/-   ##
===========================================
- Coverage    92.90%   92.90%   -0.00%     
===========================================
  Files          603      603              
  Lines        32526    32528       +2     
===========================================
+ Hits         30217    30218       +1     
- Misses        2309     2310       +1

see 2 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

pfultz2 · 2026-07-17T20:13:53Z

This improves perf for skinny gemms with 2 rows for some cases almost 2x faster:

Shape	N	Outputs	MB	Baseline (ms)	Time (ms)	GB/s	Change
2x2x2	2	4	0.0	0.0128	0.0125	0.0	+2.4%
2x8x8	8	16	0.0	0.0129	0.0128	0.0	+0.8%
2x256x8	8	512	0.01	0.0127	0.0127	0.6	+0.0%
2x1024x4	4	2048	0.02	0.0128	0.0126	1.3	+1.6%
2x4096x2	2	8192	0.03	0.0126	0.0126	2.6	+0.0%
2x655360x2	2	1310720	5.24	0.0202	0.0200	262.0	+1.0%
2x262144x4	4	524288	4.19	0.0138	0.0126	333.5	+9.5%
2x131072x8	8	262144	4.19	0.0127	0.0128	327.3	-0.8%
2x65536x16	16	131072	4.19	0.0128	0.0127	330.6	+0.8%
2x32768x32	32	65536	4.19	0.0127	0.0125	336.2	+1.6%
2x16384x64	64	32768	4.19	0.0129	0.0127	331.1	+1.6%
2x8192x128	128	16384	4.19	0.0128	0.0126	332.0	+1.6%
2x256x256	256	512	0.26	0.0148	0.0148	17.7	+0.0%
2x512x512	512	1024	1.05	0.0132	0.0141	74.3	-6.4%
2x1024x1024	1024	2048	4.19	0.0126	0.0124	339.6	+1.6%
2x2048x2048	2048	4096	16.78	0.0129	0.0122	1378.6	+5.7%
2x4096x768	768	8192	12.58	0.0128	0.0126	999.3	+1.6%
2x8192x768	768	16384	25.17	0.0168	0.0154	1634.5	+9.1%
2x8192x1024	1024	16384	33.55	0.0212	0.0182	1845.7	+16.5%
2x16384x256	256	32768	16.78	0.0129	0.0122	1380.1	+5.7%
2x32768x128	128	65536	16.78	0.0179	0.0178	942.6	+0.6%
2x2048x4096	4096	4096	33.55	0.0192	0.0192	1751.6	+0.0%
2x4096x4096	4096	8192	67.11	0.0354	0.0360	1864.7	-1.7%
2x1024x8192	8192	2048	33.55	0.0193	0.0159	2105.0	+21.4%
2x8192x512	512	16384	16.78	0.0126	0.0125	1337.8	+0.8%
2x16384x1024	1024	32768	67.11	0.0375	0.0309	2173.4	+21.4%
2x512x2048	2048	1024	4.19	0.0129	0.0126	332.2	+2.4%
2x256x4096	4096	512	4.19	0.0124	0.0122	343.1	+1.6%
2x128x8192	8192	256	4.19	0.0122	0.0121	346.1	+0.8%
2x64x16384	16384	128	4.19	0.0124	0.0128	328.4	-3.1%
2x2048x32000	32000	4096	262.14	0.4131	0.2132	1229.6	+93.8%
2x1024x50272	50272	2048	205.91	0.3258	0.1719	1197.7	+89.5%
2x8192x8192	8192	16384	268.44	0.4231	0.2156	1244.9	+96.2%
2x4096x16384	16384	8192	268.44	0.4228	0.2159	1243.6	+95.8%

It also improves the perf for strided reductions(where the reduction is not along the fastest axis). Here are the numbers for layernorm:

Shape	N	Outputs	MB	Baseline (ms)	Time (ms)	GB/s	Change
2x2	2	2	0.0	0.0135	0.0139	0.0	-2.9%
8x8	8	8	0.0	0.0129	0.0120	0.0	+7.5%
256x8	8	256	0.01	0.0138	0.0143	0.6	-3.5%
1024x4	4	1024	0.02	0.0137	0.0130	1.3	+5.4%
4096x2	2	4096	0.03	0.0137	0.0131	2.5	+4.6%
655360x2	2	655360	5.24	0.0340	0.0254	206.0	+33.9%
262144x4	4	262144	4.19	0.0291	0.0239	175.7	+21.8%
131072x8	8	131072	4.19	0.0326	0.0244	172.2	+33.6%
65536x16	16	65536	4.19	0.0486	0.0170	246.1	+185.9%
32768x32	32	32768	4.19	0.0442	0.0160	262.4	+176.3%
16384x64	64	16384	4.19	0.0405	0.0238	176.1	+70.2%
8192x128	128	8192	4.19	0.0487	0.0206	203.9	+136.4%
256x256	256	256	0.26	0.0183	0.0132	19.9	+38.6%
512x512	512	512	1.05	0.0236	0.0238	44.1	-0.8%
1024x1024	1024	1024	4.19	0.0477	0.0280	149.9	+70.4%
2048x2048	2048	2048	16.78	0.2809	0.1554	108.0	+80.8%
4096x768	768	4096	12.58	0.2325	0.0517	243.5	+349.7%
8192x768	768	8192	25.17	0.6693	0.0948	265.5	+606.0%
8192x1024	1024	8192	33.55	0.9450	0.1297	258.7	+628.6%
16384x256	256	16384	16.78	0.3116	0.0560	299.9	+456.4%
32768x128	128	32768	16.78	0.2821	0.0552	303.9	+411.1%
2048x4096	4096	2048	33.55	0.6807	0.2062	162.7	+230.1%
4096x4096	4096	4096	67.11	1.7123	0.7364	91.1	+132.5%
1024x8192	8192	1024	33.55	0.4759	0.3484	96.3	+36.6%
8192x512	512	8192	16.78	0.3075	0.0652	257.5	+371.6%
16384x1024	1024	16384	67.11	2.0872	0.3787	177.2	+451.1%
512x2048	2048	512	4.19	0.0421	0.0443	94.7	-5.0%
256x4096	4096	256	4.19	0.0319	0.0325	129.1	-1.8%
128x8192	8192	128	4.19	0.0325	0.0350	119.9	-7.1%
64x16384	16384	64	4.19	0.0292	0.0310	135.2	-5.8%
32x65536	65536	32	8.39	0.0596	0.0518	162.0	+15.1%
16x131072	131072	16	8.39	0.0551	0.0508	165.3	+8.5%
8x262144	262144	8	8.39	0.0582	0.0525	159.7	+10.9%
4x655360	655360	4	10.49	0.0595	0.0423	247.8	+40.7%
2x655360	655360	2	5.24	0.0508	0.0318	164.9	+59.7%
2048x32000	32000	2048	262.14	13.3292	2.5995	100.8	+412.8%
1024x50272	50272	1024	205.91	2.8011	1.4885	138.3	+88.2%
8192x8192	8192	8192	268.44	48.9343	4.2355	63.4	+1055.3%
4096x16384	16384	4096	268.44	30.9192	9.1689	29.3	+237.2%

And also helps with strided skinny gemms(although may not be likely for now):

Shape	N	Outputs	MB	Baseline (ms)	Time (ms)	GB/s	Change
2x2x2	2	4	0.0	0.0125	0.0126	0.0	-0.8%
2x8x8	8	16	0.0	0.0127	0.0121	0.0	+5.0%
2x256x8	8	512	0.01	0.0128	0.0126	0.7	+1.6%
2x1024x4	4	2048	0.02	0.0126	0.0125	1.3	+0.8%
2x4096x2	2	8192	0.03	0.0128	0.0131	2.5	-2.3%
2x655360x2	2	1310720	5.24	0.0363	0.0128	408.5	+183.6%
2x262144x4	4	524288	4.19	0.0292	0.0126	333.6	+131.7%
2x131072x8	8	262144	4.19	0.0373	0.0127	330.9	+193.7%
2x65536x16	16	131072	4.19	0.0801	0.0126	331.6	+535.7%
2x32768x32	32	65536	4.19	0.0707	0.0147	285.8	+381.0%
2x16384x64	64	32768	4.19	0.0472	0.0126	333.3	+274.6%
2x8192x128	128	16384	4.19	0.0599	0.0122	343.1	+391.0%
2x1024x1024	1024	2048	4.19	0.0423	0.0207	203.1	+104.3%
2x2048x2048	2048	4096	16.78	0.4782	0.0216	777.6	+2113.9%
2x4096x768	768	8192	12.58	0.2929	0.0204	616.4	+1335.8%
2x8192x768	768	16384	25.17	0.8322	0.0215	1169.6	+3770.7%
2x8192x1024	1024	16384	33.55	1.0792	0.0193	1737.0	+5491.7%
2x16384x256	256	32768	16.78	0.6379	0.0271	619.7	+2253.9%
2x32768x128	128	65536	16.78	0.3403	0.0215	781.0	+1482.8%
2x2048x4096	4096	4096	33.55	1.0203	0.0243	1382.2	+4098.8%
2x4096x4096	4096	8192	67.11	2.4739	0.0459	1461.0	+5289.8%
2x1024x8192	8192	2048	33.55	0.6790	0.0512	655.7	+1226.2%
2x8192x512	512	16384	16.78	0.5698	0.0210	800.2	+2613.3%
2x16384x1024	1024	32768	67.11	2.1766	0.0216	3111.5	+9976.9%
2x512x2048	2048	1024	4.19	0.0511	0.0198	212.1	+158.1%
2x256x4096	4096	512	4.19	0.0356	0.0208	201.3	+71.2%
2x128x8192	8192	256	4.19	0.0443	0.0438	95.9	+1.1%
2x64x16384	16384	128	4.19	0.0318	0.0344	122.1	-7.6%
2x2048x32000	32000	4096	262.14	26.0640	0.4706	557.0	+5438.5%
2x1024x50272	50272	2048	205.91	12.8998	0.3710	555.1	+3377.0%
2x8192x8192	8192	16384	268.44	97.0449	0.2136	1256.9	+45333.0%
2x4096x16384	16384	8192	268.44	53.4646	0.2944	911.8	+18060.5%

TedThemistokleous

Minor comments - Overall this is all green and makes sense. Just curious on a few things reguarding the algorithm and sizes used

pfultz2 added 3 commits June 10, 2026 12:59

Add block_tile algo

7092ca3

Add block_stride

ebfdccc

Improve picking block_tile

2f2a9ac

pfultz2 added 3 commits July 15, 2026 18:29

Merge

a27de37

Format

7a56824

Tune block stride

797c467

pfultz2 added 2 commits July 17, 2026 14:55

Fix bug in address calculation

3e3617d

Format

bc0773d

pfultz2 marked this pull request as ready for review July 17, 2026 20:13

pfultz2 requested a review from causten as a code owner July 17, 2026 20:13

pfultz2 requested review from CharlieL7, TedThemistokleous and kahmed10 July 17, 2026 20:14

Merge branch 'develop' into reduce-block-tile

567afaa

TedThemistokleous reviewed Jul 21, 2026

View reviewed changes

Comment thread test/verify/test_reduce_broadcast_tiled.cpp

TedThemistokleous reviewed Jul 21, 2026

View reviewed changes

Comment thread test/verify/test_reduce_strided.cpp

Comment thread src/targets/gpu/jit/reduce.cpp Outdated

Add 512

ff1613f

pfultz2 requested a review from TedThemistokleous July 23, 2026 16:20

TedThemistokleous approved these changes Jul 23, 2026

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Add tiling for reductions#4954

Add tiling for reductions#4954
pfultz2 wants to merge 10 commits into
developfrom
reduce-block-tile

pfultz2 commented Jun 10, 2026

Uh oh!

gh-app-migraphx-bot-pr-write Bot commented Jun 11, 2026 •

edited

Loading

Uh oh!

gh-app-migraphx-bot-pr-write Bot commented Jun 11, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Jul 17, 2026 •

edited

Loading

Uh oh!

pfultz2 commented Jul 17, 2026

Uh oh!

Uh oh!

TedThemistokleous left a comment

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

pfultz2 commented Jun 10, 2026

Motivation

Technical Details

Changelog Category

Uh oh!

gh-app-migraphx-bot-pr-write Bot commented Jun 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

gh-app-migraphx-bot-pr-write Bot commented Jun 11, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov Bot commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

pfultz2 commented Jul 17, 2026

Uh oh!

Uh oh!

TedThemistokleous left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

gh-app-migraphx-bot-pr-write Bot commented Jun 11, 2026 •

edited

Loading

gh-app-migraphx-bot-pr-write Bot commented Jun 11, 2026 •

edited

Loading

codecov Bot commented Jul 17, 2026 •

edited

Loading