v0.55.0

Changes since v0.54.0:

Performance

• real-text H100 board — e2e 4/6 wins, decode 5/6, q35 decode flipped; q9 f16 mirrors per-model gated
• q35 shexp fused sigmoid-dot (+13.7% H100 decode), E4B head-last prime (111x prefill), matmul_pre empty-fallback guard
• g26 expert-down through the int8-MMA expert GEMM — H100 prefill 2901 -> 5042 tok/s (1.74x)

Fixes

• board prompt fox-repeat -> real text; recover uncommitted board evidence
• segment exec-update captures at kernel-class boundaries — q35 graph gate green
• route gemma4 through its dc/graph twins on the device-counter entry points
• decode-batch gate1 multi-seed recalibration; round 45 ledger

Documentation

• H100 board shows e2e tok/s per model — one number per arm, no pp/decode split

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

v0.54.0

Changes since v0.53.0:

Documentation

• memra first screen + full H100 per-model board vs vLLM — honest wins, losses, and movement

Other

• rename: bw24 -> memra (public identity)
• refactor(server): single-class serving — lane subsystem moved out of the public tree

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

v0.53.0

Changes since v0.52.0:

Performance

• 31B joins the Q4_0 f16 mirror lane — H100 prefill 4765 -> 7589 tok/s (1.59x, full-trunk mirror)

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

v0.52.0

Changes since v0.51.0:

Performance

• Q4_0 -> fp16 prefill mirrors — gemma-12B H100 prefill 2.15x (campaign A step 1)

Documentation

• hd512 plan MQA correction + goal-order priority note

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

v0.51.0

Changes since v0.50.0:

Performance

• 8-warp router GEMV twin — q35 H100 decode +8.8% (router was 14.8% of the step)

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

v0.50.0

Changes since v0.49.0:

Performance

• batched prefill FA arms — the per-row decode-kernel loop retired for fresh primes

Fixes

• q4_0 stream-K corrupts ragged-k shapes — kernel contract fix + Hopper battery coverage (#23)
• BW24_REUSE_POOL — parked-cache pool cap becomes a VRAM-budget env (default 2 unchanged)

Other

• Delete .claude/workflows directory

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

v0.49.0

Changes since v0.48.0:

Performance

• continuation dark batch-prime — carried sessions join the phase (d) concat (increment b)

Documentation

• concat increment (b) continuation-primes plan (task #21)
• hd512 gqa-mma decode kernel campaign plan (g12-plain-d1736 mechanism receipt)

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

v0.48.0

Changes since v0.47.1:

Performance

• dark-lane batch-prime — concat FRESH same-lane prefills in phase (d)

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

v0.47.1

Changes since v0.47.0:

Features

• FA4-class causal grid swizzle (BW24_FA_REVX build seam, bit-identical tile remap) + verify-ckpt clone-bytes instrumentation for the ReplaySSM assessment (BW24_SPEC_STATS)
• feat(spec)+test: DSpark marginal-rate verify-window door; serving smoke battery; full-board bench harness

Fixes

• full-board harness — q35 llama arm rides self-MTP (llama rejects trimmed draft heads), qwen plain cells feed real text via BW24_PROMPT_FILE (synthetic-id knife-edge prompts retired from evidence, ledger has the kill-list), per-cell server logs

Documentation

• research-item verdicts (all five measured/assessed) + full-board summary in the next-targets ledger

Other

• test+fix: serve-smoke ALL GREEN wired into local-ci; load_draft panics fixed
• feat(spec)+test: DSpark marginal-rate verify-window door; serving smoke battery; full-board bench harness

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

v0.47.0

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...