Zyphra Inference Wave 2：在 MI355X 上加速 MoE 与长上下文服务

Speculative decoding. At low concurrency, the time for a decode step grows slowly with the number of tokens in the step. For GLM-5.3 in the vLLM baseline with MTP disabled, a step of 1 token takes 17.7 ms and a step with 32 tokens (one for each of 32 requests) takes 27.0 ms (Table A7). The step reads each weight it needs from HBM once and applies it to every token in the step, and it runs the same sequence of kernels and all-reduces for any number of tokens. Speculative decoding (Leviathan et al.; Chen et al.) uses this property. A small drafter proposes the next few tokens. The full model (the target) then processes the current token and all proposed tokens in one forward pass. The target checks the proposed tokens in order with a rejection-sampling rule. It keeps the tokens before the first rejected token and then samples one token of its own. This rule keeps the output distribution identical to decoding with the target alone. The speedup depends on how many proposed tokens the target accepts, and the acceptance rate depends on the prompt (see the Decode section).

The methods differ in how they draft. Multi-Token Prediction (MTP) adds a transformer layer that is trained with the model. It takes the target’s last hidden state and the embedding of the newest token. It uses the target’s LM head and proposes one token per drafter pass. EAGLE trains a separate small head on the target’s hidden states. It proposes a tree of candidate continuations, and the target checks all branches in one pass. DSpark proposes a block of tokens in one drafter pass, and a small sequential head conditions each proposed token on the tokens before it. A confidence head estimates how likely each proposed token is to be accepted. The scheduler uses these estimates and the server load to set how many proposed tokens each request sends to the target. We use MTP for GLM-5.3 and DeepSeek-V4 and DSpark for Kimi-K3 (for each model this is the method with the highest decode throughput among those we tested on MI355X).

The speedup decreases as concurrency increases (2.7x at 1 concurrent request and 2.0x at 32 for GLM-5.3 in the vLLM baseline, Table A7). A straight line through the two MTP-disabled step times gives about 17.4 ms per step plus 0.3 ms per token ((27.0 − 17.7) / 31). At 1 request each proposed token adds about 2% to the step time. At 32 requests each proposed position adds 32 tokens (one per request) and about 36% to the step time.

Quantization. Quantization stores weights in fewer bits. Fewer bits mean less memory for the weights and fewer bytes read from HBM in each forward pass, and on MI355X the FP4 MFMA peak is twice the FP8 peak (Table 1). The cost is lower precision, so each quantized model has to be checked for lost accuracy. MXFP4 is one of the microscaling (MX) formats in the Open Compute Project MX specification (Rouhani et al.). Each block of 32 values shares one 8-bit power-of-two scale, and each value is a 4-bit float with one of eight magnitudes (0, 0.5, 1, 1.5, 2, 3, 4, or 6). This is 4.25 bits per value, about half the size of FP8. The MI355X scaled MFMA instruction takes MXFP4 values and their block scales as inputs. We evaluate the MXFP4 quantization of GLM-5.3’s experts that AMD produced with its Quark toolkit, for retained reasoning ability and for decode throughput.

Kernel Results

The kernel benchmarks use the shapes of Kimi-K2.6, DeepSeek-V3.2, and GLM-5.1. The serving results use the models we serve now. GLM-5.1 and GLM-5.3 have the same MoE and indexer shapes, so the GLM kernel results carry over directly. 

Our fused MoE baselines are the AMD AITER kernels for FP8 and the vLLM Triton WNA16 kernel for a16w4. A dense BF16 PyTorch implementation serves as our correctness reference.

Figure 10. Fused MoE kernel speedup over the baseline kernel as a function of M, the number of tokens in one forward pass. At decode M is the number of tokens decoded in one step[^8], and at prefill it is the number of prompt tokens in the pass. Values above 1 mean our kernel is faster[^9]. Left: Kimi-K2.6 a16w4 (H=7168, E=384, top-k=8) against the vLLM Triton WNA16 kernel. Center: DeepSeek-V3.2 FP8 (H=7168, E=256, top-k=8) against AITER. In these two panels Full-N uses the full expert width and TP8 uses one eighth of it, which is the shape each GPU runs under 8-way tensor parallelism. Right: ZAYA1 FP8 (E=24, top-k=1) against AITER. All timings use HIP graphs, as in production.

Our kernel is faster than the baselines at most values of M up to 256. Small values of M correspond to decode steps and short prefill chunks. On Kimi-K2.6 a16w4 the baseline is a Triton kernel, and our kernel is faster at every M except M = 1 at full expert width, by up to 2.0x. On the FP8 shapes AITER is faster at higher M. The crossover is between M = 512 and M = 1,024 at full expert width and between M = 256 and M = 512 under TP8. We designed our kernel for decode, where each expert receives a few token rows (M × k / E is 8 or less for M up to 256 on DeepSeek-V3.2). ZOE selects the MoE kernel by M from these measurements. It runs our kernel below the crossover and AITER above it.

We also measured several variants that ran slower than the kernel described in the Optimizations section. We list them so that others can skip them. A single kernel that fuses both stages produced correct output and ran 1.7x to 2x slower for M ≥ 256 (the fused kernel needs the registers of both stages at once and spills the excess to scratch memory). Reordering the grid so that the most-used experts run first ran slower at every M. That reordering aimed to reuse expert weights from L2, and it lost the reuse of activations from L2 that the default column-major order provides. For the MX formats, an epilogue that sums each token’s k expert outputs with atomic adds ran slower than our default epilogue. Three more variants also ran slower. These were wider M-row weight reuse, an LDS cache for activations at M = 1 for a16w4, and the AITER-style INT4 weight preshuffle.

Figure 11. Lightning indexer throughput in TFLOP/s on MI355X with a causal window, for our kernel and two AITER baselines. AITER Gluon is AITER’s kernel written in Gluon, a lower-level language in the Triton project. AITER (best per shape) is the fastest of the other AITER kernels on each shape. H is the number of indexer heads and D is the head dimension. Shapes are query rows × KV positions. A square shape has as many query rows as KV positions (prefill of one prompt). A skinny shape has 2,048 or 4,096 query rows and 16K to 64K KV positions (prefill of a short chunk after a long cached prefix). The label above each of our bars is our throughput divided by AITER Gluon’s.

Our kernel is 1.9x to 2.4x faster than AITER (best per shape) on the square shapes and 1.15x to 2.1x faster on the skinny shapes. AITER’s throughput stays between 445 and 492 TFLOP/s at H=32 and between 670 and 782 TFLOP/s at H=64 on every shape. These shapes range from 2,048 to 65,536 query rows. Their KV arrays range from 0.5 MiB, which fits in the 4 MiB L2 of each XCD, to 8 MiB, which does not. AITER’s throughput changes little across these shapes, so neither memory nor the number of rows limits it. Its speed is set by the work it does for each block of KV positions.

On the square shapes our kernel is 1.07x to 1.25x faster than AITER Gluon at H=32 and runs at 0.98x to 1.11x of its throughput at H=64. Gluon lets the kernel author set data layouts, allocate LDS, and issue asynchronous copies. AITER’s Gluon kernel uses these to copy KV tiles into LDS while it computes scores with MFMA instructions, which is the same structure as our kernel (Optimizations section). Moving the causal bounds check into the buffer-load descriptor raised our throughput by about 10% on every shape.

AITER Gluon is 1.12x to 1.37x faster than our kernel on the skinny shapes with 2,048 query rows and 32K or 64K KV positions. With 2,048 rows and 16K positions the two kernels are within 3%. Our kernel divides the work across the GPU by query row (one wave per row), so the number of rows sets how much work runs in parallel. At 64K positions and H=32, going from 2,048 to 4,096 rows raises our throughput from 630 to 1,006 TFLOP/s, and our kernel becomes 1.37x faster than Gluon. Each row does about the same work in both shapes, so at 2,048 rows our kernel leaves part of the GPU unused. Gluon’s throughput rises by 4% for the same change. At H=64 the change raises our throughput from 1,045 to 1,176 TFLOP/s, and Gluon stays 1.10x faster. We fall back to the Gluon kernel on most of the skinny shapes.

Two extensions target these shapes. A BLOCK_Q > 1 variant groups several query rows and reads the union of their KV windows once, which reduces KV reads by up to a factor of BLOCK_Q. This variant is implemented and under evaluation. The prefetch pipeline currently keeps two tiles in flight, and we can increase it to three or four.

Figure 12. DeepSeek-V4 sparse attention kernels on MI355X against the SGLang baseline. Left: top-K selection kernel time for one request. Center: end-to-end single-user decode throughput on DeepSeek-V4-Pro with only our top-K kernel added. Right: dual-scope decode kernel time per call under HIP graphs, against SGLang’s Triton kernel.

Table 4. Additional DeepSeek-V4-Pro measurements on 8x MI355X (TP=8). Speedups are relative to the same configuration without the listed change.

The top-K selection scores every compressed position, so its work grows linearly with context. The SGLang kernel runs the selection for one request on one compute unit, and our kernel splits long contexts across up to 32 compute units, so our speedup grows with context (Figure 12, left). At 4K tokens our kernel is slightly slower. At 1M tokens it is 11.1x faster for one request and 9.9x faster for a batch of 32 (Table 4).

In end-to-end decode the selection is one of many steps, so its effect follows its share of the step time. That share grows with context. Adding only the top-K kernel raises DeepSeek-V4-Pro single-user throughput by 1.40x at 1M tokens, while at 32K and 64K the change is within run-to-run noise (Figure 12, center).

The dual-scope decode kernel splits each call across the GPU, because single-user decode has little work per call (Optimizations section). It is 1.30x faster per call than SGLang’s Triton kernel at batch 1. As the batch grows each call has more work, and the advantage shrinks until our kernel is slightly slower at batch 32 (Figure 12, right). The kernel raises end-to-end decode throughput by about 5% at batch 1 to 8, and MTP on top of both kernels adds 4% to 15% at 256K to 1M tokens (Table 4).

Serving Results

Microbenchmarks show whether an individual kernel is faster. Our customers see end-to-end prefill TTFT and decode throughput in tokens per second, which also depend on scheduling, parallelism, caching, and the rest of the serving stack.

To measure these, we use fixed prefill and decode benchmarks with a fixed input sequence length (ISL), output sequence length (OSL), and request concurrency. These benchmarks do not represent production traffic. They isolate the effect of our changes across models and serving recipes. Our benchmarking scripts are modeled on the vllm bench and sglang bench tools. We wrote our own so we could control prefix-cache hit rates in the prefill benchmark and use identical settings across both inference engines.

Prefill

Our prefill benchmark runs a single request at a time and generates a single output token. It sweeps input sequence length and records TTFT averaged over 10 iterations after 1 warmup iteration. We also sweep prefix-cache hit rates that correspond to different types of workload. We use 0% (no cached prefix), 60% to represent common chat, and 90% to represent agentic workloads.

Figure 13. Kimi-K3 single-request TTFT on MI355X with the best SGLang configuration we found. Random input, OSL=1, 1 concurrent request, mean ± std over 10 runs.

Figure 13 shows how prefix caching changes TTFT for Kimi-K3. Without cache hits, TTFT grows faster than the input length at long context (about 2.5x per doubling above 128K tokens), because in Kimi-K3’s full-attention layers each token attends over all earlier tokens. A cached prefix removes the per-token work for the cached tokens, but each new token still attends over the full context, cached tokens included. With 90% of the prompt cached, per-token work drops 10x while the work in those attention layers drops only about 5x[^10], and the fixed costs of each request do not drop at all. The TTFT reduction is therefore largest at medium lengths (8x at 128K tokens). It is smaller at short context, where fixed costs dominate (2.6x at 8K), and at long context, where attention dominates (6x at 512K). 

For GLM-5.3 we compare vllm/vllm-openai-rocm:v0.29.0 out of the box against the same image with ZOE, which adds our DSA lightning indexer and fused MoE kernels.

Figure 14. GLM-5.3 single-request TTFT on MI355X, vLLM v0.29.0 baseline against the same image with ZOE. Random input, OSL=1, 1 concurrent request, mean ± std over 10 runs. ZOE always matches or outperforms the vLLM implementation and pulls away strongly at long context lengths.

ZOE adds two kernels to GLM-5.3, the lightning indexer and the fused MoE kernel. In a DSA model the indexer is the only part of prefill whose work per token grows with the context. It scores each new token against every earlier position, while main attention reads a fixed 2,048 entries per token and the MoE and projection layers do the same work for every token. At 16K tokens and below the indexer is a small part of prefill and the speedup is close to 1x, which also shows that the MoE kernel changes prefill time little. As the context grows the indexer takes most of the prefill time, and the speedup grows with it, to 2.8x at 512K tokens. Prefix-cache hits raise the speedup further. Cached tokens skip the per-token work, but each new token is still scored against the full context, so the indexer is a larger share of the work that remains (3.6x at 512K with a 90% hit rate). The one result below 1x (0.80x at 32K with a 90% hit rate) comes from high run-to-run variance in the ZOE runs (a standard deviation of 0.16 s against a mean of 0.30 s).

Decode

Our decode benchmark sweeps request concurrency. Low concurrency produces the best per-user throughput, and high concurrency produces the best per-node throughput. Each decode step reads the weights once for every request in it, so more requests share each read, but each step also takes longer. We exclude TTFT and report only per-user and per-node throughput. Decode performance depends substantially on speculative decoding and is thus sensitive to the prompt/output distribution. Our benchmarks use temperature=1 and top_p=0.95, which can produce large variance between runs with the same prompt because the speculator acceptance rate changes. To show this sensitivity we report Kimi-K3 results for two prompts, one from Apex-Shortlist (heavy reasoning) and one from GSM8K (simple math).

Apex. "Determine the sum of the smallest 50 possible values of \(gcd(a2+b2+c2,abc)\) across all triples of positive integers \((a,b,c)\). Let's think step by step and output the final answer within \boxed{}."

GSM8K. "Tobias is buying a new pair of shoes that costs $95. He has been saving up his money each month for the past three months. He gets a $5 allowance a month. He also mows lawns and shovels driveways. He charges $15 to mow a lawn and $7 to shovel. After buying the shoes, he has $15 in change. If he mows 4 lawns, how many driveways did he shovel?"

Figure 15. Kimi-K3 decode throughput on MI355X with the best SGLang configuration we found (DSpark included), OSL=1024, mean ± std over 10 runs.

The model and kernels are the same for both prompts, so the gap between the two curves comes from DSpark acceptance. The target accepts more of the drafter’s proposed tokens on the GSM8K answer than on the Apex reasoning, so each step produces more tokens. Speculative decoding results therefore have to be read together with the prompt distribution.

For GLM-5.3 we report two settings. With MTP disabled the benchmark has close to zero variance, which isolates the throughput change from the kernels. With MTP enabled the configuration matches what we would use in production, and the variance from speculative decoding is larger than most of the kernel gains.

Figure 16. GLM-5.3 per-node decode throughput on MI355X, vLLM v0.29.0 baseline against the same image with ZOE. Apex prompt, OSL=1024, mean of 100 runs. Error bars show one standard deviation.

ZOE changes GLM-5.3 decode less than prefill. The Apex runs generate 1,024 tokens from a short prompt, so the context stays short and the indexer does little work. The remaining change is the fused MoE kernel, which speeds up only the MoE part of each decode step. With MTP disabled ZOE raises throughput by 1% to 7%. With MTP enabled the changes range from −3% to +13%, and most are within one standard deviation of the run-to-run variation. At short context, speculative decoding changes decode throughput far more than the kernels in this post (MTP raises GLM-5.3 throughput by 2x to 2.7x), and the long-context kernels matter once the context is long (Figure 12). 

Quantization

Kimi-K3 already ships with MXFP4 experts, which leaves little room for further reduction, so we focused on GLM-5.3, which ships with FP8 weights. AMD has released amd/GLM-5.3-Quark-MXFP4-AttnFP8, which stores the routed and shared experts in MXFP4 and keeps attention in FP8. This reduces the model size from 756 GB to 410 GB. The Hugging Face model card reports a GSM8K score of 91.28 against 91.89 for the FP8 model (99.3% recovery). We have found that GSM8K alone does not show whether a model keeps its reasoning ability through post-training quantization, so we ran both models through our own eval harness on four reasoning benchmarks (Table 5).

Table 5. GLM-5.3 quantization eval comparison with the Zyphra internal eval harness. 8 repeats per prompt, max-tokens 131072, temperature=1.0, top_p=0.95. Change is in percentage points. Image: vllm/vllm-openai-rocm:nightly-rocm100-9ea8f3ffc354901b740f0b31988900897b7221d7.

The MXFP4 model stays within 3 points of the FP8 model on all four benchmarks. Each benchmark has 30 problems (47 for Apex-Shortlist) sampled 8 times, so one problem that goes from always solved to never solved moves the score by 3.3 points (2.1 points for Apex-Shortlist). Each difference in Table 5 is smaller than that.

For serving throughput we disable speculative decoding and use our fixed decode benchmark. At 410 GB the MXFP4 model fits on four MI355X GPUs (1,152 GB of HBM) with room left for the KV cache, so we also tried two TP=4 replicas on the node (DP=2). These runs use the nightly image listed under Table 5, so the absolute numbers are not directly comparable to the v0.29.0 numbers in the Decode section.

Figure 17. GLM-5.3 per-node decode throughput on MI355X with FP8 experts, with and without our fused MoE kernel, and with MXFP4 experts. All runs use the best vLLM configuration we found and the vLLM ROCm nightly image listed under Table 5. The hatched bars split the node into two TP=4 replicas. MTP disabled, Apex prompt, OSL=1024, mean of 10 runs. 

As we described in our roofline experiments, quantization to FP8 or FP4 both reduces the HBM transfer time and increases the possible compute-bound ceiling. There are implementation caveats to both of these effects. For HBM transfer time, achievable HBM bandwidth depends on the size of the buffer being transferred, and requires a large enough buffer to saturate bandwidth[^11]. For compute speeds, the relative achievable throughput tends to decrease as one steps down the ladder to FP8 and FP4, despite the absolute throughput ceiling increasing. This is due to a few implementation-specific reasons:

• FP8 and FP4 quantization isn’t simple truncation and rounding like between BF16/FP16 and FP32. It requires quantization operations[^12] (often fused into the core op kernels) that have overhead. Those overheads take up a larger proportion of the overall operation because compute and data transfer times have decreased, and we significantly push up the required problem size to hide these operations.

• Overlap is much more fine-grained and difficult to achieve, since compute latency has decreased and we have added more quantization ops that don’t run on matrix cores and have strict ordering.

• Shapes and occupancy often introduce more constraints for FP8 and FP4 matrix cores, namely the divisor of your model dims are larger so there are fewer models that can be cleanly divided, and those models are much larger. Take for example MI355X, whose FP8 and FP4 matrix cores have a 4X increase in depth in the K GEMM dimension compared to BF16. Take the 16x16 MFMA instructions as an example, where we have a K depth of 128 for FP8 and FP4, but only 32 for BF16. That means, for example, that head dims of 96 or 192 which were divisible in BF16 now aren’t, which requires custom handling (e.g. custom padding or overlapping multiple workstreams) in the FP4/FP8 kernels to achieve speed-of-light performance. This kind of doubling or halving shows up a lot, like how low-precision kernels use larger tiles to keep the faster matrix cores fed, so a given GEMM splits into fewer tiles and fewer waves, and the partial last wave is a larger fraction of the overall run time. This wasn’t such a big deal for older GPUs like Hopper because it was a small increase that still divides basically anything, 16 for BF16 and 32 for FP8, but for FP8/FP4 on MI355X, FP4 on Blackwell, etc, this is no longer the case.

In short, MFMA gets significantly faster but most other hardware operations don’t get proportionally faster, which is why, for example, Flash Attention 4 spends a lot of its design time on removing exp2f() as a bottleneck, and kernel developers are moving to fused epilogues (e.g. QuACK in CuTe DSL) and megakernels to better hide non-GEMM work. Hardware continues to increase low-precision FLOPs much faster than all memory bandwidths. Speculative decoding and diffusion language models both process several tokens per weight read, which raises the arithmetic intensity of decoding and moves it toward the compute-bound region. For BF16 this is straightforward, since we get the ratio of achieved to advertised TFLOP/s that people are used to, and virtually all model shapes benefit. If you put quantization on top of anything memory-bandwidth-bound, you can achieve speedup assuming that quantization overhead is hidden (a linear speedup if your quantized buffer is still past the bandwidth saturation point). If you put quantization on top of anything compute-bound, you take on the implementation complexity described above, which can usually be overcome.

From here, we are equipped to understand our quantized results. Quantization speeds up decoding largely by reducing the bytes read from HBM, so its benefit depends on how much of each decode step goes to reading weights. In an MoE model at concurrency 1 each token uses only its Top-K routed experts plus the number of shared experts per layer, so only these are transferred from HBM. For example, at 1 concurrent request each GLM-5.3 MoE layer uses 9 of its 257 experts (8 routed and 1 shared). Over GLM-5.3’s 75 MoE layers these experts hold 25.5 GB in FP8, which the node’s 8 GPUs read in about 0.5 ms at the measured HBM bandwidth (Table 1). The decode step takes about 16 ms (Figure 17), and MXFP4 halves the 0.5 ms.  Most of the step goes to work that the expert format does not change, such as attention, communication between GPUs, and the fixed cost of each kernel. Halving the expert bytes can shorten such a step by only about 2%. As concurrency grows, the tokens in a step use more distinct experts, the step reads more of the expert weights, and the bytes that quantization removes become a larger share of the end-to-end time.

The MXFP4 experts run the same operations as the FP8 experts (quantize the activations, run the two expert GEMMs, and combine the results), with half the weight bytes and twice the matrix-core peak. Slower MXFP4 compute time therefore comes down to kernel implementation and how well one optimizes around the challenges described above. Our Fused MoE implementation’s FP4 support (in particular for low-batch) is ongoing work. In vLLM the MXFP4 model is slower than the FP8 model at 1 and 8 concurrent requests and faster at 32 (Figure 17). Specifically at low concurrency the MXFP4 kernels add more time per layer than the smaller weights save, and at 32 requests the bytes saved outweigh that cost. Splitting the node into two TP=4 replicas was slower at every concurrency (each GPU in a TP=4 replica reads twice as many weight bytes per step). With our fused MoE kernel the FP8 model stays faster than the MXFP4 model at low concurrency and within about 1% at 32 requests, so we continue to serve the FP8 model in production in low-concurrency cases.

Agentic Performance Benchmarks

Fixed prefill and decode benchmarks do not reliably indicate performance in production. Real traffic is multi-turn and long-context, and it includes sub-agent bursts, KV cache offloading, and tool calling. At Zyphra we use the AgentX benchmark from SemiAnalysis, which is an open-source, multi-turn agentic coding inference benchmark at 1 million tokens of context. It is available in the InferenceX repository, which describes it in more detail. We will publish agentic benchmark results soon.

Future Work

This post presents our second round of full-stack inference optimization on Zyphra Inference. We will continue to improve our inference platform with a focus on the following areas: 

• Closing the large-M gap on the FP8 frontier MoE shapes, so that prefill no longer falls back to vendor kernels through the hybrid dispatcher.

• Finishing the BLOCK_Q > 1 lightning indexer variant and a deeper prefetch pipeline for the skinny shapes.

• MXFP4 MoE kernels that are faster at low concurrency, and MXFP4 GEMMs that reach a larger fraction of the MI355X FP4 peak (Figure 5), so that the MXFP4 model decodes faster than the FP8 model at every concurrency.

• Quantization recipes informed by our pretraining experience, and custom speculators that take advantage of the MI355X's HBM capacity.

• Agentic benchmark results with AgentX at 1 million tokens of context.

Zyphra Inference is available now at cloud.zyphra.com.

Methodology and Limitations

• As in our first post, we describe our methodology and its limits, since there are many regimes in inference and we do not want to make apples-to-oranges comparisons. The limitations described at the top of our previous blog apply here as well. We restate the relevant ones below, along with those specific to our kernel benchmarking.

• All results are taken on a single node with no live traffic. Requests arrive at fixed intervals instead of randomly, and we do not test scheduling in real settings because it is outside the focus of this work.

• The performance of speculative decoding methods is sensitive to the prompts used in the benchmark, because draft acceptance rates vary across input distributions. A proposed token is accepted with probability equal to the overlap between the drafter’s and the target’s next-token distributions, and drafters match the target more closely on predictable text, such as arithmetic or code, than on open-ended text. Reading the target’s hidden states, as MTP and EAGLE do, improves the match but does not remove this dependence. EAGLE-2 averages between 3.5 and 4.9 tokens per target pass across six benchmarks at temperature 1 (Li et al.), and Figure 15 shows the same effect for DSpark.

• We hold input prompts fixed across configurations so that KV and prefix cache behavior is as similar as possible, and we report decode results for more than one prompt where the sensitivity is large.

• Kernel benchmarks flush the L2 cache with a 256 MB buffer between iterations and report medians over 200 iterations after 20 warmup iterations.

• We report timings under HIP-graph capture and replay, which we believe reflects production behavior, since per-launch (eager) timing includes launch overhead that graphs remove. Speedup is median(baseline) / median(ours).

• Both our inference software stack and the external baselines change quickly, so any results we show here are valid only for a point in time. This point in time may differ across models depending on when the experiments were run. These blogs are sliding windows that summarize the work done since the previous blog. The serving results use vLLM v0.29.0 and the SGLang ROCm image v0.5.19 (lmsysorg/sglang-rocm:v0.5.19-rocm720-mi35x-20260910), and the quantization results use a vLLM ROCm nightly image.

Appendix A: Data Tables

The tables below contain the values plotted in the figures.

[wide]A1. Fused MoE speedup (Figure 10)

[wide]Kimi-K2.6 a16w4 (H=7168, E=384, top-k=8), speedup vs vLLM Triton WNA16 under HIP graphs.

[wide]DeepSeek-V3.2 FP8 (H=7168, E=256, topk=8), speedup vs AITER under HIP graphs.

[wide]Zyphra ZAYA1 FP8 (E=24, topk=1), speedup vs AITER.

[wide]A2. Lightning indexer throughput (Figure 11)

[wide]TFLOP/s, causal window, MI355X. Ratios above 1.00x mean our kernel is faster. H=32, D=128

[wide]H=64, D=128

[wide]A3. DeepSeek-V4 sparse attention (Figure 12)

[wide]Top-K kernel time, single request, MI355X.

[wide]End-to-end single-user decode, DeepSeek-V4-Pro, SGLang, tok/s.

[wide]Dual-scope per-call decode latency under HIP graphs, vs SGLang's production Triton path.

[wide]A4. Kimi-K3 prefill TTFT (Figure 13)

[wide]MI355X, Kimi-K3, best SGLang config, random input, OSL=1, fixed prefix to simulate cache hit rate, with warmup iteration, 1 concurrent request. TTFT in seconds.

[wide]A5. GLM-5.3 prefill TTFT (Figure 14)

[wide]MI355X, GLM-5.3, best vLLM config (v0.29.0), random input, OSL=1, fixed prefix to simulate cache hit rate, with warmup iteration, 1 concurrent request. Baseline against ZOE (DSA indexer and fused MoE). TTFT in seconds.

[wide]A6. Kimi-K3 decode throughput (Figure 15)

[wide]MI355X, Kimi-K3, best SGLang config (DSpark included), OSL=1024, average of 10 runs, tok/s.

[wide]A7. GLM-5.3 decode throughput (Figure 16)

[wide]MI355X, GLM-5.3, best vLLM config (v0.29.0), Apex prompt, OSL=1024, DSA + fused MoE, average of 100 runs, per-node tok/s.

[wide]MTP disabled

[wide]MTP enabled (mean ± std)

[wide]A8. GLM-5.3 quantization throughput (Figure 17)

[wide]MI355X, GLM-5.3, best vLLM config (MTP disabled), Apex prompt, OSL=1024, average of 10 runs, per-node tok/s. vLLM ROCm nightly image (see Table 5).

Appendix B: Roofline Measurement Setup

MI355X runs FP8 and FP4 matrix math at full rate only through the new CDNA4 f8f6f4 instructions. Our FP8 and FP4 loops use one of them, v_mfma_scale_f32_32x32x64_f8f6f4, and the BF16 loop uses v_mfma_f32_32x32x16_bf16. The upstream benchmark (rocm-amdgpu-bench) used v_mfma_f32_32x32x16_fp8_fp8, the FP8 instruction carried over from CDNA3, which runs at about half the FP8 rate on MI355X. The MFMA and GEMM benchmarks set all quantization scales to 1. Before each timed GEMM we flush L2 and run a short GPU busy-wait, so the timing events measure only the GEMM kernel. TunableOp is disabled for the FP4 GEMM runs because its scaled-GEMM dispatcher has no FP4 path. For FP8 and BF16, results with and without TunableOp differ by less than 2%. For attention decode, pa_fwd_asm requires a head dimension of 128 and supports GQA ratios up to 16. We sweep batch size (1 to 256), sequence length (1K to 64K tokens), number of KV heads (4, 8, 16), and GQA ratio (1, 4, 8, 16). All sequences in a batch have the same length (production batches contain sequences of different lengths). 

Footnotes

[^1]: For an 8K-token input, TTFT with ZOE is 0.41 s with no prefix-cache hits and 0.13 s at a 90% hit rate. The 512K results run without context parallelism, so they isolate the effect of our kernels (the indexer’s share of prefill time grows with input length, and so does the speedup from our indexer kernel). Context-parallel schemes such as Zyphra’s TSP divide the input tokens across GPUs and significantly reduce TTFT further. 

[^2]: Each expert GEMM is small, and a small GEMM has too few output tiles to occupy all of the GPU’s compute units (CUs) on its own.

[^3]: Such models are often trained to balance across the global training batch size, but during inference it is difficult to achieve similar diversity of requests on all inference replicas and across time. This means that it is very common to have wide distributions of token count per expert at a given time and on a given inference replica.

[^4]: Peak matrix compute throughput generally has grown faster than HBM bandwidth, on-chip bandwidth, and scalar or transcendental op throughput. This raises the user workload’s required arithmetic intensity to reach the compute roof, and increases the fraction of total execution time spent on non-matmul ops like softmax. Lower precisions add non-matmul quantization operations such as scale processing, packing, layout conversion, and outlier transforms. Overlapping data movement with core computation also becomes more difficult when that core computation’s latency decreases.

[^5]: A roofline model gives the highest throughput a kernel can reach on a given chip. A kernel performs some number of FLOPs and moves some number of bytes between HBM and the compute units. Its arithmetic intensity is the ratio of the two (FLOPs per byte moved). HBM bandwidth caps how fast the kernel can read data, so its throughput is at most the HBM bandwidth multiplied by its arithmetic intensity. The matrix cores cap how fast it can compute, so its throughput is also at most the MFMA peak (the peak rate of the matrix fused multiply-add instructions that matrix multiplications run on). The attainable throughput is the lower of these two limits. On a log-log plot the first limit is a memory-bandwidth line and the second is a horizontal line for the compute ceiling. The two lines meet at the ridge point. A kernel with arithmetic intensity below the ridge point is memory-bound, and its run time depends on the bytes it moves. A kernel above the ridge point is compute-bound, and its run time depends on its FLOPs.

[^6]: On MI355X we measure 6.07 TB/s of HBM bandwidth with a copy kernel and 2,363 TFLOP/s for BF16 MFMA (Table 1). The BF16 ridge point is therefore 2,363 / 6.07 = 389 FLOP/byte. The FP8 MFMA peak is about twice the BF16 peak, and the FP4 peak is about twice the FP8 peak, so the ridge point is 813 FLOP/byte for FP8 and 1,620 FLOP/byte for FP4. Lower precision also halves the bytes per element. If every operand uses the lower precision, the same kernel has twice the arithmetic intensity. A memory-bound kernel then moves up the bandwidth line (it moves half the bytes for the same FLOPs). A compute-bound kernel gains from the higher compute peak. 

[^7]: Tile quantization is the extra work from padding a matrix dimension up to a multiple of the tile size. Wave quantization is the idle time when the number of tiles is not a multiple of the number of tiles the GPU runs at once (NVIDIA, Matrix Multiplication Background).

[^8]: Without speculative decoding, M at decode equals the number of requests in the batch. Speculative decoding adds each request’s proposed tokens to the step, and a batch that mixes prefill and decode adds prompt tokens.

[^9]: Speedup is the median time of the baseline kernel divided by the median time of our kernel. Each kernel runs 200 timed iterations after 20 warmup iterations, with an L2 flush using a 256 MB buffer between iterations.

[^10]: Kimi-K3 interleaves linear-attention layers (Kimi Delta Attention), whose work per token is constant, with full-attention layers (MLA). In a full-attention layer the work over a prompt of length L grows as L²/2, since each token attends to all earlier tokens. With a fraction h of the prompt cached, only the tokens from hL to L are computed, which costs (1 − h²)L²/2. For h = 0.9 this is 19% of the full work, about a 5x reduction.

[^11]: Every data transfer, whether a read from HBM, a copy over xGMI/NVLink, or a message sent over Ethernet/InfiniBand, has fixed costs (launching, synchronizing, address translation on the first access to each page, handshakes between sender and receiver, etc). Increasing the buffer size amortizes these costs away and provides more implementation flexibility for strategies like pipelining. Most works model this as a transfer of S bytes taking (α + S/β) time to transfer, where α is the fixed cost and β the peak bandwidth. Our earlier report measured this saturation for MI300X’s HBM, xGMI, and inter-node collectives (Section III). On MI300X, for example, an HBM copy needed about 16 MB to reach full bandwidth. When one decreases the buffer size by quantizing, if you’re not still above the bandwidth saturation point, you will not achieve linear speedup in transfer time.

[^12]: E.g. Hadamard rotate the matrix to spread out outliers across dimensions so they’re not destroyed, pack/unpack into the matrix core's expected bit format if applicable, apply scales per block, transpose so that data is consecutive over the reduction dim, etc.