[RadeonFlow] MXFP4 (a4w4) MoE backend for gfx950 - #3832
Merged
Conversation
* feat(mxfp4_moe): MoE sort/quant/scatter-reduce aux kernels
Add module_moe_mxfp4_aux for the MXFP4 (a4w4) MoE backend: fused token
sort+quant, 3-stage sort, scale sort/shuffle, and weighted bf16
scatter-reduce (gfx950, Kimi-K2.5 / DSR shapes).
* feat(mxfp4_moe): a4w4 gemm1/gemm2 kernels + codegen dispatch
Add module_moe_mxfp4_gemm: the two MXFP4 (a4w4) MoE GEMMs (gemm1 fused
SwiGLU, gemm2 down-proj) with codegen'd kernel-name dispatch
(gen_instances.py), MFMA f4f4 microkernels, epilogs, non-temporal
B-load dispatch, and XCD-swizzle interface (gfx950, Kimi-K2.5 / DSR).
* feat(mxfp4_moe): wire mxfp4 backend into fused_moe + tuned CSV
Route mxfp4_moe_* kernel names to the _mxfp4_moe_run pipeline
(sort/quant -> gemm1 -> gemm2 -> scatter_reduce) via a MOEMetadata
pipeline hook, with shuffle_kind-tagged CSV lookup so the backend ships
its own tuned rows alongside the default set. Add Kimi-K2.5 tuned CSV.
* feat(mxfp4_moe): mxfp4-intermediate gemm2->reduce path (CSV-gated)
Add the optional mxfp4-intermediate path: gemm2 stages flat_out as
packed fp4 + e8m0 (mxfp4-out epilog) and scatter_reduce_q reads it back
as mxfp4, cutting the reduce's HBM traffic ~3.8x. Gated by a
_MXFP4OUT tag in the tuned CSV (large-M buckets only, where the reduce
win beats the gemm2 epilog overhead). bf16-intermediate stays default.
* refactor(mxfp4_moe): remove MAX_M template parameter from gemm1 and gemm2 kernels
Eliminate the MAX_M template parameter from the gemm1 and gemm2 kernels, replacing it with a runtime-calculated max_sorted value. This change enhances flexibility and correctness in buffer resource calculations, ensuring proper handling of varying input sizes. Update related launch functions and code generation scripts accordingly.
* test(mxfp4_moe): add 128k chunk-invariance test for the max_sorted buffer bound
Validate the runtime max_sorted fix at 128k tokens via the per-token MoE invariant moe(128k) == concat(16 x moe(8k)): the 8k chunks stay within any bound (trusted reference) while the 128k run exercises the large sorted-row count (>655360) that the old compile-time MAX_M clipped. Uses random weights (self-consistency needs only identical weights + routing across runs) and masks the kernel's inherent FP-reduce nondeterminism by counting a break only where the kernel is self-deterministic yet full != chunked.
* fix(mxfp4_moe): correct gemm2 atomic A_scale buffer bound for BM=64
The A_scale buffer-resource bound divided by `kAtomic ? BM_GRID : 32`, but the read side addresses A_scale in fixed 32-row chunks (chunk_base = (BM==16) ? m_row/BM_GRID : m_row/32), i.e. the divisor is 16 only for BM=16 and 32 otherwise, never BM. For BM=64 this under-sized the descriptor 2x: addressing reaches max_sorted/32 chunks while the bound only covers max_sorted/64, so every sorted row with m_row >= max_sorted/2 fails the hardware bounds check, the scale load returns 0 (e8m0=0 -> 2^-127), and those output rows silently collapse toward zero. Match the bound divisor to the addressing: (BM==16) ? BM_GRID : 32. This was masked under the old fixed MAX_M (large enough to cover everything) and only surfaced once the bound was tightened to the runtime max_sorted.
* test(mxfp4_moe): add gemm2 BM=64 A_scale bound unit test
Drives mxfp4_moe_gemm2_a4w4 directly and compares the BM=32 atomic kernel (correct bound) against BM=64 (the kernel whose A_scale bound was 2x too small) on identical inputs, which must agree since both address A_scale in 32-row chunks. M_logical/cumsum are chosen so the BM=64 sorted length exceeds the buggy max_sorted/2, placing the upper rows in the clipped region. fp4 partials are forced non-negative so the per-token atomic sum has no cancellation (noise floor ~0), making the clipped-row signal unambiguous. Fails on the buggy bound (7688/8192 rows), passes after the fix (0/8192).
* feat(mxfp4_moe): gate fp4-intermediate gemm2 output behind AITER_MXFP4_INTERMEDIATE
The fp4 gemm2-output path quantizes each expert's output to 4-bit before the topk reduce. That before-sum quant is lossy: fine for gsm8k but it degrades other evals. Make it opt-in via AITER_MXFP4_INTERMEDIATE=1 (off by default), so real evals run the accurate bf16 reduce unless explicitly enabled.
* refactor(mxfp4_moe): generate aux kernel instances from the gemm codegen
Replace the eight hand-written template-dispatch switches in mxfp4_moe_aux.cu
with codegen'd extern "C" instances + string-keyed lookup tables, mirroring the
gemm1/gemm2 dispatch. gen_instances.py gains a --target {gemm,aux} flag so a
single shared SHAPES list drives both: adding a shape regenerates gemm and aux
together. New aux/codegen/mxfp4_moe_aux_dispatch.h holds the per-entry function
pointer types plus launch-config constants; module_moe_mxfp4_aux now runs the
codegen via blob_gen_cmd (was none).
Covers all eight aux entries: sort_quant, sort (threestage / inline_quant /
inline_quant+zero_init), quant, sort_scales, scatter_reduce, scatter_reduce_q.
Generated instances are byte-identical template instantiations of the prior
launches; verified numerically identical to the hand-switch build (old and new
show the same FP nondeterminism on the same cancellation rows) and all
MB16/MB32/MB128 keys resolve e2e on the Kimi NE=385 shape.
* test(mxfp4_moe): harden the 128k chunk-invariance noise floor
The single full-vs-full noise sample sporadically missed catastrophic-
cancellation rows (whose full-vs-chunk difference comes from the differing
reduction order, not a bug), causing ~1-in-4 false failures. Estimate the noise
floor as a union over several full re-runs (--noise-reps, default 4) and, since
full/chunked use different token counts and a same-order probe can never fully
capture cross-order cancellation, gate on the COUNT of residual breaks: the
max_sorted bug corrupts ~half the rows (~NT/2) while cancellation noise is a
handful, so tolerate <0.1% of rows and fail only on a bug-scale fraction.
Reliable PASS across repeated runs; still catches the bug with a ~500x margin.
* test(mxfp4_moe): cleanup test code
* refactor(mxfp4_moe): respect guinterleave shuffle
* refactor: gate by GateMode.INTERLEAVE, not by custom tag.
---------
Co-authored-by: Zesen Liu <ftyghome@gmail.com>
Co-authored-by: GnSight <ftyg@live.com>
Pass bare data_ptr() device addresses (fx.Int64) for the global buffer args instead of full memref descriptors. The kernels only need base pointers (contiguity + compile-time sizes), so the addresses pack contiguously into kernargs -> coalesced s_load prologue. ~7% faster at decode (M=4/8), converging to parity by M=64. cos preserved.
- backend-managed K-loop waitcnt: drop the hand-tuned inline-asm vmcnt that the LLVM waitcnt pass double-inserted (the large-M scheduling bubble); let rocdl.barrier() own the load->ds_read wait. - epilog output-address strength reduction: hoist the loop-invariant row*N_OUT i64 multiply out of the per-element store loop; per-element offsets become compile-time constants folded into the store address. - unsigned index division (divui/remui) for the non-negative grid/tile/count index math, dropping signed-division sign-correction SALU. M=4096 1.11 -> 1.00 (parity with HIP), M=16384 1.10 -> 1.03; faster/parity through M=2048 unchanged. cos preserved.
Use rocdl.exp2 (v_exp_f32) in silu_mul instead of the software math.exp2 expansion (matches HIP __expf). Removes ~31% of non-MFMA VALU (the v_ldexp + range-clamp v_cmp/v_cndmask the polynomial emitted); large-M ~2-6% faster (M=4096 1.07 -> 1.05). cos preserved.
…the HIP backend (#3828) * flydsl mxfp4 gemm1/gemm2: raw-pointer (data_ptr i64) ABI Pass bare data_ptr() device addresses (fx.Int64) for the global buffer args instead of full memref descriptors. The kernels only need base pointers (contiguity + compile-time sizes), so the addresses pack contiguously into kernargs -> coalesced s_load prologue. ~7% faster at decode (M=4/8), converging to parity by M=64. cos preserved. * flydsl mxfp4 gemm2: large-M (BM128) optimizations - backend-managed K-loop waitcnt: drop the hand-tuned inline-asm vmcnt that the LLVM waitcnt pass double-inserted (the large-M scheduling bubble); let rocdl.barrier() own the load->ds_read wait. - epilog output-address strength reduction: hoist the loop-invariant row*N_OUT i64 multiply out of the per-element store loop; per-element offsets become compile-time constants folded into the store address. - unsigned index division (divui/remui) for the non-negative grid/tile/count index math, dropping signed-division sign-correction SALU. M=4096 1.11 -> 1.00 (parity with HIP), M=16384 1.10 -> 1.03; faster/parity through M=2048 unchanged. cos preserved. * flydsl mxfp4 gemm1: hardware exp2 in silu Use rocdl.exp2 (v_exp_f32) in silu_mul instead of the software math.exp2 expansion (matches HIP __expf). Removes ~31% of non-MFMA VALU (the v_ldexp + range-clamp v_cmp/v_cndmask the polynomial emitted); large-M ~2-6% faster (M=4096 1.07 -> 1.05). cos preserved.
* refactor(mxfp4_moe): drop shuffle_kind; dispatch a4w4 by kernel name + gemm_backend
* improv: remove K=7168 constraint
* feat(mxfp4_moe): parameterize BN/BK tiles in a4w4 gemm
Make the N tile (BN) and gemm2 K tile (BK) template parameters instead of
hardcoded 256, and generalize the shapes the a4w4 gemm kernels accept:
- BN in {128,256}; N_OUT only needs %16 (MFMA-N). N_OUT not divisible by BN
is ceil-tiled and the ragged tail is dropped on write (gemm2 atomic) /
skipped per inter-block (gemm1 cshuffle epilog).
- gemm2 K=256 single-tile path, and BK=512 processed as two 128-byte units
per tile (load-bearing 8-lane load + LDS swizzle unchanged).
- gemm2 K padded to a 256-multiple for non-aligned D_INTER (scale layout /
MFMA-K); gemm1 epilog writes inter at the padded row stride.
codegen picks BN/BK per shape (pick_bn/pad_k/pick_bk) and threads them
through the launch templates. Kimi/DSR (BN=256, BK=256, K=512) is
byte-identical (numerically and in measured kernel time).
* feat(mxfp4_moe): K-pad gemm2 weights/inter for non-256 D_INTER
When D_INTER is not a multiple of 256, zero-pad w2 / w2_scale and the inter
bridge buffer along K so gemm2 runs on a clean 256-multiple K (BK / scale
layout / MFMA-K=128 require it). No-op for 256-multiples (e.g. D_INTER=512).
* test(mxfp4_moe): D_INTER=192 chunk-invariance test
Exercises gemm1 N_OUT=384 (BN=128) + gemm2 K=192 padded to 256; checks
full-batch == chunked output (MoE is per-token independent).
* fix: mxfp4 moe dispatch shuffle
* feat(mxfp4_moe): add AITER_MXFP4_MOE_BACKEND=flydsl toggle
* feat(mxfp4_moe): log resolved a4w4 gemm backend once
* refactor: integrate new sorting to fused_moe pipeline
Contributor
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
|
…gemm2 846->778us @ M=16384) (#3840)
# Conflicts: # aiter/fused_moe.py
…with token registry Append _FLYDSL to the Kimi-K2.5 tuned CSV kernelName1/2 so the a4w4 MoE selects the FlyDSL gemm port (_mxfp4_moe_run) instead of the HIP backend. Replace the _MXFP4_G1/G2_KNAME_RE regexes with a token-registry parser: names are "_"-joined tokens, classified via _MXFP4_NUMERIC_TOKENS (LETTERS+digits fields) and per-stage flag-token sets. Adding a kernel variant now needs one registry entry instead of re-deriving an optional-group regex. Parse results are field-for-field identical to the old regex (verified across 38 variants incl. XCD/SK/INLINEQUANT_CACHED/epilogs/_FLYDSL); ATOMIC+MXFP4OUT/CSHUFFLE still rejected. Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
…SL port The a4w4 MoE gemm1/gemm2 now run only on the FlyDSL port; remove the parallel HIP gemm backend entirely. moe_aux (sort/quant/sort_scales/scatter_reduce) stays in HIP and is unchanged functionally. - delete csrc/kernels/mxfp4_moe/gemm_a4w4/ (gemm cu/cuh/h/pybind, gemm dispatch header, gemm-only common/ headers, and the combined gemm+aux gen_instances.py) - move the aux codegen to csrc/kernels/mxfp4_moe/moe_aux/codegen/gen_instances.py (aux-only; byte-identical generated instances + lookup header, verified) - optCompilerConfig.json: remove module_moe_mxfp4_gemm; point module_moe_mxfp4_aux at the new aux codegen path - mxfp4_moe.py: drop the 3 module_moe_mxfp4_gemm @compile_ops bindings (gemm1/ gemm2/mxfp4out); keep aux bindings - fused_moe.py: collapse the per-stage HIP-vs-flydsl branches to flydsl-only, drop gemm_backend plumbing, the _FLYDSL kernel-name token / AITER_MXFP4_MOE_BACKEND env resolution, and the HIP-only AITER_MXFP4_INTERMEDIATE==2 path (the ==1 flydsl mxfp4out epilog supersedes it); update _mxfp4_sort_shapes to the new codegen path - Kimi tuned CSV: drop the now-meaningless _FLYDSL suffix from kernelName1/2 - csrc/include/mxfp4_moe.h: remove the HIP gemm kernel declarations Verified: 128k chunk-invariance test PASS, only module_moe_mxfp4_aux builds, gemm1/gemm2 dispatch to the flydsl port. Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
Collaborator
Author
|
@Bernard-Liu conflict. |
…IATE Drop the CSV `_MXFP4OUT` kernel-name flag from the dispatch gate so the mxfp4-out gemm2 path is controlled purely by AITER_MXFP4_INTERMEDIATE=1 (plus the Kimi/DSR shape guard). The CSV no longer needs a _MXFP4OUT row.
…conflicts) Base (dev/randomflow_pr) rewrote fused_moe.py (~1340 lines); this branch's only fused_moe.py change is the "gate gemm2 mxfp4-out solely on AITER_MXFP4_INTERMEDIATE" commit. Resolution: - aiter/fused_moe.py: take base's rewritten structure, re-apply the single semantic change -- drop the `mxfp4out and` term from the mxfp4-out gate in _mxfp4_a4w4_stage2 so it gates on AITER_MXFP4_INTERMEDIATE=="1" alone. - aiter/ops/flydsl/kernels/mxfp4_gemm2.py: keep this branch's abs-fold epilog (llvm.fabs.f32 + maxnumf amax); drop base's now-dead `i7fff` integer-mask declaration. The #3840 LDS pipelining is identical on both sides.
* flydsl mxfp4 gemm1/gemm2: raw-pointer (data_ptr i64) ABI Pass bare data_ptr() device addresses (fx.Int64) for the global buffer args instead of full memref descriptors. The kernels only need base pointers (contiguity + compile-time sizes), so the addresses pack contiguously into kernargs -> coalesced s_load prologue. ~7% faster at decode (M=4/8), converging to parity by M=64. cos preserved. * flydsl mxfp4 gemm2: large-M (BM128) optimizations - backend-managed K-loop waitcnt: drop the hand-tuned inline-asm vmcnt that the LLVM waitcnt pass double-inserted (the large-M scheduling bubble); let rocdl.barrier() own the load->ds_read wait. - epilog output-address strength reduction: hoist the loop-invariant row*N_OUT i64 multiply out of the per-element store loop; per-element offsets become compile-time constants folded into the store address. - unsigned index division (divui/remui) for the non-negative grid/tile/count index math, dropping signed-division sign-correction SALU. M=4096 1.11 -> 1.00 (parity with HIP), M=16384 1.10 -> 1.03; faster/parity through M=2048 unchanged. cos preserved. * flydsl mxfp4 gemm1: hardware exp2 in silu Use rocdl.exp2 (v_exp_f32) in silu_mul instead of the software math.exp2 expansion (matches HIP __expf). Removes ~31% of non-MFMA VALU (the v_ldexp + range-clamp v_cmp/v_cndmask the polynomial emitted); large-M ~2-6% faster (M=4096 1.07 -> 1.05). cos preserved. * lgkmcnt * flydsl mxfp4 gemm1: prologue vmcnt relax + BM128 scale-first ds_read * flydsl mxfp4 gemm2: float fabs/maxnum amax in mxfp4out epilog * flydsl mxfp4 moe: gate gemm2 mxfp4-out solely on AITER_MXFP4_INTERMEDIATE Drop the CSV `_MXFP4OUT` kernel-name flag from the dispatch gate so the mxfp4-out gemm2 path is controlled purely by AITER_MXFP4_INTERMEDIATE=1 (plus the Kimi/DSR shape guard). The CSV no longer needs a _MXFP4OUT row. --------- Co-authored-by: zhutaoyu <zhutaoyu97@gmail.com> Co-authored-by: Zesen Liu <ftyghome@gmail.com>
Add the mxfp4_moe_g*_a4w4 port to the moe 2-stage tuner alongside the generic
flydsl_moe* engine, so both FlyDSL implementations can be tuned and PK'd.
- gen_mxfp4_port_2stages_task: enumerates port candidates (BM x epilog) from the
port's _SUPPORTED variant sets (aiter/ops/flydsl/mxfp4_gemm{1,2}_kernels.py),
for a4w4 (per_1x32, fp4 act + fp4 weight) only; hooked into tune()'s main loop
- run_mxfp4_port_stage{1,2}_out: timing shells that drive the port via its
production stage entries (_mxfp4_a4w4_stage{1,2}_fw) + moe_sorting(fused_sort=True),
with a16w4 weight/scale shuffles
- _mxfp4_port_g{1,2}_kname: build kernel names matching the codegen/CSV grammar
(round-trips through _parse_mxfp4_g{1,2}_kname)
Shells take the low-level generate_data output (raw bf16 hidden + clean fp4/e8m0
weights), not generate_data_2stages (whose a4w4 a1_qt is pre-quantized to fp4).
Candidates are timing-only (ref_func=None under fast_mode); correctness is
covered by the standalone e2e tests. Verified: shells run end-to-end on a real
Kimi shape and the 16 enumerated knames round-trip through the parser.
Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
CI ruff check flagged the mxfp4 moe files: - F401 unused imports (torch in mxfp4_moe.py; aiter/ActivationType/QuantType in the chunk tests) - F841 unused local (A_SCALE_COLS in test_mxfp4_flydsl_gemm1.py) - E731 lambda assignments -> rewritten as def (test helpers u8/e8/eb/pos4) Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
Reformat the mxfp4 moe files flagged by the CI black check (psf/black@stable). Formatting-only; no behavior change. Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
moe_sort_quant's inline quant uses __builtin_amdgcn_cvt_scalef32_pk_fp4_bf16, which needs fp4-cvt-scale-insts (gfx950-only). In a multi-arch build the aux TU is also compiled for gfx942, where that builtin fails to compile. Gate it behind __gfx950__; other targets get a __builtin_trap() fallback so the TU still compiles but never silently runs (this mxfp4 MoE is gfx950-only at runtime). Verified: aux instance TU compiles for both gfx942 (trap path) and gfx950 (builtin path). Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
- remove test_mxfp4_moe_gemm2_bm64_ascale_bound.py: it tested an off-by-2x A_scale bound in the HIP gemm2 (aiter.mxfp4_moe_gemm2_a4w4), which no longer exists after the HIP backend was deleted. - remove test_mxfp4_moe_e192_chunk.py + drop (385,7168,192,9) from the aux codegen SHAPES: D_INTER=192 isn't reachable for Kimi-K2.5 (moe_intermediate _size=2048, so TP-sharded D_INTER is 512/256/128, all 256-aligned). 192 was a HIP-era shape that relied on host K-pad to 256; the FlyDSL port requires 256-aligned N_OUT/K and never implemented that pad path, so the shape has no working backend. Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
|
Accuracy: Run 1 (num_concurrent=5)
Run 2 (num_concurrent=9)
Run 3 (num_concurrent=17)
Run 4 (num_concurrent=33)
Run 5 (num_concurrent=65)
Run 6 (num_concurrent=129)
|
valarLip
reviewed
Jul 6, 2026
| # below skips, so emit its job here. The cache key needs only | ||
| # (inter_dim, topk)/(inter_dim), which the CSV shape covers regardless | ||
| # of runtime split_k. | ||
| if stage1_name.startswith("cktile_"): |
Contributor
There was a problem hiding this comment.
aiter.moe_cktile_2stages_gemm1(...)
if split_k>1 & interleaved:
flydsl_silu_and_mul_interleaved(...)
It's a ck-tile moe but will run to a FlyDSL reduce kernel.
Added this so that the Cl can passed (which require all FlyDSL kernels must be AOT ones).
| rows, | ||
| sorted_token_ids.shape[0], | ||
| float("inf"), | ||
| 0, |
Contributor
There was a problem hiding this comment.
Added，that's the dummy default stream only for aot to use.
Clarify the trailing positional args in _precompile_epilogue_to_cache:
float("inf") is swiglu_limit (unused on the silu path) and the trailing 0 is
the CUDA stream, which is null/default here because this is the compile-only
AOT path (the kernel is compiled and persisted but never launched).
Co-Authored-By: Claude Opus 4 (1M context) <noreply@anthropic.com>
# Conflicts: # aiter/jit/optCompilerConfig.json
# Conflicts: # aiter/configs/model_configs/kimik2_fp4_tuned_fmoe.csv
Contributor
|
Before the last aa65f59 commit to resolve the conflict. All test have passed, expect the kimi one dont have machine to run. |
This was referenced Jul 10, 2026
Merged
Fangzhou-Ai
added a commit
to SemiAnalysisAI/InferenceX
that referenced
this pull request
Jul 21, 2026
… vLLM Replace image: TBD with rocm/vllm-dev:nightly_cdna4, the latest gfx950 serving image. It ships AITER 0.1.19.dev (>= v0.1.16.post5 from vllm-project/vllm#48683) including the ROCm/aiter#3832 gfx950 MXFP4 MoE backend required by this recipe. 中文：将 image: TBD 替换为最新的 gfx950 服务镜像 rocm/vllm-dev:nightly_cdna4。 该镜像内置 AITER 0.1.19.dev（不低于 vllm-project/vllm#48683 引入的 v0.1.16.post5）， 包含本配方所需的 ROCm/aiter#3832 gfx950 MXFP4 MoE 后端。 Co-authored-by: Cursor <cursoragent@cursor.com>
2 tasks
Fangzhou-Ai
added a commit
to Fangzhou-Ai/recipes
that referenced
this pull request
Jul 21, 2026
The Kimi-K2.5 MXFP4 (MI355X) section documents the base AITER MXFP4 path + fp8 KV cache. Add a "Tuned AITER MXFP4 MoE (MI355X)" subsection covering the higher-throughput MoE configuration: the AITER MXFP4 (RadeonFlow) intermediate GEMM path with fused shared experts. VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS=1 VLLM_ROCM_USE_SKINNY_GEMM=0 AITER_MXFP4_INTERMEDIATE=1 # selects the RadeonFlow MXFP4 intermediate path AITER_BYPASS_TUNE_CONFIG=0 AITER_MOE_SORT_BACKEND=auto VLLM_ROCM_USE_AITER_RMSNORM=0 # TP < 8 only plus batching/scheduling flags (--block-size 16, --max-num-batched-tokens 16384, --max-num-seqs 512, --async-scheduling, --no-enable-prefix-caching). Runs at TP4 or TP8. Requires a nightly built after vLLM #48683 (AITER v0.1.16.post5, which includes ROCm/aiter#3832). These are kernel-selection choices, not precision reductions. Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com>
Fangzhou-Ai
added a commit
to Fangzhou-Ai/recipes
that referenced
this pull request
Jul 21, 2026
The Kimi-K2.5 MXFP4 (MI355X) section documents the base AITER MXFP4 path + fp8 KV cache. Add a "Tuned AITER MXFP4 MoE (MI355X)" subsection covering the higher-throughput MoE configuration: the AITER MXFP4 (RadeonFlow) intermediate GEMM path with fused shared experts. VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS=1 VLLM_ROCM_USE_SKINNY_GEMM=0 AITER_MXFP4_INTERMEDIATE=1 # selects the RadeonFlow MXFP4 intermediate path AITER_BYPASS_TUNE_CONFIG=0 AITER_MOE_SORT_BACKEND=auto VLLM_ROCM_USE_AITER_RMSNORM=0 # TP < 8 only plus batching/scheduling flags (--block-size 16, --max-num-batched-tokens 16384, --max-num-seqs 512, --async-scheduling, --no-enable-prefix-caching). Runs at TP4 or TP8. Requires a nightly built after vLLM #48683 (AITER v0.1.16.post5, which includes ROCm/aiter#3832). These are kernel-selection choices, not precision reductions. Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com>
Fangzhou-Ai
added a commit
to SemiAnalysisAI/InferenceX
that referenced
this pull request
Jul 22, 2026
The AITER MXFP4 MoE bump (vllm-project/vllm#48683 / ROCm/aiter#3832) has landed in the floating rocm/vllm-dev:nightly_cdna4 tag, so drop the immutable dated pin and track nightly_cdna4 directly. 中文：改为跟踪浮动的 rocm/vllm-dev:nightly_cdna4 标签。AITER MXFP4 MoE 的版本提升（vllm-project/vllm#48683 / ROCm/aiter#3832）已合入 nightly_cdna4 浮动标签，因此去掉固定日期的镜像 pin，直接跟踪 nightly_cdna4。 Co-authored-by: Cursor <cursoragent@cursor.com>
Fangzhou-Ai
added a commit
to SemiAnalysisAI/InferenceX
that referenced
this pull request
Jul 22, 2026
…MI355X Switch the kimik2.5-fp4-mi355x-vllm image from the rocm/vllm-dev CDNA4 nightly to the official vllm/vllm-openai-rocm nightly, commit-pinned to the immutable nightly-387189c42997b27e2c04b5d97ef8190ffa2bf909 tag (latest as of 2026-07-22). It carries the AITER MXFP4 MoE bump (vllm-project/vllm#48683 / ROCm/aiter#3832) while staying reproducible. 中文：将 kimik2.5-fp4-mi355x-vllm 的镜像从 rocm/vllm-dev CDNA4 nightly 切换为官方 vllm/vllm-openai-rocm nightly，并按提交固定到不可变的 nightly-387189c42997b27e2c04b5d97ef8190ffa2bf909 标签（截至 2026-07-22 最新）。该镜像已包含 AITER MXFP4 MoE 的版本提升 （vllm-project/vllm#48683 / ROCm/aiter#3832），同时保持可复现。 Co-authored-by: Cursor <cursoragent@cursor.com>
functionstackx
pushed a commit
to SemiAnalysisAI/InferenceX
that referenced
this pull request
Jul 30, 2026
…2.5 MXFP4 MI355X vLLM，缩小与 ATOM 的性能差距 (#2213) * perf(amd): tune Kimi K2.5 MXFP4 on MI355X [skip-sweep] Apply the accuracy-gated serving settings, preserve the existing TP4 and TP8 matrix dimensions, extend only their concurrency upper bounds to 128, and keep the image as a draft placeholder pending the required AITER bump. 中文：应用已通过精度验证的服务配置，保留现有 TP4 和 TP8 矩阵维度，仅将并发上限扩展到 128，并暂时保留镜像占位符，等待所需的 AITER 版本升级。 * Update kimik2.5_fp4_mi355x.sh * Update kimik2.5_fp4_mi355x.sh * perf(amd): pin rocm/vllm-dev:nightly_cdna4 for Kimi-K2.5 MXFP4 MI355X vLLM Replace image: TBD with rocm/vllm-dev:nightly_cdna4, the latest gfx950 serving image. It ships AITER 0.1.19.dev (>= v0.1.16.post5 from vllm-project/vllm#48683) including the ROCm/aiter#3832 gfx950 MXFP4 MoE backend required by this recipe. 中文：将 image: TBD 替换为最新的 gfx950 服务镜像 rocm/vllm-dev:nightly_cdna4。 该镜像内置 AITER 0.1.19.dev（不低于 vllm-project/vllm#48683 引入的 v0.1.16.post5）， 包含本配方所需的 ROCm/aiter#3832 gfx950 MXFP4 MoE 后端。 Co-authored-by: Cursor <cursoragent@cursor.com> * perf(amd): pin immutable nightly_cdna4 build tag for Kimi-K2.5 MXFP4 MI355X vLLM Use the reproducible dated build tag rocm/vllm-dev:nightly_cdna4_main_torch_2.11.0_rocm7.14.0a20260623_0721_b101 (same digest as floating nightly_cdna4) instead of the moving nightly_cdna4 tag. 中文：改用可复现的带日期构建标签 rocm/vllm-dev:nightly_cdna4_main_torch_2.11.0_rocm7.14.0a20260623_0721_b101 （与浮动的 nightly_cdna4 同一 digest），替换会被覆盖的 nightly_cdna4 标签。 Co-authored-by: Cursor <cursoragent@cursor.com> * fix(kimik2.5): lower MI355X gpu-memory-utilization 0.90 -> 0.85 TP4/TP8 runs OOM during CUDA-graph capture: MI355X nodes hold ~30 GiB per GPU outside the vLLM process, so a 0.90 budget (259 GiB) leaves 0 bytes physically free and capture buffers cannot allocate. 0.85 keeps weights + KV inside 288 - 30 GiB with headroom for graph capture. 中文：将 MI355X 的 gpu-memory-utilization 从 0.90 下调到 0.85。TP4/TP8 在 CUDA graph 捕获阶段发生显存溢出——MI355X 节点每张卡有约 30 GiB 被 vLLM 进程之外占用，0.90 的预算（259 GiB）导致物理显存无剩余，捕获缓冲无法 分配。0.85 使权重与 KV 缓存控制在 288 - 30 GiB 内，为 graph 捕获留出余量。 Co-authored-by: Cursor <cursoragent@cursor.com> * perf(amd): track floating nightly_cdna4 tag for Kimi-K2.5 MXFP4 MI355X The AITER MXFP4 MoE bump (vllm-project/vllm#48683 / ROCm/aiter#3832) has landed in the floating rocm/vllm-dev:nightly_cdna4 tag, so drop the immutable dated pin and track nightly_cdna4 directly. 中文：改为跟踪浮动的 rocm/vllm-dev:nightly_cdna4 标签。AITER MXFP4 MoE 的版本提升（vllm-project/vllm#48683 / ROCm/aiter#3832）已合入 nightly_cdna4 浮动标签，因此去掉固定日期的镜像 pin，直接跟踪 nightly_cdna4。 Co-authored-by: Cursor <cursoragent@cursor.com> * perf(amd): pin official vllm-openai-rocm nightly for Kimi-K2.5 MXFP4 MI355X Switch the kimik2.5-fp4-mi355x-vllm image from the rocm/vllm-dev CDNA4 nightly to the official vllm/vllm-openai-rocm nightly, commit-pinned to the immutable nightly-387189c42997b27e2c04b5d97ef8190ffa2bf909 tag (latest as of 2026-07-22). It carries the AITER MXFP4 MoE bump (vllm-project/vllm#48683 / ROCm/aiter#3832) while staying reproducible. 中文：将 kimik2.5-fp4-mi355x-vllm 的镜像从 rocm/vllm-dev CDNA4 nightly 切换为官方 vllm/vllm-openai-rocm nightly，并按提交固定到不可变的 nightly-387189c42997b27e2c04b5d97ef8190ffa2bf909 标签（截至 2026-07-22 最新）。该镜像已包含 AITER MXFP4 MoE 的版本提升 （vllm-project/vllm#48683 / ROCm/aiter#3832），同时保持可复现。 Co-authored-by: Cursor <cursoragent@cursor.com> --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Chun Fang <chun.fang@amd.com>
8 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
9 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Technical Details
Test Plan
Test Result
Submission Checklist