[vgokhale]

Qwen3.5 GDN layers on gfx1250: new prefill kernels for the conv, the chunk-local WY step, the recurrence and the output (no runtime autotune left in the GDN prefill path), and the recipe's DP points on unified attention with 64-token KV blocks.

Needs aiter alizaidy/gfx1250-qwen-support from 3d3228c1 (pushed): the chunk-form fused_gdn_verify decode kernel, ROCm/aiter#7100's current configs and routing, and ROCm/aiter#6999 (the 2D unified-attention kernel takes the block-table stride at runtime; before it, unified-attention prefill recompiled for every prefill batch).

Commits

1. conv1d rows kernel with the q / k L2 norm in its epilogue (causal_conv1d.py). The GDN prefill conv used 8-token x 256-channel tiles; the new kernel unrolls the chunk's tokens so each row is one independent load. qk_l2norm_head_dim normalizes each head after the bf16 rounding of the conv output, so chunk_gated_delta_rule skips its two l2norm_fwd launches (use_qk_l2norm_in_kernel=not qk_normed_by_conv).
2. chunk_h: a Gluon gfx1250 recurrence kernel, no autotuning (chunk_delta_h_gfx1250.py). The state stays fp32 in registers for a sequence's whole chunk walk, and the next chunks' w / k / u / gate tiles are prefetched into LDS. The launch config is a fixed rule on sequences x heads. The @triton.autotune (12 configs, timed on the first launch in the serving loop, and not even the same pick across processes) and the scratch-copy pre-tune are gone; the Triton kernel stays for other archs with one fixed config.
3. chunk_o: a Gluon gfx1250 output kernel with the gated RMSNorm in its epilogue (chunk_o_gfx1250.py). chunk_gated_delta_rule(..., output_gate=(z, weight, eps)) writes norm(o) * weight * silu(z) straight into the out_proj input; the pure-prefill call sites in attention_gdn.py pass it and skip the separate norm.
4. Chunk-local prefill kernels: kkt + inverse / WY in two kernels (fused_cumsum_kkt.py, new fused_inverse_recompute.py). k kᵀ once a key head on bf16 WMMA, coalesced gate / beta loads, and the 16x16 inverses, their merge and w / u in one kernel. This replaces solve_tril_16x16 + fused_merge_recompute, whose single configs were tuned on MI308. The short-prefill (< 64 tokens) kernels (cumsum, chunk_scaled_dot_kkt, solve_tril, wy_fast) take fixed configs instead of runtime autotune.
5. conv1d: fp32 FMA taps, hardware exp / reciprocal SiLU. The reference kernel rounds each tap product to bf16. Against an fp64 conv of the same bf16 inputs (the checkpoint's conv weights, layers 0 / 1 / 12 / 28), that puts its outputs at 1.43-1.48x the bf16 rounding error, with 66-70% correctly rounded. fp32 taps give the correctly rounded value for every output, as the decode verify kernel's conv already does. The fast SiLU moves about 1 output in 10⁵ by one bf16 step.
6. Recipe: DP points set ATOM_USE_UNIFIED_ATTN=1 and --block-size 64; the recipe doc gains the c2048 results and the exact server command.
7. Lint.

Results

c2048 DP4 + EP4, MTP3, 512 sequences an engine. Generation tok/s is the mean over the 2 minutes before a 15 s profiler window at minute 3; step times are DP rank 0's:

	branch tip 6cdc02f + aiter a0df93d2	+ GDN kernels (1-4)	+ 5, unified attention, 64-token blocks (aiter 3d3228c1)
generation tok/s	27.4k	28.9k	33.1k
prefill step	169 ms	159 ms	147 ms
decode step (median)	18.5 ms	17.5 ms	15.5 ms
MTP tokens a forward	2.73	2.72	2.72

Per GDN layer at 16k prefill tokens, 32 value heads (cold-cache microbenchmarks; the trace agrees):

Stage	Before	After
conv1d + q / k L2 norm	134 us	62-66 us
cumsum + k kᵀ, solve_tril, merge / recompute	241-246 us	147-154 us
chunk_h	208-396 us (by sequence mix)	98-202 us
chunk_o + gated RMSNorm	135-141 us	99-106 us

TP4 heads (8 value heads) gain about the same factors.

GSM8K, DEP4, 3-shot, 256 concurrent, on the last column's stack: 0.8635 / 0.8499 (0.861 / 0.851 before).

Testing

• New: tests/test_causal_conv1d_rows_gpu.py and tests/test_gdn_chunk_gfx1250.py.
  • The conv test compares q / k / v with an fp64 conv: at most one bf16 step, on 4e-6 to 2e-5 of values. The conv-state pool matches the reference kernel bit for bit.
  • The chunk test runs the eager slot path, the token-graph plan path and fixed-length batches against the Triton kernels.
• With them, the GDN tests (test_gdn_midstep_state_gpu, test_gdn_midstep_checkpoints, test_gdn_page_checkpoint, test_state_checkpoint, test_gdn_fused_decode, test_gdn_fused_verify, test_fused_sigmoid_gate) give 401 passed on gfx1250.
• Bitwise status:
  • chunk_h and chunk_o + norm are bitwise identical to the base tree on every path.
  • The w / u stage is at the fp64 floor (relative L2 1.66e-3, as before), with 99.8-100% of values bit-identical.
  • The end-to-end chunk output differs from the base tree by about 1e-3 relative.
• test_aiter_flydsl_gdn.py is MI308-only and fails identically on the base tree.
• ruff / black are clean for new code; existing files have no more findings than before.
• The tables above ran on 6cdc02f (a DSv4.1-flow harness, server command in the recipe doc). This branch is rebased on 9c6a7e7 ([Recipe] Qwen3.5 MLPerf Endpoints: DP sampling defaults (c128 +11%) #2824, [Qwen3.5] Shared expert side stream off by default (ATOM_QWEN_SHARED_EXPERT_STREAM) #2825, [Qwen3.5] TP1 mono decode: fused MTP verify layers on ReplaySSM (ATOM_QWEN35_MONO=1, -tp 1) #2827, Alizaidy/qwen35 dp prefill graph #2828, no file overlap). A c2048 point on this branch with the recipe's current DP defaults (which add the prefill token graphs and the fast launch paths) is running next; I'll add it here.