huangpustar/metainfer/int8-w8a8-gemm-prefill

INT8 W8A8 GEMM — Prefill (M > 32) — gfx928 / SGLang TP4 + TP8

Worker-29 lineage, DTK 26.04; K500SM_AI/gfx928, 120 CUs, wavefront 64, 64 KiB LDS/CU, 65,536 VGPRs/CU. Validated evidence covers M=3072 (TP4) and M=4096 in three lineages: TP4 (hy3_tp4_*_m4096, §2), TP8 (hy3_tp8_*_m4096, §3) and GLM-5.2 TP8 (glm5-2-dsh-tp8-m4096-1, §4). M in (32, 3072) [TP4] / (32, 4096) [TP8 / GLM-5.2] are measured-boundary gaps — reuse the staged recipes below and measure. This skill is the prefill half of the phase split; load int8-w8a8-gemm-foundations for the operator/layout contract, DUMMA and epilogue rules, Graph-safe interface, SGLang integration, and shared guardrails. For each M=4096 lineage the source of truth is its authoritative control-plane fact ledger (TP4: 2026-08-22 re-verified values; TP8: 2026-08-27; GLM-5.2: glm5-2-dsh-tp8-m4096-1-e6a280a2): pending worker SKILL.md tables may be stale and must not override the ledger. Baselines are fixed user-supplied Triton CUDA-graph numbers that were never re-measured — do not treat any baseline as measured or optimal; "speedup" is always vs that fixed table.

1. M = 3072 prefill (validated)

For M=3072,K=4096,N=1536, use the validated large-M path:

block tile      = 64x64
K stage         = 64
threads         = 256 = 4 wavefronts
wave ownership  = one 32x32 quadrant per wave
DUMMA unit      = m16n16k32
LDS             = two buffers of A[64,64] + B[64,64] = 16 KB total
prefetch         = two 16-byte vectors/thread into VGPR
epilogue        = direct fragment -> scale -> bf16 store
tail M          = zero-filled A loads + masked stores

Launch grid:

grid.x = N / 64
grid.y = ceil(M / 64)

Issue stage i+1 global loads before the stage i MMAC sequence. Keep the prefetched vectors in VGPR, execute the current eight MMAC instructions, wait for VMEM immediately before writing the alternate LDS buffer, then use an LDS-ready barrier. Do not infer this overlap from source order: require the gfx928 ISA to show global_load_dwordx4 -> v_mmac_i32_* -> vmcnt(0) -> ds_write_b128 in that order.

Current optimized prefill constraints are M>32, K%64==0, N%64==0, contiguous inputs/scales/output, bf16 output, and no bias.

Do not use split-K for large M unless profiling proves grid parallelism is insufficient. MxN already supplies many blocks, and reduction overhead is normally unnecessary.

For the full M=3072 experiment table, ISA procedure, resource evidence, bandwidth accounting, and rejected stage sizes, read references/wqkv-a-m3072-prefetch-isa.md.

2. M = 4096 chunked prefill — TP4 (hy3_tp4_*_m4096)

Validated on the worker-29 lineage, one worker/GPU per shape ID, DTK 26.04. Operator contract is identical to the foundations skill; weight packing/preprocessing runs outside the timed region and outside Graph capture. Source of truth = the authoritative control-plane fact ledger (2026-08-22 re-verified values); pending worker SKILL.md tables may be stale and must not override the ledger.

The 2026-08-22 re-verified values: each accepted kernel object was

re-benchmarked with correctness ON (mismatch=0, max_abs_error=0.0) using the

same protocol (100 warmup / 30 samples × 100 replays, CUDA-graph replay, hot

cache). TOPS = 2MNK/median; bandwidth = algorithmic bytes / median (not

measured HBM traffic). Speedup = fixed user-supplied Triton baseline /

median. Iteration-chain values in §6 are the historical in-run measurements

of the same accepted kernels.

• All four shapes have accepted candidates: **no GPU worker is unavailable and

no shape in this lineage is unoptimized. Rule for any future shape: a shape whose worker timed out or produced no accepted candidate stays unoptimized** — never infer success, a speedup, or optimality from a failed/timed-out worker, and never call its baseline optimal when no candidate was accepted.

• M=4096 prefill: no split-K in any accepted candidate — the M×N

output-tile grid already dwarfs 120 CUs, and the 16 MiB workspace has no useful int32 partial-plane capacity.

• Routing: exact (m,n,k) guard per shape → tuned DUMMA kernel; every other

shape falls back. Place the exact-shape guard before generic tiled paths and keep it shape-exact so decode (M≤32), M=3072, and generic arms are untouched. Each source carries a generic scalar int8/int32 fallback that decodes the packed layout when the (k,n) matches: o_proj (2048,4096), qkv (4096,2560), down_proj n-major, gate_up (4096,768); gate_up also keeps generic single-buffer tiled arms (<64,128,128>, <128,64,128>, <64,64,128>) that were never re-tuned. Untuned paths must remain byte-identical to the accepted source and must still build/pass.

• All speedups are vs the fixed user-supplied Triton CUDA-graph baseline table

above; those baselines were never re-measured and must not be called optimal.

• Decode (M≤32), M=3072, and any other shape are outside this lineage and

must keep their own validated routes/fallbacks.

3. M = 4096 chunked prefill — TP8 (hy3_tp8_*_m4096)

Validated on the worker-29 lineage (hy3-dsh-tp8-m4096-1), one worker/GPU per shape ID, DTK 26.04. Source of truth = the authoritative control-plane fact ledger (2026-08-27); pending worker SKILL.md tables may be stale and must not override the ledger (e.g. down_proj's draft names round 8 as final while the ledger also accepted rounds 10/12/13; gate_up's draft calls round 11 "not accepted" though the ledger accepts it). All four TP8 M=4096 prefill shapes have accepted candidates (4/4); no GPU worker is unavailable. ~30 rounds in this lineage ended as compile_or_agent_failure (killed ~900 s, missing proposal, restore/resume failures): these are candidate/infra failures and prove nothing — never infer success, a speedup, or optimality from a timed-out/failed worker, and a compile error is never evidence that DUMMA, INT8, a HIP API, templates, or inline asm are unsupported. Recipes do NOT transfer across TP size: the TP4 numbers in §2 are a different lineage with different shapes.

3.1 Routing and accepted evidence (authoritative ledger)

All four: correctness passed (mismatch 0 / max_abs_error 0.0), Graph capture

• changed-contents replay passed. No proven optimum on any shape (no plateau

established): accepted numbers are the best found, not ceilings. No split-K in any accepted candidate — the MxN tile grid dwarfs 120 CUs; the only measured split (gate_up S=2) cost ~40-70 us of combine.

3.2 TP8 pack layouts and fallback guards (lineage-specific)

Foundations records the TP4 layouts; the TP8 layouts below are distinct and shape-specific. Checkpoint weight is [N,K] (stride K); SGLang Triton view is W.t() — preserve it for fallback and register one one-time contiguous HIP buffer; never replace layer.weight.data with the contiguous copy. Packing is one-time, out-of-timed-region, out-of-Graph, same byte count, same graph-stable buffer; correctness is always checked against the raw logical [K,N] weight:

• qkv (k=4096,n=1280): n-major packed[n][k]
• o_proj (k=1024,n=4096): n-major packed[n*K+kk] == raw[kk*N+n]
• gate_up (k=4096,n=384): [N,K] transpose (exact (K,N)==(4096,384) only)
• down_proj (k=192,n=4096): n-major packed[n][k]

Validate every pack kernel with a byte-exact simulation of the index relation. o_proj iter 10's mis-vectorized transpose silently left most of the buffer unwritten (mismatch 16,775,923 / max_abs_error 86.5): a fast but wrong measurement is not a win; the iter-11 repair (per-thread 16-byte gather from 16 raw rows + aligned int4 store) restored mismatch 0.

Fallback routing (must stay byte-identical and building): exact-shape guard first -> tuned DUMMA kernel; every other shape -> generic tiled arms / scalar fallback (w8a8_gemm_scalar_fallback_kernel), which decodes the packed layout when the (k,n) matches and is otherwise identity. Guards (unchanged across the lineage): qkv (m >= 128 && n == 1280 && k == 4096); o_proj (m,n,k) == (4096,4096,1024); gate_up (m == 4096 && k == 4096 && n == 384); down_proj (m >= 128 && k == 192 && n % 128 == 0). M tails and paired shapes must still pass. Occupancy: verify from the exact code object (vgpr_count, sgpr, LDS bytes, scratch/spills); __launch_bounds__ minBlocks is a hint only. Math: blockDim multiple of 64; blocks x LDS <= 64 KiB; VGPR x threads x blocks <= 65,536.

3.3 Accepted architectures (per shape)