[explore] W4A8 sdot4 ranked shapes + test plan (not dest) by BlivionIaG · Pull Request #9 · opengfx1030/vllm-rdna

BlivionIaG · 2026-09-14T13:09:30Z

Purpose

Explore-only research scaffold for W4A8 = W i4 × A i8 → sdot4 (__builtin_amdgcn_sdot4 / v_dot4c_i32_i8) on gfx1030.

This is not dest, not UNC-26, and not a default path. It exists so other maintainers can research optimal tile shape.

Project 4 card: Explore: W4A8 sdot4 — ranked shapes (nibble→i8) (V620 inference).

Why this is not a duplicate

gh pr list / GitHub search on opengfx1030/vllm-rdna and vllm-project/vllm: no open PR for W4A8 sdot4 / ranked shapes.
UNC-26 stays dest (EXL3 HIP). This PR does not touch EXL3/FA or bump dest.

What landed

docs/explore/w4a8-sdot4/README.md — contract, ranked shapes, leave list, dependency.
docs/explore/w4a8-sdot4/TESTPLAN.md — V620 cells to fill (no fabricated numbers).
csrc/rocm/explore/w4a8_sdot4_shape1.{cu,cuh} — shape (1) sketch only. Not in VLLM_ROCM_EXT_SRC. Not in torch_bindings.cpp. Intended gate remains VLLM_RDNA2_W4A8_SDOT4=0.

Ranked shapes (do not skip)

ConfigA-class LDS=0, nibble→signed i8 in VGPR, K_STEP ∈ {16,32} only, two sdot4 per W dword; inner loop must cover full K_STEP (ConfigH lesson).
Recipe-9 LDS 64×64×64 i8 WG=256 ~8 KiB — only if (1) loses large-M.
Decode skinny M≤4 — A i8 in LDS, W stream; only if act quant already paid.
Later: unpack W to i8 workspace then reuse a W8A8 INT8 tile.

Leave list

sdot8 / V_DOT8_I32_I4, sudot4, FP8-bit sdot4, permanent TP≤2, unpack→fdot2 sold as sdot4.

A8 is GEMM activations only (routed expert / dense FFN already on W4 packs). Do not requant leftovers / embed / head / norms / GDN state / KV / indexer.

Dependency

Explore after W8A8 INT8 sdot4 exists. UNC-26 EXL3 stays first on dest.

Test Plan

Confirm default HIP source list unchanged: csrc/rocm/explore/ is not in VLLM_ROCM_EXT_SRC.
Confirm no can_implement / serve-script / EXL3 / FA edits.
On V620, fill cells in docs/explore/w4a8-sdot4/TESTPLAN.md (do not invent tok/s).

git diff origin/rdna_extras -- CMakeLists.txt csrc/rocm/torch_bindings.cpp
# expected: empty

Test Result

Default build/serve path: unchanged (docs + unwired stub only).
No silicon numbers in this PR (test plan is placeholders).
No model-eval delta (output path not wired).

AI assistance was used to draft the explore docs and the unwired HIP sketch. A human submitter must review every line before this leaves draft.

Do not merge. Keep draft.

Scaffold Explore-only research docs and an unwired HIP shape-(1) sketch for W i4 × A i8 → sdot4 on gfx1030. Not dest, not UNC-26, not in VLLM_ROCM_EXT_SRC / can_implement / serve scripts. Co-authored-by: Cursor Grok 4.6 Signed-off-by: Blivion <kev29lt@gmail.com>

Reed-Schimmel · 2026-09-19T06:43:20Z

I've been thinking about this too.

…ness Rework the Explore-only W4A8 (W uint4 x A int8 -> v_dot4_i32_i8) scaffold for gfx1030. Still not dest, not UNC-26, not in VLLM_ROCM_EXT_SRC, torch_bindings.cpp, can_implement or serve scripts. Contract fixes (README "What changed vs the first scaffold"): - ConfigA's "LDS=0" is LDS_PAD=0; ConfigA stages A in LDS. Shape (1) now follows ConfigA; the no-LDS variant is (1b), A via the scalar cache. - Packs are uint4 with zero points and exllama-shuffled by gptq_shuffle: zero-extend with a 3-VALU SWAR unpack, fold z*sum(a) per group, write A in the matching byte order (W stays shared with W4A16 decode). - Grouped scales make "i32 through K" impossible: i32 per group, f32 across groups. Ceiling ~1.5x main-loop VALU at G=32, ~1.9x at G=128. Draft and tooling: - csrc/rocm/explore/w4a8_sdot4.cuh: GEMM (7 sweep configs x G 32/64/128), tile-interleaved act quant, G0 issue-rate probes; compiles for gfx1030 with plain clang via w4a8_sdot4_isa_shim.h. - w4a8_sdot4_capi.cu: C ABI .so for the harness (hipcc, not CMake). - benchmarks/kernels/w4a8_sdot4_explore: NumPy oracle with error bounds, ISA audit (exact v_dot4 count per group, no spills, scratch or quarter-rate multiplies), V620 harness (G0 peak, G2 check, G3 bench vs the production W4A16 selector), G1 fake-quant accuracy script, CPU tests. - DESIGN.md (draft of the kernel); TESTPLAN.md split into agent-side checks and the V620 run, with blank result tables. Not run on a GPU. Agent-side: 56 tests pass with torch (50 + 1 skipped module without), isa_check clean with clang 18.1.3, pre-commit clean. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_012RQWdM7SmvDHYCzew7dJ1c

Survey of JartX/vllm perf/rdna3_full_stack (dafcde3) and the related upstream work (vllm-project#57925, vllm-project#55522, vllm-project#54706, 1e48f8c). Each finding that applies to gfx1030 becomes a check here: - The exllama fp16 dequant stores s*(-1024 - z) as fp16, and gfx1030's qdq_4_rdna2.cuh does the same. reference.w4a16_rdna2_weights models it bit for bit: 2.3-3.3% rel-L2 against exact on random data, several times the per-token A8 error on Gaussian rows. G1 is therefore read as pessimistic, and G2 compares the W4A16 op with both exact and baked. - Per-(token, group) activation scales (llama.cpp's Q8_1 idea) as the *_ag configs (a16_lds_k32_ag, a8_lds_k32_ag), with a per-group act quant that needs no block reduction and --act-group-size in the G1 script. Outlier A8 error goes from 5.3-6.0% to 1.0% at G=32 and 1.9% at G=128; the compiled VALU ratio is 1.36 instead of 1.47 at G=32. - Deterministic split-K: G2 runs fp16 split-K twice and records whether the bits repeat. Graduation needs f32 partials plus a fixed-order reduce, or split 1. G3 gains --split-k to time the no-split case (JartX splits only below ~2x wave oversubscription). - G3 checks every timed output against the W4A16 op (rel-L2 < 0.1) and exits non-zero otherwise. On gfx1100 a corrupted engine benchmarked faster. - Graduation plan: the M threshold goes in the C++ op entry; W4A8 gets its own TU with a W4A16 ISA diff; soak runs use a checker proven on a known-bad sample. Shape (4) waits for a fails-first harness, since dense dequant + rocBLAS wrote "!!!!" in production on gfx1100. - isa_check: fall-through "; %bb.N" blocks are parsed separately, so preheader and loop-exit code no longer counts as loop body. Not run on a GPU. Agent-side: 62 tests pass with torch (54 pass and 1 module is skipped without it), isa_check is clean with clang 18.1.3 (non-_ag rows identical to ba8ce98), and pre-commit is clean. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_012RQWdM7SmvDHYCzew7dJ1c

…pect the hard LDS cap - isa_check: clang 23 omits some asm keys (NumSgprs) for act_quant kernels, so the audit crashed with KeyError; use .get() with placeholders. - pick_split_k: the soft 16/64/32 KiB budget ignored the hard 64 KiB cap that launch_gemm enforces via the A_SRC-aware ex::lds_bytes, so it could return a split the launch rejected (kLdsTooBig). Clamp on the launch contract. - bench check: print each cell and record an explicit SKIP for ineligible cells instead of aborting the whole run. Measured after the fix on V620 (this branch): G0 sdot4/fdot2 MAC ratio 2.01 (PASS, implied 2.05 GHz); ISA audit all 27 configs ok, best a8_lds_k32 at 1.53x (G=32) / 1.77x (G=64) / 1.92x (G=128) VALU ceiling; G3 hot speedups 1.75-2.9x vs the production W4A16 op with rel-L2 ~2.5e-2; G2 180/192 pass (remaining = LDS-ineligible config/cell pairs, now explicit skips).

…eval Passing gpu_memory_utilization inside --llm-kwargs raised TypeError (duplicate kwarg), so production flags (kv_cache_memory_bytes, VL mm limits) could not be passed through. Merge harness defaults with llm_kwargs instead. G1 with this on a dense 1.5B: per-token int8 FFN activations FAIL (+10.77% PPL vs the +2% limit); per-(token, G=32) PASS (+0.75%; prefill-only +0.56%). The 27B target (cyankiwi/Qwen3.8-27B-AWQ-INT4, qwen3_5 hybrid) loads only with the production VL limits + KV cap and then yields NaN PPL under a bare LLM(), so its G1 still needs the fork's model env.

…ion groups - --clip-frac saturates the int8 grid at clip_frac*absmax, so the per-token outlier hypothesis can be tested from the harness. - fake_quant_int8 now reports a clear error when the activation group does not divide the per-rank (sharded) K, instead of torch's opaque unflatten failure. G1 investigation (dense 1.5B, prefill-only): per-token FAILS (+10.77%), and clipping makes it worse (0.9x -> +19.16%, 0.8x -> +59.94%), so the extremes carry signal; the fix is group granularity: per-(token, G=32) +0.75% PASS and per-(token, G=64) +0.72% PASS (G=128 invalid at TP=4: 128 does not divide the 2240-wide per-rank K).

…0 pass, next round Review of the first V620 pass (a1193e8, 1d601a9, 93caa86): - G2's 180/192 never covered a production-shaped cell. The f32 check runs at split 1, which needs M_TILE*K bytes of LDS: over 64 KiB for every LDS config at K=8704, and for a32n2 at K=2560. The first config that could not launch aborted the whole cell, so the 12 "skips" were 128x6144x2560, 624x6144x2560 and 2048x2560x8704 in full, and the 180 passes were the five EDGE cells. Now lib raises LdsTooBig; check_cell marks the f32 check n/a for that config and still bound-checks its fp16 path. It skips a config only when no group-aligned split fits (a32n2 at G=128, K >= 8704), and skips no longer count as failures. - G3 would have aborted the same way under --split-k 1 or --group-size 128. It now validates each launch once and records a skip. --configs ag selects the per-group configs, the only kind that passed G1. - The pick_split_k clamp from a1193e8 cannot change a split: the loops before it already stop within the <= 64 KiB soft budget or at the last split. An exhaustive sweep of 20,000+ cases found none it changes. Left in place, documented in RESULTS. Built on the results: - New config a8_smem_k32_ag (id 9): per-group scales, A through scalar loads, no LDS, so it is the one per-group config that can run split 1 at K=8704. The ISA audit (clang 18) is clean: 104-105 VGPRs, occupancy 9, 1.39x/1.68x/1.86x. a16_smem_k16_ag was tried too and spills SGPRs at G=128, so it is not added. - RESULTS-2026-09-28.md: the measured gates, the review, and a prioritized next round with each item tied to a result. P0 covers full G2, G3 on the _ag configs, counters, and G1 on the 27B. P1 covers mixed per-layer granularity, one A scale per R weight groups, K-chunked A staging and act-quant fusion. P2 covers G1 on the #31 exact build and G3 against the op serving runs (#30). - README status per gate. Shape (1d) becomes the main line. TESTPLAN and DESIGN get the new config and the skip semantics. On CPU: 54 passed and 1 skipped without torch, 64 passed with torch (including new tests for the K=8704 skip behaviour in check and bench). isa_check is clean on all 30 kernels. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_012RQWdM7SmvDHYCzew7dJ1c

BlivionIaG · 2026-09-28T21:06:37Z

Next actions after the first V620 pass (head 96b2a05; details in docs/explore/w4a8-sdot4/RESULTS-2026-09-28.md)

Updated 2026-09-29 after rechecking rdna_extras (e0112c5): added "Before the next pass", and the 27B G1 step now pins the tree and caps the KV cache.

Where it stands:

G0 PASS: sdot4/fdot2 MAC ratio 2.01.
G1 (dense 1.5B, TP=4): per-token FAIL at +10.77 % PPL, and clipping makes it worse. Per-(token, group) PASS: +0.75 % at G=32, +0.72 % at G=64. So the *_ag configs are the main line; the per-token configs stay only as timing references.
G2: the 180/192 covered only the five EDGE cells.
- The split-1 f32 check can't fit in LDS at K=8704, or at K=2560 for M_TILE 32.
- One such config aborted its whole cell, so no config was bound-checked on a production shape.
- Fixed in 96b2a05: that config's f32 check reads n/a and its fp16 path is still checked. It is skipped only where no split fits.
G3: 1.75–2.9× hot is above the VALU ceiling (≤ 1.53× at G=32), and the run doesn't say which configs produced it. The number that decides is the best _ag config's × total.

Before the next pass

Sync with rdna_extras.
- The branch is 173 commits behind and merges cleanly.
- None of those commits touch the dense W4A16 kernels, their dispatch, or the _custom_ops wrappers the harness calls, so the G3 baselines stay comparable.
- They bring the FA-RDNA2 split decode (d94e220, e0112c5) and the prefill step cadence (6ed3909). Those matter for G4 end to end, not for G1–G3.
Pin the tree for every run. From the checkout root, export PYTHONPATH=$PWD, then record python -c "import vllm; print(vllm.__file__)" and the git SHA in RESULTS.
- ae5bedf found that a copied venv's editable finder imports the deployment tree instead of the launcher's. When that happens, the W4A16 baseline op and the G1 hooks come from a different tree than the one under test.
- scripts/serve_gfx1030_full.sh, the 27B launcher, does not pin PYTHONPATH yet.
- For the first pass, note which tree provided _rocm_C.

P0: runs on the V620 (harness from `96b2a05`; the `.so` rebuilds on the new source hash)

Full G2: python -m benchmarks.kernels.w4a8_sdot4_explore.bench check --json g2-full.json, without --quick. Pass: every checked config is within bound on the prefill and decode cells, and the only skip is a32n2_lds_k32 at G=128, K ≥ 8704.

G3 on the _ag configs, including the new a8_smem_k32_ag (id 9):

M=benchmarks.kernels.w4a8_sdot4_explore.bench
python -m $M bench --configs ag --cells prefill --group-size 32 --weight-type uint4 --json g3-ag-g32.json
python -m $M bench --configs ag --cells prefill --group-size 32 --weight-type uint4 --cold --json g3-ag-g32-cold.json
python -m $M bench --configs ag --cells prefill --group-size 128 --weight-type uint4b8 --json g3-ag-g128.json
python -m $M bench --configs ag --cells prefill --group-size 32 --weight-type uint4 --split-k 1 --json g3-ag-split1.json

Stop criterion, unchanged: × total ≥ 1.25 on at least 3 of the 4 cells with M ≥ 624 and N ≥ 6144 at G=32 (≥ 1.4 at G=128), or the explore stops.
With --split-k 1, the LDS configs show as skipped at K=8704; a8_smem_k32_ag is the per-group config that can run there.

Counters (TESTPLAN B2/B6), W4A16 op vs the best _ag config on 2048×6144×2560. They explain the above-ceiling speedups and decide whether the VALU-side P1 items can pay:
rocprofv3 --pmc SQ_WAVES SQ_INSTS_VALU SQ_INSTS_LDS SQ_WAVE_CYCLES --output-format csv -d prof-g05 -- python -m $M bench --configs ag --iters 3 --warmup 1
G1 on the 27B: first get a finite PPL from the baseline arm.
- Run under the scripts/serve_gfx1030_full.sh environment with PYTHONPATH pinned as above.
- Pass the VL limits and a KV cap in --llm-kwargs, plus "dtype": "float16". Keep the cap ≤ 6e9 at 0.90 utilization: 7e9 does not fit a V620 with this tree's memory check (ae5bedf).
- If the W4A16 baseline is NaN too, the fault is in the engine path, not the fake quant.
- Then compare per-token against --act-group-size 32 and 64, and per layer (--layers mlp.gate_up_proj, --layers mlp.down_proj).
Commit the JSON records next to RESULTS-2026-09-28.md. The per-cell G3 tables in TESTPLAN stay blank until then.

P1: build only if P0 says so

Mixed granularity: per-token for gate_up_proj, _ag for down_proj. Only if the 27B's gate_up passes per-token.
One A scale per R weight groups: G=64 cost the same accuracy as G=32. The model gives 1.52× → 1.58–1.61× at M_TILE 8, at the price of a second f32 accumulator set (occupancy 10 → 8). Only if the counters show the kernel is VALU-bound.
Stream A through LDS in K chunks: this allows split 1 (deterministic, no zero-fill or atomics) with LDS configs at any K. Only if --split-k 1 beats the split runs.
Fuse per-group act quant into RMSNorm / SiLU·mul: it is thread-local, with no block reduction. First measure act quant's share of × total.

P2: with the other explores

G1 on the exact-dequant build of opengfx1030/vllm-rdna#31. Exact W4A16 weights plus fake A8 is what the W4A8 kernel computes; today's G1 also carries W4A16's baked fp16 bias.
G3 against the W4A16 op that compiled serving actually runs, per opengfx1030/vllm-rdna#30.
Both need the PYTHONPATH pin too. Without it, [explore] RDNA2 W4A16: does torch.compile freeze the M-dispatch? #30's env flag and [explore] RDNA2 W4A16: is the baked fp16 dequant bias real, and what does exact cost? #31's exact build may never reach the process that runs.

Generated by Claude Code

BlivionIaG · 2026-09-30T10:09:39Z

Superseded by landed work on rdna_extras (merge 3a0786e): the W4A8 dense + MoE kernels ship opt-in with op-level tests, ISA gates and in-model validation; the explore's accuracy conclusion (per-(token,group) scales only) and config choices are encoded in the shipped kernels. Remaining P0/PL items were perf-headroom gated on counters that are not being pursued. Closing.

…efill replacement Adds the production wiring for the validated explore kernel (PR opengfx1030#9): - csrc/rocm/w4a8_sdot4_rdna2.cu / .cuh (own TU; no existing TU touched) with two ops returning int64_t status (0 ok, positive hipError_t, negative not-eligible code so Python falls back): * w4a8_act_quant_rdna2(x, group, a_i8, a_scale, a_asum) - per-(token, group) int8 quant + f32 scales + int32 group sums in the [T][K/8][MT][8] layout * w4a8_gemm_rdna2(a_i8, w_packed, qzeros, scales, a_scale, asum, out, k, group, zero_offset, config_id, split_k) - over the SAME packed W4 buffer RDNA2W4A16LinearKernel already leaves behind. Default config id 8 (a8_lds_k32_ag) with GROUP=64 (the measured accuracy point); groups 32/64/128 supported; the hard 64 KiB LDS cap is honoured when picking a split. - CMakeLists.txt: new TU in VLLM_ROCM_EXT_SRC for gfx1030; ops.h declarations; torch_bindings.cpp registrations (single def+impl pair). - rdna2_w4a16.py: opt-in prefill branch (VLLM_RDNA2_W4A8_SDOT4=1 + on_gfx10x + ops present + shape/quant eligibility), falling back to gptq_gemm_rdna2_prefill otherwise. Default OFF - byte-identical behaviour when unset. Existing dispatch chain untouched. - tests/kernels/quantization/test_rdna2_w4a8.py: act_quant vs numpy, W4A8 vs the W4A16 prefill op (rel-L2 < 0.1) for G=32/64/128, dispatcher default-off. Needs a single-rank TP group (default_vllm_config + probe-and-init helper). - docs/rdna2/w4a8-wiring.md. Validated on 4x V620 (venv-7.14.0_0.28.0): _rocm_C builds clean, both ops registered, 17/17 tests pass. Wired-path A/B vs the production W4A16 prefill op (act_quant + gemm, G=64): M=128 1.10x, M=512 1.55x, M=2048 1.78x, M=2048/K=8704 2.09x, rel-L2 ~1.3e-2. Still open: in-model validation needs a dense W4A16 checkpoint (the 27B AWQ bring-up currently resets the host); determinism at split_k>1 (pk4 CAS).

This was referenced Sep 14, 2026

[explore] W4A4 sdot8 ranked shapes + test plan (not dest) #10

Draft

silicon: lock W4A4 sdot8 Explore ConfigA LDS=0 first BlivionIaG/rdna-hip-wiki#4

Merged

GeorgeMA-Strong mentioned this pull request Sep 16, 2026

fix(rocm): restore Intel CPU PLE and V620 MTP2 startup #12

Closed

wsantos mentioned this pull request Sep 24, 2026

Fixed Piecewise #20

Merged

claude and others added 6 commits September 28, 2026 01:40

BlivionIaG closed this Sep 30, 2026

This was referenced Oct 11, 2026

fix(rdna): launcher stops re-enabling GQA decode (VLLM_FA_RDNA2_GQA_DECODE default 0) #53

Draft

feat(rdna2): opt-in W8A8 int8 sdot4 dense and MoE kernels for gfx1030 #93

Draft

Navigation Menu

Repository navigation

[explore] W4A8 sdot4 ranked shapes + test plan (not dest) - #9

[explore] W4A8 sdot4 ranked shapes + test plan (not dest)#9
BlivionIaG wants to merge 7 commits into
rdna_extrasfrom
explore/w4a8-sdot4

BlivionIaG commented Sep 14, 2026

Uh oh!

Reed-Schimmel commented Sep 19, 2026

Uh oh!

BlivionIaG commented Sep 28, 2026 •

edited

Loading

Uh oh!

BlivionIaG commented Sep 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

BlivionIaG commented Sep 14, 2026

Purpose

Why this is not a duplicate

What landed

Ranked shapes (do not skip)

Leave list

Dependency

Test Plan

Test Result

Uh oh!

Reed-Schimmel commented Sep 19, 2026

Uh oh!

BlivionIaG commented Sep 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Before the next pass

P0: runs on the V620 (harness from 96b2a05; the .so rebuilds on the new source hash)

P1: build only if P0 says so

P2: with the other explores

Uh oh!

BlivionIaG commented Sep 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

BlivionIaG commented Sep 28, 2026 •

edited

Loading

P0: runs on the V620 (harness from `96b2a05`; the `.so` rebuilds on the new source hash)