diff --git a/.agents/NOW.md b/.agents/NOW.md index 2a1f95a3c..3ac09ebee 100644 --- a/.agents/NOW.md +++ b/.agents/NOW.md @@ -26,7 +26,7 @@ Work: exact-chunks on main `1ce0d662b`; sm_120 measured at `3d2581551`. | CPU levers (`QUANT-GGUF-CIQ-GEMM`) | Profile DONE: decode **47% threadpool sync**, prefill **~39% paged attn**. **G5 not next** | Parakeet encoder; attn dtype hoist | | Supported-models list | **LANDED**: FEATURES arch table CI-bound (33 archs) | — | | `/v1/videos` OpenAI shape | **MERGED** (#71): Sora `model`/`size`/`seconds` + `GET /{id}/content` | `row/SERVE-VIDEOS-REFS` PR open: reference conditioning | -| Vulkan 27B decode | **4.285 vs 4.35**; host ~3 ms/tok, pipeline -1.41 ms (5/5), spin rejected | NOT met: 900 barriers/tok | +| Vulkan 27B decode | **MET: 4.36 vs llama.cpp 4.35** (7 legs, main). Barriers -19.8%/tok, GPU -1.09 ms, e2e 8/12, OFF | Re-measure before flipping `VT_VULKAN_SMART_BARRIERS` | | `BACKEND-ROCM` | **(b) fix in; #140 gfx1201 hipBLAS + Gemma-4 MoE landed (contributor, authorship-preserved); W0 green 4 archs** | compile + M2 ([spec](specs/rocm-unified-memory-b.md)) | | TP spike #287 (PR #143) | **TP-W1 LANDED**: rank-group table + TP handle (6/6); DSR leak FIXED (unblocks #127/#154/#155) | TP-W2 (linears + loader) | | Release | **ACTIVE; W5 19/19+10/10; contract 30/30** | #141; artifacts pending | diff --git a/.agents/benchmark-record.md b/.agents/benchmark-record.md index 07d702813..010efb485 100644 --- a/.agents/benchmark-record.md +++ b/.agents/benchmark-record.md @@ -16781,3 +16781,179 @@ on the page and the history is preserved here byte-for-byte, per POL-EVIDENCE-PRESERVE. Nothing was rewritten or dropped. **BINDING 2026-08-04: 87% of vLLM (37.55 vs 43.10, SAME-TOOL nsys both engines); the whole +3.1 ms/step is the bf16 M=1 GEMV bucket (2/3 o_proj, ~196-204 vs 139 us/call, identical `gemvx` kernel); attention/MoE/glue tied or ours-ahead. Invocation match (bf16-out `cublasGemmEx`) A/B'd = WASH, ruled out; ROOT CAUSE FOUND 2026-08-04 (`VT_LAGUNA_RESIDENT_BF16W`): the bf16 projections read UNIFIED/ATS host memory, not `cudaMalloc`'d device memory — staging them device-resident (byte-exact ids) gives 38.8→44.6 tok/s (o_proj 194→131, lm_head 2410→1620 us/call), parity+ vs vLLM 43.1; **default-ON** (flip smoke-verified: canonical byte-exact ids, 44.6 clean-median). Earlier ceiling/diffuse verdicts below were cross-tool artifacts.** **REAL vLLM BAR ESTABLISHED (2026-07-31, `CLAIM-LAGUNA-VLLM-NVFP4`): FIRST-EVER vLLM Laguna run** — prior numbers (incl. the correctness oracle) were all llama.cpp, never vLLM. vLLM on official `poolside/Laguna-S-2.1-NVFP4` (single GB10, greedy, eager, MARLIN backend forced via `VLLM_TEST_FORCE_FP8_MARLIN=1` because the auto-default `FLASHINFER_CUTLASS` needs an absent `nvcc`): **~18.8 tok/s** (64-tok steady) — a LOWER bound. Our GGUF-Q4_K engine = 7.7 tok/s (vLLM ~2.4×); llama.cpp GGUF = 27.8 (still fastest at batch-1). llama.cpp is now a labeled SECONDARY "beat best-in-class GGUF" note; vLLM-NVFP4 is the headline bar. TRUE apples-to-apple still owes OUR NVFP4 Laguna forward arm (same tensor-core path as 27B/35B) — bring-up W-plan SPEC'D in `.agents/specs/laguna-nvfp4-arm-2026-07-31.md` (~85% reuse of the 35B NVFP4 W4A4 MoE infra + a name-map; bf16 attn/dense + fp4 experts; N1-N5 bricks, DGX-gated). **N1-scaffold LANDED (2026-07-31):** additive `LagunaMoeWeights.experts_{gate,up,down}_fp4` + `shared_{gate,up,down}_fp4` (`Nvfp4Weight`, mirror qwen3_5), dead until the N1 loader; CPU build clean + `test_laguna_scaffold` 8/8·167 unchanged. **N1b loader IMPLEMENTED (2026-07-31, build-verified):** `LoadLagunaForCausalLMWeights` (`laguna_weights.cpp`) replaces the `VT_CHECK(false)` stub — resolver + per-layer `LoadBf16Direct` (attn/dense/norms/embed/lm_head/router/shared-expert) + F32 `e_score_correction_bias` + `LnLoadCtNvfp4Raw` W4A4 experts. Name-map + dtypes VERIFIED against the real `poolside/Laguna-S-2.1-NVFP4` index (router `mlp.gate` BF16, bias F32, experts W4A4, shared-expert BF16). **N1b RUN-VERIFIED (2026-07-31):** loader round-trips a synthetic NVFP4 checkpoint byte-identically (`test_laguna_nvfp4_loader` 2/2·29; full detail in the benchmark record). **N2 FORWARD-BRANCH LANDED + CPU-GATED (2026-07-31):** `LqGemmNvfp4Fp4` (per-expert TRUE-W4A4: `ScaledFp4Quant(input_global_scale_inv)`→`MatmulNvfp4Fp4(alpha)`, unified-memory pattern like `LqGemm`) + `LagunaFfnBlock` branches on `fp4=!experts_gate_fp4.empty()` (routed experts fp4; keep-quant grouped fast-path gated off `!fp4`; bf16 attn/dense/router/shared-expert/lm_head unchanged) + both `LagunaForwardGguf{,Cached}` guards relaxed to `has_gguf_weights||has_nvfp4_weights`. **CORRECTION:** routed experts are W4A4 ⇒ per-expert `MatmulNvfp4Fp4`, NOT the grouped W4A16 `MoeGroupedGemmNvfp4` (grouped W4A4 deferred to N5 speed). `test_laguna_nvfp4_loader` 3/3·61 (added a forward run-gate: fp4 MoE branch runs through the real `LagunaForwardGguf` → finite+deterministic logits + routed-experts-consumed); `test_laguna_scaffold` 8/8 unchanged (GGUF byte-identical). **N3 DRIVER LANDED + CPU-SMOKE-VERIFIED (2026-07-31):** `examples/laguna_gen` auto-detects a safetensors DIRECTORY (→ NVFP4: `LoadHfConfig(config.json)` + `LoadLagunaForCausalLMWeights` + `LagunaForwardGguf{,Cached}`) vs a `.gguf` FILE (→ keep-quant), sharing the greedy loop; `--token-ids` bypass the tokenizer for the id-vs-golden gate. Verified on a synthetic NVFP4 dir with a REAL config.json (exercises the `LoadHfConfig`→`ParseLagunaParams` seam the loader test bypassed) → `has_nvfp4=1`, KV-cache decode runs finite. **N4 RAN on GB10 (2026-08-01) — the arm works end-to-end; correctness coherent+near-tie, speed 120× off.** git-archived `84fab587` → clean CUDA build (`121a`) → `laguna-gen --gpu` on the real 67 GiB `ckpt` with vLLM's exact prompt ids injected (`2,785,9626,377,15360,395`, captured via the HF tokenizer). Two GB10 memory fixes landed to run: release the mmap'd shards after the loader's memcpy-copy (114→67 GiB RSS), and create the CUDA context BEFORE the load (the 67 GiB reclaimable page cache otherwise starves `cudaStreamCreate`). **Correctness:** ours `22345 83 350 71070 395 340 9626 372 1703 …` vs golden `22345 83 290 350 674 330 5541 966 340 9626 377 15360 …` — **first 2 tokens match vLLM exactly**, then near-tie divergence; coherent ("France is" = 9626/377/15360; shares golden vocab). EXPECTED: our TRUE-W4A4 (fp4 activations) vs the MARLIN golden's W4A16 (bf16 activations) — different precision, not a bug. **Speed: 6.34 s/tok (0.16 tok/s), prefill 17.3s — ~120× slower than vLLM 18.8.** ROOT CAUSE (source-confirmed): `LqGemmNvfp4Fp4` uses the generic `vt::MatmulNvfp4Fp4` = the hand-written EMULATION CUDA kernel, NOT the cutlass sm120a fp4 tensor-core path the 27B/35B W4A4 use (`MatmulNvfp4Fp4DirectD`); + per-expert loop + per-GEMM host sync + no device residency. **nsys (2026-08-01) trace-confirmed + refined:** only 2 GPU kernels — `MatmulNvfp4Fp4Naive` = 99.3% of GPU time + fp4-quant 0.7%; GPU busy only ~18% of wall. NO bf16 GEMM on the GPU ⇒ `LqGemm`'s bf16 branch runs the host `MatmulNK` reference on the CUDA queue (attention/dense/router/shared/lm_head are CPU-bound, ~4.8 s/tok) — a second lever the source scan missed. **N5 LEVER #2 LANDED (2026-08-01) — 16× decode.** Routed the bf16 tower (attention/dense/router/shared/lm_head) off the host `MatmulNK` onto the GPU (`LqGemm` bf16 branch: `vt::CastBf16` the small activation + `vt::MatmulBT` bf16×bf16→f32, weight stays bf16 — no per-token `ReadF32` of `lm_head [100352,H]`): **decode 6.34 → 0.39 s/tok (16.3×; 0.16 → 2.56 tok/s), prefill 17.3 → 2.24s**; coherence preserved (near-tie). CPU path unchanged (run-gate byte-identical). **N5 LEVER #1 LANDED (2026-08-01) — native fp4 tensor-core, another ~2×.** The engine's native sm120a fp4 tensor-core MMA (`MatmulNvfp4Fp4Native`, `mma.sync kind::mxf4nvf4`) reads the same linear scale layout `LqGemmNvfp4Fp4` produces — it was gated OFF behind `VT_NVFP4_FP4_NATIVE`; the Laguna driver now defaults it ON (scoped; 27B/35B untouched). **decode 0.39 → ~0.20-0.24 s/tok (~2×; ~4.2-5.0 tok/s)**; coherent (byte-identical ids to the emulation path — numerically equivalent), first token matches the golden. **Cumulative N5: 0.16 → ~4.5 tok/s (~28×), now ~4× from vLLM 18.8.** **Device-resident MoE block LANDED + MEASURED (2026-08-01, `LagunaMoeResidentFp4`, `VT_LAGUNA_RESIDENT_MOE` default-ON):** the whole token's routed experts as ONE async device chain (fp4-quant→GEMM gate/up, `MoeSiluMul`, →down stacked, ONE `MoeCombine`), draining once vs ~Pk×3 syncs. **Speed EAGER-NEUTRAL (0.20 s/tok)** — empirically confirms the ds4 precedent (per-op syncs overlap GPU compute; wall is GPU-serial-bound; the graph is the payoff). **CORRECTNESS WIN: golden-token match 2 → 13** (the device `MoeSiluMul`/`MoeCombine` mirror vLLM's fused MoE faithfully). Lands default-ON (better correctness, no speed cost, graph prerequisite). **CORRECTED CEILING (from the measured state): a perfect decode graph caps at ~5.9 tok/s** (GPU already ~87% busy at 0.20 s/tok), still 3.3× short of vLLM 18.8 — the graph is necessary but NOT sufficient; the remaining 3.3× is KERNEL EFFICIENCY (native fp4 MMA ~302µs/M=1 expert GEMM vs vLLM's tuned cutlass sm120a fp4 + fused norm/quant/silu). Parity = TWO campaigns: (A) device-resident+graph → ~5.9; (B) cutlass DirectD experts + fused ops + M=1-tuned GEMV → the rest. **CAMPAIGN-B FIRST BRICK LANDED (2026-08-01): coalesced M=1 fp4 GEMV** (`MatmulNvfp4Fp4Gemv`, one warp/column, coalesced weight-row reads, `VT_NVFP4_FP4_GEMV` default-ON) — same-binary A/B: **decode 0.20 → 0.15 s/tok (1.33×; → ~6.7 tok/s), prefill 1.14 → 0.86s**, coherent+near-tie. **Cumulative this session: 0.16 → ~6.7 tok/s (~42×), now ~2.8× from vLLM 18.8.** (ILP variant `kCpw=4` measured SLOWER — 0.21 s/tok, occupancy loss > activation-reuse gain — reverted to `kCpw=1`; kernel kept templated as a re-measurable knob.) **ncu of the GEMV (sudo): sm__throughput 35-71%, DRAM n/a — COMPUTE/LATENCY-bound, not BW-bound.** Corrects the earlier "~6× BW → ~16-17 tok/s" estimate: the next GEMV lever is HARDWARE fp4 dequant (`cvt.e2m1x2`), not more bandwidth. Parity (18.8) is a multi-brick campaign (decode graph + fused norm/quant + hardware-dequant GEMV), not one more kernel. **B0 hw-fp8 SCALE-decode: MEASURED NEGATIVE, reverted (2026-08-01, `ab7a1c1e`).** Replacing the GEMV's per-byte software fp8-e4m3 group-scale decode (`F8E4M3ToF32Dev`/`ldexpf`) with hardware `cvt.rn.f16.e4m3` (`__nv_fp8_e4m3`→float) is bit-exact (ids byte-identical on the real ckpt) but paging-immune ncu shows it NEUTRAL-to-slightly-WORSE (grid768 41.2 vs 41.9µs tie; mean 53.6 vs 49.4µs) — GPU `ldexpf` is a cheap exponent-bit add, not a libcall. NOTE this is the fp8 SCALE decode, NOT the fp4-e2m1 WEIGHT dequant (the `kE2M1` `__constant__` LUT); the LUT→arithmetic/`cvt.e2m1x2` weight-dequant is a SEPARATE still-open lever (spec brick B1). Also: end-to-end wall-clock is unusable for kernel A/B here (67 GiB unified reload swings TPOT 0.16↔1.08 s/tok run-to-run) — kernel-duration ncu is the only honest anchor. **★ B2 SCOPED + DE-RISKED (2026-08-01, zero-DGX) — the real 18.8 lever:** vLLM's 18.8 bar is MARLIN W4A16 (`VLLM_TEST_FORCE_FP8_MARLIN=1`), which is LOW-M-optimized (decode-correct, unlike a tensor-core W4A4 GEMM that wastes M=1 tile rows). The engine already ships the EXACT kernel `vt::MoeGroupedGemmNvfp4Marlin` (1:1 lift of vLLM `moe_wna16_marlin_gemm`) + shared `MarlinRepackExpertWeight`, and qwen3_5 (27B/35B) already routes its NVFP4 experts through it (default-ON `VT_NVFP4_MARLIN`, 16/16-vs-oracle, +22% gate/+80% decode) via `BuildMoeMarlinResident`. So B2 = mirror that for `LagunaMoeWeights.experts_*_fp4` (a `BuildLagunaMoeMarlinResident` reusing the shared repack + route `LagunaFfnBlock`'s fp4 branch to the Marlin grouped GEMM, GEMV kept as the `=0` escape hatch) — pure reuse, no new kernel, matches vLLM's exact W4A16 numerics. **B2 IMPLEMENTED (2026-08-01, `3c49ef37`) — COMPILES CLEAN on GB10 sm_121a, runtime bug pending.** `LagunaMoeResidentMarlin` + `BuildLagunaMoeMarlinResident` (laguna.cpp, `#ifdef VT_MARLIN_NVFP4`) reconstruct the MoE Marlin path over the SHARED `dense_nvfp4::Dev`/`DBuf`/`ResidentNvfp4` + shared `vt::cuda` Marlin repack/align ops + `vt::MoeGroupedGemmNvfp4Marlin`; SACRED 27B/35B path BYTE-UNTOUCHED; gated `VT_LAGUNA_MARLIN_MOE=1` **default-OFF** (zero regression to the default GEMV path). Compiles clean on the full CUDA build. RUN: loads OK (48 layers, 256 experts) but the FIRST FORWARD device-faults silently on the Marlin path — a layout/param bug (suspects: `MoeCombine` bf16-in/f32-out dtype, the down-GEMM reusing the gate/up align, or the fp4-original free omitted → mem ~doubles). NEXT: `compute-sanitizer` localize → fix → near-tie vs the vLLM-Marlin golden + kernel-duration ncu → flip default-ON. Default path unaffected. **UPDATE (`22d6e146`): added the qwen3_5-style fp4-original free after repack** (device transients + host bytes; peak was ~3× the expert tower → past the 119 GiB pool → null-alloc → silent fault the likely cause); compiles clean. The runtime gate stayed INCONCLUSIVE this session (contended/orphaned processes on the shared box, no captured ids) — rerun on a clean uncontended session, compute-sanitizer if it still faults. **★★ B2 VALIDATED on GB10 (2026-08-01, with the mem-free fix): RUN_EXIT=0, coherent, first 13 generated tokens MATCH the vLLM-Marlin golden EXACTLY** (`22345 83 290 350 674 330 5541 966 340 9626 377 15360 81` — the best Laguna-NVFP4 correctness yet, W4A16 matching vLLM's config). **Steady-state decode 0.10 s/tok = ~10 tok/s** (steps 10-17 all 0.10; the TPOT-0.56 average is warmup-polluted — the DevicePool warms over ~9 decode steps then reuses). vs the GEMV path's 6.7 tok/s = **~1.5× faster; the gap to vLLM 18.8 closes from ~3× to ~1.9×.** Memory flat (7.9 GiB host RSS — the fp4-original free worked; it also fixed the first-forward fault). Still `VT_LAGUNA_MARLIN_MOE=1` default-OFF. TO DONE: move the lazy Marlin-resident build (216s first-forward, 48L×256E repack) to model-LOAD time → clean warm A/B + ncu → flip default-ON → matrix/roadmap. Remaining ~1.9×: vLLM graphs its decode (ours still eager) — decode CUDA-graph is the next lever. **REPRODUCED 3× (reproduction gate MET): GB10 runs deterministic — first 18-20 tokens byte-identical, steady-state 0.10 s/tok confirmed each — so the ~10 tok/s + golden-match is gated, not a single sample.** **#234 item (1) — load-time resident-build LANDED (`LagunaBuildMarlinResidents`, called from the example after load; mirrors vLLM process_weights_after_loading): builds all 48L×256E Marlin residents at LOAD so the repack is not a first-token TTFT spike. Fixed an anon-namespace linkage bug (public fn was defined with internal linkage → moved outside the anon namespace); BUILD CLEAN + links on GB10 sm_121a, default-OFF. Runtime prewarm-fires-at-load timing UNVERIFIED this session (repeated ssh-drops ate the run capture); the forward's lazy build is the validated fallback so it cannot regress. Owed: one clean run to confirm the build moved to load + then flip default-ON.** **★★ DONE (2026-08-01): Marlin is now the UNCONDITIONAL DEFAULT (`LagunaMarlinMoeEnabled` default-ON; `=0` is a code-level A/B opt-out no user needs) — "it just works" with NO env. Confirmed in a no-env GB10 run captured via tmux: `MARLIN residents built at load in 238.4s`, prefill 14.78s (build moved OUT of first-forward), golden-matching ids, steady-state 0.10 s/tok = ~10 tok/s (4th reproduction), RSS ~5-8 GiB. So a default Laguna-NVFP4 load on GB10 gets vLLM's own W4A16 Marlin decode (~10 tok/s, ~1.9× from vLLM 18.8) with zero flags. The 238s load-time repack is a one-time cost (mirrors vLLM process_weights_after_loading); optimizing its 48×256 per-expert sync count is a follow-up. Residual to 18.8 = decode CUDA-graph (deferred; user refocusing on DeepSeek next).** Post-lever-1 nsys: the remaining ~4× is HOST-SYNC-bound — 22,115 `cudaStreamSynchronize` (78.6% of API time, ~2,760/token, the per-GEMM `DrainQueue`), GPU kernels fast. Remaining levers: grouped W4A4 MoE (design input: `vt::MoeGroupedGemmNvfp4` is W4A16, so true-W4A4 grouped needs a new fp4×fp4 op or the `use_a16` mode + expert-stacking — needs a spike), device-resident decode (RECOMMENDED — the current forward is host-style so every GEMM drains; keep activations on-device, drain once/step; reuse qwen3_5's `Dev`/`Nvfp4Dev`/`ResidentNvfp4`/device-SwiGLU machinery; kills the 22k syncs; converges with the pending GGUF #228 and lifts both quant paths), decode CUDA-graph. Binding number needs a clean 2-3× re-run. See `docs/BENCHMARKS.md` + the spec N5 plan. See `docs/BENCHMARKS.md` `CLAIM-LAGUNA-VLLM-NVFP4`. Prior W7 nsys attribution: host-orchestration-bound, levers ranked (spec `laguna-s21-w7-speed-2026-07-31.md`, ledger `CLAIM-LAGUNA-W7-SPEED`). Prior RUNNABLE + FAST DECODE (W6, 2026-07-31): a per-layer K/V cache + single-token incremental decode replaces W5's O(n²) STATELESS full-recompute — TOKEN-IDENTICAL (byte-equal ids, md5 match, == the W5 golden) and 5.05× faster per token: decode 3.33 → 0.66 s/tok on the real 3-shard UD-Q4_K_XL GGUF (GB10, `--gpu`, keep-quant), same "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\"…". `LagunaKvCache` (mirrors `DeepseekV4KvCache`, MLA-latent → GQA multi-head K/V) caches post-QK-RMSNorm/post-RoPE K + raw V at f32 (bit-exact by construction: RoPE/QK-norm are position-only and attention is causal). MIXED attention handled per-layer: 12 GLOBAL layers grow the cache unbounded (full causal); 36 SLIDING-WINDOW-512 layers EVICT the oldest rows beyond the 512 window (gemma2/3 `is_sliding`), capping their K/V. `LagunaForwardGgufCached` + shared `LagunaAttention`/`LagunaFfnBlock` helpers used by BOTH forwards (identical float ops — the recompute path's ids are unchanged after the refactor); `examples/laguna_gen --stateless` forces the W5 recompute for the A/B gate. No cache bug: bit-exact on the first run. Next speed: grouped-expert GEMM + device-resident decode (both in-tree from ds4). See `.agents/specs/laguna-s21-w6-2026-07-31.md`. Prior RUNNABLE (W5, 2026-07-31): our engine greedy-generates COHERENT text on the REAL 3-shard UD-Q4_K_XL GGUF (GB10, keep-quant). `laguna-gen` "The capital of France is" → " Paris.\n\nThe user is seeking a detailed explanation of the concept of \"cultural capital\" as developed by French soci…" — the FIRST token is "Paris.", matching the llama.cpp-Poolside reference on the identical bytes. Multi-shard GGUF reader (LagunaGgufCtx routes each of 814 tensors to its shard; shard-1 = header only) + keep-quant tower (attn/dense/shared/experts/lm_head stay Q8_0/Q4_K/Q5_K COMPRESSED, consumed via `vt::MatmulBT`; norms/router/bias/embed → f32) + `LagunaForwardGguf` (the f32 composition with the ~9 GEMM sites swapped to keep-quant Gemm/GemmRowSlice, ds4 precedent) + `examples/laguna_gen`. Real GGUF metadata verified: dual-RoPE freq_base 500000/10000, dims 64/128, YaRN factor 32, sigmoid ungrouped-noaux router (scale 2.5), per-layer Q-head [48 global/72 sliding], per-head softplus out-gate, QK-RMSNorm. Load 20.6s, peak 71 GiB (fits 119 pool). Prior W4 IN PROGRESS (2026-07-31): 73.4 GiB UD-Q4_K_XL GGUF FETCHED + read authoritatively (814 tensors); 3 CPU-verified fidelity corrections grounded in the real GGUF + llama.cpp — per-head QK-RMSNorm (`attn_q/k_norm`, the scope MISSED it), GGUF-authoritative dual-RoPE mscale (llama.cpp `yarn_attn_factor·(1+0.1·ln(factor))`, factor 32 not HF 128), separate `ffn_gate/up_exps`. Keep-quant tower materialization + `ForwardGguf` + the real-model greedy run vs llama.cpp-laguna same-quant oracle = W5 close. Prior: W3 REAL host-reference forward + 3 new ops (`laguna_ops.cpp`, CPU `-Werror` clean, `test_laguna_scaffold` unit-gated)** + +## BACKEND-VULKAN-BARRIERS — the llama.cpp barrier lever, MEASURED: 19.8% of the ~900 per-token barriers are provably unnecessary and removing them is worth ~1 ms/token by the GPU-side instrument, unresolved e2e at 12 pairs; the llama.cpp goal is ALREADY MET on current main WITHOUT it (2026-08-09, GB10, base `c1716fd0`) + +**THE HEADLINE IS THE BASELINE, NOT THE LEVER.** No absolute tok/s number existed +for main after `#191` landed pipelined submission. Eight legs of current main +(`c1716fd0`, page cache dropped before each, `flock $HOME/gpu.lock`, Qwen3.6-27B +bf16, 1 prompt, 32-in/32-out, c1) read TPOT **228.66, 229.16, 229.52, 229.54, +229.61, 229.83, 230.25** ms and one 242.84 ms leg that is DISCARDED for a named +cause: a compile job this session started overlapped it, which is contention, not +the ~1-in-16 `vt_matmul` collapse. **Median 229.54 ms = 4.36 tok/s**, clean-leg +spread 1.59 ms (0.69%). llama.cpp Vulkan on the same box and weights is **4.35**. +**The gate is MET on current main, before this row.** The previous binding number +(4.285 at `81ea01f0`) predates `#191`. + +**THE LEVER.** `vulkan_context.cpp` recorded an unconditional COMPUTE->COMPUTE +`VkMemoryBarrier` before EVERY batched dispatch. llama.cpp instead syncs at 42 +explicit `ggml_vk_sync_buffers` call sites with per-scratch-buffer `need_sync` +flags (`ggml-vulkan.cpp:3193`, `:2108`, `:8174`, `:8687-8748` @ `237ad9b96`). +`VT_VULKAN_SMART_BARRIERS=1` generalizes that: a barrier is recorded only on a +true RAW, WAW or WAR hazard against everything recorded since the previous +barrier. + +**HOW THE READ/WRITE SETS WERE OBTAINED, WHICH WAS THE WHOLE PROBLEM.** +`Dispatch()` receives one flat array of VkBuffers with NO roles attached, and +guessing a role from the binding index is exactly the unverified assumption whose +failure mode is a missing barrier. The GLSL already states it and glslang +ENFORCES it: writing a `readonly buffer` is a compile error, and the promise is +recorded as a `NonWritable` decoration in the emitted SPIR-V. So +`gen-vulkan-spirv.py` now reflects each module's binding count and +`writable_mask` OUT OF THE COMPILED BLOB (the same technique the existing +`spec_ids` reflection uses), and `GetPipeline` hard-checks the declared binding +count against what the host binds. No op-layer plumbing, no hand-maintained +table, and the source of truth is machine-checked. + +**COMPLETENESS ARGUMENT.** The only device commands this backend records that +touch a storage buffer are `vkCmdDispatch` — there is no `vkCmdCopyBuffer`, +`vkCmdFillBuffer` or `vkCmdUpdateBuffer` anywhere in `src/vt/vulkan/` +(`Copy`/`Memset` are host memcpy over the persistently mapped allocation, and +they already drain). So every device-side reader and writer passes through the +analysis. `AllocBuffer` binds each VkBuffer to its OWN VkDeviceMemory at offset +0, so two distinct VkBuffers cannot alias; two tensors sharing one buffer are +reported as colliding even when their byte ranges do not, which costs a barrier +and never misses one. The access sets are cleared ONLY when a barrier is +recorded — never at a flush, never at a drain — which is what carries the +invariant across the command-buffer boundary that pipelined submission creates, +using the same submission-order scope the unconditional barrier relied on. + +**MEASURED, two-length GPU-timestamp diff (output-len 36 minus 4, 32 decode +tokens), same binary, one env variable:** + +| per decode token | always-barrier | smart | delta | +|---|---:|---:|---:| +| barriers recorded | 900.0 | 722.0 | **-178.0 (-19.8%)** | +| barriers skipped | 0 | 178.0 | - | +| GPU span (per-cmdbuf first-TOP to last-BOTTOM) | 231.38 ms | 230.29 ms | **-1.09** | +| sum of per-dispatch intervals | 230.83 ms | 229.73 ms | -1.10 | +| span - sum (the inter-dispatch GAP) | 0.556 ms | 0.563 ms | +0.007 | + +**A PRIOR CLAIM IS CORRECTED. The "~2.3 ms of GPU-timeline time OUTSIDE any +timestamped dispatch" is NOT what this instrument reads.** The gap outside the +timestamped intervals is **0.56 ms/token**, and it does NOT move when 178 +barriers per token are removed. What moves is the sum of the intervals +themselves. `vkCmdWriteTimestamp` at `TOP_OF_PIPE` is written when the command +reaches the top of the pipe, which a preceding barrier does not hold back, so the +barrier's drain is billed INSIDE the following dispatch's interval, not between +them. Span-minus-sum therefore UNDERSTATES barrier cost and is the wrong probe +for it; the per-command-buffer SPAN is the honest one, and it is also the only +one that stays meaningful once dispatches may overlap. + +**MEASURED, e2e, 8 order-alternated AB/BA pairs, same binary, cache dropped +before every leg, `flock` held (TPOT ms):** + +| pair | order | always | smart | delta | +|---|---|---:|---:|---:| +| 0 | AB | 231.16 | 228.61 | **-2.55** | +| 1 | BA | 231.15 | 230.51 | **-0.64** | +| 2 | AB | 229.14 | 229.39 | +0.25 | +| 3 | BA | 230.39 | 228.15 | **-2.24** | +| 4 | AB | 229.98 | 231.41 | +1.43 | +| 5 | BA | 230.86 | 230.03 | **-0.83** | +| 6 | AB | 230.56 | 228.59 | **-1.97** | +| 7 | BA | 230.30 | 228.29 | **-2.01** | + +**6 of 8 pairs to smart. Median TPOT 230.48 -> 229.00 (-1.48 ms, 4.339 -> 4.367 +tok/s); mean paired delta -1.07 ms.** SAY WHAT IS RESOLVABLE: a 6-of-8 sign test +alone is p = 0.145 and does NOT establish the effect. What raises it above +suggestion is that a SECOND, INDEPENDENT instrument — the GPU-side timestamp span, +which is paging-immune and does not see the host at all — reads -1.09 ms/token on +the same tree. The two agree to within 0.4 ms. The mechanism is also measured +directly rather than inferred: 178 fewer barriers per token. + +**THE BOX DRIFTS MORE THAN THE EFFECT, WHICH IS WHY ONLY THE PAIRING IS +QUOTABLE.** The always-barrier arm of the AB/BA block medians 230.48 ms while the +standalone current-main block medians 229.54 ms — 0.94 ms apart on +BEHAVIOURALLY IDENTICAL code, in two sessions an hour apart. A later block on a +freshly rebooted box read 225.87 ms for the same arm. No cross-session absolute +comparison of a ~1 ms effect on this box is meaningful. + +**A SECOND BLOCK OF 4 PAIRS RAN AFTER THE BOX REBOOTED, IN A DIFFERENT REGIME, +AND IS REPORTED SEPARATELY RATHER THAN POOLED.** Post-reboot the same +always-barrier arm runs 218-227 ms instead of 229-231, so this is not the same +operating point: pair 40 218.39/218.29 (-0.10), pair 41 224.98/221.28 (**-3.70**), +pair 42 224.31/225.26 (+0.95), pair 43 226.90/227.00 (+0.10). **2 of 4 to smart, +mean paired delta -0.69 ms.** The block's own base arm drifts 8.6 ms across four +pairs, which is six times the effect being looked for, so it constrains little on +its own. **Across BOTH blocks, 12 pairs: 8 to smart, 4 to always; mean paired +delta -0.94 ms/token (~0.4%); sign test p = 0.19.** The e2e instrument therefore +does NOT establish this effect at 12 pairs. What does stand is the mechanism (178 +fewer barriers per token, counted) and the GPU-side span (-1.09 ms/token), and +those two are what the default decision below rests on. + +**AN ATTEMPT TO EXTEND TO 16 PAIRS WAS LOST AND IS RECORDED BECAUSE THE CAUSE IS +REUSABLE.** After 17 back-to-back legs, each preceded by `drop_caches` and a cold +50 GB reload, the box hit `NVRM ... Out of memory [NV_ERR_NO_MEMORY]` repeatedly +and REBOOTED (boot -1 ended 12:04:29, boot 0 started 12:07:45). The 27B is ~50 GB +on a 119 GB UNIFIED box and the dropped page cache is refilled from zero every +leg; this is the known GB10 unified-memory reboot pathology, not this change. All +earlier evidence survived on disk. Two relaunch attempts also died (one to a shell +quoting bug in a `sed`-generated script that silently dropped the pair index, one +to the user systemd manager going away after the reboot), and the 8-pair block +above is what stands. + +**CORRECTNESS, ALL ON GB10, RUN BEFORE ANY TIMING.** 12 combinations — both +barrier arms x both `VT_VULKAN_INFLIGHT` depths x three gates: +`test_vulkan_backend` **35/35 (2647 assertions at INFLIGHT=1, 2650 at 2)**, +`test_backend_cross_device` **11/11 (132)**, `VLLM_CPP_DEVICE=vulkan +test_opt_paged_engine` **1/1 (63)** reporting *"6/6 prompts token-exact (96/96 +tokens) vs the vLLM 0.25.0 oracle"* and *"all 9 OPT ops dispatched on device type +3 with 0 declines"*. Every one exit 0. (35 test cases, up from 33: this row adds +two.) + +**LONG DECODE, 256 OUTPUT TOKENS, BYTE-IDENTICAL.** `--output-token-ids` md5 over +all four arms — {always, smart} x {INFLIGHT 1, 2} — is the SAME value +`02f7606ec96c9a21775210c2ab33a2c5`. 32 tokens is not enough to expose a rare +ordering hazard; this is. + +**THE MUTATION, AND IT IS THE MOST INFORMATIVE RESULT HERE.** Two scratch +mutations, each built and run IN PLACE on the gated tree (a `cp -a` of the build +directory does NOT work — CMake caches absolute source paths, so the copy rebuilt +the ORIGINAL sources and produced a byte-identical binary and a false green; the +tree was restored and both source and binary md5s verified back to +`c6f4d518e1588dc7233fb5ab83b6a8ce` / `92a0334910be5faa3d0e73b2c689f20d`): + +* **M1, hazard predicate always reports INDEPENDENT** -> `test_vulkan_backend` + **32/35, 330 assertions failed**, and they are VALUE failures, not counters: + `CHECK( back[i] == ... )` reading `0 == -150`, i.e. a dispatch observing memory + before its producer's writes. **The unconditional barrier is load-bearing and + this proves it on the real driver.** +* **M2, only the WRITE-AFTER-READ half removed** -> **34/35, exactly 2 assertions + failed**, both of them the new structural counters (`barriers == 4` read 3, + `skips == 1` read 2). **Every value assertion in the file still passed.** A + plausible partial implementation — RAW and WAW handled, WAR forgotten — + computes correct numbers on this hardware today and is caught ONLY by asserting + the barrier COUNT against a dependency structure known in advance. That is why + the gate asserts counters. + +**DEFAULT: OFF, DELIBERATELY.** The lever is gate-clean on GB10 in every +combination tried and its mechanism is directly measured, but (a) the e2e effect +is ~0.6% and the e2e instrument alone does not resolve it past this box's +session drift, (b) the llama.cpp goal is already MET on current main without it, +and (c) this is the same hazard class that shipped a measured-faster arm which +computed garbage on GB10 two days ago. Flipping it is one environment variable +with all of the above attached, and it should be flipped on an operator's +re-measurement, not on this one. + +**OUT-OF-SCOPE OBSERVATIONS, NEITHER TOUCHED, NEITHER CAUSED BY THIS ROW.** +(1) `vt_fused_chain`'s step list is memcpy'd by the HOST into ONE shared +`ScratchBuffer` per dispatch while batching may leave earlier fused-chain +dispatches unexecuted; a pipeline barrier does not and never did protect against +a host overwrite of host-visible memory, so this is orthogonal to the barrier +policy, but the "safe because dispatch is synchronous" comment on +`ScratchBuffer()` predates batching. (2) `FreeBuffer` destroys a VkBuffer without +draining. The hazard sets hold raw handle VALUES and are never dereferenced, so a +destroyed handle in them is harmless (handle reuse can only manufacture an extra +barrier). (3) `docs/ENVIRONMENT.md` documents `VT_VULKAN_RING` as defaulting to +128; `kDescriptorRing` is 256. diff --git a/.agents/coordination.md b/.agents/coordination.md index 6e1386c1e..8e80d874d 100644 --- a/.agents/coordination.md +++ b/.agents/coordination.md @@ -118,6 +118,22 @@ without the selected contention proof for their entire run are discarded. ## Active claims +**Vulkan dependency-driven barriers (`BACKEND-VULKAN`, 2026-08-09, `CLAIM-VULKAN-FULL-1`).** Claude Code (opus-5), isolated worktree +`/home/mudler/_git/vllm.cpp/.claude/worktrees/agent-aab93052e5a3c7025`, branch +`row/BACKEND-VULKAN-BARRIERS`, base `origin/main` `c1716fd0` (rebased forward; +the commits main gained carry NO `src/`, `examples/` or gate-test change, so the +GB10 evidence stands). NOT PUSHED, no PR, by instruction. Scope: the batched +dispatch path's barrier policy in `src/vt/vulkan/vulkan_context.{h,cpp}`, the +SPIR-V binding-access reflection in `scripts/gen-vulkan-spirv.py` and its two +generated artifacts, two new cases in `tests/vt/test_vulkan_backend.cpp`, and one +`docs/ENVIRONMENT.md` row. NO kernel, shader, model, scheduler or op-table change. +The row's headline result is that the llama.cpp Vulkan decode gate is ALREADY MET +on current main (4.36 vs 4.35, 7 clean legs) independently of the lever; +`VT_VULKAN_SMART_BARRIERS` is implemented, GB10-gated in 12 combinations, and +ships DEFAULT-OFF pending an operator re-measurement. dgx build/gate roots +`dgx.casa:~/vkbase` (main control) and `~/vkbar` (this branch), both +`-DVLLM_CPP_VULKAN=ON`, every GPU stage under `flock $HOME/gpu.lock`. + **Server binary release manifest W5 (`ENG-RELEASE-BINARIES`, 2026-08-08, `CLAIM-ENG-RELEASE-BINARIES-W5`).** Codex (GPT-5), isolated worktree `/home/mudler/_git/vllm.cpp-release-binaries`, branch diff --git a/.agents/state.md b/.agents/state.md index e79acd955..e266a8892 100644 --- a/.agents/state.md +++ b/.agents/state.md @@ -43725,3 +43725,92 @@ ggml-vulkan.cpp:3193). That is the largest identified residual and the next row. FIRST RESUME COMMAND: git log --oneline -1 # expect 66275b1f on row/BACKEND-VULKAN-HOSTDISPATCH ssh dgx.casa 'grep -E "^VLLM_CPP_VULKAN" ~/vk-hostdispatch/build-vk/CMakeCache.txt' # must be ON + +## 2026-08-09 — BACKEND-VULKAN-BARRIERS: llama.cpp goal already MET on main (4.36 vs 4.35); the barrier lever is real (19.8% skipped, GPU -1.09 ms/tok) but e2e-unresolved at 12 pairs, ships default-OFF + + +HEAD `row/BACKEND-VULKAN-BARRIERS`, rebased onto `origin/main`. NOT PUSHED, no +PR, by instruction. Built and gated on GB10 from `c1716fd0`; the two commits main +gained afterwards (`dc2139b3` docs/CI/release-tooling, `7534da65` a C-ABI symbol +hoist in `src/capi/vllm_c.cpp`) leave `src/vt/`, `examples/` and the gate tests +BYTE-IDENTICAL, verified by `git diff --stat`, so the measurements below stand on +the rebased tree. + +THE GOAL WAS ALREADY MET BEFORE THIS ROW, AND NOBODY KNEW. No absolute tok/s +number existed for main after `#191`. Eight legs of `c1716fd0` (drop_caches per +leg, `flock`, 27B bf16 32-in/32-out c1) give a clean-leg median of 229.54 ms = +**4.36 tok/s** against llama.cpp Vulkan's 4.35 on the same box and weights. Seven +clean legs spread 1.59 ms (0.69%); the eighth (242.84) is discarded for a NAMED +cause, a compile job this session started overlapping it. The 4.285 of record +predates `#191`. + +THE BARRIER LEVER IS IMPLEMENTED, GATED AND MEASURED, AND SHIPS DEFAULT-OFF. +`VT_VULKAN_SMART_BARRIERS=1` records a barrier only on a true RAW/WAW/WAR hazard +against everything recorded since the previous barrier. The read and write sets +are NOT plumbed from the op layer and NOT guessed from binding index: they are +reflected out of the committed SPIR-V's `NonWritable` decorations, which glslang +enforces (writing a `readonly buffer` is a compile error), and `GetPipeline` now +hard-checks the module's declared binding count against what the host binds. +MEASURED per decode token, two-length GPU-timestamp diff: barriers 900.0 -> 722.0 +(178 skipped, 19.8%), GPU span 231.38 -> 230.29 ms. e2e, 12 order-alternated AB/BA +pairs across two regimes (the box rebooted mid-campaign and came back ~10 ms +faster): 8 to smart, mean paired delta -0.94 ms, sign test p = 0.19, i.e. NOT +established by the e2e instrument alone. The pre-reboot 8-pair block alone is +6/8 and median 230.48 -> 229.00 ms. + +WHAT THE MEASUREMENT CORRECTS. The handoff named "~2.3 ms of GPU-timeline time +OUTSIDE any timestamped dispatch" as this lever's size. Direct instrumentation +(new `gpu_span_ms()`: per-command-buffer first-TOP-of-pipe to last-BOTTOM) reads +**0.56 ms/token** outside the dispatch intervals, and that number does NOT move +when 178 barriers per token are removed. `TOP_OF_PIPE` is timestamped when the +command reaches the top of the pipe, which a preceding barrier does not hold +back, so the drain is billed INSIDE the following dispatch's interval. +Span-minus-sum understates barrier cost; the SPAN is the honest probe, and it is +the only one that stays meaningful once dispatches may overlap. + +WHAT THE UNIT GATE CATCHES THAT VALUES DO NOT, MEASURED BY MUTATION ON GB10. +Forcing the hazard predicate to always report "independent" fails +`test_vulkan_backend` 32/35 with 330 VALUE failures (`0 == -150`, a dispatch +reading memory before its producer wrote it) -- the unconditional barrier is +load-bearing and this is the proof. Removing ONLY the write-after-read half fails +34/35 with exactly TWO assertions, both of them the new structural barrier +counters; every value assertion in the file still passes. A partial +implementation that handles RAW and WAW and forgets WAR computes correct numbers +on this hardware today. Both mutants were built IN PLACE: a `cp -a` of the build +directory silently rebuilds the ORIGINAL sources (CMake caches absolute source +paths) and produced a byte-identical binary and a false green before this was +noticed; the tree was restored with source and binary md5s verified. + +CORRECTNESS RUN BEFORE ANY TIMING, ALL ON GB10, 12 COMBINATIONS (both barrier +arms x `VT_VULKAN_INFLIGHT` 1 and 2): `test_vulkan_backend` 35/35 (2647/2650 +assertions), `test_backend_cross_device` 11/11 (132), `VLLM_CPP_DEVICE=vulkan +test_opt_paged_engine` 6/6 prompts token-exact (96/96 tokens), 0 declines. A +256-token decode gives BYTE-IDENTICAL token ids across all four +{arm} x {INFLIGHT} combinations, md5 `02f7606ec96c9a21775210c2ab33a2c5`. + +WHY THE DEFAULT STAYS OFF. The e2e effect is ~0.6% and the e2e instrument alone +does not resolve it past this box's session drift (the always-barrier arm medians +230.48 ms in the AB block and 229.54 ms in the standalone main block, on +behaviourally identical code); the goal is already met without it; and this is +the same hazard class as the fence spin that was measured FASTER and computed +garbage on GB10 two days ago. It is one environment variable away with the +evidence attached. + +PROHIBITIONS: never two 27B processes on dgx at once; take `flock $HOME/gpu.lock`; +`local-ai-worker` stopped before GPU work and restored to `--restart=always` +after (restored at the end of this session). + +BLOCKER: the box OOM-REBOOTED partway through an attempt to extend the AB/BA +block to 16 pairs (`NVRM ... NV_ERR_NO_MEMORY`, boot -1 ended 12:04:29). 17 +back-to-back legs each preceded by `drop_caches` refill 50 GB from cold on a +119 GB unified box. The 8-pair block stands; a longer block needs either fewer +cold reloads or a smaller model. + +NEXT LEVER: with the barrier gap measured at 0.56 ms/token and unmoved, the +remaining decode time is 84% `vt_matmul_vec` at ~248 GB/s and 8.3% +`vt_matmul_coopmat`. The named residuals are the lm_head shape's unexplained +7% shortfall against its own measured streaming ceiling, and prefill. + +FIRST RESUME COMMAND: + git log --oneline -1 # row/BACKEND-VULKAN-BARRIERS + ssh dgx.casa 'grep -E "^VLLM_CPP_VULKAN" ~/vkbar/build-vk/CMakeCache.txt' # must be ON diff --git a/docs/BENCHMARKS.md b/docs/BENCHMARKS.md index df5935493..b03fd675a 100644 --- a/docs/BENCHMARKS.md +++ b/docs/BENCHMARKS.md @@ -363,7 +363,7 @@ built on it rather than keeping the flattering one. | Memory footprint vs declared workload (`ROAD-V1-MEM`, #83) | **Never measured, and not measurable today**: there is no auto-sizing to compare against, because the KV pool is a hand-typed `--num-blocks`, so "what the run actually needed" has no number | Once M1's `MemoryBudget` lands: predicted-vs-actual bytes per allocation class, then peak footprint ours-auto vs vLLM at its 0.9 default on the same model and config | | Startup latency (cold to first `/health`) | **36.51 s vs vLLM 0.25.0's 221.51 s = 6.07x** (medians of 3, 27B-NVFP4, GB10). PROVISIONAL: 3 of 6 legs contended, repeat killed by a host reboot. [Detail](../.agents/benchmark-record.md) | Uncontended 3-rep re-run on a quiet box | | Speculation depth (`ROAD-V1-D3-SPEC-K`, #81) | **Never measured, MTP is k=1** (our port covers vLLM's k=1 branch only), so no acceptance-vs-depth curve exists | k=2..4 three-way greedy gate, then the c1/c>1 A/B + the per-workload (prose vs code) acceptance-vs-depth curve any dynamic or adaptive depth policy needs | -| Vulkan vs llama.cpp Vulkan (`BENCH-VK-LLAMA`) | 25 NATIVE (+8 GDN). **27B prefill 21.5x**, decode **4.285**/4.35. Host MEASURED ~3 ms/tok not 5.3; pipelined submit **-1.41 ms/tok 5/5**; fence spin REJECTED. [source](../benchmarks/demo/vulkan_27b_llamacpp.json) | `VK-C` coopmat A/B on Thor (`VT_VULKAN_COOPMAT=0` A/Bs it): **11.1x-32.9x** vs our UNTILED scalar kernel, not vs a competent GEMM. `VK-E`: llama.cpp `-DGGML_VULKAN=ON` at `237ad9b96` on dgx, same GGUF, three columns | +| Vulkan vs llama.cpp Vulkan (`BENCH-VK-LLAMA`) | 25 NATIVE (+8 GDN). **27B prefill 21.5x**; decode **4.36 vs 4.35, MET** (7 clean legs). Smart barriers skip 19.8%/tok, GPU -1.09 ms; e2e 8/12, unresolved. OFF. [source](../benchmarks/demo/vulkan_27b_llamacpp.json) | `VK-C` coopmat A/B on Thor (`VT_VULKAN_COOPMAT=0` A/Bs it): **11.1x-32.9x** vs our UNTILED scalar kernel, not vs a competent GEMM. `VK-E`: llama.cpp `-DGGML_VULKAN=ON` at `237ad9b96` on dgx, same GGUF, three columns | | ROCm (`BACKEND-GATE-ROCM-VLLM` / `-SGLANG`) | **NOT APPLICABLE: no number measured, claimed or owed.** W0 ctest-green on 4 gfx archs (#41); gfx1201 hipBLAS + Gemma-4 MoE (#140, contributor) ran M0/M1 on 2× R9700, our side CPU-link-verified only. No AMD HW here | The approach-(b) fix (PENDING community) unblocks the first APU model run (M2); the gate becomes a same-box vLLM-ROCm oracle once a model runs ([#41](https://github.com/mudler/vllm.cpp/issues/41)); floor: vLLM | | SGLang floor arms | Never ran | Both arms of the SGLang comparison | | Embeddings on the ONE surface (ROW 6, `LlamaModel` + `vllm_embed` + `/v1/embeddings`) | **NO number measured, claimed or owed.** Correctness-gated only, CPU: the 2026-08-08 fold (engine path == direct registry path, f64 LAST+normalize reference on the committed fixture) is plumbing, no speed claim | A REAL embedding checkpoint (e5-mistral class) + a same-box `vllm.LLM(task="embed")` oracle; only then does an embed-throughput bar exist | diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md index c3071bb90..3690a7bd5 100644 --- a/docs/ENVIRONMENT.md +++ b/docs/ENVIRONMENT.md @@ -114,6 +114,7 @@ portable/reference path. In normal operation leave them unset. | `VT_VULKAN_INFLIGHT` | 2 | How many command-buffer submissions may be in flight at once. At 1 a flush submits and BLOCKS on its fence, so all host recording is serialized against an idle GPU; at 2 the flush returns and the next token records while the previous one runs, worth a MEASURED -1.41 ms/token on 27B (5/5 paired legs). Each slot owns a disjoint descriptor-ring slice, so raising this raises descriptor memory | | `VT_VULKAN_HOST_PROFILE` | off | `=1` prints a per-phase HOST time breakdown at exit (submit, record, descriptor update, bookkeeping, pipeline lookup, and time blocked on the GPU). Diagnostic companion to `VT_VULKAN_DISPATCH_STATS`, which covers the DEVICE side; use it to separate host cost from GPU wait | | `VT_VULKAN_BATCH` | on | `=0` forces the per-dispatch submit-and-wait path instead of recording many dispatches into ONE command buffer (with a pipeline barrier between them) and submitting once. MEASURED on GB10: decode **2.62x**, faster in **8 of 8** interleaved pairs (24.8 -> 64.5 tok/s), 40-46 dispatches per submit, opt-125m still 6/6 token-exact. Batching is sound only because all three host-read paths drain the batch first: `Backend::Copy` and `Memset` (host memcpy over the mapped allocation), `Synchronize`, and `Backend::FlushPending`, which op_provider.cpp calls before running a PORTABLE REFERENCE TIER kernel directly over device memory. Without that third one a host kernel would read stale bytes silently. It exists for the same-binary A/B and as the bisect lever. Vulkan-only | +| `VT_VULKAN_SMART_BARRIERS` | off | `=1` records a `vkCmdPipelineBarrier` before a batched dispatch ONLY where a real dependency exists, instead of before every one. The batched path otherwise emits ~900 full COMPUTE->COMPUTE memory barriers per 27B decode token; llama.cpp instead calls `ggml_vk_sync_buffers` at 42 explicit call sites and tracks per-scratch-buffer `need_sync` flags (`ggml-vulkan.cpp:3193`, `:8687-8748` @ pin `237ad9b96`). This is that idea made general: each dispatch's READ set and WRITE set come from the committed SPIR-V's `NonWritable` decorations (`vulkan_spirv.h` § `writable_mask`), which glslang enforces, and a barrier is skipped only when none of read-after-write, write-after-write or write-after-read holds against everything recorded since the previous barrier. It DEFAULTS TO EMITTING: buffers are compared whole rather than by byte range, an operand that is both read and written counts as a write, and an unrecognised binding counts as writable, so every uncertainty costs a barrier rather than correctness. Both arms live in one binary because a cross-BUILD comparison of two barrier policies has already produced a false 1.2x in this campaign -- and because a policy that drops a real dependency computes WRONG NUMBERS, which is invisible on llvmpipe and was invisible for a whole benchmark block the last time it happened here. Vulkan-only | | `VT_VULKAN_GEMV` | on | `=0` forces the decode GEMV onto the portable scalar GEMM instead of `vt_matmul_vec`. The vec tactic gives each output element a WORKGROUP whose lanes stride K, so adjacent lanes read adjacent addresses; the scalar kernel puts one invocation on each output element and loops K there, which for the `MatmulBT` weight layout leaves adjacent lanes `K*2` bytes apart, each pulling its own cache line for 2 bytes of it. Selected only for `MatmulBT` at M=1 with K at least one workgroup wide -- the other orientation is already coalesced and the vec shape would make it worse. MEASURED on GB10: GEMM `ms/call` improves in 7 of 8 interleaved pairs, median 1.8x. Its K reduction is a tree, so unlike the scalar kernel it does NOT share the CPU's accumulation order; the opt-125m STRICT gate (6/6 token-exact) is what clears it. Vulkan-only | | `VT_VULKAN_RMSNORM` | auto | Which `vt::RmsNorm` SPIR-V module runs: `wide` forces the 1024-invocation subgroup-reducing one, `base` forces the portable 128-invocation one, unset lets the device capability decide (1024 invocations on the X axis plus compute subgroup BASIC and ARITHMETIC). The wide module exists because `RmsNorm` dispatches ONE WORKGROUP PER ROW and a batch-1 decode step has exactly one row: on Qwen3.6-27B that put 128 invocations on a 5120-wide row, four warps of one SM, with the rest of the GPU idle. MEASURED on GB10 by the two-length GPU-timestamp diff: **0.0611 -> 0.0123 ms/call, 7.88 -> 1.59 ms/token**, and paired decode **241.9 -> 235.6 ms** median TPOT. The tell that it was OCCUPANCY and not the reduction is that the SAME shader costs 0.066 ms/call during PREFILL, where 32 rows give it 32 workgroups and 32x the data. It exists for the same-binary A/B and so the unit gate can exercise the fallback on hardware that would always pick the wide arm. Vulkan-only | | `VT_VULKAN_MATMUL_NCOLS` | 4 | Output columns each lane of the portable scalar GEMM computes, in the `[K,N]` (non-transposed) orientation only. At 1 the kernel is the flat one-invocation-per-output-element body; above 1 a workgroup takes `128*NCOLS` CONSECUTIVE output columns of one row, so at each step of K it reads a contiguous run of that many elements instead of 128. This is the ONE decode GEMM that cannot reach the `vt_matmul_vec` tactic, because in `[K,N]` the lanes are already coalesced and the GEMV shape would make them strided; on the 27B it is the lm_head, `m=1 k=5120 n=248320`, 2.54 GB moved per token. MEASURED on GB10, 27B decode, `ms/call` medians over interleaved replicates: NCOLS 1 = 12.48, 2 = 12.46, **4 = 11.54**, 8 = 12.81, with 4 winning **6 of 6** interleaved pairs against 1. Blocking is a TRADE, not a monotone win: at 8 the dispatch falls to 243 workgroups (~31k threads) and the device runs out of work to hide memory latency with faster than the longer contiguous run buys back. It rides a specialization constant, so every arm is the same committed module and they A/B in one binary. Every arm is BIT-IDENTICAL -- each accumulator owns one output element and sums the whole K sequentially, which is the CPU kernel's order -- so this kernel keeps the byte-exact tier that the coopmat and GEMV tactics gave up; a memcmp gates that. Vulkan-only | diff --git a/docs/STATUS.md b/docs/STATUS.md index 9a8b630e1..a39b558d0 100644 --- a/docs/STATUS.md +++ b/docs/STATUS.md @@ -422,7 +422,7 @@ Parakeet ASR (2026-08-07): *CPU-correct, ON THE ONE SURFACE (ROW 1)*. Ids exact LoRA (W1 CPU runtime brick landed; not yet usable end-to-end), multi-GPU, Vulkan (opt-125m exact; 25 native +8 GDN, both recurrences + fused attn preamble; 27B prefill 21.5x, decode -4.28/4.35; submit pipelined -1.41 ms/tok 5/5; #125 +4.36/4.35 MET; barriers -19.8%, GPU -1.09 ms; #125 [campaign](../.agents/specs/vulkan-full-support.md)), ROCm (W0 community-green on 4 gfx archs (#41); the ratified (b) APU unified-memory fix is in (**blind-written, unverified**); M2 needs verification; gfx1201 hipBLAS + diff --git a/scripts/gen-vulkan-spirv.py b/scripts/gen-vulkan-spirv.py index e7207e544..4238e2ec4 100755 --- a/scripts/gen-vulkan-spirv.py +++ b/scripts/gen-vulkan-spirv.py @@ -114,6 +114,95 @@ def compile_one(cc: pathlib.Path, src: pathlib.Path) -> bytes: SPIRV_DECORATION_SPEC_ID = 1 SPIRV_HEADER_WORDS = 5 +# The rest of the annotation/type opcodes the BINDING ACCESS reflection below +# needs (SPIR-V core spec §3.52 Instructions, §3.20 Decoration). +SPIRV_OP_TYPE_POINTER = 32 +SPIRV_OP_VARIABLE = 59 +SPIRV_OP_MEMBER_DECORATE = 72 +SPIRV_DECORATION_NON_WRITABLE = 24 +SPIRV_DECORATION_BINDING = 33 + +# Mirrors kMaxDispatchBindings in src/vt/vulkan/vulkan_context.cpp. The mask +# below is one bit per binding, so a shader with more bindings than this could +# not be described and must fail LOUDLY rather than silently reporting the +# overflow bindings as read-only. +MAX_BINDINGS = 32 + + +def _instructions(blob: bytes): + """Yield (opcode, [operand words]) for one SPIR-V module.""" + words = [int.from_bytes(blob[i:i + 4], "little") for i in range(0, len(blob), 4)] + i = SPIRV_HEADER_WORDS + while i < len(words): + word_count = words[i] >> 16 + opcode = words[i] & 0xFFFF + if word_count == 0: + sys.exit("malformed SPIR-V: zero-length instruction") + yield opcode, words[i + 1:i + word_count] + i += word_count + + +def binding_access(blob: bytes, name: str) -> tuple[int, int]: + """Return (binding_count, writable_mask) for one SPIR-V module. + + WHY THIS IS PARSED OUT OF THE COMPILED MODULE AND NOT WRITTEN DOWN BY HAND. + The Vulkan backend can only skip a pipeline barrier between two dispatches if + it can prove they do not share a buffer in a hazardous way, and that proof + needs each dispatch's READ set and WRITE set. Nothing on the host side knows + them: `VulkanContext::Dispatch` receives one flat array of VkBuffers with no + roles attached, and guessing a role from the binding index would be exactly + the unverified assumption that turns a missing barrier into silently wrong + numbers. + + The GLSL already states it, and glslang ENFORCES it: a `readonly buffer` + block is a compile error to write to, and the compiler records the promise as + a `NonWritable` decoration in the emitted module. So the committed SPIR-V is + an authoritative, machine-checked statement of which bindings the shader can + write, and the freshness check on this generator keeps it from drifting from + the .comp source. + + A binding is reported WRITABLE unless a NonWritable decoration is positively + found for it, so an unrecognised encoding degrades to "assume it is written", + which costs a barrier rather than correctness. + """ + binding_of: dict[int, int] = {} # variable id -> binding number + var_non_writable: set[int] = set() # variable ids decorated NonWritable + member_non_writable: set[tuple[int, int]] = set() + pointee_of: dict[int, int] = {} # pointer type id -> pointee type id + ptr_type_of: dict[int, int] = {} # variable id -> pointer type id + for opcode, ops in _instructions(blob): + if opcode == SPIRV_OP_DECORATE and len(ops) >= 3: + if ops[1] == SPIRV_DECORATION_BINDING: + binding_of[ops[0]] = ops[2] + elif opcode == SPIRV_OP_DECORATE and len(ops) == 2: + if ops[1] == SPIRV_DECORATION_NON_WRITABLE: + var_non_writable.add(ops[0]) + elif opcode == SPIRV_OP_MEMBER_DECORATE and len(ops) >= 3: + if ops[2] == SPIRV_DECORATION_NON_WRITABLE: + member_non_writable.add((ops[0], ops[1])) + elif opcode == SPIRV_OP_TYPE_POINTER and len(ops) == 3: + pointee_of[ops[0]] = ops[2] + elif opcode == SPIRV_OP_VARIABLE and len(ops) >= 3: + ptr_type_of[ops[1]] = ops[0] + + if not binding_of: + sys.exit(f"{name}: no descriptor bindings found in the SPIR-V") + mask = 0 + for var, binding in binding_of.items(): + if binding >= MAX_BINDINGS: + sys.exit(f"{name}: binding {binding} is at or above the backend's " + f"{MAX_BINDINGS}-binding limit") + block = pointee_of.get(ptr_type_of.get(var, -1), -1) + read_only = var in var_non_writable or (block, 0) in member_non_writable + if not read_only: + mask |= 1 << binding + count = max(binding_of.values()) + 1 + # The host binds descriptors 0..count-1 densely, so a hole would mean the + # mask's bit positions no longer line up with the dispatch's buffer array. + if len(set(binding_of.values())) != count: + sys.exit(f"{name}: descriptor bindings are not a dense 0..{count - 1} range") + return count, mask + def spec_ids(blob: bytes) -> list[int]: """Return the SpecId values decorated in one SPIR-V module, sorted ascending.""" @@ -178,12 +267,31 @@ def render_header(blobs: dict[str, bytes], version: str) -> str: add("// passes specialization values BY ID, and Vulkan SILENTLY IGNORES a map entry") add("// whose ID the module does not declare - so without this table a host/shader") add("// drift produces WRONG NUMBERS instead of a clean failure.") + add("//") + add("// binding_count is the number of descriptor bindings the module declares, and") + add("// writable_mask has bit i set iff binding i is NOT decorated NonWritable, i.e.") + add("// iff the shader is permitted to WRITE that buffer. Both are parsed out of the") + add("// compiled module rather than written down beside it.") + add("//") + add("// This is the READ/WRITE SET the dispatch path needs to decide whether two") + add("// dispatches are genuinely independent. Dispatch() is handed one flat array of") + add("// VkBuffers with no roles attached, and inferring a role from the binding index") + add("// would be an unverified assumption whose failure mode is a MISSING BARRIER and") + add("// therefore silently wrong numbers. `readonly` in the GLSL is enforced by") + add("// glslang (writing such a block is a compile error) and recorded as NonWritable") + add("// in the SPIR-V, so this mask is a machine-checked fact about the shader.") + add("//") + add("// A bit is SET unless read-only is positively proven, so anything the") + add("// reflection does not understand degrades to an extra barrier, never a") + add("// missing one.") add("struct SpirvModule {") add(" const char* name;") add(" const uint32_t* words;") add(" size_t word_count;") add(" const uint32_t* spec_ids;") add(" size_t spec_id_count;") + add(" uint32_t binding_count;") + add(" uint32_t writable_mask;") add("};") add("") add("// DEFINED IN vulkan_spirv.cpp. The array is `extern` and therefore of unknown") @@ -231,8 +339,9 @@ def render_source(blobs: dict[str, bytes], version: str) -> str: for name in sorted(blobs): ids = spec_ids(blobs[name]) idp = f"kSpecIds_{name}" if ids else "nullptr" + count, mask = binding_access(blobs[name], name) add(f' {{"{name}", kSpv_{name}, sizeof(kSpv_{name}) / sizeof(uint32_t), ' - f'{idp}, {len(ids)}}},') + f'{idp}, {len(ids)}, {count}u, 0x{mask:08x}u}},') add("};") add("const size_t kSpirvModuleCount = sizeof(kSpirvModules) / sizeof(kSpirvModules[0]);") add("") diff --git a/src/vt/vulkan/vulkan_context.cpp b/src/vt/vulkan/vulkan_context.cpp index 67d55b6aa..5821ca974 100644 --- a/src/vt/vulkan/vulkan_context.cpp +++ b/src/vt/vulkan/vulkan_context.cpp @@ -435,6 +435,30 @@ const uint32_t kInFlight = [] { // checked at dispatch so an overrun names itself instead of smashing the stack. constexpr uint32_t kMaxDispatchBindings = 32; +// SMART BARRIERS (BACKEND-VULKAN-BARRIERS), VT_VULKAN_SMART_BARRIERS. +// +// The batched dispatch path records a full COMPUTE->COMPUTE memory barrier before +// EVERY dispatch -- about 900 per 27B decode token -- because the ops in a decode +// step are sequentially dependent and a dispatch that misses its producer's writes +// computes garbage. llama.cpp does not: it calls `ggml_vk_sync_buffers` +// (ggml-vulkan.cpp:3193 @ pin 237ad9b96) at 42 explicit call sites and tracks +// per-scratch-buffer `prealloc_{x,y,split_k}_need_sync` flags (:2108, :8174, +// :8687-8748) so a sync happens only where a real dependency exists. +// +// This is the same idea made GENERAL rather than per-scratch-buffer: the buffers +// each dispatch reads and writes are known exactly (the committed SPIR-V records +// which bindings are NonWritable, see vulkan_spirv.h), so the barrier can be +// driven by an actual hazard test instead of by a call site the author remembered. +// +// DEFAULT OFF. The failure mode of a wrong answer here is not a crash, it is +// silently wrong numbers, and this backend has already shipped one measured-faster +// arm that computed garbage on the real driver while passing on llvmpipe (see the +// FENCE SPIN note above). The default moves only on GB10 evidence. +const bool kSmartBarriersEnv = [] { + const char* v = std::getenv("VT_VULKAN_SMART_BARRIERS"); + return v != nullptr && std::strcmp(v, "0") != 0; +}(); + // Descriptor-ring sets available to ONE slot. The ring is partitioned rather // than shared, so slot s owns [s*kRingSlice, (s+1)*kRingSlice) and no set can be // rewritten by a batch while an earlier, still-executing batch reads it. @@ -485,6 +509,10 @@ struct VulkanContext::Pipeline { // the counter lives here and DispatchHistogram aggregates on demand. uint64_t dispatches = 0; const char* module_name = nullptr; // points into the committed SPIR-V table + // Bit i set iff binding i is one the SHADER MAY WRITE, copied from the + // committed SPIR-V table (vulkan_spirv.h § writable_mask). The dispatch path + // splits its buffer array into a read set and a write set with it. + uint32_t writable_mask = 0; }; // The set of buffers bound by any batch that has not yet been drained. @@ -818,6 +846,8 @@ VulkanContext::VulkanContext() { pipelines_ = new std::map(); dispatch_ms_ = new std::map(); batch_buffers_ = new BufferSet(); + hazard_written_ = new BufferSet(); + hazard_read_ = new BufferSet(); // TWO timestamps per dispatch (before and after), for a whole batch, and ONE // RANGE PER SLOT: a slot's queries are read back only when that slot retires, // so two in-flight batches must not share query indices. Created only under the @@ -895,6 +925,24 @@ VulkanContext::VulkanContext() { } std::fprintf(stderr, "[vt vulkan] %-24s %8llu %10.1f\n", "TOTAL", static_cast(ctx.dispatch_count()), total_ms); + // THE BARRIER LINE. The span is the GPU's own wall over each command + // buffer; total_ms above is the sum of the individual dispatch intervals. + // Their difference is the GPU time spent BETWEEN dispatches -- the barrier + // drains and the launch setup -- which is the only direct measurement of + // what BACKEND-VULKAN-BARRIERS is trying to remove, and it is invisible in + // every per-shader row above. Read the gap, not the arms' per-shader sums: + // once barriers are skipped, adjacent dispatches may OVERLAP and their + // intervals double-count. + const double span = ctx.gpu_span_ms(); + std::fprintf(stderr, + "[vt vulkan] BARRIERS recorded=%llu skipped=%llu smart=%d\n" + "[vt vulkan] GPU span=%.1f ms over %llu cmdbufs, " + "sum-of-dispatch=%.1f ms, GAP=%.1f ms (%.1f%%)\n", + static_cast(ctx.barrier_count()), + static_cast(ctx.barrier_skip_count()), + ctx.smart_barriers() ? 1 : 0, span, + static_cast(ctx.gpu_span_batches()), total_ms, + span - total_ms, span > 0 ? 100.0 * (span - total_ms) / span : 0.0); }); } mutex_ = new std::mutex(); @@ -1037,9 +1085,24 @@ VulkanContext::Pipeline& VulkanContext::GetPipeline(const std::string& name, " value(s) were supplied — host and committed SPIR-V have drifted;" " regenerate with scripts/gen-vulkan-spirv.py"); + // The binding COUNT must agree too, and for a second reason beyond the layout: + // the writable mask below is indexed by binding, so a host that bound more + // buffers than the module declares would get mask bit 0 -- "read-only" -- for + // the extra ones, and the hazard analysis would then be free to skip a barrier + // for a buffer the shader can write. Bound LOUDLY rather than conservatively, + // because a mismatch is a host/shader drift that also breaks the descriptor set + // layout and should never be tolerated silently. + VT_CHECK(module->binding_count == buffer_count, + "vulkan: kernel '" + name + "' declares " + + std::to_string(module->binding_count) + " descriptor binding(s) but " + + std::to_string(buffer_count) + + " buffer(s) were bound — host and committed SPIR-V have drifted;" + " regenerate with scripts/gen-vulkan-spirv.py"); + Pipeline p; p.buffer_count = buffer_count; p.push_size = push_size; + p.writable_mask = module->writable_mask; // The committed SPIR-V table has static storage duration, so this pointer // outlives the pipeline and the histogram can key on it without copying. p.module_name = module->name; @@ -1240,6 +1303,41 @@ uint64_t VulkanContext::barrier_count() const { return barrier_count_; } +uint64_t VulkanContext::barrier_skip_count() const { + std::lock_guard guard(*static_cast(mutex_)); + return barrier_skipped_; +} + +bool VulkanContext::smart_barriers() const { + if (smart_barriers_override_ != 0) return smart_barriers_override_ > 0; + return kSmartBarriersEnv; +} + +void VulkanContext::set_smart_barriers_override(int v) { + // Drain FIRST. hazard_written_/hazard_read_ describe commands already recorded + // under the OLD policy; switching while a batch is open or in flight would let + // the new policy reason about a history it did not build. After a drain the GPU + // is idle and the next dispatch starts from a clean slate either way. + // + // The sets are deliberately NOT cleared here even so -- see the header for why + // they are cleared only by an emitted barrier. Keeping them merely costs the + // next dispatch a barrier it may not have needed. + FlushBatch("smart-barrier-policy-change"); + std::lock_guard guard(*static_cast(mutex_)); + smart_barriers_override_ = v; + force_barrier_next_ = true; +} + +double VulkanContext::gpu_span_ms() const { + std::lock_guard guard(*static_cast(mutex_)); + return static_cast(gpu_span_ns_) / 1.0e6; +} + +uint64_t VulkanContext::gpu_span_batches() const { + std::lock_guard guard(*static_cast(mutex_)); + return gpu_span_batches_; +} + void VulkanContext::FlushIfBatchTouches(void* buffer, const char* why) { std::lock_guard guard(*static_cast(mutex_)); // A host pointer (nullptr) cannot alias a bound VkBuffer, so it never forces a @@ -1296,9 +1394,23 @@ void VulkanContext::RetireSlotLocked(uint32_t s) { VK_QUERY_RESULT_64_BIT | VK_QUERY_RESULT_WAIT_BIT) == VK_SUCCESS) { auto& acc = *static_cast*>(dispatch_ms_); + // Earliest top-of-pipe and latest bottom-of-pipe over the whole command + // buffer. Taken as a MIN and a MAX rather than "first and last", because + // once barriers are skipped the dispatches may overlap and need not retire + // in the order they were recorded. Zero ticks mean the query was never + // written and are excluded, or they would drag the span back to the epoch. + uint64_t span_lo = 0, span_hi = 0; for (uint32_t i = 0; i < n; ++i) { const uint64_t t0 = ticks[i * 2], t1 = ticks[i * 2 + 1]; if (t1 > t0) acc[(*names)[i]] += double(t1 - t0) * timestamp_period_ns_ / 1.0e6; + if (t0 == 0 || t1 == 0) continue; + if (span_lo == 0 || t0 < span_lo) span_lo = t0; + if (t1 > span_hi) span_hi = t1; + } + if (span_hi > span_lo && span_lo != 0) { + gpu_span_ns_ += static_cast(double(span_hi - span_lo) * + timestamp_period_ns_); + ++gpu_span_batches_; } } names->clear(); @@ -1509,14 +1621,76 @@ void VulkanContext::Dispatch(const std::string& name, const void* const* buffers // executing, and this barrier -- which orders against everything earlier in // SUBMISSION ORDER on the same queue, not merely earlier in this buffer -- is // what carries the dependency across the command-buffer boundary. - VkMemoryBarrier mb{}; - mb.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER; - mb.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT; - mb.dstAccessMask = VK_ACCESS_SHADER_READ_BIT | VK_ACCESS_SHADER_WRITE_BIT; - vk.vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, - VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &mb, 0, nullptr, 0, - nullptr); - ++barrier_count_; + // + // SMART BARRIERS (BACKEND-VULKAN-BARRIERS) narrow "every dispatch" to "every + // dispatch that could actually observe or clobber an earlier one". The + // decision DEFAULTS TO EMITTING: a barrier is skipped only when this + // dispatch's operands are proven disjoint, in the hazardous direction, from + // every operand recorded since the previous barrier. The three sets and the + // invariant they maintain are documented on hazard_written_ in the header, + // and the cross-command-buffer guarantee above is unchanged -- when the test + // says hazard, the barrier recorded is exactly the one that was always + // recorded, with exactly the same submission-order scope. + auto& written = *static_cast(hazard_written_); + auto& read = *static_cast(hazard_read_); + const bool analyse = smart_barriers_override_ != 0 ? smart_barriers_override_ > 0 + : kSmartBarriersEnv; + // The always-barrier arm does not maintain the access history -- it does not + // need it, and paying for it there would make the control arm slower than the + // tree it is the control FOR. That leaves the history empty on the first + // dispatch after the policy is switched ON, which would read as "no hazard" + // when the truth is "no history". force_barrier_next_ makes that transition + // emit one barrier, after which the invariant is established honestly. + bool hazard = true; + if (analyse && !force_barrier_next_) { + hazard = false; + for (uint32_t i = 0; i < buffer_count && !hazard; ++i) { + void* b = const_cast(buffers[i]); + if ((p.writable_mask >> i) & 1u) { + // A WRITE collides with an earlier write (WAW) and with an earlier read + // (WAR); the latter is why the read set has to exist at all. + hazard = written.Contains(b) || read.Contains(b); + } else { + // A READ collides only with an earlier WRITE (RAW). Two dispatches that + // merely read the same buffer -- the gate and up projections sharing one + // normalized activation, say -- need no ordering between them. + hazard = written.Contains(b); + } + } + } + if (hazard) { + VkMemoryBarrier mb{}; + mb.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER; + mb.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT; + mb.dstAccessMask = VK_ACCESS_SHADER_READ_BIT | VK_ACCESS_SHADER_WRITE_BIT; + vk.vkCmdPipelineBarrier(cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &mb, 0, nullptr, + 0, nullptr); + ++barrier_count_; + // The barrier orders this dispatch after every command submitted earlier on + // this queue, so nothing before it can still be a hazard and the history + // restarts empty. Cleared HERE and nowhere else -- not at a flush, not at a + // drain -- which is what makes the invariant hold across command buffers. + written.Clear(); + read.Clear(); + force_barrier_next_ = false; + } else { + ++barrier_skipped_; + } + // Record this dispatch's accesses AFTER the decision, so a dispatch is never + // tested against itself. A buffer bound at both a readable and a writable + // binding is entered in BOTH sets, which is a superset of the truth and can + // only ever produce an extra barrier. + if (analyse) { + for (uint32_t i = 0; i < buffer_count; ++i) { + void* b = const_cast(buffers[i]); + if ((p.writable_mask >> i) & 1u) { + written.Insert(b); + } else { + read.Insert(b); + } + } + } } vk.vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, p.pipeline); diff --git a/src/vt/vulkan/vulkan_context.h b/src/vt/vulkan/vulkan_context.h index 2e2169c80..8c80decf1 100644 --- a/src/vt/vulkan/vulkan_context.h +++ b/src/vt/vulkan/vulkan_context.h @@ -205,6 +205,54 @@ class VulkanContext { // dropping it is likewise invisible in llvmpipe's numbers. uint64_t barrier_count() const; + // --- SMART BARRIERS (BACKEND-VULKAN-BARRIERS). Dispatches whose buffers do not + // collide with anything recorded since the last barrier get NO barrier. + // + // Barriers NOT recorded because the incoming dispatch was proven independent of + // everything since the previous barrier. barrier_count() + barrier_skip_count() + // is the dispatch count in a batched run, which is what the always-barrier arm + // would have recorded. This is the ONLY direct evidence that the analysis is + // doing anything: the two arms compute identical numbers by construction, so a + // value check cannot tell them apart, and neither can the wall clock when the + // skip rate is small. + uint64_t barrier_skip_count() const; + // Whether the hazard analysis is active (VT_VULKAN_SMART_BARRIERS, or the + // override below). Asked rather than re-derived from the environment, for the + // reason batching_enabled() exists. + bool smart_barriers() const; + // A/B LEVER over that decision, inside ONE binary. -1 forces the unconditional + // barrier before every dispatch (main's behaviour, byte for byte), +1 forces the + // hazard analysis, 0 lets VT_VULKAN_SMART_BARRIERS decide. + // + // A cross-BUILD comparison of two barrier policies is exactly the shape that + // produced a false 1.2x reading earlier in this campaign, and -- far worse here + // -- a barrier policy that drops a REAL dependency computes wrong numbers, so + // the correctness gate has to be able to run both arms against the same inputs + // in one process and compare them directly. + // + // Drains first: the analysis state describes work already recorded, so changing + // the policy underneath an open or in-flight batch would reason about the wrong + // history. + void set_smart_barriers_override(int v); + int smart_barriers_override() const { return smart_barriers_override_; } + + // GPU-TIMELINE SPAN, milliseconds: summed over command buffers, the interval + // from the FIRST dispatch's top-of-pipe timestamp to the LAST one's + // bottom-of-pipe. Only collected when VT_VULKAN_DISPATCH_STATS is set, like the + // per-dispatch timestamps it is derived from. + // + // WHY IT IS SEPARATE FROM DispatchTimeMs. Those sum each dispatch's OWN + // interval; this measures the wall the GPU spent on the whole command buffer. + // The difference between them is the time the GPU spent BETWEEN dispatches -- + // barrier drains and launch setup -- which is the quantity this row exists to + // move and which no per-shader number can show. It is also the honest metric + // once barriers are skipped: without a barrier between them two dispatches may + // OVERLAP, so their individual intervals double-count and only the span stays + // meaningful. + double gpu_span_ms() const; + // Command buffers the span above was accumulated over. + uint64_t gpu_span_batches() const; + // Number of distinct pipelines currently cached. Exposed for the unit gate: it // is how a test proves a new specialization produced a NEW pipeline rather than // silently reusing an existing one — which would look identical in the results. @@ -350,6 +398,41 @@ class VulkanContext { uint32_t batch_count_ = 0; // dispatches recorded into it void* batch_buffers_ = nullptr; // BufferSet*, buffers any UNRETIRED batch bound + // --- SMART BARRIERS (BACKEND-VULKAN-BARRIERS). The two halves of the access + // history the hazard test consults, both BufferSet*. + // + // THE INVARIANT, which is the whole correctness argument. At every point, + // hazard_written_ contains every buffer WRITTEN, and hazard_read_ every buffer + // READ, by any command recorded AFTER the most recent vkCmdPipelineBarrier this + // context emitted -- across command-buffer and submission boundaries, because + // they are cleared ONLY when a barrier is recorded and by nothing else. + // + // A dispatch is independent of that history iff none of its reads is in + // hazard_written_ (read-after-write), none of its writes is in hazard_written_ + // (write-after-write) and none of its writes is in hazard_read_ + // (write-after-read). If any of the three holds, a barrier is recorded, which + // by Vulkan's submission-order scope orders the dispatch after EVERY command + // submitted earlier on this queue -- so the sets may then be emptied and the + // invariant restarts. That scope is also what makes the analysis sound across + // the command-buffer boundary a pipelined submission creates. + // + // Both are keyed on the whole VkBuffer, never on the tensor's byte range: two + // tensors that share an allocation are reported as colliding even when their + // ranges do not, which costs a barrier and never misses one. An operand that is + // both read and written by the same shader lands in the WRITE set only, which is + // strictly stronger -- a write collides with everything a read collides with. + void* hazard_written_ = nullptr; + void* hazard_read_ = nullptr; + uint64_t barrier_skipped_ = 0; // barriers the analysis proved unnecessary + int smart_barriers_override_ = 0; // -1 always barrier, +1 analyse, 0 = env + // "The access history is not trustworthy, barrier unconditionally once." Set at + // construction and whenever the barrier policy changes, because the + // always-barrier arm deliberately does not maintain the history and an empty + // history must never be mistaken for an absence of hazards. + bool force_barrier_next_ = true; + uint64_t gpu_span_ns_ = 0; // summed per-command-buffer GPU spans + uint64_t gpu_span_batches_ = 0; + // PIPELINED SUBMISSION (BACKEND-VULKAN-HOSTDISPATCH). Ported from llama.cpp // `ggml_backend_vk_graph_compute` (ggml-vulkan.cpp:16192-16195 and // :16417-16423 @ pin 237ad9b96), whose own comment states the mechanism: diff --git a/src/vt/vulkan/vulkan_spirv.cpp b/src/vt/vulkan/vulkan_spirv.cpp index bc634bcee..201b1c1c0 100644 --- a/src/vt/vulkan/vulkan_spirv.cpp +++ b/src/vt/vulkan/vulkan_spirv.cpp @@ -14187,32 +14187,32 @@ constexpr uint32_t kSpecIds_vt_sigmoid_gate_bf16[] = { } // namespace const SpirvModule kSpirvModules[] = { - {"vt_add", kSpv_vt_add, sizeof(kSpv_vt_add) / sizeof(uint32_t), nullptr, 0}, - {"vt_attn_qk_norm_rope_gate", kSpv_vt_attn_qk_norm_rope_gate, sizeof(kSpv_vt_attn_qk_norm_rope_gate) / sizeof(uint32_t), kSpecIds_vt_attn_qk_norm_rope_gate, 3}, - {"vt_cast", kSpv_vt_cast, sizeof(kSpv_vt_cast) / sizeof(uint32_t), kSpecIds_vt_cast, 2}, - {"vt_causal_conv1d_update", kSpv_vt_causal_conv1d_update, sizeof(kSpv_vt_causal_conv1d_update) / sizeof(uint32_t), nullptr, 0}, - {"vt_embedding", kSpv_vt_embedding, sizeof(kSpv_vt_embedding) / sizeof(uint32_t), kSpecIds_vt_embedding, 3}, - {"vt_fused_chain", kSpv_vt_fused_chain, sizeof(kSpv_vt_fused_chain) / sizeof(uint32_t), nullptr, 0}, - {"vt_gdn_decode", kSpv_vt_gdn_decode, sizeof(kSpv_vt_gdn_decode) / sizeof(uint32_t), kSpecIds_vt_gdn_decode, 2}, - {"vt_gdn_post_conv", kSpv_vt_gdn_post_conv, sizeof(kSpv_vt_gdn_post_conv) / sizeof(uint32_t), kSpecIds_vt_gdn_post_conv, 3}, - {"vt_gdn_prefill", kSpv_vt_gdn_prefill, sizeof(kSpv_vt_gdn_prefill) / sizeof(uint32_t), kSpecIds_vt_gdn_prefill, 2}, - {"vt_gdn_state_gather", kSpv_vt_gdn_state_gather, sizeof(kSpv_vt_gdn_state_gather) / sizeof(uint32_t), nullptr, 0}, - {"vt_gdn_state_scatter", kSpv_vt_gdn_state_scatter, sizeof(kSpv_vt_gdn_state_scatter) / sizeof(uint32_t), nullptr, 0}, - {"vt_greedy_argmax", kSpv_vt_greedy_argmax, sizeof(kSpv_vt_greedy_argmax) / sizeof(uint32_t), nullptr, 0}, - {"vt_layer_norm", kSpv_vt_layer_norm, sizeof(kSpv_vt_layer_norm) / sizeof(uint32_t), nullptr, 0}, - {"vt_matmul", kSpv_vt_matmul, sizeof(kSpv_vt_matmul) / sizeof(uint32_t), kSpecIds_vt_matmul, 5}, - {"vt_matmul_coopmat", kSpv_vt_matmul_coopmat, sizeof(kSpv_vt_matmul_coopmat) / sizeof(uint32_t), kSpecIds_vt_matmul_coopmat, 2}, - {"vt_matmul_vec", kSpv_vt_matmul_vec, sizeof(kSpv_vt_matmul_vec) / sizeof(uint32_t), kSpecIds_vt_matmul_vec, 6}, - {"vt_paged_attn", kSpv_vt_paged_attn, sizeof(kSpv_vt_paged_attn) / sizeof(uint32_t), kSpecIds_vt_paged_attn, 4}, - {"vt_qkv_split", kSpv_vt_qkv_split, sizeof(kSpv_vt_qkv_split) / sizeof(uint32_t), kSpecIds_vt_qkv_split, 2}, - {"vt_relu", kSpv_vt_relu, sizeof(kSpv_vt_relu) / sizeof(uint32_t), nullptr, 0}, - {"vt_reshape_and_cache", kSpv_vt_reshape_and_cache, sizeof(kSpv_vt_reshape_and_cache) / sizeof(uint32_t), kSpecIds_vt_reshape_and_cache, 1}, - {"vt_rms_norm", kSpv_vt_rms_norm, sizeof(kSpv_vt_rms_norm) / sizeof(uint32_t), nullptr, 0}, - {"vt_rms_norm_gated", kSpv_vt_rms_norm_gated, sizeof(kSpv_vt_rms_norm_gated) / sizeof(uint32_t), nullptr, 0}, - {"vt_rms_norm_wide", kSpv_vt_rms_norm_wide, sizeof(kSpv_vt_rms_norm_wide) / sizeof(uint32_t), nullptr, 0}, - {"vt_rope_from_cache", kSpv_vt_rope_from_cache, sizeof(kSpv_vt_rope_from_cache) / sizeof(uint32_t), kSpecIds_vt_rope_from_cache, 5}, - {"vt_sigmoid_gate_bf16", kSpv_vt_sigmoid_gate_bf16, sizeof(kSpv_vt_sigmoid_gate_bf16) / sizeof(uint32_t), kSpecIds_vt_sigmoid_gate_bf16, 1}, - {"vt_silu_and_mul", kSpv_vt_silu_and_mul, sizeof(kSpv_vt_silu_and_mul) / sizeof(uint32_t), nullptr, 0}, + {"vt_add", kSpv_vt_add, sizeof(kSpv_vt_add) / sizeof(uint32_t), nullptr, 0, 6u, 0x00000030u}, + {"vt_attn_qk_norm_rope_gate", kSpv_vt_attn_qk_norm_rope_gate, sizeof(kSpv_vt_attn_qk_norm_rope_gate) / sizeof(uint32_t), kSpecIds_vt_attn_qk_norm_rope_gate, 3, 13u, 0x000003f0u}, + {"vt_cast", kSpv_vt_cast, sizeof(kSpv_vt_cast) / sizeof(uint32_t), kSpecIds_vt_cast, 2, 4u, 0x0000000cu}, + {"vt_causal_conv1d_update", kSpv_vt_causal_conv1d_update, sizeof(kSpv_vt_causal_conv1d_update) / sizeof(uint32_t), nullptr, 0, 11u, 0x00000303u}, + {"vt_embedding", kSpv_vt_embedding, sizeof(kSpv_vt_embedding) / sizeof(uint32_t), kSpecIds_vt_embedding, 3, 6u, 0x00000030u}, + {"vt_fused_chain", kSpv_vt_fused_chain, sizeof(kSpv_vt_fused_chain) / sizeof(uint32_t), nullptr, 0, 9u, 0x000000f0u}, + {"vt_gdn_decode", kSpv_vt_gdn_decode, sizeof(kSpv_vt_gdn_decode) / sizeof(uint32_t), kSpecIds_vt_gdn_decode, 2, 12u, 0x000004c0u}, + {"vt_gdn_post_conv", kSpv_vt_gdn_post_conv, sizeof(kSpv_vt_gdn_post_conv) / sizeof(uint32_t), kSpecIds_vt_gdn_post_conv, 3, 16u, 0x000030fcu}, + {"vt_gdn_prefill", kSpv_vt_gdn_prefill, sizeof(kSpv_vt_gdn_prefill) / sizeof(uint32_t), kSpecIds_vt_gdn_prefill, 2, 12u, 0x000004c0u}, + {"vt_gdn_state_gather", kSpv_vt_gdn_state_gather, sizeof(kSpv_vt_gdn_state_gather) / sizeof(uint32_t), nullptr, 0, 6u, 0x00000003u}, + {"vt_gdn_state_scatter", kSpv_vt_gdn_state_scatter, sizeof(kSpv_vt_gdn_state_scatter) / sizeof(uint32_t), nullptr, 0, 5u, 0x00000003u}, + {"vt_greedy_argmax", kSpv_vt_greedy_argmax, sizeof(kSpv_vt_greedy_argmax) / sizeof(uint32_t), nullptr, 0, 2u, 0x00000002u}, + {"vt_layer_norm", kSpv_vt_layer_norm, sizeof(kSpv_vt_layer_norm) / sizeof(uint32_t), nullptr, 0, 8u, 0x000000c0u}, + {"vt_matmul", kSpv_vt_matmul, sizeof(kSpv_vt_matmul) / sizeof(uint32_t), kSpecIds_vt_matmul, 5, 6u, 0x00000030u}, + {"vt_matmul_coopmat", kSpv_vt_matmul_coopmat, sizeof(kSpv_vt_matmul_coopmat) / sizeof(uint32_t), kSpecIds_vt_matmul_coopmat, 2, 6u, 0x00000030u}, + {"vt_matmul_vec", kSpv_vt_matmul_vec, sizeof(kSpv_vt_matmul_vec) / sizeof(uint32_t), kSpecIds_vt_matmul_vec, 6, 8u, 0x00000030u}, + {"vt_paged_attn", kSpv_vt_paged_attn, sizeof(kSpv_vt_paged_attn) / sizeof(uint32_t), kSpecIds_vt_paged_attn, 4, 11u, 0x000000c0u}, + {"vt_qkv_split", kSpv_vt_qkv_split, sizeof(kSpv_vt_qkv_split) / sizeof(uint32_t), kSpecIds_vt_qkv_split, 2, 8u, 0x000000fcu}, + {"vt_relu", kSpv_vt_relu, sizeof(kSpv_vt_relu) / sizeof(uint32_t), nullptr, 0, 4u, 0x0000000cu}, + {"vt_reshape_and_cache", kSpv_vt_reshape_and_cache, sizeof(kSpv_vt_reshape_and_cache) / sizeof(uint32_t), kSpecIds_vt_reshape_and_cache, 1, 9u, 0x000000f0u}, + {"vt_rms_norm", kSpv_vt_rms_norm, sizeof(kSpv_vt_rms_norm) / sizeof(uint32_t), nullptr, 0, 8u, 0x000000f0u}, + {"vt_rms_norm_gated", kSpv_vt_rms_norm_gated, sizeof(kSpv_vt_rms_norm_gated) / sizeof(uint32_t), nullptr, 0, 8u, 0x000000c0u}, + {"vt_rms_norm_wide", kSpv_vt_rms_norm_wide, sizeof(kSpv_vt_rms_norm_wide) / sizeof(uint32_t), nullptr, 0, 8u, 0x000000f0u}, + {"vt_rope_from_cache", kSpv_vt_rope_from_cache, sizeof(kSpv_vt_rope_from_cache) / sizeof(uint32_t), kSpecIds_vt_rope_from_cache, 5, 7u, 0x0000000fu}, + {"vt_sigmoid_gate_bf16", kSpv_vt_sigmoid_gate_bf16, sizeof(kSpv_vt_sigmoid_gate_bf16) / sizeof(uint32_t), kSpecIds_vt_sigmoid_gate_bf16, 1, 6u, 0x00000030u}, + {"vt_silu_and_mul", kSpv_vt_silu_and_mul, sizeof(kSpv_vt_silu_and_mul) / sizeof(uint32_t), nullptr, 0, 4u, 0x0000000cu}, }; const size_t kSpirvModuleCount = sizeof(kSpirvModules) / sizeof(kSpirvModules[0]); diff --git a/src/vt/vulkan/vulkan_spirv.h b/src/vt/vulkan/vulkan_spirv.h index b7623bf90..e08481558 100644 --- a/src/vt/vulkan/vulkan_spirv.h +++ b/src/vt/vulkan/vulkan_spirv.h @@ -36,12 +36,31 @@ namespace vt::vulkan { // passes specialization values BY ID, and Vulkan SILENTLY IGNORES a map entry // whose ID the module does not declare - so without this table a host/shader // drift produces WRONG NUMBERS instead of a clean failure. +// +// binding_count is the number of descriptor bindings the module declares, and +// writable_mask has bit i set iff binding i is NOT decorated NonWritable, i.e. +// iff the shader is permitted to WRITE that buffer. Both are parsed out of the +// compiled module rather than written down beside it. +// +// This is the READ/WRITE SET the dispatch path needs to decide whether two +// dispatches are genuinely independent. Dispatch() is handed one flat array of +// VkBuffers with no roles attached, and inferring a role from the binding index +// would be an unverified assumption whose failure mode is a MISSING BARRIER and +// therefore silently wrong numbers. `readonly` in the GLSL is enforced by +// glslang (writing such a block is a compile error) and recorded as NonWritable +// in the SPIR-V, so this mask is a machine-checked fact about the shader. +// +// A bit is SET unless read-only is positively proven, so anything the +// reflection does not understand degrades to an extra barrier, never a +// missing one. struct SpirvModule { const char* name; const uint32_t* words; size_t word_count; const uint32_t* spec_ids; size_t spec_id_count; + uint32_t binding_count; + uint32_t writable_mask; }; // DEFINED IN vulkan_spirv.cpp. The array is `extern` and therefore of unknown diff --git a/tests/vt/test_vulkan_backend.cpp b/tests/vt/test_vulkan_backend.cpp index 8f92898af..1cec70450 100644 --- a/tests/vt/test_vulkan_backend.cpp +++ b/tests/vt/test_vulkan_backend.cpp @@ -93,6 +93,72 @@ TEST_CASE("the committed SPIR-V table is present and well-formed") { } } +TEST_CASE("the committed SPIR-V table records each module's WRITABLE bindings") { + // Device-independent: a property of the checked-in artifact. + // + // WHY THIS IS THE MOST LOAD-BEARING ASSERTION IN THE FILE + // (BACKEND-VULKAN-BARRIERS). writable_mask is what lets the dispatch path skip + // a pipeline barrier, and its DANGEROUS failure mode is silent: a mask that + // came back all-zero would describe every binding as read-only, no dispatch + // would ever appear to write anything, no hazard would ever be detected, and + // every barrier in the batch would be dropped. That produces no error and no + // crash -- only wrong numbers, on real hardware, in a way a software + // rasterizer's effectively serial execution hides. So the mask is asserted as a + // STRUCTURE here rather than trusted because the numbers came out right. + for (size_t mi = 0; mi < vt::vulkan::kSpirvModuleCount; ++mi) { + const auto& m = vt::vulkan::kSpirvModules[mi]; + CAPTURE(m.name); + // Every compute shader in this backend produces an output. A module with NO + // writable binding is a reflection failure, not a legitimate shader. + CHECK(m.writable_mask != 0u); + REQUIRE(m.binding_count >= 1u); + // One bit per binding, and the dispatch path's stack arrays are 32 wide. + CHECK(m.binding_count <= 32u); + // No bit may be set above the declared binding count, or the mask and the + // dispatch's buffer array have drifted out of alignment. + const uint32_t above = + m.binding_count >= 32u ? 0u : (m.writable_mask >> m.binding_count); + CHECK(above == 0u); + // Not EVERY binding writable either: each of these shaders reads at least one + // operand, so an all-ones mask is the other reflection failure (it would cost + // barriers rather than correctness, but it would mean nothing was parsed). + const uint32_t all = m.binding_count >= 32u + ? 0xffffffffu + : ((1u << m.binding_count) - 1u); + CHECK(m.writable_mask != all); + } + // Spot checks against the GLSL, read directly from src/vt/vulkan/shaders. These + // pin the reflection to specific known-correct answers, so a generator change + // that starts reporting plausible-but-wrong masks fails here. + struct Expect { const char* name; uint32_t bindings; uint32_t mask; }; + // vt_add.comp: A at 0/1 readonly, B at 2/3 readonly, D (out) at 4/5 writable. + // vt_silu_and_mul.comp: x at 0/1 readonly, out at 2/3 writable. + // vt_greedy_argmax.comp: logits at 0 readonly, out at 1 writable. + // vt_rms_norm.comp: x/weight readonly at 0..3, out at 4/5 and the in-place + // residual stream at 6/7 both writable. + // vt_matmul_vec.comp: a/b at 0..3, out at 4/5, and the two 64-bit ALIASES of a + // and b at 6/7 -- aliases of READ operands, so still read-only. + const Expect kExpect[] = { + {"vt_add", 6u, 0x30u}, + {"vt_silu_and_mul", 4u, 0x0cu}, + {"vt_greedy_argmax", 2u, 0x02u}, + {"vt_rms_norm", 8u, 0xf0u}, + {"vt_matmul_vec", 8u, 0x30u}, + }; + for (const auto& e : kExpect) { + CAPTURE(e.name); + bool found = false; + for (size_t mi = 0; mi < vt::vulkan::kSpirvModuleCount; ++mi) { + const auto& m = vt::vulkan::kSpirvModules[mi]; + if (std::strcmp(m.name, e.name) != 0) continue; + found = true; + CHECK(m.binding_count == e.bindings); + CHECK(m.writable_mask == e.mask); + } + CHECK(found); + } +} + TEST_CASE("the committed SPIR-V table records each module's specialization constants") { // Device-independent: a property of the checked-in artifact, so this also gates // the generator on a box with no Vulkan. @@ -1369,6 +1435,10 @@ TEST_CASE("submission is PIPELINED: a flush does not block, and a drain retires // still-executing batch; dropping it computes correct numbers on a software // rasterizer and races on real hardware, so it is asserted by count and not by // value. + // + // The hazard analysis reaches the SAME count on this loop rather than being + // excused from it: every iteration rewrites `to`, which is a write-after-write + // against the previous iteration, so each one is a genuine dependency. CHECK(ctx.barrier_count() - barriers0 == static_cast(slice) + 1u); // THE SLICES PARTITION THE RING. Without this a mutation collapsing every slot @@ -1445,6 +1515,197 @@ TEST_CASE("submission is PIPELINED: a flush does not block, and a drain retires vk.DestroyQueue(q); } +// A scoped restore for the barrier-policy lever, so a failing assertion inside a +// case below cannot leave the whole rest of the file running under a policy it +// did not ask for. +namespace { +class SmartBarrierArm { + public: + explicit SmartBarrierArm(int v) : prev_(vt::vulkan::VulkanContext::Get().smart_barriers_override()) { + vt::vulkan::VulkanContext::Get().set_smart_barriers_override(v); + } + ~SmartBarrierArm() { vt::vulkan::VulkanContext::Get().set_smart_barriers_override(prev_); } + SmartBarrierArm(const SmartBarrierArm&) = delete; + SmartBarrierArm& operator=(const SmartBarrierArm&) = delete; + + private: + int prev_; +}; +} // namespace + +TEST_CASE("SMART BARRIERS: a real hazard barriers, a proven-independent pair does not") { + if (!VulkanPresent()) return; + auto& ctx = vt::vulkan::VulkanContext::Get(); + if (!ctx.batching_enabled()) return; // no barriers are recorded at all without it + Backend& vk = vt::GetBackend(DeviceType::kVULKAN); + Queue q = vk.CreateQueue(); + const Device d{DeviceType::kVULKAN, 0}; + + // WHY THIS GATE EXISTS (BACKEND-VULKAN-BARRIERS), AND WHY IT ASSERTS COUNTERS. + // Skipping a barrier changes NO number when it is correct and changes numbers + // only on real hardware when it is wrong -- a software rasterizer runs the + // dispatches serially enough to hide a dropped dependency, which is how this + // backend's fence-spin arm passed 33/33 on llvmpipe while producing "the the + // capital capital of of" on GB10. So the mechanism is asserted directly: the + // pairs below are constructed so that the CORRECT number of barriers is known + // in advance from the dependency structure alone, and any predicate that stops + // detecting hazards moves that number. + constexpr int64_t kN = 1024; + const size_t bytes = static_cast(kN) * sizeof(float); + std::vector a(static_cast(kN), 1.5f), b(static_cast(kN), 2.25f); + // FIVE SEPARATE ALLOCATIONS, which is what makes the independence real: this + // backend gives every allocation its own VkBuffer and its own VkDeviceMemory, + // so distinct pointers here cannot alias and "different buffer" is a fact + // rather than an assumption about offsets. + void* da = vk.Alloc(bytes); + void* db = vk.Alloc(bytes); + void* d1 = vk.Alloc(bytes); + void* d2 = vk.Alloc(bytes); + void* d3 = vk.Alloc(bytes); + vk.Copy(q, da, a.data(), bytes); + vk.Copy(q, db, b.data(), bytes); + vk.Synchronize(q); + + Tensor ta = Tensor::Contiguous(da, vt::DType::kF32, d, {kN}); + Tensor tb = Tensor::Contiguous(db, vt::DType::kF32, d, {kN}); + Tensor t1 = Tensor::Contiguous(d1, vt::DType::kF32, d, {kN}); + Tensor t2 = Tensor::Contiguous(d2, vt::DType::kF32, d, {kN}); + Tensor t3 = Tensor::Contiguous(d3, vt::DType::kF32, d, {kN}); + + // vt_add reads both inputs and writes the output -- the writable mask asserted + // by the SPIR-V table case above -- so the five dispatches below have exactly + // one dependency structure and it is visible by inspection. + // + // d0 t1 <- a + b (the policy switch forces this one to barrier) + // d1 t2 <- a + b INDEPENDENT of d0: shares only the READ operands + // d2 t3 <- t1 + b READ-AFTER-WRITE on t1 + // d3 t1 <- a + b WRITE-AFTER-READ on t1 (d2 read it) + // d4 t1 <- a + b WRITE-AFTER-WRITE on t1 (d3 wrote it) + // + // So: 4 barriers and exactly 1 skip. Each of the three hazard KINDS appears + // once, because a predicate that only implements read-after-write is a real + // and plausible mistake that produces correct numbers most of the time. + SUBCASE("the analysis skips exactly the independent pair") { + SmartBarrierArm arm(+1); + REQUIRE(ctx.smart_barriers()); + const uint64_t barriers0 = ctx.barrier_count(); + const uint64_t skips0 = ctx.barrier_skip_count(); + + vt::Add(q, t1, ta, tb); + vt::Add(q, t2, ta, tb); + vt::Add(q, t3, t1, tb); + vt::Add(q, t1, ta, tb); + vt::Add(q, t1, ta, tb); + + const uint64_t barriers = ctx.barrier_count() - barriers0; + const uint64_t skips = ctx.barrier_skip_count() - skips0; + CAPTURE(barriers); + CAPTURE(skips); + // THE MUTATION TARGET. Forcing the hazard predicate to report "independent" + // takes this to 1 and 4 (only the policy-switch barrier survives); dropping + // the write-after-read half alone takes it to 3 and 2. Neither changes any + // value this file checks. + CHECK(barriers == 4u); + CHECK(skips == 1u); + // Every dispatch is accounted for exactly once. + CHECK(barriers + skips == 5u); + vk.Synchronize(q); + } + + // THE CONTROL ARM, in the SAME BINARY. The two arms run the same kernels in the + // same order on the same inputs, so a cross-BUILD comparison could not tell + // them apart at all -- and this campaign has already been given a false 1.2x by + // exactly that shape. Here the lever is proven to MOVE. + SUBCASE("the always-barrier arm records one per dispatch and skips none") { + SmartBarrierArm arm(-1); + REQUIRE(!ctx.smart_barriers()); + const uint64_t barriers0 = ctx.barrier_count(); + const uint64_t skips0 = ctx.barrier_skip_count(); + + vt::Add(q, t1, ta, tb); + vt::Add(q, t2, ta, tb); + vt::Add(q, t3, t1, tb); + vt::Add(q, t1, ta, tb); + vt::Add(q, t1, ta, tb); + + CHECK(ctx.barrier_count() - barriers0 == 5u); + CHECK(ctx.barrier_skip_count() - skips0 == 0u); + vk.Synchronize(q); + } + + // A DEPENDENT CHAIN LONG ENOUGH TO CROSS SEVERAL COMMAND BUFFERS, with the + // analysis on. This is the value half, and it is the half only real hardware + // can fail: each step must observe exactly its predecessor's output, and the + // chain wraps the slot rotation repeatedly so the dependency has to survive the + // command-buffer boundary that a pipelined submission creates. + SUBCASE("a long dependent chain is still exact with the analysis on") { + SmartBarrierArm arm(+1); + const uint32_t steps = ctx.ring_slice() * (ctx.in_flight_limit() + 1) * 2 + 3; + CAPTURE(steps); + std::vector zero(static_cast(kN), 0.0f), one(static_cast(kN), 1.0f); + vk.Copy(q, d1, zero.data(), bytes); + vk.Copy(q, db, one.data(), bytes); + vk.Synchronize(q); + const uint64_t barriers0 = ctx.barrier_count(); + for (uint32_t i = 0; i < steps; ++i) vt::Add(q, t1, t1, tb); + // Every step is a read-after-write AND a write-after-write on t1, so the + // analysis must find a hazard at every single one. A chain like this is + // exactly where a dropped barrier turns into a wrong number. + CHECK(ctx.barrier_count() - barriers0 == static_cast(steps)); + std::vector got(static_cast(kN), -1.0f); + vk.Copy(q, got.data(), d1, bytes); + vk.Synchronize(q); + for (int64_t i = 0; i < kN; i += 128) { + CAPTURE(i); + CHECK(got[static_cast(i)] == doctest::Approx(static_cast(steps))); + } + } + + // INDEPENDENT WRITES FOLLOWED BY A READER. The writes may all skip; the reader + // must not, and its value proves the skipped ones really did complete. This is + // the pattern the decode step would actually exploit -- several projections off + // one normalized activation, then a consumer -- so it is checked for VALUE and + // not only for counts. + SUBCASE("independent writers then a consumer: the consumer barriers, values hold") { + SmartBarrierArm arm(+1); + std::vector three(static_cast(kN), 3.0f); + vk.Copy(q, da, a.data(), bytes); + vk.Copy(q, db, three.data(), bytes); + vk.Synchronize(q); + // Re-assert the policy to force the history to a KNOWN state: the setter + // drains and marks the next dispatch as unconditionally barriered, so the + // dispatch below clears the access sets and the counts that follow start from + // a history containing exactly it. Without this the sets would still carry + // whatever an earlier subcase left, and t2/t3 could look written. + ctx.set_smart_barriers_override(+1); + vt::Add(q, t1, ta, tb); + const uint64_t barriers0 = ctx.barrier_count(); + const uint64_t skips0 = ctx.barrier_skip_count(); + vt::Add(q, t2, ta, tb); // independent of t1: skip + vt::Add(q, t3, ta, tb); // independent of both: skip + vt::Add(q, t1, t2, t3); // reads t2 AND t3: read-after-write, barrier + CHECK(ctx.barrier_count() - barriers0 == 1u); + CHECK(ctx.barrier_skip_count() - skips0 == 2u); + std::vector got(static_cast(kN), -1.0f); + vk.Copy(q, got.data(), d1, bytes); + vk.Synchronize(q); + // t2 = t3 = 1.5 + 3.0 = 4.5, so t1 = 9.0. A dropped barrier before the + // consumer reads whatever t2/t3 held before -- which on this box is the 4.5 + // of an earlier subcase or uninitialised memory, either way not 9.0. + for (int64_t i = 0; i < kN; i += 128) { + CAPTURE(i); + CHECK(got[static_cast(i)] == doctest::Approx(9.0f)); + } + } + + vk.Free(da); + vk.Free(db); + vk.Free(d1); + vk.Free(d2); + vk.Free(d3); + vk.DestroyQueue(q); +} + TEST_CASE("a REFERENCE-TIER op drains the batch before it touches device memory") { if (!VulkanPresent()) return; auto& ctx = vt::vulkan::VulkanContext::Get();