Serve Kimi-K2.5-MXFP4 on MI355X with ATOM

In our previous Kimi-K2.5 posts, we moved from fused MoE optimization with FlyDSL [8] to W4A8 and W8A8 quantization with AMD Quark [9]. Those posts focused on kernel and quantization work: how to make the dominant MoE path faster, and how to trade off INT4, INT8, and FP8 formats on AMD Instinct™ MI300X and MI325X GPUs.

Serve Kimi-K2.5-MXFP4 on MI355X with ATOM#

20 min read. | 5017 total words.

In our previous Kimi-K2.5 posts, we moved from fused MoE optimization with FlyDSL [8] to W4A8 and W8A8 quantization with AMD Quark [9]. Those posts focused on kernel and quantization work: how to make the dominant MoE path faster, and how to trade off INT4, INT8, and FP8 formats on AMD Instinct™ MI300X and MI325X GPUs.

This post shifts the focus to the serving stack. We show how to serve the pre-quantized amd/Kimi-K2.5-MXFP4 checkpoint [2] on AMD Instinct MI355X with ATOM [3], so the runtime can recognize MXFP4 metadata, dispatch optimized gfx950 FP4 kernels, and expose an OpenAI-compatible endpoint without manually wiring every model and kernel component.

ATOM connects model-specific execution, AITER [4] kernels, torch.compile, CUDA graph capture, distributed execution, and serving interfaces into a deployable stack. In this post, we use ATOM’s standalone OpenAI-compatible server for measurement and describe the vLLM out-of-tree plugin path for users who want to keep vLLM [5] as their serving framework.

ATOM, short for AiTer Optimized Model, is a lightweight vLLM-like implementation focused on integrating and optimizing AMD GPU model execution with AITER kernels. In practice, ATOM provides the framework layer around the optimized kernels:

ROCm and AITER kernel integration: ATOM routes model operations to AITER kernels, including ASM, CK, Triton, and FlyDSL-backed paths.

OpenAI-compatible serving: ATOM provides /v1/completions and /v1/chat/completions endpoints for standalone serving.

Compilation and graph capture: ATOM supports piecewise torch.compile and CUDA graph capture for low-latency decode.

Multi-GPU execution: ATOM supports tensor parallelism, data parallelism, and expert-parallel execution paths for large MoE models.

Quantization-aware dispatch: ATOM reads HuggingFace quantization metadata and maps supported formats, including MXFP4, to runtime quantization and kernel choices.

For Kimi-K2.5-MXFP4, that last point is key. The checkpoint already contains packed FP4 weights and E8M0 scale tensors. ATOM maps those metadata fields to the per_1x32 FP4 runtime path and dispatches AITER gfx950 kernels that execute scaled MFMA instructions.

ATOM can run as its own OpenAI-compatible server. This mode is useful for benchmarking and for validating new model support because the optimized ATOM model implementation and AITER kernels are directly visible in the serving stack.

For many deployments, vLLM is the serving framework that users already use in production. ATOM can also work as a vLLM out-of-tree (OOT) plugin backend: ATOM is installed as a separate Python package and plugs into vLLM through vLLM’s official plugin interfaces.

ATOM registers a vLLM platform plugin so vLLM can resolve the ATOM platform.

ATOM supplies AITER-backed attention backends for supported models.

vLLM continues to drive request scheduling and serving, while the hot model execution path runs through ATOM model code and AITER kernels.

This split matters for adoption. The standalone server is the shortest path to validate a new optimized model stack; the vLLM plugin path lets users keep the vLLM interface and operational model while benefiting from ATOM’s AMD-specific model and kernel optimizations.

The W4A8 post used AMD Quark [6] to produce a checkpoint compatible with the INT8 MFMA runtime path on gfx942. That required a different weight layout from W4A16: INT4 weights with per-channel scales, plus dynamic activation quantization at runtime. The reason was structural: the INT8 MFMA inner loop could not apply a different floating-point weight scale every 32 K values, so W4A8 moved weight scales to a per-channel layout that could be applied in the epilogue.

For MXFP4 on gfx950, the config changes in a different direction. Instead of coarsening the weight scale to make an INT8 epilogue work, MXFP4 keeps a fine-grained 32-value block scale and uses a hardware path that can consume that block scale in the matrix instruction.

For MXFP4, both weights and input activations use dtype=fp4, qscheme=per_group, group_size=32, and scale_format=e8m0. The weights are stored statically in the checkpoint as packed FP4 payloads plus E8M0 weight_scale tensors; activations are quantized dynamically at runtime into the same 1x32 block-scaled representation.

That makes the checkpoint and runtime contract different from W4A8. W4A8 changes scale granularity so the INT8 path can stay efficient on gfx942. MXFP4 keeps per-32 block scaling because gfx950 adds the scaled MFMA family needed to consume those scales directly.

The Kimi-K2.5-MXFP4 Model on gfx950#

4-bit element storage: the value payload uses FP4 E2M1, packed two elements per byte.

Microscaling: each block of 32 K-dimension values carries an E8M0 scale, giving the block a shared exponent.

The Kimi-K2.5-MXFP4 model configuration uses this scheme for both activations and weights:

Value observed in config.json

The W4A8 post focused on a gfx942 question: how do we move the MoE path from BF16 MFMA to an 8-bit MFMA path without losing too much accuracy? MXFP4 asks the next-generation gfx950 version of that question: if the hardware can consume block-scaled 4-bit floating-point operands directly, can we keep both the compact payload and a fine-grained scale layout in the matrix core?

That is the important architectural change. On CDNA3/gfx942, the common choices were BF16, INT8, FP8, and BF8 MFMA forms. On CDNA4/gfx950, the MFMA family adds block-scaled low-precision formats, including FP4, FP6, and BF6, with an E8M0 block exponent shared by 32 K-dimension elements [7].

Compared with FP8, the main MXFP4 advantage is not just “a smaller number.” It is a different memory/computation trade-off:

Half the element bytes of FP8: FP4 payloads are 4 bits instead of 8 bits, reducing weight and activation traffic for quantized matrix paths.

Hardware-visible block scales: gfx950 can consume the FP4 payload and the E8M0 scale operands in scaled MFMA instructions, rather than requiring a full software dequantization back to BF16 first.

Better range than raw FP4: the shared E8M0 exponent lets a 32-element block shift its dynamic range, which is essential for using FP4 in large MoE models.

There is still a trade-off. FP8 keeps more mantissa/exponent information per element and is often simpler for accuracy-sensitive paths. MXFP4 wins when the model can tolerate the lower per-element precision and when memory bandwidth, model footprint, or MoE weight traffic dominate.

For a representative Kimi-K2.5 expert gate_proj or up_proj shard, the logical weight matrix has 2,048 output channels and 7,168 K-dimension values per row. In the MXFP4 checkpoint, this appears as:

The scale overhead is small compared with the payload. A 32-value block needs 16 bytes of FP4 payload plus 1 byte of E8M0 scale, or about 0.53125 bytes per original value before additional tensor metadata and alignment effects. The key point is that this scale is not a dequantization side table that must disappear before compute. On gfx950, it is part of the block-scaled operand stream consumed by the FP4 matrix instruction.

Figure 1. MXFP4 dataflow on gfx950. Both activations and weights are represented as packed FP4 payloads plus E8M0 block scales. The important difference from a storage-only FP4 path is that the block scales stay visible to the matrix instruction through the gfx950 scaled MFMA family.

The previous W4A8 post framed quantization as a layout problem as much as a datatype problem: the target MFMA instruction determines where scales can be applied, and that determines which checkpoint layout is efficient at runtime. MXFP4 follows the same logic, but with a different answer on gfx950.

The key MXFP4 difference is that the scale is not merely checkpoint metadata. The E8M0 scale is part of the block-scaled compute contract: ATOM maps the model metadata to 1x32 FP4 quantization, and AITER can dispatch gfx950 kernels that keep the FP4 payload and scale operands close to the matrix instruction.

When Is the Scale Applied? The gfx950 Scaled-MFMA Difference#

The most important question is not only “how many bits does each value use?” It is “where does the scale enter the GEMM?”

That placement determines whether a scale layout is cheap, expensive, or structurally incompatible with the target MFMA instruction.

W4A16: scale is absorbed before MFMA.

After dequantization, the scale has been baked into the BF16 value. The BF16 MFMA instruction sees only BF16 operands; it does not know whether the original checkpoint had one scale per row or one scale per 32 values. This makes fine-grained per-group scaling accuracy-friendly, but it also keeps the compute path in BF16.

W4A8 on gfx942: no block-scaled MFMA, so scale is deferred until the epilogue.

The weight scale cannot be applied before INT8 MFMA, because multiplying by a floating-point scale would turn the operand into a floating-point value. The gfx942 W4A8 path also does not have a block-scaled MFMA instruction that can consume INT4 or INT8 payloads together with per-32 scale operands. The scale therefore moves to the epilogue, after the dot product is complete. That is why W4A8 prefers per-channel weight scales: one row scale can be applied once at the end. A per-group scale every 32 K values would force the dot product into many separately scaled partial sums, which defeats the clean INT8 inner loop.

MXFP4 on gfx950: scaled MFMA makes per-32 scaling a hardware compute path.

MXFP4 takes a different route. Its E8M0 scale is neither fully baked into a BF16 value before MFMA nor delayed as a single correction after accumulation. The scale is part of the block-scaled FP4 operand representation:

This is the key gfx950 distinction. CDNA4 adds the V_MFMA_SCALE_F32_* family for block-scaled low-precision inputs, including FP4 [7]. The scaled MFMA path can consume low-precision FP4 payloads together with their block scales, so a 1x32 scale layout remains compatible with the high-throughput matrix instruction. In other words, MXFP4 keeps the local dynamic-range benefit of per-32 scaling without reverting the main compute path to BF16 dequantization or forcing all scale application into a late epilogue.

This is why MXFP4 is more than “FP8 but smaller.” The format changes where scale arithmetic lives, and gfx950 is the hardware generation that turns that scale placement into a first-class MFMA path. On gfx950, the serving stack can keep FP4 values compact while still giving each 32-value block its own exponent-like scale at compute time.

The MXFP4 checkpoint carries the runtime contract in config.json, not only in tensor names. Both weights and activations use the same high-level quantization scheme:

Weights are static: the packed FP4 payload and E8M0 weight_scale tensors are stored in the checkpoint.

Activations are dynamic: ATOM quantizes live BF16 activation blocks into FP4 payloads and E8M0 scales at runtime.

Some layers remain outside MXFP4: the checkpoint excludes attention projections, MoE router gates, the language model head, and vision/projector components from this global FP4 quantization path. That keeps precision-sensitive or non-MoE components on their own kernels while concentrating MXFP4 on the matrix-heavy MLP and expert projections.

In the checkpoint index, quantized modules appear as paired weight and weight_scale tensors. For Kimi-K2.5-MXFP4, the index contains 69,303 weight_scale tensors, matching 69,303 modules with stored block-scale metadata. That is the checkpoint-side reason ATOM can map the model to QuantType.per_1x32 rather than treating the FP4 payload as a generic packed byte tensor.

The MXFP4 serving path matters because it keeps FP4 payloads close to the matrix core instead of expanding the model back to a wider format before the main compute. At runtime, the path has three connected stages: parse the checkpoint metadata, quantize live activations into the same 1x32 block-scaled format, then dispatch GEMM and MoE kernels that use gfx950 scaled MFMA.

ATOM maps per_group quantization to QuantType.per_1x32. The same parser recognizes fp4, float4, and mxfp4 as fp4x2, then forces MXFP4 into per_1x32 block scaling when needed.

This matches the model config: FP4 values with group size 32 and E8M0 scales.

For weight tensors, the FP4 payload and E8M0 scales are already stored in the checkpoint. Activations are different: they are produced live in BF16 by the surrounding transformer layers, so ATOM dynamically quantizes activation blocks at inference time.

For the MXFP4 linear path, ATOM calls get_hip_quant(QuantType.per_1x32), views activations as torch.float4_e2m1fn_x2, views scales as torch.float8_e8m0fnu, and passes both into the low-precision GEMM path. Conceptually:

This is the activation-side counterpart to the checkpoint’s packed weight and weight_scale tensors.

This is the key gfx950 instruction family for the FP4 path: the kernel is not simply converting FP4 weights to BF16 and using BF16 MFMA; it is using the scaled MFMA form that accepts low-precision block-scaled operands.

ATOM’s Mxfp4MoEMethod creates FP4 expert weights and uint8 scale tensors with a hard-coded MXFP4 block size of 32. For MoE execution, it passes:

quant_type=self.quant_type

w1_scale=layer.w13_weight_scale

w2_scale=layer.w2_weight_scale

into AITER fused_moe.

Together, these linear and MoE paths show the core benefit of MXFP4 on gfx950: FP4 payloads and 1x32 scaling can be carried into the compute kernels instead of serving only as a compressed checkpoint format.

The scaled-MFMA claim was verified by disassembling representative gfx950 FP4 code objects used by the ATOM/AITER stack. The FP4 GEMM kernels contain v_mfma_scale_f32_16x16x128_f8f6f4 instructions, for example:

Representative f4gemm_bf16_per1x32Fp4_* code objects contained hundreds of v_mfma_scale_f32_16x16x128_f8f6f4 instructions. The MXFP4 MoE code objects also contained the same scaled MFMA family, with typical counts in the hundreds per code object. This confirms that the serving path is not merely loading a compressed FP4 checkpoint and expanding it to BF16 before the matrix core; the gfx950 kernels really do execute the block-scaled MFMA path.

mla_a8w8_qh16_qseqlen1_gqaratio16_ps.co for MLA decode

BF16 GEMM fallback kernels such as bf16gemm_fp32bf16_tn_80x64_splitk_clean.co

This is expected. The model is not “FP4 everywhere.” The MXFP4 path is concentrated in quantized linear/MoE compute, while attention, routing, normalization, logits, and other framework components still use their own precision choices.

We benchmarked the ATOM OpenAI-compatible server using the same input/output shape used in the previous Kimi-K2.5 ROCm blog posts:

random_range_ratio = 1.0

This keeps the benchmark shape aligned with the earlier MI300X and MI325X posts. The hardware and serving stack are different, so the numbers below should be read as the MI355X/gfx950 MXFP4 result under the same workload shape, not as a same-system speedup over W4A16, W4A8, or W8A8.

amd/Kimi-K2.5-MXFP4

The following command shows the serving setup used for the 8-GPU run. Replace /models/Kimi-K2.5-MXFP4 with the location where the MXFP4 checkpoint is mounted inside the container.

We use the same random input/output lengths as the earlier Kimi-K2.5 posts: 10,240 input tokens and 512 output tokens.

The low-concurrency case highlights first-token latency and per-token decode cost when only two requests are active. The high-concurrency case stresses parallel decoding, where aggregate output throughput rises to 1,055.95 tok/s while mean TPOT increases to 29.34 ms/token.

These are end-to-end serving measurements, not isolated GEMM or MoE microbenchmarks. They include request scheduling, attention, MoE, sampling, KV cache management, and framework overhead around the MXFP4 kernels.

The strongest claim supported by this experiment is that Kimi-K2.5-MXFP4 can be served end-to-end on 8x MI355X with ATOM while exercising gfx950 scaled-MFMA FP4 kernels in the matrix-heavy paths. The benchmark includes the full serving stack, so it is more relevant to deployment than an isolated GEMM microbenchmark.

The comparison below is useful for context, but it is not a controlled same-system speedup experiment. The previous W4A16, W4A8, and W8A8 rows were measured on MI325X with SGLang/AITER, while the MXFP4 row was measured on MI355X with ATOM/AITER. That means several variables change at once: hardware generation, quantization format, serving framework, kernel dispatch path, and runtime scheduler behavior.

For this reason, the MXFP4 numbers should be read as a gfx950 enablement result and a method-aligned workload comparison. A strict speedup claim would require a same-system baseline, such as another Kimi-K2.5 format on the same 8x MI355X ATOM stack, or an MXFP4 fallback path on the same hardware.

Results may vary based on system configuration, software versions, model settings, and workload characteristics.

The table below mirrors the benchmark comparison style from the W4A8/W8A8 post, but the interpretation is different. The W4A16, W4A8, and W8A8 rows were reported on an 8x MI325X SGLang/AITER setup; the MXFP4 row here was measured on an 8x MI355X ATOM/AITER setup. The workload shape and GPU count are aligned, but the hardware generation, quantization format, and serving stack are not identical, so this is context rather than a same-system speedup claim.

The comparison highlights why MXFP4 is interesting on gfx950: it combines the small payload width of 4-bit weights with a scale layout that the hardware can consume directly in the FP4 matrix path. That is the key distinction from both W4A8, which changes scale granularity to make INT8 MFMA efficient, and W8A8, which uses a wider FP8 payload to keep the compute pipeline simple.

When Should You Choose MXFP4?#

MXFP4 is attractive when the deployment target has gfx950 block-scaled MFMA support and the model is available in a checkpoint layout that preserves FP4 payloads plus E8M0 scales. In that setting, MXFP4 gives the serving stack a compact 4-bit payload without forcing per-32 scales to be baked into BF16 before compute or delayed into a single epilogue scale.

In practice, MXFP4 is a good candidate when weight bandwidth, MoE expert traffic, or model footprint are important constraints, and when the workload can tolerate FP4 element precision with E8M0 block scaling. FP8 remains a simpler choice when maximum precision margin or broader kernel availability is more important than the smallest payload. W4A8 remains useful for gfx942 paths where INT8 MFMA is the available high-throughput option.

To validate that the ATOM MXFP4 path preserves model quality on a downstream reasoning benchmark, we ran GSM8K 10-shot evaluation against the same ATOM server. This matches the GSM8K protocol used in the previous W4A8/W8A8 post. The run used greedy generation (do_sample=False, temperature=0.0) and evaluated the full GSM8K test set of 1,319 samples.

The strict-match filter follows the canonical GSM8K answer marker, while the flexible-extract filter extracts the final numeric answer more permissively. Reporting both filters makes the result easier to compare with the earlier Kimi-K2.5 posts.

The result shows that the MXFP4 serving path maintains strong downstream reasoning quality while using the lower-precision FP4 payload and gfx950 block-scaled compute path.

This third Kimi-K2.5 optimization note extends the earlier MI300X/MI325X story to MI355X and gfx950 by moving from kernel and quantization work to optimized serving:

ATOM makes the optimized path deployable: users can run a standalone OpenAI-compatible server, while the vLLM OOT plugin path lets vLLM remain the serving framework and ATOM provide AMD-optimized model execution.

The checkpoint is truly MXFP4: the model configuration records FP4 tensors with group_size=32 and E8M0 scales, and stores packed FP4 payloads plus weight_scale tensors for quantized modules.

ATOM parses MXFP4 into a 1x32 FP4 execution path: per_group maps to QuantType.per_1x32, and MXFP4 is represented as torch.float4_e2m1fn_x2 plus torch.float8_e8m0fnu scales at runtime.

AITER dispatches gfx950 FP4 kernels: the serving path uses f4gemm kernels and FlyDSL FP4 fused MoE kernels built around gfx950 scaled MFMA support.

End-to-end serving is strong on 8x MI355X: under the aligned 10,240/512 workload, MXFP4 reached 165.30 output tok/s at concurrency 2 and 1,055.95 output tok/s at concurrency 40.

Accuracy is maintained on GSM8K: the same ATOM server scored 94.01% flexible-extract exact match on GSM8K 10-shot (0.9401 +/- 0.0065) on the full 1,319-sample test set.

The broader takeaway is that gfx950 turns MXFP4 into a first-class compute format, and ATOM turns that hardware capability into a serving path. Compared with FP8, MXFP4 reduces data movement through 4-bit payloads while preserving dynamic range through block scales; compared with older FP4-as-storage approaches, gfx950 scaled MFMA lets the kernel keep the low-precision representation much closer to the matrix core. ATOM is the framework layer that makes this usable: it loads the HuggingFace checkpoint, maps quantization metadata to runtime operators, chooses the AITER/FlyDSL kernels, exposes a serving API, and provides benchmark tooling for the full stack rather than isolated kernels.

[1] Kimi-K2.5 on HuggingFace

[2] Kimi-K2.5-MXFP4 on HuggingFace

[8] Previous ROCm blog: Accelerating Kimi-K2.5 on AMD Instinct MI300X: Optimizing Fused MoE with FlyDSL

[9] Previous ROCm blog: Further Accelerating Kimi-K2.5 on AMD Instinct MI325X: W4A8 and W8A8 Quantization with AMD Quark

The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions, and typographical errors. The information contained herein is subject to change and may be rendered inaccurate for many reasons, including but not limited to product and roadmap changes, component and motherboard version changes, new model and/or product releases, product differences between differing manufacturers, software changes, BIOS flashes, firmware upgrades, or the like. Any computer system has risks of security vulnerabilities that cannot be completely prevented or mitigated. AMD assumes no obligation to update or otherwise correct or revise this information. However, AMD reserves the right to revise this information and to make changes from time to time to the content hereof without obligation of AMD to notify any person of such revisions or changes. THIS INFORMATION IS PROVIDED ‘AS IS.” AMD MAKES NO REPRESENTATIONS OR WARRANTIES WITH RESPECT TO THE CONTENTS HEREOF AND ASSUMES NO RESPONSIBILITY FOR ANY INACCURACIES, ERRORS, OR OMISSIONS THAT MAY APPEAR IN THIS INFORMATION. AMD SPECIFICALLY DISCLAIMS ANY IMPLIED WARRANTIES OF NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR ANY PARTICULAR PURPOSE. IN NO EVENT WILL AMD BE LIABLE TO ANY PERSON FOR ANY RELIANCE, DIRECT, INDIRECT, SPECIAL, OR OTHER CONSEQUENTIAL DAMAGES ARISING FROM THE USE OF ANY INFORMATION CONTAINED HEREIN, EVEN IF AMD IS EXPRESSLY ADVISED OF THE POSSIBILITY OF SUCH DAMAGES. AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies. © 2026 Advanced Micro Devices, Inc. All rights reserved

Full text under the publisher's open licence · 4006 words

aiqbench analysis

Serve Kimi-K2.5-MXFP4 on MI355X with ATOM

Related coverage