art7_src/README.md

# int3 g128 MLP on A10G — sub-4-bit feasibility study (v0-v3), lane parked

Goal: spend the idle PPL budget (2.027 vs 2.42 cap) on the verify path's
biggest line item — MLP gate_up+down = 1.70 GB of the 2.44 GB step (70%).
int3 bit-plane packing (32 weights -> 3 int32 + bf16 g128 scale) = 1.29 GB.

## What works (reusable)
- In-boot requantization from the served int4 checkpoint via identity-probe
  extraction (no quant-format unpacking), 42 layers in ~6s
- MSE scale search: relerr 0.232 -> 0.189
- Bit-plane pack/unpack with exact roundtrip; composed selftests pass at
  rel 0.005 (decode) / 0.004 (prefill)
- fullgraph-safe integration via torch.library.custom_op (dynamo.disable is
  FATAL under vLLM fullgraph compile — v0 finding)
- Single-numerics discipline: decode kernels + prefill dequant-scratch serve
  the SAME values; PPL measures the real model. Partial mode (per-component
  int3/int4) keeps that property.
- Gate architecture: microbench >= threshold else stock fallback — 5 runs,
  zero leaderboard damage.

## The kernel-efficiency ladder (the actual finding)
  v1 row-major packing:                 55.7 GB/s  (uncoalesced lane loads)
  v1 tiled-coalesced + stages:          55.7 -> no change (NOT a load problem)
  v2 BLOCK_K=128, trans-free unpack:    179 GB/s   (3.2x: per-iteration
                                         tl.trans/barrier latency was the wall)
  v3 per-kernel layout+config autotune: gateup 224.8 (bo=64,nw=4,ns=3)
                                         down  125.4 (config-INSENSITIVE)
Full per-config tables in kernel_ladder.txt.

## Conclusion
Beating int4-Marlin time (needs ~>=400 GB/s on packed bytes at M=8) is not
reachable with single-CTA-per-tile Triton GEMV on sm_86: the down projection
(K=10240, 80 CTAs x 80 chained dot-iterations) is dependency-latency-bound
and insensitive to launch configs. The remaining 2x needs split-K partial
accumulation, cp.async double-buffering, and warp specialization — a
Marlin-class CUDA kernel. "No sub-4-bit Ampere kernel" is now a measured
engineering bill, not folklore. PPL cost of int3 RTN (relerr 0.189) remains
unmeasured — no run got past the speed gate, by design.