Fused GDN prefill produces inaccurate output on gfx1151

[randomvariable]

The fused Gated DeltaNet prefill kernel (vllm/model_executor/layers/mamba/gdn/chunk_fused.py) is materially less accurate than the two-kernel path it replaces when run on gfx1151 (RDNA3.5). Measured on real hardware, not inferred.

Evidence

Both GPU paths were compared against a float64 CPU per-token recurrence used as ground truth, so the result does not depend on a chosen tolerance:

seq_len	two-kernel max err	fused max err	ratio
63	4.61e-04	4.38e-03	~9.5x
64	4.20e-04	5.18e-03	~12x
128	5.95e-04	3.78e-02	~64x
192	7.19e-04	5.47e-02	~76x

Error grows with chunk count, so long prefills — the case the kernel exists for — are worst affected. It is not a tiling artifact: at seq_len 128, BV=16 gives 3.33e-02, BV=32 3.14e-02, BV=64 3.10e-02.

Where the error is

The recurrence itself is fine. final_state is bit-exact in every configuration, so the register-resident state carry is correct. The output is only wrong once the chunk state is consumed through the inter-chunk term:

case	output error	final_state error
1 chunk, no initial state (inter-chunk term is zero)	0.000e+00	0.000e+00
1 chunk, initial state supplied	4.76e-03	0.000e+00
2 chunks, no initial state	2.88e-02	0.000e+00

That isolates it to b_o += tl.dot(b_q, b_h_q). The suspect is chunk_fused.py:207-217, which downcasts the fp32 register state to bf16 before transposing, deliberately imitating the reference path's HBM round-trip. Whatever numerical regime that reproduces on gfx9, it does not reproduce on gfx1151.

No fix attempted — that needs the kernel author's intent about the downcast.

Also: it does not compile as shipped

Before the accuracy question can even be reached, the shipped autotune list fails codegen:

error: no matching matrix core intrinsic for wmma version 1 with instruction
shape [0, 0, 64] and element types A='bf16', B='bf16', C='f32'

chunk_fused.py:79 autotunes BV in [8, 16, 32, 64]. BV=8 produces an 8-wide tl.dot, below RDNA3.5 WMMA's 16-element minimum. The comment at lines 72-78 attributes the small-BV configs to MI300X grid occupancy; gfx9 MFMA accepts them and WMMA does not. Restricting to BV >= 16 compiles, which is how the accuracy numbers above were obtained.

Reproducing

A test file exists that establishes this: 23 cases covering chunk-aligned and unaligned lengths, padded batches, ragged varlen batches, and both state dtypes. It uses the tolerances of the equivalent CuteDSL backend test unchanged, so the bar cannot be said to have been set to suit the outcome.

Note this defect may not be gfx1151-specific — the downcast is architecture-independent, though MFMA and WMMA differ in accumulation behaviour. Untested on gfx9 for lack of hardware.

[randomvariable]

Resolved by 3c9a8d2 ([ROCm] Remove the fused GDN prefill backend).

That commit deletes the fused kernel and removes fla_fused from the GDN prefill resolver and CLI configuration. Current supported selections are flashinfer, cutedsl, and triton; unsupported values fall back to Triton/FLA. Therefore the inaccurate fused-GDN output path described here can no longer be selected or run.