ggml-cuda: make mul_mat_q tile selection tunable per-type and per-arch by aviallon · Pull Request #21849 · ggml-org/llama.cpp

aviallon · 2026-04-13T10:47:10Z

Overview

Modifies mmq.cuh tile selection logic to somewhat imitate what fattn-mma-f16.cuh does.
This allows tuning tile selection logic per-architecture, which is then used in this PR to avoid some tile configurations on CDNA2 (gfx90a) that make poor use of the hardware.

This is my first time touching actual computation code in llama.cpp, so it may not be of the highest quality, and I apologize for that.
I really tried to carefully benchmark my changes though.

This work was sponsored by JustAI

Additional information

Benchmarks:

Comparison: Tuned vs Baseline — Qwen3.5-122B-A10B IQ3_XXS pp throughput (ubatch=2048)
Tuned: IQ3_XXS tile=128 disabled (forces tile=64 for n>=128)
Baseline: No MMQ_CONFIG_CASE (pre-config-table, commit 1e9d771)
Device: AMD Instinct MI210 gfx90a, ROCm 7.2.1, locked at 800MHz in Performance Determinism mode.
Flags: -ctk q8_0 -ctv q8_0 -fa 1 -dio 1 -ngl 99 -r 3 -ub 2048
Date: 2026-04-13

Before (baseline)

model	size	params	backend	ngl	n_ubatch	type_k	type_v	fa	dio	test	t/s
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp128	119.08 ± 3.12
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp256	197.96 ± 1.53
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp512	295.08 ± 0.94
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp1024	398.65 ± 0.34
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp2048	491.17 ± 1.83
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp4096	479.88 ± 0.72
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp8192	462.77 ± 0.72

After (tuned: IQ3_XXS tile=128 disabled)

model	size	params	backend	ngl	n_ubatch	type_k	type_v	fa	dio	test	t/s
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp128	132.48 ± 3.05
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp256	212.58 ± 1.64
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp512	312.88 ± 0.92
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp1024	415.14 ± 0.16
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp2048	502.88 ± 1.76
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp4096	491.14 ± 0.42
qwen35moe 122B.A10B IQ3_XXS - 3.0625 bpw	41.66 GiB	122.11 B	ROCm	99	2048	q8_0	q8_0	1	1	pp8192	473.72 ± 0.84

Summary

test	before (t/s)	after (t/s)	delta
pp128	119.08	132.48	+11.3%
pp256	197.96	212.58	+7.4%
pp512	295.08	312.88	+6.0%
pp1024	398.65	415.14	+4.1%
pp2048	491.17	502.88	+2.4%
pp4096	479.88	491.14	+2.3%
pp8192	462.77	473.72	+2.4%

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES, for tests (see the test-backend-ops commit) + benchmarks (because repetitive rsync + ninja + llama-bench was too boring for me).

Add mmq_config struct with per-(type, mmq_x, arch) min_blocks and enabled fields, following the fattn_mma_config pattern. Replace the __launch_bounds__ #ifdef maze with a single constexpr config lookup. This will allow easier per-arch tuning later-on. Separate config functions per arch family: - cdna4: defaults (min_blocks=2), placeholder for future tuning - cdna3: defaults (min_blocks=2), placeholder for future tuning - cdna2: gfx90a (min_blocks=2), placeholder for tuning - cdna1: gfx908 (min_blocks=2), placeholder for tuning - nvidia_volta_up: Volta and newer (min_blocks=1) - rdna1: gfx1010/gfx1012 (min_blocks=1, preserves old no-arg behavior since we changed the default) - default: everything else (min_blocks=2) Add if constexpr device-side early-exit for disabled tiles (via the enabled field) and host-side filter in the tile selection loop. Behavior should be exactly the same as before, unless I made a mistake.

Use the mmq_config enabled/min_blocks fields to tune IQ3_XXS MMQ tile selection on CDNA2 (gfx90a). Tested on MI210. IQ3_XXS overrides: - mmq_x=112 disabled (260 VGPRs, 1 wave/SIMD at 10% occupancy; falling through to mmq_x=128 gives the same tile count with 2 waves) - mmq_x=128 relaxed to min_blocks=1 (eliminates 412B scratch spill)

Assisted-By: github-copilot/claude-sonnet-4.6

aviallon · 2026-04-13T10:47:53Z

Pinging @IMbackK and @JohannesGaessler for review. Please don't yell at me too loud, I'm scared 🫠

JohannesGaessler · 2026-04-13T11:13:24Z

If something like this is done at all it should be a table that can be easily looked up and modified like for fattn-mma-f16.cuh.

aviallon · 2026-04-13T11:16:13Z

Ok. So constexpr functions are not ok here? In fattn-mma-f16.cuh, I saw that there were some sanity checks in the macro too. Perhaps I could replicate that entirely?

JohannesGaessler · 2026-04-13T11:17:01Z

To clarify what I mean: for the mma FA kernel all relevant kernel parameters are packaged inside of a single table and explicitly listed for all GPUs and kernel configurations. If something like this is added it should be done in that format and that format should then encompass all relevant paramters. I don't want to make the way MMQ configuration is done more convoluted for a few % on some niche GPUs.

IMbackK

In general A centralized place for mmq tuneing parameters would be great.

IMbackK · 2026-04-13T10:55:47Z

+    return mmq_config{2, true};
+}
+
+// CDNA2 (gfx90a, MI210): wavefront64, 512 VGPRs/SIMD.


Kind of misleading, cdan2+ dosent have 512 VGPRS despite what amd has in the datasheet.
All GCN/CDNA gpus have only 256 vector registers.
starting cdan2 there are two banks per smid, if the second bank is not used for AGPRs it can be used to increase the occupancy of the kernel, but no kernel can ever have more than 256 VGPRs allocated.

Ah, indeed. You are better informed than me. Finding good doc on the matter was so hard.

What should I write on those headers instead?

IMbackK · 2026-04-13T10:58:10Z

+}
+
+// CDNA1 (gfx908, MI100): wavefront64, (256 VGPRs + 256 AccGPRs)/SIMD.
+// Default min_blocks=2 targets <=256 VGPRs/wave (2 waves/SIMD).


Misleading, there are only 256 VGPRs a kernel with 256 registers allocated has occupancy 1 unless the kernel happens to be allocating exactly half as AGPRS

aviallon · 2026-04-13T12:22:17Z

In general A centralized place for mmq tuneing parameters would be great.

Centralized with both fattn-mma-f16 and mmq? Or you are you just saying that the general idea I tried to implement makes sense?

JohannesGaessler · 2026-04-13T12:58:07Z

Historically I wrote MMQ first and after that FA. For MMQ I implemented functions per kernel parameter, for FA I eventually ended up changing the configuration to the current table-like setup. That table-like setup should also be used for MMQ and only after that would I consider approving a PR that extends the configuration for MMQ.

JohannesGaessler · 2026-04-13T13:01:25Z

To make it clear: the configuration for FA and MMQ should still be separate, but the pattern by which it is done should be the same.

aviallon · 2026-04-14T00:02:47Z

What kind of tunables would you want to see? Is the current list of tunables enough, or do you expect something much more comprehensive?

JohannesGaessler · 2026-04-14T08:58:43Z

If you want to work on this, please first make a PR that does nothing but change how the kernels are orchestrated using the existing tuning parameters and doesn't yet add any new ones. Currently there are a bunch of functions like get_mmq_x_max_host, their logic should be consolidated as was done for FA.

iacopPBK · 2026-04-16T08:37:21Z

@aviallon hello, we can put our efforts together. I propose to start with a minimal

struct mmq_config {
int min_blocks;
bool enabled; //yours
int threads_per_row; //mine
};

Then making some get_config as in your pr. We can parametrize these three things without changing any wrt upstream. I can update the vec dot q8 to use it then.

If this struct gets approved will be much easier to implement our speedups, to expand to other params if necessary.

Read again last response from maintaibers. We may need to start even simpler without new params.

@JohannesGaessler

Consolidates per-arch MMQ tuning into a single table, per @JohannesGaessler request in ggml-org#21849 to mirror the FA orchestration pattern. No behavior change, verified on gfx1151 with test-backend-ops and llama-bench parity. Follows up PR ggml-org#21849, to allow tuning MMQ params per arch and perf PR with the gfx1151 tuning from ggml-org#21344 will sit on top of this orchestration, hopefully Co-authored-by: Antoine Viallon <antoine@lesviallon.fr> Co-authored-by: Sunil Pedapudi <424940+pedapudi@users.noreply.github.com>

@JohannesGaessler

Consolidates per-arch MMQ tuning into a single table, per @JohannesGaessler request in ggml-org#21849 to mirror the FA orchestration pattern. No behavior change, verified on gfx1151 with test-backend-ops and llama-bench parity. Follows up PR ggml-org#21849, to allow tuning MMQ params per arch and perf PR with the gfx1151 tuning from ggml-org#21344 will sit on top of this orchestration, hopefully Co-authored-by: Antoine Viallon <antoine@lesviallon.fr> Co-authored-by: Sunil Pedapudi <424940+pedapudi@users.noreply.github.com>

aviallon added 3 commits April 13, 2026 12:38

test : add Qwen3.5-122B MUL_MAT_ID perf tests (IQ2_S, IQ3_XXS, Q6_K)

d58c297

Assisted-By: github-copilot/claude-sonnet-4.6

aviallon requested review from a team and ggerganov as code owners April 13, 2026 10:47

aviallon changed the title ~~Upstream/mmq config table~~ [HIP/CUDA] make mul_mat_q tile selection tunable per-type and per-arch Apr 13, 2026

aviallon changed the title ~~[HIP/CUDA] make mul_mat_q tile selection tunable per-type and per-arch~~ ggml-cuda: make mul_mat_q tile selection tunable per-type and per-arch Apr 13, 2026

IMbackK reviewed Apr 13, 2026

View reviewed changes

github-actions Bot added testing Everything test related Nvidia GPU Issues specific to Nvidia GPUs ggml changes relating to the ggml tensor library for machine learning labels Apr 13, 2026

JohannesGaessler mentioned this pull request Apr 14, 2026

ggml-cuda: better vram to lds loading pipeline in load_tiles_q8_0 #21698

Open

chrismcmacken mentioned this pull request May 14, 2026

ggml-cuda: per-arch MMQ config variable refactor into tabular structure #23031

Closed

Uh oh!

Conversation

aviallon commented Apr 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Overview

Additional information

Before (baseline)

After (tuned: IQ3_XXS tile=128 disabled)

Summary

Requirements

Uh oh!

aviallon commented Apr 13, 2026

Uh oh!

JohannesGaessler commented Apr 13, 2026

Uh oh!

aviallon commented Apr 13, 2026

Uh oh!

JohannesGaessler commented Apr 13, 2026

Uh oh!

IMbackK left a comment

Choose a reason for hiding this comment

Uh oh!

IMbackK Apr 13, 2026

Choose a reason for hiding this comment

Uh oh!

aviallon Apr 13, 2026

Choose a reason for hiding this comment

Uh oh!

aviallon Apr 13, 2026

Choose a reason for hiding this comment

Uh oh!

IMbackK Apr 13, 2026

Choose a reason for hiding this comment

Uh oh!

aviallon commented Apr 13, 2026

Uh oh!

JohannesGaessler commented Apr 13, 2026

Uh oh!

JohannesGaessler commented Apr 13, 2026

Uh oh!

aviallon commented Apr 14, 2026

Uh oh!

JohannesGaessler commented Apr 14, 2026

Uh oh!

iacopPBK commented Apr 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

aviallon commented Apr 13, 2026 •

edited

Loading

iacopPBK commented Apr 16, 2026 •

edited

Loading