Skip to content

v0.47.0

Choose a tag to compare

View all tags
@github-actions github-actions released this 30 Jul 00:23
v0.47.0
19d4dd4

Changes since v0.46.0:

Performance

  • cache serve-mode flags, throttle metrics publish to every 32nd tick

Features

  • arch auto-detection + runtime arch guard; post-merge semantic fixes
  • wgmma int8 block-scale GEMM CORRECT at 3.84x the MMQ class
  • wgmma v0 RUNS — 169us = 4.1x MMQ class raw (correctness mapping = last bug)
  • wgmma prefill-GEMM dev harness + v0 kernel — arc opened
  • single-session CUDA-graph serving lane (+34% at B=1)
  • GraphSession step decomposition prof (BW24_GS_PROF)
  • GraphSession — step-wise CUDA-graph decode for serving
  • Q8_0 rp g2 twin — grid fill for sub-wave shapes
  • Q8_0 m=1 rpca — cp.async-staged weight ring (H100 latency fix)
  • Q8_0 mr2 rp m=1 twin — the latency lever (2 rows/warp)
  • Q8_0 split-plane decode mirrors — the H100 coalescing fix
  • BW24_BVAR variant-override seam + serving-regime bench prompts
  • batched GDN state kernels via device pointer arrays
  • lane-aware batched scheduler + native yield gate
  • decode_step_batch v2 — batched GDN projections
  • batched decode step v1 + exactness battery + B-scaling bench
  • Phase A — re-admit portable-PTX tensor cores on Hopper
  • sm_90a boot arch — Hopper rides the portable-CUDA seam

Fixes

  • link libcuda via toolkit stubs dir — GPU-less release/CI machines
  • hoist gdnseq_t/gdnvl_t typedefs out of the mma guard — sm_89 portable build
  • [[maybe_unused]] on dot_i8_16's right_sum — the param is consumed only by the AVX-VNNI arm, and a runner without AVXVNNI fails -Werror=unused-parameter (surfaced on the v0.47.0 push; companion rebuilt clean under the exact CI flags locally)
  • restore the sm_120a graph-door default lost to the branch arch-gate; unified-tree battery evidence
  • graph-lane promotion gated on generation length
  • rpca opt-in — measured -2% on H100 Q8_0 m=1; ledger updated
  • Q8_0 rp defaults to mr1 — mr2 measured negative on H100
  • adaptive dark-prefill chunks + full-tick TPOT estimator
  • interactive cap follows batched capacity, not the round-robin bound
  • judge-prefill starvation — the two demo-battery findings

Documentation

  • unified dual-arch README + flag catalog; recover protocol bench into repo
  • README standing + mechanism prose synced to v0.46 state
  • tree-draft lane design — root-anchored paths, v1 path-duplicated verify, zero new kernels
  • H100 lane shipped — README chapter, FLAGS §7 catalog, CLAUDE.md lane laws; private-repo separation scrub
  • docs+tools: wgmma B-transpose direction confirmed; exact layout needs the ISA table — guess-loop stopped
  • wgmma layout search verdict — B transpose belongs in the smem write, not the descriptor
  • wgmma v0 debug hint — core-matrix smem order, first LBO/SBO candidates
  • task-8 rescoped by clean prime profile — wgmma target is the GEMM (60%), not FA (9%)
  • task-8 prime-profile attempt — tokenwise-path gotcha recorded, FA share still open
  • graph-serving live milestone — 217 tok/s long-gen over HTTP
  • step decomposition — 99% GPU-bound, multi-step replay killed pre-build; in-graph fusion is the lane
  • docs+fix: GraphSession +30.4% measured (233.8 tok/s, gate PASS); embd hold annotated
  • graph-serving design pinned — whole-gen bursts disqualified, GraphSession step API is the arc
  • g2 kept on microbench evidence; graph-serving is the last big m=1 lever
  • docs+tools: per-shape microbench — kernels AT wall; m=1 gap = launch/D2H overhead
  • m=1 ncu stall verdict — rpca port is the next kernel arc
  • tuning-lane ledger — measured verdicts + evidence-ranked frontier
  • Q8_0 split-plane measured results
  • lane-battery demo record — four defects fixed, contract validated
  • serving-regime B-curve (487 tok/s aggregate at B=8, 3.27x)
  • record batched-engine + live lane-server measurements
  • record Phase A measured results (22x prime)
  • H100 serving-lane architecture (sm_90a)

Other

  • arch-gate the H100 promotions off the naked sm_120a build (rig-divergence law)
  • V1 exact-rescale triple-refuted: C7517 compiler-serializes cross-bank GMMA register reads (17x); owner-gate verdict final
  • s8 prefill GEMM verdict: exact route refuted (5.37x rescale overhead), w8a8-class route blocked on owner accuracy law
  • s8 wgmma pairing cracked EXACT (probe_s8.cu): canonical family law complete (bf16/tf32/s8, one byte-geometry)
  • same-session showdown: decode 122.8% (beaten), prefill 73% (stale-denominator correction); s8-wgmma GEMM arc opened
  • serving: graph promotion captures over the primed cache (fixes 3x solo long-prompt regression)
  • hybrid graph door in generate_with: official decode 190.3 -> 220.7 tok/s (+16%, 5/5 x5, bit-identical)
  • graph-lane gate rot: emission off-by-one in graph-decode-gate's eager arm (engine was bit-correct)
  • qb16 mirror-fold: l2-v2 epilogue emits q's bf16 twin (kb16 pattern); standalone bulk cvt dies on both paths
  • ledger: FA3-TMA-ring k45 rebuild refuted by composition; task #22 kernel frontier closed
  • extract cu/wgmma_common.cuh: shared Hopper wgmma toolkit (verbatim hoist from hybrid.cu + tf32 kind)
  • serving: raise prime-batch T cap 320 -> 2048 (old crossover refuted on wgmma vl cores)
  • k45 refutation #12 (q prefetch, 91.0->92.2us): kernel micro-frontier closed on measurement; next frontier = cross-request prefill batching scheduler
  • K3 solve Route A refuted on perf (51.0 vs 18.9us f32, correct at 4.2e-4): f32 substitution is K3's structural optimum
  • tf32 wgmma canonical pairing cracked: bf16-analog layout w/ 4-elem kk groups, same desc constants (probe MATCH 7.9e-5)
  • v6 exchange-free k45 refuted (91.0 -> 114.1us): duplicate full-k doubles wgmma operand traffic; ledger updated
  • ledger: solve32 ILP split refuted (55.1 vs 48.8us + instantiation bit-identity break); mechanism recorded
  • ledger: K3 solve32 arc design note (inverse-product route vs SIMT latency route)
  • wgmma-fused vl twins: batched prime +3.8% interleaved x5, all gates green
  • K2-wgmma: A/P GEMMs on tensor cores inside BW24_GDN_WGMMA, lane +2.56% cumulative x5
  • BW24_GDN_WGMMA: K4+K5 fused wgmma engine seam, promoted default-ON hopper
  • K4+K5 fusion v5: in-band (Y 1.07e-2, state 1.03e-2, O 1.08e-2), 91.3us fused
  • ledger: K4-wgmma v3/v4 arc verdict - 70.3us floor, wash vs mma; pivot to K4+K5 fusion
  • K4-wgmma v4: cp.async W staging (real rows only, pad zeroed once): 73.7 -> 70.3us, in-band
  • K4-wgmma v4: U register prefetch under staging/step-A shadow: 82.3 -> 73.7us, in-band
  • K4-wgmma v4: 2 warpgroups x i-halves, 116.6 -> 82.3us, in-band (Y 1.04e-2, state 7.8e-3)
  • K4-wgmma v3: col-32 grid (128 CTAs) + 16B staging vectorization: 210.9 -> 116.6us, in-band unchanged
  • K4-wgmma v2: full chain in-band (Y 1.04e-2, state 7.8e-3, band 3e-2) at 210.9us vs mma 68.3us
  • GDN-wgmma harness v1: K4 step-A tile (Y = U - W.M) MATCH 9.2e-5 first try - the FA3 descriptor toolkit transfers; chunk chain (M accumulator + smem round-trip + step B) is the next assembly
  • K4-wgmma design resolved: M smem round-trip per chunk (on-SM, noise), m64 with pad, per-head B blocks head packing
  • GDN chunk stack wgmma arc opened: 12.4ms at single-digit %-of-peak, FA3 toolkit maps 1:1, target = lane crossing (~62ms vs vLLM 66)
  • K4 Ssnap column-block layout: contiguous 8KB CTA writes, streams identical, honest-small; lane map exhausted above the ms floor
  • FA3 batched vl twin: per-seq tensor maps by value + grid.z, +0.7% batched 5/5; serving FA fully wgmma/TMA
  • de-broadcast q/k: [T,16,128] compact layout, +1.4% lane / +1.7% batched, byte-identical
  • Lt m=2048 sweep: GEMM ceiling reconfirmed at lane shape; parity map pinned (de-broadcast + gaps + Ssnap ~= 3ms of the 4ms distance)
  • device embed gather: +12.4% official lane (29,050 tok/s = 93.6% vLLM, TTFT 70ms), streams MATCH, battery green
  • FA3 v11 SHIPPED: full-TMA staging + trans_b PV, 205us harness (4.8x), official lane +7.6% -> 25,924 tok/s (83.5% vLLM), TTFT 79ms
  • FA3 v10 engine-integrated + PROMOTED default-ON hopper: official lane 24,675 tok/s (79.5% vLLM), TTFT 83ms
  • FA3 v10 WINS: 883us vs shipped 993us (+12.5%), correct at all T
  • swizzle pairing cracked: TMA SWIZZLE_128B + wgmma desc mode=1/SBO=1024/LBO-ignored, k-slice = base+j*32; all v10 primitives proven
  • TMA foundation proven first-run (probe_tma.cu: tensor map + bulk-tensor 2D + mbarrier expect-tx byte-exact); FA3 v10 = swizzle pairing inside the v5 shape, all primitives individually proven
  • add+norm f16out fusion: 3 trunk sites, byte-identical streams, honest-small (+0.17% batched 4/5, ~+0.6% arithmetic single-seq); kept as strict-less-work
  • FA3 coda: setmaxnreg-C7515 interaction pinned (its presence costs more than the spills it prevents); 6 producer variants refuted; v5 definitive for this generation
  • FA3 v9: mbarrier producer correct (cp.async.mbarrier.arrive.noinc lesson) but == v8; real wall = C7515 wgmma serialization in the 240-reg region; arc ledger complete at 9 versions, v5 parity floor
  • FA3 v8b/c: setmaxnreg contract cracked (probe-verified: entry cap + pool redistribution + region allocation; placement matters); producer shape refuted at this geometry (barrier serialization 1829us) - TMA/mbarrier is the remaining piece
  • FA3 v8a: producer-WG-without-setmaxnreg refuted (launch_bounds 170-reg ceiling spills the 254-reg consumers); barrier choreography proven correct, setmaxnreg compile contract pinned as the next unknown
  • fp8-for-Q8_0 refuted (per-block-scale wall, F8-origin only); FA3 v8 design pinned build-ready (TMA+mbarrier+producer+setmaxnreg choreography)
  • FA3 harness v7: split-D refuted (WG1 starvation + smem cap); arc state = v5 parity with all foundations proven, remaining = full producer/consumer+TMA+setmaxnreg choreography (priced)
  • FA3 harness v6 refuted (overlap reorder loses to natural 2-WG interleave); ncu(v5): 254 regs = the same register disease -> split-D + setmaxnreg is the canonical next shape
  • FA3 harness v3-v5: full wgmma FA kernel correct at all T (online-semantics ref); v5 (2 warpgroups, shared K/V ring, int4 cp.async) = 999us at T=2048 -- PARITY with the shipped mma kernel (993us) before overlap tuning
  • FA3 harness v2: full FA tile on wgmma MATCH 1.2e-5 (QK^T -> causal softmax -> P bf16 restage -> PV 4x m64n64k16 -> normalize); ref aligned to kernel P-rounding semantics
  • FA3 arc: wgmma bf16 QK^T tile proven (harness v1 MATCH 1.75e-4)
  • mirror folds: producers emit the K4 bf16 operands (l2 v2 -> kb16, K3 store -> wb16); +1.2% official / +2.8% batched
  • l2 prefill v2: warp-per-row float4, +2.3% official lane / +2.1% batched (gated numeric config)
  • conv-fuse: carried-ring conv+SiLU+repack in one pass, +9.7% official prefill / +7.8% batched, bit-identical
  • varlen attn pre-FA: split/norm/rope/append in 4 vl launches, +2.8% batched prime; attn split copies gone
  • varlen FA: one launch for all fresh seqs, +5.7% batched prime interleaved 5/5
  • batch-trunk f16out + batched lm_head: +4.7% batched prime interleaved 5/5
  • serving policy: BW24_PRIME_BATCH default 4 -> 6 (varlen core moved the concat sweet spot; 16501 vs 15950 tok/s at T=152)
  • varlen GDN increment 3: full varlen core (prep+mirrors+tail), +13.1% batched prime interleaved 5/5
  • varlen GDN increment 2: K1-K3 varlen (full K1-K5 vl chain), +5.5% batched prime interleaved 5/5
  • varlen GDN increment 1: K4/K5 vl pair, +2.3% batched prime interleaved 5/5
  • batch-prime scatter/gather elimination: +6.1% batched prime interleaved-verified 5/5
  • f16out epilogue fusion: +1.6% pp512 interleaved-verified 5/5, bit-identical; big-segment arc closed by gap anatomy + launch probe
  • piecewise 3-5 corrected verdict: small segments REFUTED under interleaved A/B (earlier +0.9% was clock drift); segments opt-in, core-split foundation kept
  • piecewise increment 4 probe: S-mid-via-copy refuted (-0.6%), reverted to S-glue baseline; core-contract refactor pinned as the gating increment
  • piecewise increment 3: first segment LIVE, pp512 crosses 20k (+0.9%, bit-identical)
  • piecewise increment 2: ffn boundary slabs (gate/up/ffn_out) + _into wiring
  • piecewise increment 1: try_f16_gemm_pre_into (GEMM into resident slab views; FFI unchanged)
  • piecewise prime graphs: build-ready segmentation design (cache-free segments, eager Lt/K4/append/ring, slab inventory, _into variants, bit-gate plan)
  • prime activation slabs: landed (piecewise foundation, frozen Lt addresses), third launch-diet neutral, all gates green
  • vLLM decomposed on-box: prefill 31.0k today (bw24 = 64%), decode lead widened to 107%; PIECEWISE graphs = the unblocked gap-floor route
  • cutlass int8 probe: 569-780 TOP default config, deterministic; vLLM-35k arithmetic now suspect - nsys decomposition running
  • CUTLASS deterministic-GEMM probe: determinism CONFIRMED, rate ceiling 0.7x Lt -> graph reclaim refuted at current rates (7-config sweep)
  • task #14 closing verdict: prime graphs blocked by Lt address-variant numerics (mechanism identified via 7-experiment kill-list)
  • task #14 defect hunt: five hypotheses killed by experiment, live lead = graph-baked shared f16 scratch
  • task #14 gate: padded replays GREEN end-to-end, copy-out byte-exact; ONE pinned defect at exact bucket length
  • task #14 PAD-PROOF GREEN: padded prime graph bit-identical to eager (bucket 512/true 300, maxdiff 0.0)
  • task #14 pad-proofing pieces 1+3: gdn_pad_mask + row_gather_dev kernels + launchers (design-v3 insertion pending)
  • task #14 design v3: copy-out replaces pointer table (25-50us vs 2.3ms saved, zero kernel/node surgery); pad-proofing insertion points pinned
  • task #14 SMOKE GREEN: prime graph captures (13ms) + replays bit-identical at +10% (23.3ms vs 25.7)
  • task #14 implementation state: capture-smoke constraints audited, increment order pinned
  • task #14 design v2: prime graphs tractable - pad-to-bucket absorbs dynamics except 4 device-length spots; ~1.3KB pointer table
  • task #14: norm+cvt fusion landed (bit-identical, neutral) - gap floor proven launch-count-insensitive
  • task #13 increment 3: worker batch-prime (+21.6% serving throughput, 98% batch coverage)
  • task #13 increment 2: prime_cache_batch + gate ALL GREEN, regime mapped (+80% at B=8 T=64, crossover ~T=320)
  • task #13 increment 1: mixer proj/core split + matmul_group_multi (bit-neutral, gates green)
  • task #13 design: cross-request prefill batching (concat GEMMs, per-seq views, gate plan)
  • task-9 kernel sweep CLOSED: wall audit complete, remaining gap attributed to scheduler arcs (tracked)
  • round-9 wall audit: conv float4 neutral (reverted, scatter-bound), de-broadcast option mapped, graphs = the remaining arc
  • FA cp.async 2-stage ring on bf16 K/V tiles: pp512 19886, bit-identical
  • FA BF16-KV staging: +11% pp512 (19718), bit-identical, default ON
  • GDN K5-MMA coupled pair: pp512 17786 (+105% night), fp16 coupling channel after bf16 pin failure
  • promote GDN K4-MMA default on Hopper lane: state-carry battery green, pp512 17240 (+99% cumulative)
  • GDN K4-MMA engine integration: opt-in seam, dual kernel-check pins, +3.5% pp512 available
  • GDN K4 bf16-mma kernel PROVEN in harness: 68.3us vs 119.4 f32 (1.75x), correct
  • W8A8 per-row arc refuted by probe: int8 IMMA nets 0.87-1.04x vs shipped fp16 mirror
  • FA W2 probe: 2-warp CTA variant refuted at T=512 (-7%), seam kept for the small-chunk serving hypothesis
  • ledger: Phase D extraction closed (4 crates live, cuda-util refuted as speculative)
  • extract bw24-kv (Phase D): dual cache + KV format policy behind the KvDev seam
  • extract bw24-validate (Phase D): shared validation-protocol core + FA occupancy seam verdict
  • ledger: FA prefill ncu diagnosis - 6.25% occupancy (255 regs by design), 67% long-scoreboard on sync K/V staging; fix = pipelined staging (bit-identical)
  • prime grind round 2 ledger: +7% landed (uninit diet, float4 wave), probes recorded, remaining arcs ranked
  • elementwise float4 wave: silu_mul, add, f16 convert (+2.7% pp512 combined)
  • prime launch diet: 47 zeros->uninit on full-overwrite prefill buffers (+4.5% pp512)
  • matmul_group: convert-once f16 dispatch for shared-activation projection groups
  • FP16-mirror prefill: +80% pp512 (8674->15626 tok/s), promoted default on Hopper lane
  • wgmma prefill GEMM: engine integration + honest per-shape verdict (exact path refuted for default)
  • docs+tools: wgmma B-transpose direction confirmed; exact layout needs the ISA table — guess-loop stopped
  • docs+fix: GraphSession +30.4% measured (233.8 tok/s, gate PASS); embd hold annotated
  • tools+docs: persistent-CTA refuted; m=1 matvec per-shape survey complete
  • tools+docs: ldcs-vs-ldg per-shape A/B — streaming exonerated; wave-tail is the wide-shape suspect
  • docs+tools: per-shape microbench — kernels AT wall; m=1 gap = launch/D2H overhead
  • refactor: extract host sampler into shared bw24-sampling crate
  • tools: validate-h100.sh — one-command gate battery + perf record
  • refactor: extract lane machinery into shared bw24-lanes crate

Boards + reproduction artifacts: https://huggingface.co/Avifenesh/bw24-bench · full experiment log in research/tune-data/

Assets 6
Loading