[Kernel] Add tuned W8A8 block-FP8 GEMM configs for gfx950 (MI350X)

[mustafayildirim]

Purpose

Add tuned W8A8 block-FP8 triton GEMM configs for gfx950 (AMD Instinct MI350X, device_name=0x75b0). These cover the four MLA projection GEMM shapes exercised by GLM-5.3-Flash (Glm5NextForConditionalGeneration); the configs directory currently has zero entries for gfx950, so the kernel falls back to default tile configs with a "Performance might be sub-optimal!" warning at boot.

Shapes added (N, K): (3072, 4096), (4096, 1536), (4096, 256), (512, 4096) — each with 12 M buckets spanning decode (M=1..256) and prefill (M=1024..16384).

Methodology

Swept triton launch configs (BLOCK_SIZE_M/N/K, GROUP_SIZE_M, num_warps, num_stages) per M bucket against _w8a8_triton_block_scaled_mm on MI350X (ROCm 7.2.3, vllm nightly 0bfc7a1), triton.testing.do_bench timing, selecting the fastest valid config per bucket. JSONs follow the existing get_w8a8_block_fp8_configs() lookup format.

Validation

Serving GLM-5.3-Flash on 8x MI350X (TP=8, breakable cudagraphs, aiter MoE/MLA), identical vllm bench serve runs before/after:

Profile	Metric	default configs	tuned configs	delta
in=2048 out=512 conc=32	mean TTFT	5214 ms	3178 ms	-39%
in=2048 out=512 conc=32	mean TPOT	22.98 ms	17.99 ms	-22%
in=32768 out=256 conc=8	mean TTFT	3385 ms	2934 ms	-13%
in=32768 out=256 conc=8	mean TPOT	30.88 ms	26.93 ms	-13%

Boot log confirms all 32 "Using default W8A8 Block FP8 kernel config" warnings replaced by "Using configuration from ..." lookups. Full correctness suite (mixed chunked-prefill/decode repro, 32K-524K long-context ladder) passes with zero memory-access faults.

Test Plan

N/A — data-only change (JSON config files); no code paths modified.

[claude]

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

[simondanielsson]

Thanks for adding this!

FYI: ROCm/aiter#5602

[mustafayildirim]

Both addressed in 1ad2a5b:

1. indent=4 — JSONs rewritten to match the dir's formatting convention.
2. Device name — files renamed to AMD_Instinct_MI350X, and added "0x75b0": "AMD_Instinct_MI350X" to _ROCM_DEVICE_ID_NAME_MAP so get_device_name_as_file_name() resolves deterministically instead of depending on the amd-smi market_name string.

Re aiter#5602: complementary rather than overlapping — that one tunes aiter's own GEMM backends (FlyDSL/ASM/Torch, BF16 + A8W8) on MI355X, while these JSONs feed vllm's triton _w8a8_triton_block_scaled_mm path that the fp8_utils loader picks. Both can land; whichever backend the dispatcher selects at runtime gets tuned tiles. (For what it's worth, our serving path on this nightly resolved the MLA projections to the triton kernel — the 32 'Using default W8A8' warnings — so these configs are the ones hit in practice on this stack.)

[mustafayildirim]

Thanks for the pointer! Looked at ROCm/aiter#5602 — that one covers gfx950 A8W8 block-scale + BF16 configs for the GLM-5.3 Flash projection shapes, on MI355X (256 CUs). This PR is MI350X (304 CUs) W8A8 per-tensor configs from tuning runs on our fleet, so they complement rather than overlap — different SKU, different scale scheme, and the vllm-side lookup needs the device-specific entries regardless. Happy to cross-check shape coverage against the 1,826 captured shapes from that PR if useful.