Skip to content

[Config] configs: add tuned configs for Qwen3-VL-235B MXFP4 (gfx950) - #4554

Merged
yifehuan merged 4 commits into
ROCm:mainfrom
vorapolsiloai:tune-qwen3-vl-235b-mxfp4-gfx950
Aug 27, 2026
Merged

yifehuan merged 4 commits into
ROCm:mainfrom
vorapolsiloai:tune-qwen3-vl-235b-mxfp4-gfx950

Conversation

@vorapolsiloai

@vorapolsiloai vorapolsiloai commented Aug 4, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Add gfx950 tuning coverage for amd/Qwen3-VL-235B-A22B-Instruct-MXFP4 across TP1/TP2/TP4/TP8.

Added tuning data

  • q3vl_a4w4_blockscale_tuned_gemm.csv: 7,212 A4W4 blockscale GEMM rows.
  • q3vl_bf16_tuned_gemm.csv: 945 BF16 GEMM rows.
  • q3vl_fp4_tuned_fmoe.csv: 60 FP4 fused-MoE rows.
  • Total: 8,217 rows, all for gfx950 / 256 CUs.

Data validation

All three files were parsed after rebasing and checked against the runtime merge keys used by their corresponding tuners:

  • exact expected row counts — passed;
  • no null or malformed rows — passed;
  • gfx950 and 256-CU metadata on every row — passed;
  • no duplicate runtime keys within any added file — passed;
  • no runtime-key collisions with the rebased main model configs — passed;
  • two slower current-main BF16 rows that collided with faster rows in this PR were removed, matching AITER's global best-us merge policy;
  • git diff --check — passed.
  • The corrected wheel build passed config merging; the only observed upstream failure is the same unrelated MI35X test_moe_2stage.py GPU-memory exhaustion seen on other PRs, and a maintainer rerun has been requested.

End-to-end integration and performance

The Qwen3-VL integration image selected FlyDSL stage-2 _reduce rows from this PR. At TP8, 1024 input/output tokens, concurrency 4 on MI355X:

  • selected tuned _reduce: 375.19 output tok/s, 889.47 total tok/s;
  • the same tuning rows with #4556's atomic override: 391.46 output tok/s, 928.27 total tok/s;
  • reference range: 389.48–391.05 output tok/s, 924.04–928.77 total tok/s.

This confirms that the rows are active in the intended deployment and documents their measured relationship with #4556. It does not claim an isolated tuned-versus-untuned uplift; the PR is ready for review with the existing data and integration validation.

Qwen3-VL integration stack

Rebase provenance

  • Frozen AITER main: fc2e5d57fb5b8ad8e7e23f7103071dde798ea618
  • Initial tuning commit: ed4b626fbd85e55796df41b86cf0c08b23723d13
  • Current PR head: ce421f8b38764105f7015765c136e5cbe9cf380f

This PR was developed with AI assistance. The human author reviewed the tuning provenance and validation, and owns the submitted data. Commit metadata contains only human attribution and DCO trailers.

@github-actions

github-actions Bot commented Aug 4, 2026

Copy link
Copy Markdown
Contributor

🏷️ CI Guide

Runs automatically on every PR:

  • ✅ Pre-checks (submodule verification, code formatting)
  • ✅ Aiter op tests (gfx942 + gfx950)
  • ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label Tests
ci:triton-300x Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
ci:sglang SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 4554 --add-label <label>

vorapolsiloai and others added 2 commits August 18, 2026 13:45
Signed-off-by: Vorapol Assavasangthong <Vorapol.Assavasangthong@amd.com>
Co-authored-by: Bishwo Adhikari <Bishwo.Adhikari@amd.com>
Signed-off-by: Vorapol Assavasangthong <Vorapol.Assavasangthong@amd.com>
Co-authored-by: Bishwo Adhikari <Bishwo.Adhikari@amd.com>
@vorapolsiloai
vorapolsiloai force-pushed the tune-qwen3-vl-235b-mxfp4-gfx950 branch from bb8c63d to 169f1b3 Compare August 18, 2026 14:02
@vorapolsiloai
vorapolsiloai marked this pull request as ready for review August 18, 2026 14:03
@vorapolsiloai
vorapolsiloai requested a review from a team August 18, 2026 14:03
@github-actions github-actions Bot changed the title configs: add tuned configs for Qwen3-VL-235B MXFP4 (gfx950) [Config] configs: add tuned configs for Qwen3-VL-235B MXFP4 (gfx950) Aug 18, 2026
Current main added two rows that collide with faster Qwen3-VL entries for the same runtime keys. Remove the slower rows so global config merging remains unique and deterministic.

Signed-off-by: Vorapol Assavasangthong <Vorapol.Assavasangthong@amd.com>
@vorapolsiloai

Copy link
Copy Markdown
Contributor Author

The earlier BF16 duplicate-key failure is fixed on the current head, and the corrected wheel build now passes config merging. One unrelated standard-test shard still needs a maintainer rerun:

  • linux-aiter-mi35x-1, shard 4 failed in op_tests/test_moe_2stage.py with torch.OutOfMemoryError (requested 21.05 GiB; 14.64 GiB free; 229.52 GiB already allocated).
  • This is the same test, allocation, and available-memory failure seen on [HIP] [kernel] mrope cache-quant: accept strided flash KV-cache view #4531, confirming a suite resource issue rather than a tuning-data regression.
  • Fork authors do not have permission to invoke the job rerun API.

Could a maintainer please rerun the failed shard?

@zufayu
zufayu requested a review from yifehuan August 19, 2026 01:25
Replace trace-specific rows with reusable runtime buckets and exclude fragile BF16 backends. Add matching untuned manifests while restoring shared BF16 entries.

Co-authored-by: Cursor <cursoragent@cursor.com>
@vorapolsiloai

Copy link
Copy Markdown
Contributor Author

Thanks for the detailed review. I addressed both points:

1. hipBLASLt exclusion

  • Retuned BF16 with direct libtype=hipblaslt excluded.
  • Also excluded OPUS after E2E testing exposed a packaged-kernel failure.
  • Final BF16 file contains 92 rows: 78 Torch, 9 ASM, 3 Skinny, and 2 Triton; zero hipBLASLt/OPUS rows.
  • BF16 correctness and HIP graph capture/replay passed for all 92/92 rows.

2. Runtime padded-M grid

  • Mapped the raw request shapes through runtime getPaddedM, deduplicated them by runtime key, and retuned the canonical bucket endpoints.
  • Used the power-of-two/fine-grained tier approach established by [Config] [gemm][gfx950] Tune Kimi-K3 A4W4 and FP8 projections #4873 and [Tune] Add qwen3.5-397B MXFP4 a16w16 GEMM tuning configs #3974.
  • Reduced A4W4 from 7,212 to 87 rows and BF16 from 945 to 92 rows.
  • Reused existing shared config entries instead of adding duplicate overrides.
  • Added matching untuned CSVs whose endpoint keys match the tuned files.
  • Runtime collision and dispatch validation passed, with no missing A4W4 configuration warnings during tuned TP2/TP4 E2E runs.

E2E validation completed without kernel, graph-capture, or serving errors. The revised configuration retained measurable uplift versus Qwen3-VL tuning disabled: +3.53% TP2 and +4.65% TP4 total throughput on the quick warm workload.

Comment thread aiter/configs/model_configs/q3vl_bf16_tuned_gemm.csv Outdated
Comment thread aiter/configs/model_configs/q3vl_a4w4_blockscale_tuned_gemm.csv Outdated
@yifehuan
yifehuan merged commit 23f8372 into ROCm:main Aug 27, 2026
80 of 82 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants