Skip to content

[executorch][cuda] Shorter autotune benchmark window and one compile pool for every split candidate - #23949

Merged
Gasoonjia merged 1 commit into
gh/Gasoonjia/52/origfrom
gh/Gasoonjia/53/orig
Oct 18, 2026
Merged

Gasoonjia merged 1 commit into
gh/Gasoonjia/52/origfrom
gh/Gasoonjia/53/orig

Conversation

@pytorchbot

Copy link
Copy Markdown
Collaborator

This PR was created by the merge bot to help merge the original PR into the main branch.
ghstack PR number: #23737 by @Gasoonjia
^ Please use this as the source of truth for the PR details, comments, and reviews
ghstack PR base: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/53/base
ghstack PR head: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/53/head
Merge bot PR base: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/52/orig
Merge bot PR head: https://github.com/pytorch/executorch/tree/gh/Gasoonjia/53/orig
Differential Revision: D124496888
@diff-train-skip-merge

…pool for every split candidate

Pull Request resolved: #23737

Builds on D124355785 (configs compiled in parallel before Triton benchmarks them). Two remaining costs of the compile-time quantized GEMM autotune:

1. **Benchmark window.** The family kernels' `triton.autotune` used Triton's default `do_bench(warmup=25, rep=100)`: about 125 ms per config whatever the kernel's length. These kernels run 10-100 us, so every config ran thousands of times; a bucket-64 shape with 288 (config, split) pairs spent ~36 s in it. They now use `quantized_gemm_utils.autotune_do_bench` (5 ms warmup, 20 ms timing), which still runs each config a few hundred times. Only the INT4/5/6/8 family kernels use it; SDPA, MoE and the rest keep Triton's default.

2. **One compile pool per shape.** `@autotune_launch_param` tries each split-K candidate in turn, and D124355785 opened one compile pool per candidate, one after another (e.g. 72, 72, 72, 48 and 24 configs at bucket 64). `@autotune_launch_param` now takes an optional `candidate_prepass`: a scope around one untimed call of every legal candidate before any is timed. The quantized GEMM launches pass `quantized_gemm_family._compile_every_candidate_at_once`, in which `launch_split_k_gemm` only queues its kernel's untuned configs on one shared pool and skips the launch. Leaving the scope waits for all compiles; timing then runs exactly as before. Each key compiled ahead is consumed once by its timed launch, which then does not open a pool of its own.

Neither changes the search: the same configs and split candidates are compiled and timed, and the picks are the fastest by the same rules.

**Does the shorter window still pick the best config?** For the models' linear shapes (INT4: 8448x6656, 6656x4096, 6656x19968, 256x6656, 1024x6656; INT6: 4096x5376, 5376x21504, 6656x19968, 1024x6656; INT5 and INT8: two each) at buckets 1/4/16/64 (INT5, INT8: 1/4) and every legal split, every pruned config was timed with the current window and with the new one; the union of each window's five fastest was then re-timed with a long window (100 ms warmup, 300 ms timing) as the reference.
- 232 (shape, bucket, split) cases: the new window's pick is within 1% of the reference best in 199 cases (current window: 202), within 2% in 214 (213); mean regret 1.0054 (1.0050).
- What ships is one split per shape: over the 44 (shape, bucket) pairs, with the split chosen as the fastest, the new window is more than 2% off the best in 2 pairs and the current window in 3. Worst: new 8.8% on INT4 256 x 6656 at M = 1 (an 8 us kernel, 0.7 us), current 6.0% on INT4 1024 x 6656 at M = 16. Both windows miss by a few percent on the shortest kernels, at the same rate.

**Compile time.** One INT4 W4A16 bucket-64 shape (5376 x 4096) under `autotune_launch_params()` + `cuda_graph_timing()`, A100, empty Triton cache: 334.6 s before D124355785, 77.8 s with it, 58.0 s with this diff (26.7 s with a warm cache, was 44.8 s); split pick 4 in all, the same 292 kernels compiled, per-split times within 1%.

Full Muse Glimmer 30B batching export (forward_{1,2,4,8,16,32,64} + forward_others + samplers; A100; empty Triton and Inductor caches), per method in minutes from the AOTI wrapper times:

| method | before D124355785 | D124355785 | this diff |
|---|---|---|---|
| load + trace + forward_1 | 46.5 | 51.3 | 40.8 |
| forward_2 | 21.4 | 15.3 | 14.5 |
| forward_4 | 23.4 | 15.5 | 14.4 |
| forward_8 | 22.3 | 12.2 | 10.3 |
| forward_16 | 39.3 | 15.5 | 12.9 |
| forward_32 | 52.8 | 18.8 | 14.7 |
| forward_64 | 73.3 | 24.2 | 18.7 |
| forward_others | 10.5 | 10.3 | 11.1 |
| total | 4h55m | 2h50m | 2h25m |

The export generates the same tokens as before (greedy; 512 and 4 x 2k prompts) at the same speed as with D124355785 (single_512 63.0 / 62.7 tok/s vs 62.5 / 62.9; 4 x 2k 111.3 vs 111.0).

What is left per autotuned shape is mostly compiling: during forward_64 only ~5 of the 32 compile threads are busy on average, and 87% of their time is in Triton's MLIR/LLVM passes (make_llir, make_ttgir, make_ptx), which appear to serialize on the GIL.
ghstack-source-id: 451021038
@exported-using-ghexport

Differential Revision: [D124496888](https://our.internmc.facebook.com/intern/diff/D124496888/)
@pytorch-bot

pytorch-bot Bot commented Oct 18, 2026

Copy link
Copy Markdown

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/executorch/23949

Note: Links to docs will display an error until the docs builds have been completed.

This comment was automatically generated by Dr. CI and updates every 15 minutes.

@Gasoonjia
Gasoonjia added this pull request to stack #23954 October 18, 2026 07:47
@meta-cla meta-cla Bot added the CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. label Oct 18, 2026
@github-actions

Copy link
Copy Markdown

This PR needs a release notes: label

If your change should be included in the release notes (i.e. would users of this library care about this change?), please use a label starting with release notes:. This helps us keep track and include your important work in the next release notes.

To add a label, you can comment to pytorchbot, for example
@pytorchbot label "release notes: none"

For more information, see
https://github.com/pytorch/pytorch/wiki/PyTorch-AutoLabel-Bot#why-categorize-for-release-notes-and-how-does-it-work.

@Gasoonjia
Gasoonjia merged commit cf8cfce into main Oct 18, 2026
265 of 273 checks passed
@Gasoonjia
Gasoonjia deleted the gh/Gasoonjia/53/orig branch October 18, 2026 07:48
@pytorchbot
pytorchbot deployed to upload-benchmark-results October 18, 2026 09:30 — with GitHub Actions Active

This branch was successfully deployed

2 active deployments
upload-benchmark-results — 58376aab Deployed Oct 18, 2026 by pytorchbot via upload-benchmark-results #20984
cadence — 58376aab Deployed Oct 18, 2026 by pytorchbot via hifi-op-test / hifi4 #33965
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants