Skip to content

perf(attention): add packed fp8 kv reads on gfx1201 - #897

Merged
jason-fxz merged 4 commits into
FlashML-org:mainfrom
zihaomu:fix/amd-fp8-kv-pr894-20261112
Nov 14, 2026
Merged

jason-fxz merged 4 commits into
FlashML-org:mainfrom
zihaomu:fix/amd-fp8-kv-pr894-20261112

Conversation

@zihaomu

@zihaomu zihaomu commented Nov 12, 2026 •

Copy link
Copy Markdown
Contributor

Adds a packed FP8 KV reader for the qualified gfx1201, BF16 query, D128, GQA4/8, batch-one single-row decode path. Four contiguous E4M3FN bytes are loaded as one uint32, converted with native packed instructions, and reconstructed as BF16 for the existing attention math. KV storage, per-token/head scales and BF16 WMMA remain unchanged; this is not native FP8 WMMA attention.

Merged current main 66190a9 (including #900) into the PR. This revision addresses the review with 64/1/4 (BLOCK_N/num_stages/num_warps), replacing the previous spilling 64/1/2 launch. Current-revision kernel and regression validation applies to fc3cb03 and the pinned main above. The model-context matrix below is explicitly labelled as historical.

  • The gate and launch live in attention.py, using the existing get_rocm_gfx_arch() helper. Packed dispatch requires 16-byte aligned bases and row/head strides: a merely 4-byte aligned padded layout produced 33 SGPR spills in validation, so it now uses the original generic reader. kv_rows.py contains only the packed addressing/conversion; its converter is private and the test-only PACK=1 arm is removed.
  • CUDA keeps the original inline address expressions and generic reader. INT8, BF16 KV, D256, multi-row decode and unqualified layouts keep main's dispatch. perf(attention): run up to 8 query rows per request in the triton decode kernel #900's multi-row implementation is retained.
  • gfx1200 is excluded only because no gfx1200 hardware was available for qualification, not because its ISA lacks the required instructions.

Model speed across context lengths: BF16 KV vs FP8 KV

Historical measurements from 2026-11-12, restored here for visibility; these models were not rerun for the current review revision.

These measurements compare BF16 KV versus FP8 KV on the same branch, not main versus the new native reader. Neither model enters this PR's D128 native dispatch: Qwen3.6 uses D256 full attention among its GDN layers; Qwen3.8 uses QSA sparse attention. These results qualify the exercised generic paths and document remaining performance work, not gains attributable to this patch.

Both checkpoints were already present locally; no model was downloaded or requantized:

  • Qwen/Qwen3.6-35B-A3B, original BF16 weights, revision 995ad96eacd98c81ed38be0c5b274b04031597b0.
  • Qwen/Qwen3.8-Flash-Next-NVFP4, original NVFP4/FP8 mixed weights and PLE files, revision fc694b54fb0174e0913e6adf86691ef85a4ead47.

Controls: same physical GPU0 / PCI 0000:03:00.0, TP=1, batch=1, BF16 compute, Graph replay, MoE offload, fixed 16 GiB expert cache, 35,072 KV slots, naive cache, prefill chunks of 2,048, greedy and exactly 128 output tokens. Qwen3.6 uses triton; Qwen3.8 uses qsa_sparse, moe.nvfp4=triton, and disk PLE. CPU affinity is 16-31,80-95; memory NUMA policy remains default because the container denies set_mempolicy.

Each format has one service process; each input length is warmed once, then measured three times. All 36 measured requests and 12 warmups completed, and usage/admission counters verified exact prompt/output lengths and no prefix reuse. Medians below include every completed measurement, including slow requests.

Model Input tokens BF16 KV tok/s FP8 KV tok/s Observed change BF16 / FP8 TTFT (s)
Qwen3.6-35B-A3B 512 29.07 18.53 -36.24% 2.287 / 2.233
Qwen3.6-35B-A3B 8,192 28.49 30.32 +6.39% 8.517 / 7.562
Qwen3.6-35B-A3B 32,768 28.54 26.25 -8.05% 33.131 / 52.270
Qwen3.8-Flash-Next-NVFP4 512 29.08 29.37 +1.00% 2.393 / 1.865
Qwen3.8-Flash-Next-NVFP4 8,192 30.95 29.96 -3.20% 8.165 / 7.745
Qwen3.8-Flash-Next-NVFP4 32,768 24.93 30.72 +23.24% 53.904 / 33.941

Decode timing for this matrix is (completion_tokens - 1) / (terminal finish_reason SSE timestamp - first nonempty text timestamp). It includes the terminal event even when its text is empty. Raw runner logs retain the older first/last-nonempty estimate; this table is consistently recomputed from the saved SSE events. TTFT includes full prefill.

Known regression and uncertainty

FP8 KV is slower in the measured Qwen3.6 512-token case. Its BF16 samples are 31.93 / 28.18 / 29.07 tok/s; FP8 samples are 18.53 / 16.88 / 20.33. This D256 case is not fixed by the D128 dispatch in this PR.

Short contexts offer less KV bandwidth to save while conversion and row-scale work remain. The model's D256 decode specialization reports 73 VGPR spills with FP8 versus 2 with BF16. This is a profiling lead, not an isolated root-cause result: kernel timing and controlled ablations have not established how much each factor contributes.

There is substantial runtime variation. Qwen3.6 FP8 at 32K has samples 9.03 / 26.25 / 28.63 tok/s; Qwen3.8 BF16 at 32K has 24.78 / 24.93 / 31.80. No slow sample was dropped. These are within-service repeats in sequential format runs, not independent alternating service A/B rounds. The percentages are descriptive and do not establish stable gains or a universal -36.24% penalty.

At the same KV capacity, logged K/V allocations are 0.67 -> 0.34 GiB for Qwen3.6 and 0.83 -> 0.43 GiB for Qwen3.8. GDN state and expert weights remain unchanged. Saved VRAM is not reassigned to the expert cache in these comparisons.

Why the handwritten conversion remains

The pinned Triton 3.8 ROCm build lowers plain OCP FP8-to-BF16 casts in software on RDNA4. The handwritten v_cvt_pk_f32_fp8_e64 forces VOP3: LLVM's word-0 packed intrinsic can select a VOP1 form rejected by the gfx12 assembler for VGPR inputs. Triton #11809, merged as 7e239bb, adds hardware upcasts using one byte-select FP8-to-FP32 conversion per byte, then narrowing. Its merge does not change our pinned compiler. Recheck plain casts for correctness, generated instructions and cold/Graph performance on a compiler containing that change before removing this workaround. FP8-to-BF16 casts were already numerically possible before that PR; efficient lowering is the issue.

Fresh R9700 stage-1 A/B

Radeon AI PRO R9700/gfx1201, PCI 0000:03:00.0; dual EPYC 9334, CPU affinity 16-31,80-95; Linux 7.0.0-42-generic, amdgpu 7.1.9.31600000. Locked ROCm 10.1, HIP 7.16.26385, Torch 2.14.0+rocm10.1.0, Triton 3.8.0 (669b31ac ROCm build); image sha256:a79c374ab2bc93c0e4358da6c7bda5cb84cf1308e8aff74fc99fc95c511cbd03. Physical GPU0 was reserved and monitored; no foreign GPU process was observed.

Same randomized physical slots, Q/K/V/scales, D128, 32 query heads, batch 1, 64 scratch/split cap. Independent CPU FP8 dequantization plus FP32 PyTorch attention reference, unchanged atol=0.002, rtol=0.02. Five seeded shuffled rounds per case; every sample retained. Stage 1 is timed separately from the merge kernel. Each cell is cache-evicted / warm Graph latency in microseconds, median across five rounds. Cache eviction uses triton.testing.do_bench; Graph uses do_bench_cudagraph. Main selects 32/3/4 in this matrix.

GQA / keys Main Generic 64/1/2 Packed 64/1/2 Generic 64/1/4 Packed 64/1/4
4 / 8,192 258.82 / 166.09 212.46 / 106.03 60.03 / 49.37 130.33 / 78.33 46.87 / 40.28
4 / 32,768 654.00 / 509.15 705.05 / 444.21 170.96 / 159.03 450.36 / 302.86 154.75 / 149.01
8 / 8,192 132.30 / 80.08 112.83 / 56.75 39.58 / 29.97 94.46 / 60.65 32.23 / 26.12
8 / 32,768 335.83 / 235.69 402.98 / 221.21 94.90 / 88.16 288.88 / 227.41 98.39 / 92.64

The same-launch controls isolate the packed reader from launch tuning; the reader delta includes uint32 loads, .cg policy and native conversion. The final packed reader is faster than every main and generic-64/1/4 sample at each tested length, in both cache conditions. One main GQA4/8K cold round is a long tail (max/min 1.363); it is retained, not discarded or corrected. These are kernel measurements, not model throughput multipliers.

Register/spill correction: the previous packed 2-warps kernel had 61/65 native VGPR spills and 248/264 bytes scratch for GQA4/8. Four warps remove them. At GQA8/32K, this costs about 4% cold / 5% Graph latency relative to the spilling 2-warps version; the final configuration follows the zero-spill requirement.

GQA Final packed VGPRs SGPRs VGPR spills SGPR spills Scratch bytes LDS bytes
4 223 79 0 0 0 16,384
8 229 79 0 0 0 16,384

Counts above come from native AMDGPU metadata (.vgpr_count, .vgpr_spill_count, .sgpr_spill_count, .private_segment_fixed_size), not conflated with Triton's n_spills. Final packed kernels also report n_spills=0. Zero-spill qualification covers the new packed path; it does not claim to fix existing spills in main's generic AMD kernels.

Regression and CUDA isolation

  • AMD broad regression: 193 passed, 4 failed, 0 skipped.
  • AMD final dispatch/reader regression: 58 passed, 0 failed, 0 skipped.
  • AMD forced emulation D128: 14 passed, 0 failed, 0 skipped.
  • NVIDIA native: 194 passed, 0 failed, 3 skipped.
  • 14 AMD integration cases passed, including 17/513/32768 keys, GQA4/8, odd sliding window, split caps 1/8/64, 16-byte aligned padding and 4-byte/unaligned fallback, plus Graph replay. Maximum reference absolute error: 0.000763267. All exercised final packed attention specializations have zero native spills/scratch.
  • Both K/V orientations preserve every finite E4M3FN value bit-for-bit against CPU PyTorch, including signed zero. No tolerance was relaxed. NVIDIA's three skips are ROCm-only tests.

All four AMD failures reproduce on pinned main with the same failure reasons. The BF16 D256 two-query backend case routed through #900's decode requires 65,792 bytes of shared memory on this compiler, exceeding the 65,536-byte device limit. The three e4m3 suite failures are the native-vs-emulation assertion (e4m3_native() is false on ROCm) and the sm89/sm120 foreign-target compile gates (bf16/uint8 operand mismatch). Their tests and production paths are unchanged by this PR; they are retained as failures, not skipped or made green. Main's multi-row AMD shared-memory launch needs a separate fix. The new packed reader is disabled for that path.

Fresh RTX 5090 GPU1 / PCI 0000:e1:00.0, EPYC 9554, driver 615.71.09, immutable NGC image sha256:42263b2424fc237b34c4fc4a91c30d603c57eed36e37d31ff6d9a4f1f801edee, existing Torch 2.14.1+cu130, CUDA 13.0, Triton 3.8.0. 81 main/head comparisons: BF16/FP8/INT8 KV; D128 GQA4/8 and D256 GQA8; 17/8192/32768 keys; 1/3/8 query rows. All independently checked outputs and Graph replays pass; output bits, CUBIN executable sections, normalized SASS, launch parameters and register/shared/spill counts match main. All NVIDIA cases report zero spills. This verifies the tested SM120 stack, not every NVIDIA architecture.

Fresh model A/B against current main

Local Qwen/Qwen3-4B, BF16 weights, revision 1cfa9a7208912126459214e8b04321603b3df60c; no download. Same reserved R9700, TP=1, batch=1, BF16 compute, FP8 KV, Triton attention, Graph, greedy, 32,768 prompt tokens and 128 outputs. Fixed 35,072 KV slots, page size 1, prefill chunks of 2048, memory ratio 0.7, radix cache. Four fresh services in main/head/head/main order, each warmed once then measured three times. TTFT is for the warm prefix, not full prefill. Decode timing includes the terminal finish_reason event: 127 / (terminal - first nonempty text).

The model's compiled packed specialization uses 245 VGPRs/77 SGPRs and also has zero VGPR/SGPR spills and zero scratch; its real cache/scale strides differ from the microbenchmark's separate allocations.

Service Source Decode median tok/s Warm TTFT ms
1 main-1 33.66 111.45
2 head-1 46.51 105.92
3 head-2 46.41 103.30
4 main-2 38.00 109.29

All 12 measurements and four warmups completed; no slow sample was dropped. Within-source service-median max/min: main 1.1290, head 1.0022.
The predeclared 1.10 startup-drift gate failed; these service rates are descriptive only and do not establish a stable end-to-end gain.
Output text SHA256 sets: main 00a5408efb8f9b76a4b92a0536ec84b6b1635e6ac6bb1140216cbf8a4064e78d; head 9bc101aaf7a671cc0bda7ac257ffda622359b8fd69fd24ce848f92cd76dd9901.
Main/head synthetic continuations differ; do not treat this as a model-quality or token-identity gate. Independent operator numerical checks above remain the correctness evidence.

These synthetic continuations are not a task-quality evaluation. Previous service rates of 33.65/37.6/43.7 tok/s were from earlier revisions and inadequately paired cohorts; they are superseded as evidence for this revision's performance.

Reproduction and retained evidence

Run the same commands at pinned main and this head inside the recorded hardware/runtime, after selecting an idle device by PCI identity:

PYTHONPATH=python python -m pytest tests/kernels/test_triton_attention.py \
  tests/kernels/test_e4m3_compat.py tests/kvcache/test_kv_cache_rebuild.py \
  tests/kvcache/test_pool_sizing_surface.py tests/engine/test_attention_backend_arch.py -q
FREETOKEN_FORCE_E4M3_EMU=1 TRITON_CACHE_DIR=/tmp/ft-pr897-emu PYTHONPATH=python \
  python -m pytest tests/kernels/test_triton_attention.py \
  -k 'matches_dequantized_reference and 128' -q

PYTHONPATH=python python -m freetoken.cli serve \
  --model /models/qwen3-4b --served-model-name cast-4b \
  --dtype bfloat16 --attention-backend triton --kv-cache-dtype fp8_per_token_head \
  --moe-strategy offload --num-tokens 35072 --page-size 1 \
  --max-seq-len-override 34816 --max-prefill-length 2048 \
  --max-running-requests 1 --cuda-graph-max-bs 1 --memory-ratio 0.7 \
  --text-model-only --enable-cache-report --sampling-defaults none \
  --reasoning-parser off --host 127.0.0.1 --port 19854

Send the saved exact 32K text fixture to /v1/completions with model=cast-4b, max_tokens=128, temperature=0, ignore_eos=true, streaming with usage. Prompt SHA256: 28e8514172da8e98fd27b06d31c4650eb048c3a61faf82de233957f566ffed9e. Warm once and measure three times in each ABBA service; verify 32,768 input/128 output and 32,767 cached prompt tokens per measured request.

New host artifacts: .freetoken-zihao/pr897-review-20261113/: bench.py (main, same-launch generic, packed controls; captures stage1), all individual timing samples, asm/ HSACO/native assembly, *-sources.json, runtime/ownership records, test XML, padded/unaligned/window integration checks, exact prompt/request/SSE files, and summary.json. NVIDIA artifacts: same root's nvidia/, remote test_lap/pr897-review-20261113/. Historical primitive casts, mutation checks and hybrid runs remain under .freetoken-zihao/amd-fp8-kv-pr894-20261112/. These paths are retained local artifacts, not public download URLs.

@jason-fxz

jason-fxz commented Nov 13, 2026 •

Copy link
Copy Markdown
Collaborator

@zihaomu
The packed reader looks correct to me, and the CUDA path is untouched (KV_PACK is a constexpr and the else arm is the original expression). A few things before we merge.

1. Kernel A/B against main. The only main number is one service at 33.65 tok/s, while the same HSACO measured 37.6 and 43.7 tok/s in two services, so it cannot show the gain. Please add a stage-1 A/B on the R9700 against main (D128, GQA 4/8, 8K and 32K keys), plus the generic reader with the same 64/1/2 launch, so we can see what the packed load buys apart from the launch change. Please also give the VGPR and spill count of the packed kernel; we keep spills at 0.

2. Launch in attention.py, arch from the existing helper. kv_rows.py holds the format's read and write code; decode launch geometry lives in _quantized_decode_launch (attention.py:40). Please move the gate and the 64/1/2 next to it, and read the arch with get_rocm_gfx_arch() from freetoken.utils.arch instead of a new _gfx1201:

def _use_gfx12_fp8_reader(kv_format, q, k_cache, v_cache) -> bool:
    return (
        kv_format.name == "fp8_per_token_head" and not FORCE_EMU
        and get_rocm_gfx_arch() == "gfx1201"
        and ...  # the same dtype, shape and alignment checks
    )

Is there a reason to leave out gfx1200 (RX 9060 series)? LLVM gives gfx1200 and gfx1201 the same feature set (both FeatureISAVersion12, both covered by gfx12-generic), and Triton maps both to ISAFamily::RDNA4. If it is only missing hardware, one line in the description is enough.

3. Drop the test-only arm. The kernel calls kv_tile_rows only when KV_PACK == 4, so its PACK == 1 arm (kv_rows.py:101-114) runs only in the new test. Please make kv_tile_rows the packed reader only, drop the pack=1 case of test_rdna4_kv_decode_preserves_every_finite_fp8_code, and make kv_tile_gfx12_bf16 private.

4. Comments. "Preserve the merged reader's address expressions on non-packed targets." (attention.py:320) describes the PR history; please state the reason instead, e.g. # NOTE: keep the address math inline: moving it into a helper changes CUDA codegen. Please also say why the asm is hand-written, so whoever upgrades Triton knows to A/B a plain cast against it: Triton 3.8/3.9 cast OCP fp8 in software on RDNA4, and triton-lang/triton#11809 uses one v_cvt_f32_fp8 per byte because LLVM emits word 0 of v_cvt_pk_f32_fp8 as a VOP1 the gfx12 assembler rejects, e.g. # NOTE: VOP3 (e64) by hand: LLVM emits word 0 of v_cvt_pk_f32_fp8 as a VOP1 the gfx12 assembler rejects.

@jason-fxz

Copy link
Copy Markdown
Collaborator

@zihaomu #900 just merged and touches the same decode kernel, please merge or rebase onto main.

Merge main through 66190a9, retaining the multi-row decode changes from FlashML-org#900.
Move architecture and launch selection into attention.py, reuse the shared
architecture helper, and remove the packed reader's test-only generic arm.

Use four warps and require 16-byte aligned cache bases and strides to avoid
packed-reader register spills. Keep CUDA address expressions inline and
document the gfx12 VOP3 conversion workaround.

Validate D128 GQA4/8 against main and same-launch generic controls. Confirm
zero packed spills, 14 layout/Graph cases, 58 final operator cases and 14
forced-emulation cases. All 81 RTX 5090 main/head codegen comparisons match.
Preserve four unrelated AMD suite failures reproduced on pinned main.
@TashaSkyUp

Copy link
Copy Markdown

I reproduced and narrowed the pre-existing gfx1201 BF16 D256 multi-row LDS failure mentioned in the PR description, using current main 66190a9 and Triton-ROCm 3.8.0. This is compile-only against GPUTarget("hip", "gfx1201", 64), so it does not depend on a runtime launch or model weights.

Minimal stage-1 specialization that reproduces the exact reported number:

  • ROWS=2
  • BLOCK_H=2, VALID_BLOCK_H=2 (M = 4)
  • BLOCK_D=BLOCK_DV=D=DV=256
  • BF16 Q/K/V, unquantized KV (KV_FMT="none")
  • BLOCK_N=64
  • num_warps=4
  • num_stages=3

Triton metadata:

BLOCK_N=64 stages=3 shared=65792   # +256 over gfx1201 65536 limit
BLOCK_N=32 stages=3 shared=32768   # 32768 bytes headroom
BLOCK_N=16 stages=3 shared=16384

The 65,792-byte result is therefore independently reproducible from the kernel compiler alone and exactly matches the failure reported here.

I also isolated the stage effect:

BLOCK_N=64 stages=1 shared=32768
BLOCK_N=64 stages=2 shared=65792
BLOCK_N=64 stages=3 shared=65792

One useful detail: _multi_row_decode_launch(4, 256, False, 2, 2, smem) already returns BLOCK_N=32 when smem=65536, but returns BLOCK_N=64 when smem is 0 or >=67584. So the launch heuristic itself would avoid this if its budget matched the actual gfx1201 64 KiB launch limit. The failure points to a mismatch between _optin_smem_bytes() / the budget used by the heuristic and the actual HIP launch limit, rather than a requirement for D256 to use >64 KiB.

The lowest-risk fix looks like either:

  1. cap the ROCm/gfx12 multi-row LDS budget at the actual 64 KiB limit before _multi_row_decode_launch, or
  2. conservatively force BLOCK_N <= 32 for ROCm, multi-row, BF16 D>=256.

Option 2 preserves the 3-stage path and, in this exact specialization, cuts compiled LDS from 65,792 to 32,768 bytes. I have the compile-only reproducer saved locally if any other shapes are useful to sweep.

Hardware context for the host doing the work: 2x 32 GiB gfx1201 R9700 + one gfx1201 RX 9070 XT; host ROCm 7.2.3. The compile result above is target-pinned and does not claim a runtime benchmark.

@jason-fxz
jason-fxz merged commit 140faab into FlashML-org:main Nov 14, 2026
2 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants