diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index 4192e13f2b..66faee3a92 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -241,7 +241,7 @@ this repository. State remains `ACTIVE`; no lifecycle transition is claimed. | `BACKEND-TPU` | vLLM TPU parity surface | `platforms/__init__.py:35-56,202-208`, `platforms/tpu.py:9-20` | - | - | [CUDA inventory](specs/cuda-architecture-inventory.md) | `INVENTORIED` | - | | `BACKEND-ACCEL-PROVIDER` | **The acceleration-PROVIDER seam** — two or more implementations of ONE `vt::` op on ONE `DeviceType` coexisting, selected DETERMINISTICALLY and observably. Answers the user's standing requirement "build it so we can extend acceleration easily to other platforms", which is a question about the SEAM, not about any one backend | no upstream mirror (vllm.cpp original). Mirrors the SHAPE of the runtime tactic/heuristic dispatch every provider in vLLM's chain uses instead of compile-time pinning: flashinfer's per-arch tactic registry (`flashinfer/gemm/fp4_gemm_cutlass_template_sm120.h:187-220`), cuBLASLt/CUTLASS per-call heuristics | `vt::OpProvider` + device-neutral `vt::ProviderCaps` [op_provider.h](../include/vt/op_provider.h); registry, deterministic selection, decline-and-fall-back and stats [op_provider.cpp](../src/vt/op_provider.cpp). `RegisterOp`/`GetOp`/`OpRegistered` MOVED OUT of [ops.cpp](../src/vt/ops.cpp) with **identical signatures and semantics** — all ~70 op wrappers in that file are byte-unchanged, which is what "zero call-site edits" means. GENERALIZED FROM [cuda_arch_tactics.h](../src/vt/cuda/cuda_arch_tactics.h) (capacity-bounded static storage, capability predicate, decline-by-return, selection stats), lifted out of `vt::cuda` and keyed on (OpId, DeviceType). First consumer: the MLX GEMM provider on Metal [metal_mlx_provider.mm](../src/vt/metal/metal_mlx_provider.mm) | **THE DEFECT FIXED, STATED PRECISELY:** the old table held ONE `void*` per (OpId, DeviceType) and `RegisterOp` overwrote it with no check and no warning, so two providers of one op resolved by STATIC-INIT ORDER ACROSS TUs — unspecified by the standard, i.e. a nondeterministic BUILD. Selection is now `(priority DESC, name ASC by strcmp)`, both compile-time constants of the registering TU, hence a pure function of WHICH providers are linked. **PROVEN, not asserted:** [test_op_provider.cpp](../tests/vt/test_op_provider.cpp) registers the SAME three providers in OPPOSITE orders on two slots and requires the same winner AND the same full order (11 cases / 47 assertions), plus equal-priority name tie-break, duplicate-name rejection, capability-predicate skip, caps re-publication re-resolution, decline-and-fall-back down a 3-deep stack, the `declines` counter, per-call `selections` stats, and the `VT_OP_PROVIDER_DISABLE` same-binary A/B lever. **END-TO-END on a real accelerator (M4):** MLX and the native MSL GEMM coexist on `kMatmul`/`kMatmulBT`, MLX wins by priority, and an interior-pointer activation makes MLX DECLINE exactly once and fall through to ours with the right answer ([test_metal_backend.cpp](../tests/vt/test_metal_backend.cpp), 9 cases / 108 assertions with MLX ON). clean `-Werror` 0 warnings on all three toolchains (AppleClang 21 CLT-only macOS 26.5.2 Metal ON and Metal+MLX ON; GCC Linux CPU; nvcc 13.0 sm_121a on dgx with `VLLM_CPP_TRITON=ON`). **REGRESSION-SAFE on the hottest shared file:** `GetOp` steady state is one relaxed atomic load of a resolved-selection cache (was one array load); negative resolution is memoized so `OpRegistered`, which the fused-recipe ladder calls per step for ops a backend lacks, stays O(1); the provider-disable lookup short-circuits lock-free when nothing is disabled. dgx regression set ALL UNCHANGED, each STANDALONE (see the state log entry) — anchor `tests/vt/test_op_provider.cpp:64` | [Metal/MLX reuse study §6](specs/metal-mlx-reuse-study.md) (which specced it, work row `W0b-2`); reconciled with — not rivalling — [drop-in kernel ABI](specs/dropin-kernel-abi.md), which is the ARGUMENT half for raw-C launchers while this is the SELECTION half | `ACTIVE` — the mechanism is landed and gated with two real providers on one op; it is deliberately not closed, because the CUDA (cuBLASLt/CUTLASS/flashinfer), CPU (llama.cpp `vec_dot`) and Vulkan (coopmat) rows of the §6.1 table are DESIGNED FOR but not yet POPULATED, and the `QuantTypeTraits` split (study §3.4, work row `W0b-3`) that keys on the same predicate is not landed | `CLAIM-BACKEND-ACCEL-PROVIDER-1` | | `BACKEND-METAL-MLX` | Apple Metal — **native MSL** (runtime-compiled), with MLX demoted to an ALTERNATIVE kernel source (work row M5) but PROMOTED to the named competitor/benchmark floor (`BACKEND-GATE-METAL-MLXLM`) | vllm.cpp extension through upstream seam `platforms/interface.py:134-229` | **W0 SKELETON LANDED 2026-07-22.** `vt::Backend` + registrar + allocation registry [metal_backend.mm](../src/vt/metal/metal_backend.mm); runtime-MSL device/library/pipeline context [metal_context.mm](../src/vt/metal/metal_context.mm); embedded MSL [metal_msl.h](../src/vt/metal/metal_msl.h); op encode + `RegisterOp` [metal_ops.mm](../src/vt/metal/metal_ops.mm); `Platform` + registrar [platforms/metal.cpp](../src/vllm/platforms/metal.cpp); CMake tri-state `VLLM_CPP_METAL` (AUTO-on for an Apple host with an ObjC++ compiler). All ADDITIVE — no existing source file changed to enable it. `DeviceType::kMETAL` slot [device.h:16](../include/vt/device.h#L16) needed no edit | **`ACTIVE` MEANS A GATED SKELETON, NOT A SUPPORTED BACKEND — no model runs on Metal.** REGISTERED ops (8): `kAdd` (incl. rank-1 bias broadcast), `kRelu`, `kSiluAndMul`, `kCastBf16`, `kCastF32`, `kLayerNorm`, `kRmsNorm` (incl. in-place residual stream), and ONE `kFusedChain` Tier-1 interpreter — which inherits the whole portable fusion catalog (non-Tier-1 recipes come free via the device-agnostic Tier-0 composite). **2026-07-22 ADDITION — the NATIVE MSL dense GEMM pair `kMatmul`/`kMatmulBT` is now REGISTERED** ([metal_msl.h](../src/vt/metal/metal_msl.h) `vt_matmul`, a threadgroup-tiled f32-accumulating GEMM transcribed from our own CPU `MatmulKernel`/`MatmulBTKernel`; one kernel serves both orientations via a `bt` flag and carries `lda` because `vt::MatmulBT` admits a row-strided activation), taking the registered set to **10**. It exists so the OPTIONAL MLX provider is a CONFIGURATION rather than the only way to get a GEMM here — the seam's premise is two providers of one op coexisting and being A/B-able. **MEASURED vs the CPU oracle on the M4, at real projection widths:** decode-shaped 1x2048x2048 bf16 NMSE 2.80e-06 (`kMatmul`) / 2.76e-06 (`kMatmulBT`); prefill-shaped 32x2048x6144 bf16 2.75e-06 / 2.74e-06; 128x512x512 f32 3.81e-14 / 3.74e-14 — all far inside the NMSE <= 5e-4 bar; no bit-exactness is claimed for a reducing op. **THE MLX PROVIDER IS LANDED, GATED AND DEFAULT-OFF** (`-DVLLM_CPP_MLX=ON -DMLX_ROOT=...`, [metal_mlx_provider.mm](../src/vt/metal/metal_mlx_provider.mm)): inputs are zero-copy via `array::set_data` over our own `MTLBuffer` with a no-op deleter, the OUTPUT is NOT (see the study correction below), `mlx::core::matmul` + an explicit `eval()` is the boundary, and the kernel DECLINES per call (mixed dtypes, row-strided or INTERIOR-pointer activation) and forwards to ours. MLX-vs-CPU NMSE equals the MSL-vs-CPU figure on all six arms and MLX-vs-MSL is 0 on these shapes — recorded as an OBSERVATION, NOT A PROMISE: bit-exactness across providers is explicitly not on offer, and the gate is NMSE. That MLX genuinely COMPUTED (rather than silently declining) is proven by `declines == 0` alongside `last_selected == "mlx"`, which is the Risk-4 failure mode made detectable. **ONE STUDY CORRECTION, MEASURED:** study §5.2 proposed calling `steel_matmul` directly with a pre-bound output; `nm -gU libmlx.dylib` shows `steel_matmul_axpby` is **NOT EXPORTED** by the shipped wheel (only the `get_steel_*_kernel` helpers are), and `Matmul::eval_gpu` re-`set_data`s its output from MLX's allocator anyway, so the output is a host memcpy on unified memory — O(M*N) against an O(M*N*K) GEMM, real, recorded, and never described as zero-copy. **`kPagedAttention` STAYS OURS regardless of MLX** (no paged-KV primitive exists anywhere in MLX, study §5.3). STILL UNREGISTERED/STUBBED: `kPagedAttention`, `kReshapeAndCache`, `kEmbedding`, the entire quant tier and every sampler op (`vt::GetOp` throws; a partial backend is a supported, tested state, `src/vt/ops.cpp:104-111`). `SupportsGraphCapture()` FALSE (`MTLIndirectCommandBuffer` unimplemented); dispatch SYNCHRONOUS (one command buffer per op, commit+wait) — correct, not fast; `get_attn_backend_priority()` deliberately EMPTY (naming a backend with no Metal attention kernel would hand back one whose kernels do not exist). Both the registered AND the stubbed set are asserted as executable facts in [test_metal_backend.cpp](../tests/vt/test_metal_backend.cpp). **EVIDENCE (M4, Apple M4 / macOS 26.5.2 / AppleClang 21):** clean `-Werror` build of the WHOLE tree (lib + all tests) with Metal ON; `test_backend` **7/7 (18/18)** — was 5/7 FAIL before the registrar fix; [test_metal_backend](../tests/vt/test_metal_backend.cpp) **6/6 cases (59/59)**; [test_backend_cross_device](../tests/vt/test_backend_cross_device.cpp) **5/5 cases (73/73)**. **OP NMSE vs the CPU oracle on the same host** (widths 128/100/17 = power-of-two, ragged, sub-simd): `kAdd`/`kAdd`-broadcast/`kRelu` **0 (exact)**; `kSiluAndMul` 2.50e-15; `kRmsNorm` 5.30e-15..1.90e-14; `kRmsNorm` residual stream **0 (exact)**; `kLayerNorm` 3.86e-15..9.95e-15; `kFusedChain` (Tier-0 AND Tier-1) 9.05e-15 — worst case **1.9e-14, eleven orders of magnitude inside the NMSE <= 5e-4 bar**. BIT-EXACT (`memcmp`) for `Copy`/`Memset` and the bf16<->f32 codec including NaN, +-inf, +-0 and 16 exact rounding ties. No bit-exactness is claimed for the reducing ops. Two persistent macOS ctest failures (`test_serve_low_tools`, `test_safetensors`) are PRE-EXISTING platform gaps PROVEN unrelated by a `-DVLLM_CPP_METAL=OFF` A/B on the same tree — Linux-only `os.sched_getaffinity`/`POSIX_FADV_DONTNEED` and `/proc/self/smaps` respectively. **The macOS registrar fix also unblocks `BACKEND-CPU` on macOS.** NOT DONE: GEMM/attention/KV/quant/sampling kernels, async dispatch, graph capture, any model, and any performance number (none is owed until a model runs). **REUSE MAP + FIRST-MODEL COST NOW QUANTIFIED (2026-07-22, [study](specs/metal-mlx-reuse-study.md)):** running **OPT** on Metal needs exactly **9 ops, 6 of them new** (`kEmbedding`, `kMatmulBT`, `kPagedAttention`, `kQkvSplit`, `kReshapeAndCache`, `kGreedyArgmax`) — and all four OPT TUs contain **ZERO** CUDA references, making it the cheapest correct first non-CUDA model; **Qwen3-dense** needs **10 ops, 7 new**, with its entire BF16 forward (`qwen3.cpp:106-231`) already CUDA-free. Weight loading needs **0 OpIds** (host memcpy/transpose only). Seam fixes are **M=4**, one of which was a **REAL PORTABILITY BUG**: `dense_attn_block.h:140,157` selected host-pointer aliasing on `!is_cuda()`, so a Metal/Vulkan run would hand a HOST pointer to a DEVICE kernel — latent only because no model runs on one. **FIXED 2026-07-22** (`!is_cuda()` -> `is_cpu()` in BOTH `ResidentWeight` and `ResidentWeightF32`): host-pointer aliasing is a CPU property, not a not-CUDA property. Behaviour on kCPU and kCUDA is bit-identical, which the unchanged dgx regression set and the 156/156 Linux CPU suite evidence; the three remaining seam fixes (1, 3, 4) are owed at the first Metal model. Corrected counts vs the fan-out spike: CUDA coverage is **74/75** not 73, `vt::Backend` is **6 pure + 18 defaulted** not 6+20/26, raw `kCUDA` comparisons are **54 not 43** plus **13 uncounted `is_cuda()` sites** **=== M3a LANDED 2026-07-22 — THE FIRST MODEL RUNS ON APPLE GPU, AND IT RUNS TOKEN-EXACT. ===** `OPTForCausalLM` (facebook/opt-125m, bf16) generates END TO END through the ordinary engine stack on the M4 and is **STRICT token-exact 6/6 prompts / 96/96 tokens** against the SAME committed dgx-captured vLLM 0.25.0 goldens the CUDA arm is gated on — a DEVICE-INDEPENDENT bar (they are vLLM's tokens, not ours), so Metal met the bar CUDA already met rather than one re-derived on Metal. **FIVE new MSL kernels** took the registered set from 10 to **15 of 75**: `kEmbedding`, `kQkvSplit`, `kReshapeAndCache`, `kPagedAttention`, `kGreedyArgmax` ([metal_msl.h](../src/vt/metal/metal_msl.h), [metal_ops.mm](../src/vt/metal/metal_ops.mm)); each transcribes the per-element math of our own CPU reference and takes only its dispatch shape from llama.cpp. **MEASURED vs the CPU oracle on the same M4:** `kEmbedding`/`kQkvSplit`/`kReshapeAndCache`/`kGreedyArgmax` **BIT-EXACT** (no floating-point reduction => no reordering freedom; argmax reduces but over an order-INDEPENDENT (value, lowest-index) max, pinned by a deliberate two-position tie and an all-equal row); `kPagedAttention` **NMSE 4.99e-13** vs the 5e-4 bar — and bit-exactness is explicitly NOT claimed for it, because the kernel is the algebraically identical ONLINE (flash) softmax where the CPU reference is a materialized 3-pass, i.e. a different reduction order BY CONSTRUCTION. **THE METAL PATH IS PROVEN TO HAVE EXECUTED, NOT INFERRED:** [test_opt_paged_engine.cpp](../tests/vllm/models/test_opt_paged_engine.cpp) asserts `runner().device().type == kMETAL` and, for ALL NINE ops OPT dispatches, `selections > 0` AND `declines == 0` (`kPagedAttention` selections = **1152**); the per-op unit tests additionally NaN-POISON every output buffer so an un-executed kernel cannot pass a numeric check by accident. `last_selected` alone would NOT be proof — a provider can decline INSIDE its kernel and forward down (fan-out spike Risk 4). **SEAM FIXES — the study predicted 4; reality was 3 real, 1 REFUTED, 1 NEW:** item 1 (`model_loader.cpp` hardcoded `GetBackend(kCUDA)`) CONFIRMED and fixed to ask `CurrentPlatform()` — this was the single line keeping every non-NVIDIA accelerator on the CPU reference no matter how complete it was; item 4 (empty `get_attn_backend_priority`) CONFIRMED and fixed to `{"FLASH_ATTN"}` with `FlashAttentionBackend` self-registering for `kMETAL` (a NAME registration only — its host metadata is device-agnostic and the Metal kernels read the same NHD layout), MLA deliberately still unoffered; item 3 (the unguarded `vt/cuda/` include in the runner) **REFUTED BY MEASUREMENT** — declaration-only, it COMPILES AND LINKS in a Metal-only macOS build, so no change was needed or made; and **ONE NEW BUG THE STUDY DID NOT PREDICT** — `runner.cpp:516` gated KV-cache DEVICE RESIDENCY on `is_cuda()`, so on Metal the cache fell into a host `std::vector` and `vt::ReshapeAndCache` was handed a HOST pointer. Same defect CLASS as the `dense_attn_block.h` bug and invisible for the same reason (no test can see it until a model runs on a non-NVIDIA device); fixed to `!is_cpu()`. **NEW SEAM, forced by the above:** `Platform::supports_model_architecture()` (default `true`, so CUDA/CPU are byte-unchanged) — because once `SelectQueue` asks the platform, "which device is this process on" stops being the same question as "which device can run THIS model". Metal's list is exactly `{"OPTForCausalLM"}`; anything else falls back to the CPU reference and runs correctly rather than dying in a kernel bind. **Caught by the macOS regression suite, not by design** — three green tests went red the moment SelectQueue started selecting Metal, and that is recorded as such. **TREE-FRICTION COUNTS RE-JUDGED:** the 54 `kCUDA` / 13 `is_cuda()` counts are accurate but NOT uniform in severity — OPT-on-Metal needed **2 of the 13** changed and **0 of the 5** CUDA includes. The honest headline is not "67 sites to unpick" but "**2 of 13 `is_cuda()` sites encoded *not-NVIDIA means no device memory*, and both were real bugs**". No tree-wide unpicking campaign was started and none is warranted on this evidence. **EVIDENCE:** clean `-Werror` **0 warnings** on a CLEAN FULL macOS rebuild (AppleClang 21, CLT-only, MSL compiled at RUNTIME — there is no offline `metal` compiler on this box); [test_metal_backend](../tests/vt/test_metal_backend.cpp) **12 cases / 18,535 assertions**; full macOS ctest **154/156**, both misses the documented pre-existing platform gaps. (`test_capi` was separately seen failing STANDALONE on macOS and PROVEN pre-existing by building unmodified `origin/main` on the same box and reproducing the identical `CHECK(1 == 2)`; it passes under ctest.) **NOT DONE / NOT CLAIMED: any Metal SPEED number.** The M4 could not be quieted (root LaunchDaemon needs interactive sudo; the desktop aerial wallpaper is the larger contender), so any timing is void under the standing contended-run rule — correctness and speed are separate bars and ONLY CORRECTNESS IS MET. The row therefore stays `ACTIVE`, not `DONE`. Dispatch is still one command buffer per op (commit+wait, work row `M3c`); the GEMM is a plain threadgroup tile loop with no simdgroup-matrix use. **=== M3b LANDED 2026-07-23 — QWEN3-DENSE RUNS ON APPLE GPU, THE SECOND MODEL; FORWARD CORRECT (ORACLE-CONFIRMED NEAR-TIE-ROBUST), NOT STRICT-TOKEN-EXACT (0.6B is ill-posed for strict). ===** `Qwen3ForCausalLM` (Qwen3-0.6B bf16) generates END TO END on the M4. **HONEST STATE — this SUPERSEDES the earlier "SACRED 16/16 strict token-exact / 4 near-ties all at gap-0" claim, which was UNSUBSTANTIATED (the branch test had no teeth; see the RCA + oracle ledger rows).** The Metal forward is a DIFFERENT but equally correct bf16 decoder: on a CLEAN M4 build it diverges from the CUDA golden at 4 first-divergence positions (p0 tok5 = 15344 " Italy" vs 9625 " France"; p5 tok10; p10 tok10; p11 tok1) — a GENUINE bf16 NEAR-TIE, **oracle-confirmed**, NOT a forward bug. **THE DECISIVE MEASUREMENT (vLLM 0.25.0 teacher-forced on the METAL prefix, `scripts/qwen3-neartie-gap.py`, batch=1, gpu_mem_util=0.40, enforce_eager):** every one of the 60 Metal-vs-CUDA divergent positions has Metal's token within **0.5 nats of vLLM's OWN argmax given the Metal prefix — max gap 0.125 nats (worst p5 tok10 / p10 tok12 / p11 tok1 all 0.125), none outside vLLM's top-20**. The France/Italy flip is p0 tok5: metal=15344, and vLLM's teacher-forced argmax on the identical prefix IS **15344 at gap 0.0000** — vLLM contradicts its own CUDA-capture pick of France, the literal near-tie signature. **CUDA is itself build-sensitive at this tie:** the production build (FA2+Marlin+Triton+CUTLASS) resolves p0 tok5 → France 9625; a portable-kernel-only CUDA build resolves it → Italy 15344 (same as Metal), proof it is a numerical near-tie and not a Metal defect. **THE GATE (landed, honest, oracle-backed):** device-appropriate goldens — `our_ids_metal.npy` (the Metal forward's deterministic sequence) + `neartie_gap_mnats_metal.npy` (vLLM teacher-forced on the Metal prefix) — driven through the IDENTICAL gate logic on every device (hard anchor REQUIRE + ≤0.5-nat near-tie band; NO cross-device latitude — Metal gated against Metal's OWN oracle golden, never excused against CUDA's). Metal PASSES **16/16** (10 strict token-exact vs vLLM greedy + 6 near-tie-band, max gap 0.125 nats, 0 forward-divergent). **PROVEN to have teeth:** perturbing the committed Metal anchor (p2 tok0) → hard anchor-drift FAILURE; perturbing a committed gap to 0.6 nats (>0.5 band) → FORWARD-DIVERGENCE band FAILURE; restored → PASS. Per-op correctness stays the NMSE ≤5e-4-vs-CPU proof (RCA, all 28 layers). **STRICT token-exactness on Qwen3-0.6B is ILL-POSED (a near-tie model); a strict Metal gate needs a bigger DETERMINISTIC dense model (Qwen3-4B) — not present on the M4 — DEFERRED.** **THREE new MSL RoPE kernels** took the registered set **15 -> 18 of 75**: `kRopeFromCache` (BIT-EXACT — reads the bf16 cos\|sin cache and only rotates, no reduction/transcendental), `kRopeCosSinCache` and `kRopeNeox` (NMSE bar — Metal has no double, so f32 `pow`/`precise::cos/sin`; measured **0** and **4.07e-15** vs the CPU oracle). **STUDY PREDICTION CORRECTED:** the default path is `VT_QWEN3_ROPE_CACHE`-ON, which dispatches `kRopeCosSinCache` (build the per-step cache) + `kRopeFromCache` (apply it) — not "kRopeFromCache alone"; `kRopeNeox` is the cache-off opt-out. **METAL EXECUTION PROVEN:** device==kMETAL + `selections>0 ∧ declines==0` for all 9 Qwen3-dense ops (`kRopeFromCache`/`kPagedAttention` 7168 each); per-op tests NaN-poison outputs. `Qwen3ForCausalLM` added to `supports_model_architecture`. `test_metal_backend` now **15 cases / 19,331 assertions**, clean `-Werror` 0 warn on the M4; DSR holds 86; Metal TUs `VLLM_CPP_METAL` AUTO->OFF on Linux so the dgx CUDA build is unaffected. The ours-vs-MLX benchmark is now produced (INDICATIVE / BLOCKED-ON-SUDO) — see `BACKEND-GATE-METAL-MLXLM`. **=== S5 REFERENCE TIER (2026-07-23, `CLAIM-BACKEND-SEAM-S5-1`) — Metal is now CORRECT-BY-DEFAULT BEYOND its 18 native kernels. ===** Because Metal is unified memory (StorageModeShared), the S5 portable reference tier ([op_provider.cpp](../src/vt/op_provider.cpp)) installs the CPU kernel as a negative-priority `vt-cpu-ref` fallback on the first `GetOp` miss, so an op Metal lacks a native MSL kernel for runs on the portable CPU reference instead of throwing. **HONEST STATE:** the 18 native MSL kernels remain the FAST path (they always win by priority — Metal's measured numbers above are unchanged); the reference tier only removes the "throw on the first missing op" cliff, so a model needing an op outside the native 18 now RUNS (slowly) rather than not at all. Op count on Metal became a PERFORMANCE budget, not a correctness gate. The Metal OPT run is the hardware form of the S5 zero-native-kernel proof; the hardware-free proof is `test_reference_tier`. Still owed: `M3c` (batched encoders — the named speed lever), `M2r`, `M4`. **=== DISPATCH ATTRIBUTION MEASURED 2026-07-27 — THE BACKEND IS SUBMIT BOUND, NOT KERNEL BOUND, AND THE RECORDED LEVER RANKING WAS WRONG. ===** Instruments has no Metal System Trace without a full Xcode and the M4 has CLT only (`xcrun -f xctrace` -> "not a developer tool"), so the execution trace AGENTS.md requires before any throughput claim did not exist for this backend. It does now: **`VT_METAL_PROFILE`** ([include/vt/metal_profile.h](../include/vt/metal_profile.h), implemented in [metal_ops.mm](../src/vt/metal/metal_ops.mm)) splits every dispatch into HOST encode, submit+wait wall, and REAL GPU busy (`MTLCommandBuffer.GPUEndTime - GPUStartTime`), env-gated (`VT_METAL_PROFILE=1`) with a programmatic switch for tests, costing one relaxed atomic load per dispatch when off. **MEASURED (M4, Qwen3-1.7B-bf16 p=512 g=128 b=1, one binary, arms by `VT_OP_PROVIDER_DISABLE=mlx`, whole series under the GPU lock):** native MSL arm 34.47 s / 3.71 tok/s, **50,944 dispatches**, host encode **0.213 s (0.6%)**, submit+wait **33.882 s = 98.3% of the entire run**, real GPU busy **22.566 s = 66.6% of that wall**, so **11.316 s (33% of the run) is pure round trip**. **THE ROUND-TRIP CONSTANT IS ~186 us**, agreed independently by four near-zero-work kernels as `(wait-gpu)/count`: `vt_rms_norm` 186 us, `vt_silu_and_mul` 195, `vt_reshape_and_cache` 195, `vt_rope_from_cache` 170 — a property of the dispatch model, not of any kernel. **THE SMALL ELEMENTWISE KERNELS ARE ~96% OVERHEAD:** those four are **25,216 of the 50,944 dispatches** and spend **4.69 s of wall for 0.19 s of GPU work** (13.6% of runtime for 0.8% of the GPU's work). **THE CEILING, AND IT IS THE POINT:** ~395 dispatches per decode token x 186 us = **~73 ms/token of pure round trip = a ~13.6 tok/s ceiling WITH INFINITELY FAST KERNELS** (11.4 tok/s at the measured 222 us mean), against MLX-LM's **27.9 tok/s** on the same box and model. **So the competitor floor is unreachable by kernel work alone, and `M3c` is not one lever among several but the PRECONDITION for every other Metal perf lever having anywhere to land.** The M3b narrative naming "`M3c` and a simdgroup GEMM" as co-equal is CORRECTED: the GEMM lever (`vt_matmul` 21,632 dispatches, 26.5 s wait, 20.75 s GPU, 78.3% efficient) is real but its win cannot be OBSERVED end to end while dispatch dominates. **HYPOTHESIS CLOSED, NOT LEFT OPEN:** `GetReferenceTierHits()` is **0** in both arms, so no op is silently running on the S5 portable CPU tier. **MLX ARM IS ONLY PARTIALLY ATTRIBUTED AND IS LABELLED AS SUCH:** MLX-served GEMMs bypass our `Encoder` (`vt_matmul` 21,632 -> 7,168; MLX absorbed exactly 14,464 `kMatmulBT`), leaving **10.08 s of its 20.55 s unaccounted**; closing that is row `M3c-4`. Our-side non-GPU time is still 6.92 s, which is study §5.3's predicted "sync tax" now MEASURED. **EVIDENCE:** [test_metal_backend](../tests/vt/test_metal_backend.cpp) "Metal dispatch profile attributes host encode, wait and GPU busy" (records nothing when off; exactly one row per dispatch when on; pins the `gpu_s <= wait_s` invariant the whole attribution rests on) — suite **18 cases / 19,377 assertions** green on the M4. **STILL INDICATIVE, NOT BINDING** (worker daemon + aerial wallpaper up, no passwordless sudo), though the gpu/wait RATIOS this turns on are far more contention-robust than absolute throughput, both terms being measured inside the same dispatch. Work breakdown `M3c-1`..`M3c-4`, `M3d`, `M5b` in the spec. **=== `M3c-1` LANDED 2026-07-27 — BATCHED COMMAND BUFFERS; THE BACKEND IS NOW COMPUTE BOUND. ===** Dispatches are appended to ONE shared `MTLCommandBuffer`/`MTLComputeCommandEncoder` and committed at a flush point instead of one buffer per op ([metal_ops.mm](../src/vt/metal/metal_ops.mm) `Batch`/`FlushLocked`). Ordering is free inside the encoder because `computeCommandEncoder` is `MTLDispatchTypeSerial`; the HOST is kept from seeing stale bytes by making `Synchronize`/`Copy`/`Memset`/`Free`/`DestroyQueue` flush points ([metal_backend.mm](../src/vt/metal/metal_backend.mm)); and a bind that throws now flushes rather than leaking a half-bound shared encoder into the next op. **NEW SEAM, required for correctness:** `Backend::FlushPending()` (default no-op, [backend.h](../include/vt/backend.h)) called from `GetOp` when the S5 portable CPU tier is the selection ([op_provider.cpp](../src/vt/op_provider.cpp)) — that tier is a HOST kernel about to read and write DEVICE memory, so a deferred submission must drain first. `VT_METAL_SYNC_DISPATCH=1` restores one-buffer-per-op, which is both the spec's required debug mode and the same-binary A/B lever below. **MEASURED (M4, Qwen3-1.7B-bf16 p=512 g=128, SAME binary, arms by `VT_METAL_SYNC_DISPATCH`, 2 reps, arm order alternated, runners paused, whole series under the GPU lock):** b=1 **3.68 -> 5.52 agg tok/s (1.50x)**, 34.79 s -> 23.22 s, TTFT 5369 -> 4937 ms; b=16 **19.12 -> 21.64 agg tok/s (1.13x)**, 107.13 s -> 94.65 s. **Command buffers 50,944 -> 454 at b=1** (112.2 dispatches per commit) and 52,536 -> 469 at b=16. **THE CONTROL THAT MAKES IT A CLEAN RESULT: GPU busy time is UNCHANGED** (22.566 s -> 22.490 s at b=1; 93.957 s -> 93.845 s at b=16), so this removed OVERHEAD, not work — and `gpu_busy_frac_of_wait` moves **66.6% -> 98.4%** (b=1) and **88.4% -> 99.6%** (b=16). **The backend is therefore no longer submit bound; it is compute bound, which promotes `M3d` (simdgroup GEMM) to the next lever exactly as the attribution spec predicted.** The ~11.0 s recovered at b=1 matches the 11.316 s the spec attributed to round trips, which is the prediction confirming the model rather than a new claim. b=16 gains less BECAUSE it was already 88.4% GPU-bound: larger per-dispatch work amortises the fixed ~186 us, which is self-consistent, not a disappointment. **SHIPPING CONFIG (MLX provider + batching), b=1: 9.32 / 9.52 tok/s (13.74 / 13.44 s)** against 6.23 with MLX serial, i.e. the same ~1.5x on top of the GEMM win, narrowing the MLX-LM gap at b=1 from ~4.8x to **~2.96x**. **CORRECTNESS PRECONDITION MET:** the Qwen3-dense Metal gate ran on **device type 2 (METAL)** and passes **128/128** against Metal's own oracle-backed golden; `test_metal_backend` **19 cases / 20,118 assertions**, including three new tests pinning batching (commits==1 for 8 ops + Synchronize), flush-before-host-read, and in-order chaining without intermediate sync. **The OPT gate SKIPPED** (its checkpoint is dgx-only) and is NOT evidence here. **STILL INDICATIVE, NOT BINDING** (worker daemon + wallpaper up), though rep spread was 0.4%. **=== `M3d` LANDED 2026-07-27 — DECODE GEMV, AND THE ROW'S OWN PREMISE WAS WRONG. ===** `M3d` was written as "simdgroup-matrix GEMM". Shape-class profiling after `M3c-1` refuted that: **21,464 of the 21,632 matmuls in a 128-token generation are m=1 decode GEMVs and ALL take the BT orientation**; only **168** are prefill GEMMs, so a simdgroup-matrix kernel would have optimised 0.8% of the dispatches. Implemented instead: `vt_matmul_bt_gemv` ([metal_msl.h](../src/vt/metal/metal_msl.h)) — one simdgroup per output column, streaming the CONTIGUOUS BT weight row fully coalesced, reduced with `simd_sum`. The 16x16 tile kernel wasted 15 of every 16 threadgroup rows at m=1. **MEASURED (same binary, arms by `VT_METAL_NO_GEMV`, 2 reps alternated, under the GPU lock):** b=1 **5.41 -> 10.51 agg tok/s = 1.94x**; b=16 21.55 -> 21.64 = **1.00x, which is an INERTNESS CONTROL rather than a disappointment** — at b=16 the decode m is 16, so the GEMV path is not taken and the numbers SHOULD be identical (21.62 vs 21.64). Batched decode (m=2..16) still runs the tile kernel and is the obvious follow-up. **IT IS ALSO MORE ACCURATE, WHICH IS WHAT FORCED THE GOLDEN RE-CAPTURE:** f32 NMSE vs an f64 oracle **2.68e-14 for the GEMV against 7.40e-13 for the tile kernel, 27x better** (a `simd_sum` tree reduction beats sequential tile accumulation). The bf16 parity arms CANNOT discriminate the two kernels at all — at bf16 output the NMSE is dominated by store rounding (~2.8e-06 for both, identical to six digits) — so an f32 arm was added to the suite precisely to tell a defect from a rounding-order change. **THE SACRED GATE WAS RE-CAPTURED, NOT WEAKENED.** The committed `our_ids_metal.npy` was captured with the LESS accurate tile kernel, so it is kernel-specific and went stale the moment the numerics improved; the gate correctly went red (hard anchor drift p10 tok12). Re-capture followed the documented flow: dump the new Metal sequence on the M4, then **teacher-force vLLM 0.25.0 on dgx** (`qwen3-neartie-gap.py`, oracle venv, run in a SCRATCH dir so the CUDA goldens could not be touched) to regenerate the gaps. **Only 4 of 256 tokens changed, all in prompt 10.** The oracle's verdict on the NEW sequence is BETTER than on the old: **255/256 tokens are vLLM's own argmax (old golden: 253/256)**, mean gap 0.7 vs 1.5 mnats, **0 tokens outside vLLM's top-K**. Gate now **16/16 (10 strict token-exact + 6 near-tie band, max gap 0.188 nats, 0 forward-divergent)** — the same strict/band split as before. **HARNESS FIX REQUIRED TO DO THIS AT ALL:** the anchor `REQUIRE` fired BEFORE the `VT_DUMP_IDS` dump was written, so the documented re-capture was impossible to run; the anchor assertion is now skipped in dump mode ONLY (a dump run writes no golden and asserts nothing). The gate's teeth are unchanged and were demonstrated empirically — it failed on this very change before the re-capture. **EVIDENCE:** `test_metal_backend` **20 cases / 20,125 assertions**, including a test that proves the GEMV actually ROUTED (a numeric check alone cannot: the tile kernel computes the same answer) plus ragged-K and ragged-N arms for the strided reduction tail. `VT_METAL_NO_GEMV=1` is the same-binary A/B lever and the bisect switch. **=== 2-D BLOCKED SIMDGROUP GEMM LANDED 2026-07-27 — PREFILL, +26% END TO END. ===** Re-attribution after `M3c-1`+`M3d` (per-kernel GPU column restored for sync-dispatch runs) found the bottleneck had MOVED: decode GEMV **5,788 ms (50.5%)**, **prefill tile GEMM 3,863 ms (33.7%) from only 168 dispatches**, paged attention 1,627 ms (14.2%), all six other kernels 188 ms (1.6%). Against MLX-LM that split is **prefill ~8.2x slower (3.86 s vs ~0.47 s)** but **decode only ~1.6x (7.4 s vs 4.59 s)** — the opposite of the prior assumption. `vt_matmul_bt_mm` ([metal_msl.h](../src/vt/metal/metal_msl.h)) replaces the 16x16 scalar tile loop for every m > 1: a 32x32 output tile per threadgroup, 4 simdgroups in a 2x2 grid each owning a 2x2 block of `simdgroup_float8x8`, with A and B tiles staged through THREADGROUP memory because `simdgroup_load` needs a typed pointer while our operands carry a runtime dtype code. That staging is also what supplies the A-reuse-across-columns the small-m dead-end proved was the missing property, so ONE kernel serves prefill and batched decode. **MEASURED (isolated same-binary A/B via a new `VT_METAL_NO_MM` lever added so this kernel could be measured independently of the m=1 GEMV; 2 reps, GPU lock, runners idle):** **10.55 -> 13.27 agg tok/s (+26%)**, duration 12.14 -> 9.65 s, **TTFT 4955 -> 2524 ms (halved)** — the prefill win landing exactly where the attribution predicted. **f32 NMSE 1.63e-13** vs the tile kernel's 6.61e-13. **The SACRED gate passed UNCHANGED (16/16, max gap 0.188 nats) with NO golden re-capture:** this kernel's numerics stay on the same side of every near-tie in the gate set. Unit suite **21 cases / 20,134 assertions**, including a ragged 37x333x201 arm that exercises tile edges in M, N and K simultaneously. **Toward the MLX-LM parity goal: 10.5 -> 13.3 of 27.9 tok/s.** Next ranked lever is the decode GEMV (still ~50% of GPU time; per-element dtype switch with scalar loads cannot saturate bandwidth) — anchor `tests/vt/test_metal_backend.cpp:44` | [backend fan-out](specs/backend-fanout-metal-vulkan-xpu.md); **[Metal dispatch attribution](specs/metal-dispatch-attribution.md)**; **[Metal/MLX reuse study + `vt::OpProvider` seam](specs/metal-mlx-reuse-study.md)**; [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md); [CUDA inventory](specs/cuda-architecture-inventory.md) | `ACTIVE` | `CLAIM-BACKEND-FANOUT-1` | -| `BACKEND-VULKAN` | Vulkan compute backend — portable GPU via ahead-of-time-compiled SPIR-V; llama.cpp `ggml/src/ggml-vulkan/` is the port source (vLLM has NO Vulkan path anywhere) | vllm.cpp extension through upstream seam `platforms/interface.py:134-229`; llama.cpp `ggml/src/ggml-vulkan/ggml-vulkan.cpp` @ `237ad9b96` is the maturity reference AND the cited port source | **W0/V1 SKELETON LANDED 2026-07-22.** `vt::Backend` + registrar + allocation registry [vulkan_backend.cpp](../src/vt/vulkan/vulkan_backend.cpp); dlopen entry-point loader [vulkan_loader.cpp](../src/vt/vulkan/vulkan_loader.cpp); instance/device/queue/memory/pipeline context [vulkan_context.cpp](../src/vt/vulkan/vulkan_context.cpp); GLSL compute shaders [shaders/](../src/vt/vulkan/shaders/) compiled AHEAD OF TIME into the committed [vulkan_spirv.h](../src/vt/vulkan/vulkan_spirv.h) by [gen-vulkan-spirv.py](../scripts/gen-vulkan-spirv.py); op dispatch + `RegisterOp` [vulkan_ops.cpp](../src/vt/vulkan/vulkan_ops.cpp); `Platform` + registrar [platforms/vulkan.cpp](../src/vllm/platforms/vulkan.cpp); vendored Khronos TYPE headers [third_party/vulkan/](../third_party/vulkan/) @ `vulkan-sdk-1.4.328.1` under `VK_NO_PROTOTYPES` (nothing linked); CMake tri-state `VLLM_CPP_VULKAN` whose **AUTO resolves OFF** — deliberately UNLIKE Metal's AUTO, because Vulkan overlaps CUDA on the gate box and auto-enabling a second GPU backend would perturb the CUDA regressions. All ADDITIVE — no existing source file changed to enable it. `DeviceType::kVULKAN` slot [device.h:16](../include/vt/device.h#L16) needed no edit | **`ACTIVE` MEANS A GATED SKELETON, NOT A SUPPORTED BACKEND — no model runs on Vulkan.** **AMENDED 2026-08-06 (`VK-A1`, [campaign spec](specs/vulkan-full-support.md) §1.1): the claim below that unregistered ops make `vt::GetOp` THROW is FALSE and has been since accelerator-seam row `S5` (`af0b21ba`) gave unified-memory devices the portable reference tier. MEASURED at runtime: of 87 CPU-registered ops, 8 are NATIVE on Vulkan, 79 are served by the reference tier (the CPU kernel, on shared memory — correct and arbitrarily slow), and ZERO throw. `test_vulkan_backend` asserted the throw and had been RED since S5, unseen because `VLLM_CPP_VULKAN=ON` appeared nowhere in `ci.yml`; both are repaired and a `build-test-vulkan` CI leg now runs the suite GPU-free on llvmpipe. Op resolution is NOT an end-to-end claim: `get_attn_backend_priority()` is still EMPTY and no model has been run.** **EXL3 IS NOW NATIVE HERE** (`BACKEND-VULKAN-EXL3`, [#2530](https://github.com/mudler/vllm.cpp/issues/2530), [spec](specs/backend-vulkan-exl3.md)): `kCastF16` and `kExl3Gemm` were the exactly two ops an EXL3 checkpoint ran on the portable CPU tier when the queue was Vulkan, MEASURED at two `VT_OP_PROVIDER_STATS` fallback notices and now at ZERO. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference rather than ported from `cuda_exl3.cu`, whose 90 KiB shared-memory budget alone exceeds Vulkan's 16 KiB guarantee before one reaches `mma.sync`, `ldmatrix`, `cp.async` or a grid-wide barrier Vulkan has at no version -- so the gate is BYTE equality with the CPU arm across all three codebooks and every width, not a tolerance, and it passed byte-exact on the FIRST run (8 cases / 47 assertions on llvmpipe, NO GPU and NO lease). `kExl3HadR128` is deliberately NOT registered although its shader ships as steps 1 and 3 of that GEMM: no dense forward path calls the op, and the suite ASSERTS the absence. NO speed number is claimed on any axis -- llvmpipe is a software rasteriser, so the shader runs on the same cores the reference kernel does. `kExl3MoeMlp` is owed and needs a grid barrier Vulkan does not have. REGISTERED ops (8, the SAME set as the Metal skeleton so both are comparable through one harness): `kAdd` (incl. rank-1 bias broadcast), `kRelu`, `kSiluAndMul`, `kCastBf16`, `kCastF32`, `kLayerNorm`, `kRmsNorm` (incl. in-place residual stream), and ONE `kFusedChain` Tier-1 interpreter inheriting the whole portable fusion catalog. UNREGISTERED/STUBBED: `kMatmul`/`kMatmulBT`, `kPagedAttention`, `kReshapeAndCache`, `kEmbedding`, the entire quant tier and every sampler op (`vt::GetOp` throws; a partial backend is a supported, tested state, `src/vt/ops.cpp:104-111`). `SupportsGraphCapture()` FALSE (a pre-recorded `VkCommandBuffer` is the eventual mapping); dispatch SYNCHRONOUS and mutex-serialized (record + submit + fence wait per op) — correct, not fast; no staging path for non-host-visible memory; `get_attn_backend_priority()` deliberately EMPTY. Both the registered AND the stubbed set are asserted as executable facts in [test_vulkan_backend.cpp](../tests/vt/test_vulkan_backend.cpp). **SHADER ROUTE, determined and recorded:** libshaderc would be a forbidden compiled dep, so unlike llama.cpp (which shells to `glslc` at BUILD time, `ggml-vulkan/CMakeLists.txt:21-31,141-188`) the SPIR-V is COMMITTED (now 28 modules / 478,256 bytes, glslang **16.5.0**), making the build hermetic with NO shader toolchain anywhere, CI included. **The 2026-07-22 premise that "neither box has `glslc`/`glslangValidator`/`libshaderc` ... neither grants sudo" is CORRECTED as of 2026-09-02** (`BACKEND-VULKAN-EXL3`, [#2530](https://github.com/mudler/vllm.cpp/issues/2530)), and in both directions: `/usr/bin/glslc` DOES exist on the dev box now, and it CANNOT regenerate this tree, because it is glslang 14.0.0 while `vt_matmul_coopmat.comp` requires `GL_EXT_bfloat16`. The route that works needs NO sudo and is the one CI's `vulkan-spirv-freshness` job already takes -- unpack the PINNED glslang 16.5.0 release tarball into a temp dir and put it first on `PATH` -- and it reproduces the committed blob BYTE FOR BYTE, which is the check to run BEFORE regenerating anything. **RELAXED-PRECISION KNOBS PINNED:** `1.0/sqrt` instead of llama.cpp's `inversesqrt` (`rms_norm.comp:86`, `norm.comp:39`) — the Vulkan analogue of Metal's `MTLMathModeSafe`; no `RelaxedPrecision` decorations; fp32 float-controls PROBED and reported rather than assumed (llvmpipe: denorm-preserve false, signed-zero/Inf/NaN-preserve true), which cannot move a gated result because the bf16/f16 codecs are integer transcriptions of `src/vt/dtype.cpp`. **EVIDENCE — dev box (`llvmpipe`, Vulkan 1.4, `mesa-vulkan-drivers`, NO GPU):** clean `-Werror` 0 warn; [test_vulkan_backend](../tests/vt/test_vulkan_backend.cpp) **8/8 (82/82)**; [test_backend_cross_device](../tests/vt/test_backend_cross_device.cpp) **5/5 (73/73)**; full CPU+Vulkan ctest **155/156**, the one failure `test_openai_conformance` a known `-j` parallelism flake that PASSES on rerun — **a GPU-FREE CI PATH IS PROVEN**. **EVIDENCE — GB10:** clean CUDA `-Werror` **0 warnings** on BOTH the production (Vulkan-OFF) and Vulkan-ON builds; `test_vulkan_backend` **8/8 (82/82)**; `test_backend_cross_device` **5/5 (144/144)** — DOUBLE the dev-box count because the SAME BINARY compares the CPU oracle against **both CUDA and Vulkan**, the cross-backend gate this row's spike promised; `test_backend` 7/7, `test_platform` 7/7. NOT DONE: GEMM/attention/KV/quant/sampling kernels, coopmat/coopmat2 tactics, async dispatch, graph capture, any model, and any performance number (none is owed until a model runs; llama.cpp's Vulkan backend is the eventual floor, `BACKEND-GATE-VULKAN-LLAMACPP`, still `INVENTORIED`). **=== S5 REFERENCE TIER (2026-07-23, `CLAIM-BACKEND-SEAM-S5-1`) — Vulkan is CORRECT-BY-DEFAULT BEYOND its 8 native kernels ONLY on a UNIFIED-MEMORY Vulkan device (GB10, integrated GPUs). ===** The S5 portable reference tier ([op_provider.cpp](../src/vt/op_provider.cpp)) installs the CPU kernel as a negative-priority `vt-cpu-ref` fallback on the first `GetOp` miss — but its gate is `Backend::UnifiedMemory()`, and `VulkanContext::unified_memory()` is FALSE on a DISCRETE Vulkan GPU, so a discrete Vulkan device NEVER gets a CPU fallback (a host kernel against discrete VRAM is corruption) and still throws on a missing op, exactly as before. On llvmpipe/GB10 (unified) the fallback runs, removing the throw-on-first-missing-op cliff; a discrete-Vulkan staging-copy variant is a later, separately-gated row (audit Risk 8). No model runs on Vulkan yet regardless — anchor `tests/vt/test_vulkan_backend.cpp:46` | [backend fan-out](specs/backend-fanout-metal-vulkan-xpu.md); [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md); [CUDA inventory](specs/cuda-architecture-inventory.md) | `ACTIVE` | `CLAIM-BACKEND-FANOUT-1` | +| `BACKEND-VULKAN` | Vulkan compute backend — portable GPU via ahead-of-time-compiled SPIR-V; llama.cpp `ggml/src/ggml-vulkan/` is the port source (vLLM has NO Vulkan path anywhere) | vllm.cpp extension through upstream seam `platforms/interface.py:134-229`; llama.cpp `ggml/src/ggml-vulkan/ggml-vulkan.cpp` @ `237ad9b96` is the maturity reference AND the cited port source | **W0/V1 SKELETON LANDED 2026-07-22.** `vt::Backend` + registrar + allocation registry [vulkan_backend.cpp](../src/vt/vulkan/vulkan_backend.cpp); dlopen entry-point loader [vulkan_loader.cpp](../src/vt/vulkan/vulkan_loader.cpp); instance/device/queue/memory/pipeline context [vulkan_context.cpp](../src/vt/vulkan/vulkan_context.cpp); GLSL compute shaders [shaders/](../src/vt/vulkan/shaders/) compiled AHEAD OF TIME into the committed [vulkan_spirv.h](../src/vt/vulkan/vulkan_spirv.h) by [gen-vulkan-spirv.py](../scripts/gen-vulkan-spirv.py); op dispatch + `RegisterOp` [vulkan_ops.cpp](../src/vt/vulkan/vulkan_ops.cpp); `Platform` + registrar [platforms/vulkan.cpp](../src/vllm/platforms/vulkan.cpp); vendored Khronos TYPE headers [third_party/vulkan/](../third_party/vulkan/) @ `vulkan-sdk-1.4.328.1` under `VK_NO_PROTOTYPES` (nothing linked); CMake tri-state `VLLM_CPP_VULKAN` whose **AUTO resolves OFF** — deliberately UNLIKE Metal's AUTO, because Vulkan overlaps CUDA on the gate box and auto-enabling a second GPU backend would perturb the CUDA regressions. All ADDITIVE — no existing source file changed to enable it. `DeviceType::kVULKAN` slot [device.h:16](../include/vt/device.h#L16) needed no edit | **`ACTIVE` MEANS A GATED SKELETON, NOT A SUPPORTED BACKEND — no model runs on Vulkan.** **AMENDED 2026-08-06 (`VK-A1`, [campaign spec](specs/vulkan-full-support.md) §1.1): the claim below that unregistered ops make `vt::GetOp` THROW is FALSE and has been since accelerator-seam row `S5` (`af0b21ba`) gave unified-memory devices the portable reference tier. MEASURED at runtime: of 87 CPU-registered ops, 8 are NATIVE on Vulkan, 79 are served by the reference tier (the CPU kernel, on shared memory — correct and arbitrarily slow), and ZERO throw. `test_vulkan_backend` asserted the throw and had been RED since S5, unseen because `VLLM_CPP_VULKAN=ON` appeared nowhere in `ci.yml`; both are repaired and a `build-test-vulkan` CI leg now runs the suite GPU-free on llvmpipe. Op resolution is NOT an end-to-end claim: `get_attn_backend_priority()` is still EMPTY and no model has been run.** **EXL3 IS NOW NATIVE HERE** (`BACKEND-VULKAN-EXL3`, [#2530](https://github.com/mudler/vllm.cpp/issues/2530), [spec](specs/backend-vulkan-exl3.md)): `kCastF16` and `kExl3Gemm` were the exactly two ops an EXL3 checkpoint ran on the portable CPU tier when the queue was Vulkan, MEASURED at two `VT_OP_PROVIDER_STATS` fallback notices and now at ZERO. `kExl3Gemm` is TRANSCRIBED from the portable CPU reference rather than ported from `cuda_exl3.cu`, whose 90 KiB shared-memory budget alone exceeds Vulkan's 16 KiB guarantee before one reaches `mma.sync`, `ldmatrix`, `cp.async` or a grid-wide barrier Vulkan has at no version -- so the gate is BYTE equality with the CPU arm across all three codebooks and every width, not a tolerance, and it passed byte-exact on the FIRST run (8 cases / 47 assertions on llvmpipe, NO GPU and NO lease). `kExl3HadR128` is deliberately NOT registered although its shader ships as steps 1 and 3 of that GEMM: no dense forward path calls the op, and the suite ASSERTS the absence. NO speed number is claimed on any axis -- llvmpipe is a software rasteriser, so the shader runs on the same cores the reference kernel does. `kExl3MoeMlp` is owed and needs a grid barrier Vulkan does not have. REGISTERED ops (8, the SAME set as the Metal skeleton so both are comparable through one harness): `kAdd` (incl. rank-1 bias broadcast), `kRelu`, `kSiluAndMul`, `kCastBf16`, `kCastF32`, `kLayerNorm`, `kRmsNorm` (incl. in-place residual stream), and ONE `kFusedChain` Tier-1 interpreter inheriting the whole portable fusion catalog. UNREGISTERED/STUBBED: `kMatmul`/`kMatmulBT`, `kPagedAttention`, `kReshapeAndCache`, `kEmbedding`, the entire quant tier and every sampler op (`vt::GetOp` throws; a partial backend is a supported, tested state, `src/vt/ops.cpp:104-111`). `SupportsGraphCapture()` FALSE (a pre-recorded `VkCommandBuffer` is the eventual mapping); dispatch SYNCHRONOUS and mutex-serialized (record + submit + fence wait per op) — correct, not fast; no staging path for non-host-visible memory; `get_attn_backend_priority()` deliberately EMPTY. Both the registered AND the stubbed set are asserted as executable facts in [test_vulkan_backend.cpp](../tests/vt/test_vulkan_backend.cpp). **SHADER ROUTE, determined and recorded:** libshaderc would be a forbidden compiled dep, so unlike llama.cpp (which shells to `glslc` at BUILD time, `ggml-vulkan/CMakeLists.txt:21-31,141-188`) the SPIR-V is COMMITTED (now 28 modules / 478,256 bytes, glslang **16.5.0**), making the build hermetic with NO shader toolchain anywhere, CI included. **The 2026-07-22 premise that "neither box has `glslc`/`glslangValidator`/`libshaderc` ... neither grants sudo" is CORRECTED as of 2026-09-02** (`BACKEND-VULKAN-EXL3`, [#2530](https://github.com/mudler/vllm.cpp/issues/2530)), and in both directions: `/usr/bin/glslc` DOES exist on the dev box now, and it CANNOT regenerate this tree, because it is glslang 14.0.0 while `vt_matmul_coopmat.comp` requires `GL_EXT_bfloat16`. The route that works needs NO sudo and is the one CI's `vulkan-spirv-freshness` job already takes -- unpack the PINNED glslang 16.5.0 release tarball into a temp dir and put it first on `PATH` -- and it reproduces the committed blob BYTE FOR BYTE, which is the check to run BEFORE regenerating anything. **RELAXED-PRECISION KNOBS PINNED:** `1.0/sqrt` instead of llama.cpp's `inversesqrt` (`rms_norm.comp:86`, `norm.comp:39`) — the Vulkan analogue of Metal's `MTLMathModeSafe`; no `RelaxedPrecision` decorations; fp32 float-controls PROBED and reported rather than assumed (llvmpipe: denorm-preserve false, signed-zero/Inf/NaN-preserve true), which cannot move a gated result because the bf16/f16 codecs are integer transcriptions of `src/vt/dtype.cpp`. **EVIDENCE — dev box (`llvmpipe`, Vulkan 1.4, `mesa-vulkan-drivers`, NO GPU):** clean `-Werror` 0 warn; [test_vulkan_backend](../tests/vt/test_vulkan_backend.cpp) **8/8 (82/82)**; [test_backend_cross_device](../tests/vt/test_backend_cross_device.cpp) **5/5 (73/73)**; full CPU+Vulkan ctest **155/156**, the one failure `test_openai_conformance` a known `-j` parallelism flake that PASSES on rerun — **a GPU-FREE CI PATH IS PROVEN**. **EVIDENCE — GB10:** clean CUDA `-Werror` **0 warnings** on BOTH the production (Vulkan-OFF) and Vulkan-ON builds; `test_vulkan_backend` **8/8 (82/82)**; `test_backend_cross_device` **5/5 (144/144)** — DOUBLE the dev-box count because the SAME BINARY compares the CPU oracle against **both CUDA and Vulkan**, the cross-backend gate this row's spike promised; `test_backend` 7/7, `test_platform` 7/7. NOT DONE: GEMM/attention/KV/quant/sampling kernels, coopmat/coopmat2 tactics, async dispatch, graph capture, any model, and any performance number (none is owed until a model runs; llama.cpp's Vulkan backend is the eventual floor, `BACKEND-GATE-VULKAN-LLAMACPP`, still `INVENTORIED`). **=== S5 REFERENCE TIER (2026-07-23, `CLAIM-BACKEND-SEAM-S5-1`) — Vulkan is CORRECT-BY-DEFAULT BEYOND its 8 native kernels ONLY on a UNIFIED-MEMORY Vulkan device (GB10, integrated GPUs). ===** The S5 portable reference tier ([op_provider.cpp](../src/vt/op_provider.cpp)) installs the CPU kernel as a negative-priority `vt-cpu-ref` fallback on the first `GetOp` miss — but its gate is `Backend::UnifiedMemory()`, and `VulkanContext::unified_memory()` is FALSE on a DISCRETE Vulkan GPU, so a discrete Vulkan device NEVER gets a CPU fallback (a host kernel against discrete VRAM is corruption) and still throws on a missing op, exactly as before. On llvmpipe/GB10 (unified) the fallback runs, removing the throw-on-first-missing-op cliff; a discrete-Vulkan staging-copy variant is a later, separately-gated row (audit Risk 8). No model runs on Vulkan yet regardless. **=== 2026-09-27: the two claims in the preceding sentence are STALE, and the discrete-staging half of Risk 8 is ANSWERED. ===** "A discrete-Vulkan staging-copy variant is a later row" and "no model runs on Vulkan yet" were both true of the 2026-07-22 skeleton this cell originally recorded, and neither is true now. **(a) A MODEL RUNS ON VULKAN:** `test_opt_paged_engine` (OPT-125m) is 6/6 prompts token-exact (96/96) with 0 declines on BOTH GB10 and llvmpipe, and Qwen3-4B bf16 and Qwen3.6-27B bf16 have both been run (`BACKEND-VULKAN-LOADMEM`, where `VT_ADOPT_DEVICE_BYTES` took VmHWM 100.759 -> 53.413 GiB on the 27B). Counts: `test_vulkan_backend` **35/35 (2650)**, `test_backend_cross_device` **11/11 (132)**. **THE REGISTERED SURFACE IS 35 OPS, not the 8 this cell's first paragraph lists** — the tree registers the dense set (`kMatmul`/`kMatmulBT`, `kPagedAttention`, `kEmbedding`, `kQkvSplit`, `kReshapeAndCache`, `kGreedyArgmax`), the RoPE trio, the full GDN/SSM set (`kGdnStateGather`/`kGdnStateScatter`/`kGdnPostConv`/`kGdnPrefill`/`kGdnDecode`), `kSigmoidGateBf16`/`kRmsNormGated`, `kExl3Gemm`, and the GGUF keep-quant tier (`kMatmulBTQuant`, TQ1_0/TQ2_0, #2248). **(b) THE STAGING PATH IS ANSWERED BY ReBAR, not owed as code.** An **Intel Arc Pro B60** (`garlic-clove`, `8086:e211`, `xe` driver) is on the estate — the hardware `VK-I` was scoped against, and the "acquire later" half of the 2026-08-06 decision. MEASURED 2026-09-27: `PHYSICAL_DEVICE_TYPE_DISCRETE_GPU`, but `memoryTypes[3]`/`[6]` expose `DEVICE_LOCAL \| HOST_VISIBLE \| HOST_COHERENT` (0x0007) via Resizable BAR, which is the type `vulkan_context.cpp:873` already prefers — so the "non-host-visible memory" condition a staging path would survive **does not arise on this card**. That is a property of **ReBAR, not of Arc**: on a discrete card without it, the keep-quant CPU fall-through and the portable reference tier are corruption, not slowness, so the ordered-fallback requirement STANDS. The `BACKEND-VULKAN-KEEPQUANT` comment that asserted "The B60 is integrated" was false (`vulkaninfo` says DISCRETE) and is corrected in this change. **`VK-I` is PARTIALLY answered: the gate re-run where Vulkan is the only path is still OWED** and 20.91 GiB is its binding constraint — see [full-support](specs/vulkan-full-support.md) §6.2. Also note `VLLM_CPP_DEVICE=vulkan` is a PLACEBO: it is read nowhere in the tree, so every "on Vulkan" claim here must cite the test's printed `BACKEND PROOF` device type — anchor `tests/vt/test_vulkan_backend.cpp:46` | [backend fan-out](specs/backend-fanout-metal-vulkan-xpu.md); [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md); [CUDA inventory](specs/cuda-architecture-inventory.md) | `ACTIVE` | `CLAIM-BACKEND-FANOUT-1` | | `BACKEND-ANE` | Apple Neural Engine for encoder/pooling/fixed-shape draft classes | vllm.cpp extension through upstream seam `platforms/interface.py:134-229`; not a paged decode backend | Platform seam anchored [interface.h:56](../include/vllm/platforms/interface.h#L56) (a `Platform` subclass; not a paged-decode backend) | - | [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md); [CUDA inventory](specs/cuda-architecture-inventory.md) | `INVENTORIED` | - | | `BACKEND-TENSTORRENT` | Tenstorrent Blackhole (Tensix multicore, discrete PCIe, no unified memory) — thin `vt::` adapter over ttnn's existing C++ op library rather than hand-written kernels, mirroring the Metal/MLX decision (E1); no upstream TT platform existed at W0; vLLM's `vllm-tt-plugin` (2026-09-07, [`oracles/vllm-tt-plugin.md`](oracles/vllm-tt-plugin.md)) now registers `TTQwen3_5ForConditionalGeneration` and is the primary-rank denominator where it serves | vllm.cpp extension through upstream seam `platforms/interface.py:134-229` (same pattern as Metal/Vulkan) | **ACTIVE 2026-08-10.** `vt::tenstorrent::Backend` + registrar [tenstorrent_backend.cpp](../src/vt/tenstorrent/tenstorrent_backend.cpp); shared mesh-device lifecycle [tenstorrent_device.cpp](../src/vt/tenstorrent/tenstorrent_device.cpp); 17 registered ops cover OPT-125m and the Qwen3-0.6B forward (`kMatmul`, `kMatmulBT`, `kAdd`, `kRelu`, `kEmbedding`, `kLayerNorm`, `kRmsNorm`, `kSiluAndMul`, bf16/f32 casts, three RoPE forms, `kQkvSplit`, `kReshapeAndCache`, host-oracle `kPagedAttention`, `kGreedyArgmax`) [tenstorrent_ops.cpp](../src/vt/tenstorrent/tenstorrent_ops.cpp); platform allow-list selects OPT and Qwen3 [platforms/tenstorrent.cpp](../src/vllm/platforms/tenstorrent.cpp). `DeviceType::kTENSTORRENT` [device.h](../include/vt/device.h) | [test_tenstorrent_backend.cpp](../tests/vt/test_tenstorrent_backend.cpp) carries real-Blackhole op gates; [test_qwen3_paged_engine.cpp](../tests/parity/test_qwen3_paged_engine.cpp) selects Tenstorrent device-specific anchor and teacher-forced near-tie goldens. OPT-125m STRICT 6/6 passed. Qwen3.5-27B GSQ-RCO GGUF e2e PASSED (2026-09-22): both ISTA-DASLab IQ3_XXS and IQ3_S files decode keep-quant on the P150, tokens argmax-exact vs the pinned llama.cpp b10451 batch decode (gap 0.0 mnats, 16/16, band 500; evidence 10 in the row issue). Qwen3 short warm smoke ran 4 tokens at about 0.28 tok/s; full 16x16 gate remains pending behind host paged attention | [tenstorrent-backend.md](specs/tenstorrent-backend.md) | `ACTIVE` | `CLAIM-BACKEND-TENSTORRENT-SPIKE` | | `BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | Child of `BACKEND-TENSTORRENT` — the owed op-level numerics evidence at the residual-RMS device boundary (`kDeviceResidualMinRows == 32`): device path does `ttnn::add`+`ttnn::rms_norm` in bf16; host/CPU path accumulates in f32. Bot-flagged on #289; never measured at the boundary. | vllm.cpp CPU oracle `RmsNormKernel` mirrors vLLM `fused_add_rms_norm` (add in model dtype, variance in f32); `src/vt/cpu/cpu_ops.cpp:371-398` | `src/vt/tenstorrent/tenstorrent_ops.cpp:1067-1117` (host/device split, `kDeviceResidualMinRows=32`) | [test_tenstorrent_backend.cpp](../tests/vt/test_tenstorrent_backend.cpp) `kRmsNorm residual: device vs CPU f32 oracle across the rows=32 boundary`: 22/22 cases on real Blackhole P150. **Measured 2026-08-11:** host path `rows<32` bit-identical to CPU (`max_abs=0`); device bf16 path `rows>=32` diverges by constant **0.0459 abs** (1.9–2.6× rel on near-zero outputs) — bf16 rounding signature, not accumulation. Decision pending the e2e golden tie-break | [tenstorrent-residual-golden.md](specs/tenstorrent-residual-golden.md) | `SPIKE` | `CLAIM-BACKEND-TENSTORRENT-RESIDUAL-GOLDEN` | diff --git a/.agents/environment.md b/.agents/environment.md index ab766038fd..b4eb24506e 100644 --- a/.agents/environment.md +++ b/.agents/environment.md @@ -117,6 +117,56 @@ Three consequences for anyone sizing work here: at 73.9 GiB `VmHWM`; the host side is now 31 GiB total. Run a CPU comparison on `thor` or `dgx`, or on-box against an oracle instead. +### `garlic-clove` — Intel Arc Pro B60 + +**It is NOT in the fleet table above, and it is not a fleet device.** Do not +read its absence as unavailability: it is reachable, and it is the only Intel +GPU on this estate. It is deliberately absent from the `rc` table because it is +not enrolled in resource-controller — `rc devices` could not be consulted from +the shell this was written in (no `rc` client), and nothing here may claim +membership it has not verified. **Therefore the lease rule does not apply, and +the file mutex does: take `${GPU_LOCK:-$HOME/gpu.lock}` ON THE BOX, as the +non-fleet-device clause of §"Reaching a GPU" requires.** Do not `ssh` in and +start GPU work unguarded, and do not add it to the fleet table until `rc +devices` actually reports it. + +| | | +|---|---| +| Host | `garlic-clove`, Tailscale name, `100.67.232.8`, Linux, account `yoav` (lowercase) | +| GPU | **Intel Arc Pro B60 Graphics (BMG G21)**, `8086:e211`, ASUS subsys `1849:6023`, `xe` kernel driver | +| Vulkan | device API **1.4.354**, conformance 1.4.0.0, Mesa **26.2.3** (kisak PPA), LLVM 21.1.8; `VK_KHR_cooperative_matrix` = true, `VK_KHR_shader_bfloat16`, `VK_KHR_shader_integer_dot_product` | +| Device type | **`PHYSICAL_DEVICE_TYPE_DISCRETE_GPU`** | +| Memory | 20.91 GiB `DEVICE_LOCAL` heap + 23.44 GiB host heap; `memoryTypes[3]`/`[6]` = `DEVICE_LOCAL \| HOST_VISIBLE \| HOST_COHERENT` (0x0007) via ReBAR | +| Host | Ubuntu 24.04.4 LTS, kernel 7.0.0-34, i7-10700 8c/16t, 31 GiB RAM, 465 GiB NVMe (159 GiB free) | +| Toolchain | cmake 3.28.3, gcc 13.3.0, ninja 1.11.1, git, py3.12. **No `icpx`, no `sycl-ls`, no `/dev/accel/*`.** | +| Assets | **no model weights** — `~/.cache/huggingface` is 3.5 MB; only vocab-only GGUFs under `~/llama-maple/models` | +| Checkout | `~/vllm.cpp`, was on `row/BACKEND-VULKAN-TQ1_0-finish` with `VLLM_CPP_VULKAN=ON`, `BUILD EXIT 0`; that work appears to have landed on `main` as #2248, so the branch needs a fetch/prune, not a merge | + +**Its load-bearing property is ReBAR, and that is not a property of Arc — but it +is a performance property, not a correctness one.** ReBAR is what makes the +allocation *device-local AND* host-addressable on a *discrete* card, which is +what the GGUF keep-quant CPU fall-through and the portable reference tier both +want. `VulkanContext` (`vulkan_context.cpp:872-875`) does not require +`DEVICE_LOCAL`: it prefers the device-local host-visible type and falls back to +plain `HOST_VISIBLE | HOST_COHERENT`, and `AllocBuffer` maps whichever it picks, +so `DeviceMemoryIsHostAddressable()` stays true without ReBAR. A discrete board +without ReBAR therefore runs those paths **slower**, with the GPU reaching the +allocation across the bus, not incorrectly. Host visibility is the correctness +requirement; device locality is the performance preference. See +[`specs/vulkan-full-support.md`](specs/vulkan-full-support.md) §4 and the +comment at `src/vt/vulkan/vulkan_ops.cpp` in the `BACKEND-VULKAN-KEEPQUANT` +block. + +**Two cosmetic warts, recorded because a clean enumeration is a gate here.** A +stale `dzn_icd.json` (declaring `api_version 1.1.354`) makes the loader print +`Received return code -9 from call to vkCreateInstance in ICD libvulkan_dzn.so. +Skipping this driver` on every enumeration. It is HARMLESS — the working +`intel_icd.json` provides the device, and `vulkaninfo --summary` lists the B60 +as GPU0 — but it will pollute any recorded enumeration. The loader's *instance* +version is also 1.3.275 while the *device* is 1.4.354; the backend `dlopen`s the +loader and only the device capability matters, but the two numbers should not +be quoted interchangeably. + ### `orin:gpu0` needs L4T CUDA 12.6, and the DGX recipe breaks it **Do not install `cuda-toolkit-13-*` from the generic `sbsa` repo on `orin`.** That @@ -2340,6 +2390,16 @@ inner 4096, state 128; context 262144. Enumeration and the clean-tree rule: [`specs/oracle-llamacpp-repin-stock.md`](specs/oracle-llamacpp-repin-stock.md). -- **No Intel GPU exists on any box here**, so `BACKEND-XPU` end-to-end work is - HW-BLOCKED; only policy-port, compile coverage and oneAPI CPU-device unit - numerics are available. +- **An Intel GPU now exists on this estate: `garlic-clove`, an Intel Arc Pro B60.** + This line used to read "**No Intel GPU exists on any box here**, so + `BACKEND-XPU` end-to-end work is HW-BLOCKED", and that was true when written + but is FALSE as of 2026-09-27. The board is the hardware `VK-I` was scoped + against; see [garlic-clove](#garlic-clove--intel-arc-pro-b60) below. **`BACKEND-XPU` + nonetheless stays `SPIKE`/HW-blocked for a DIFFERENT and still-true reason: the + SYCL toolchain is absent, not the GPU.** Measured on the box 2026-09-27 — no + `icpx`, no `sycl-ls`, and no `/dev/accel/*` nodes, so there is no Level Zero + driver userspace to dispatch through. The Level Zero *runtime* libraries are + installed (`libze_loader.so.1`, `libze_intel_gpu.so.1`) and are not + sufficient. So the available work remains policy-port, compile coverage and + oneAPI CPU-device unit numerics, and closing `BACKEND-XPU` now needs an + oneAPI install rather than hardware acquisition. diff --git a/.agents/issues/BACKEND-VULKAN/ISSUE-LOCAL-01M3JGPW1FN5SG506ANGBT3C0Q.md b/.agents/issues/BACKEND-VULKAN/ISSUE-LOCAL-01M3JGPW1FN5SG506ANGBT3C0Q.md new file mode 100644 index 0000000000..b5614baf15 --- /dev/null +++ b/.agents/issues/BACKEND-VULKAN/ISSUE-LOCAL-01M3JGPW1FN5SG506ANGBT3C0Q.md @@ -0,0 +1,27 @@ +ID: ISSUE-LOCAL-01M3JGPW1FN5SG506ANGBT3C0Q +Title: VK-I: B60 ReBAR retires the discrete staging path; the 'B60 is integrated' comment and the 'no Intel GPU' registry line are both false +Row: BACKEND-VULKAN +State: OPEN +Kind: bug +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-27 +Updated: 2026-10-02 +Closed: - + +## Problem + +Three records about the Intel Arc Pro B60 are wrong, and one planned deliverable is answered by hardware rather than by code. + +(1) FALSE CODE COMMENT. src/vt/vulkan/vulkan_ops.cpp:2109 asserts 'The B60 is integrated'. Measured on garlic-clove with vulkaninfo, the device reports deviceType = PHYSICAL_DEVICE_TYPE_DISCRETE_GPU, with two separate heaps (20.91 GiB DEVICE_LOCAL, 23.44 GiB host-visible). The CODE is right for the stated WRONG reason: VulkanBackend::DeviceMemoryIsHostAddressable() (vulkan_backend.cpp:135) returns true unconditionally, and vulkan_context.cpp:873 prefers a DEVICE_LOCAL|HOST_VISIBLE|HOST_COHERENT type, which memoryTypes[3] and memoryTypes[6] (propertyFlags 0x0007) satisfy. The property the code depends on holds; the stated reason does not. This is a trap rather than a typo: the comment grounds host addressability in the card being integrated, but the allocator guarantees it on any board -- vulkan_context.cpp:872-875 prefers DEVICE_LOCAL|HOST_VISIBLE|HOST_COHERENT and falls back to plain HOST_VISIBLE|HOST_COHERENT, refusing to initialize only if that fails too. A non-ReBAR Arc card is therefore still host-addressable through the second type; what it loses is device locality, so the keep-quant fall-through and the portable reference tier run slower there, not unsafe. + +(2) FALSE REGISTRY LINE. .agents/environment.md asserts 'No Intel GPU exists on any box here, so BACKEND-XPU end-to-end work is HW-BLOCKED'. garlic-clove is an Intel Arc Pro B60 (8086:e211, ASUS subsys 1849:6023) on the xe driver. + +(3) VK-I's STAGING-PATH DELIVERABLE IS ANSWERED, NOT BUILDABLE. vulkan-full-support.md scopes VK-I as 'The staging path for non-host-visible memory, and the gate re-run where Vulkan actually matters', and records the 2026-08-06 decision 'GB10 first, acquire later'. The acquire has happened, but the first half needs no code: ReBAR maps the B60's 20.91 GiB as host-visible, so the non-host-visible condition VK-I was written to handle does not arise on this card. Recording that as a measurement retires the risk; writing a staging path against it would be code for a condition this hardware does not have. + +Also stale, and the reason this was not caught: the BACKEND-VULKAN matrix cell still describes the 2026-07-22 skeleton (8 native ops, 'no model runs on Vulkan'). The tree now registers 35 Vulkan ops including the full GDN/SSM set, GGUF keep-quant/TQ1_0, EXL3 and MoE, and vulkan_ops.cpp:2109 itself names this board. + +## Resolution + +- diff --git a/.agents/issues/BACKEND-VULKAN/ISSUE-LOCAL-01M3JGXN1QTRE1D7WX3JMV0KE2.md b/.agents/issues/BACKEND-VULKAN/ISSUE-LOCAL-01M3JGXN1QTRE1D7WX3JMV0KE2.md new file mode 100644 index 0000000000..38940d5de8 --- /dev/null +++ b/.agents/issues/BACKEND-VULKAN/ISSUE-LOCAL-01M3JGXN1QTRE1D7WX3JMV0KE2.md @@ -0,0 +1,25 @@ +ID: ISSUE-LOCAL-01M3JGXN1QTRE1D7WX3JMV0KE2 +Title: There is no runtime device selector: VLLM_CPP_DEVICE is read nowhere, so every non-CUDA gate selects its device by accident +Row: BACKEND-VULKAN +State: OPEN +Kind: bug +GitHub: - +Mirror: PENDING +Availability: FULL +Created: 2026-09-27 +Updated: 2026-09-27 +Closed: - + +## Problem + +The engine picks its device in src/vllm/entrypoints/model_loader.cpp via CurrentPlatform(), which walks {kCUDA, kXPU, kVULKAN, kMETAL, kCPU} and takes the FIRST backend that probed a device. Nothing overrides that. The env var the campaign recorded, VLLM_CPP_DEVICE, is read NOWHERE: grep over the whole tree returns only record files, never src/, include/, tests/ or examples/. + +So `VLLM_CPP_DEVICE=vulkan test_opt_paged_engine` selected Vulkan only because those builds had no CUDA compiler. benchmark-record.md already MEASURED both ways on one source: with /usr/local/cuda/bin on PATH at configure time the identical command reports 'the engine selected device type 1' (kCUDA) and passes 6/6; without it, 'device type 3' (kVULKAN), 6/6, 0 declines. + +Why this is filed now rather than read and skipped: the estate gained a SECOND accelerator host (garlic-clove, Intel Arc Pro B60, no CUDA toolchain at all). A gate that selects its device by build accident is exactly the failure that a new box invites -- on a host with no CUDA the placebo is indistinguishable from a working selector, so the first Vulkan-only venue cannot tell the difference. The test already prints the truth ('the engine selected device type N') and load-direct-upload.md already uses the correct idiom, so the defect is that no selector EXISTS, not that the evidence is unobtainable. + +Scope: add a real runtime device override read at the SelectQueue seam, make it fail LOUDLY on an unknown value rather than silently falling through, and pin the selection in the test's own assertions. Then replace the remaining VLLM_CPP_DEVICE command citations in the record with the BACKEND PROOF form. Not attempted in the records-only VK-I change, because it is a code change to the device seam and needs its own row, spec and gate. + +## Resolution + +- diff --git a/.agents/specs/vulkan-full-support.md b/.agents/specs/vulkan-full-support.md index 9f2c687206..1cfffdf677 100644 --- a/.agents/specs/vulkan-full-support.md +++ b/.agents/specs/vulkan-full-support.md @@ -43,7 +43,11 @@ a win over *the Vulkan maturity floor on the box we own*, and the record must sa exactly that rather than "we beat llama.cpp". The claim that would matter to a user — Vulkan winning where CUDA/ROCm/Metal do not exist — is `VK-I`, and it is hardware-blocked until an RDNA or Arc board is acquired (user decision 2026-08-06: -GB10 first, acquire later). +GB10 first, acquire later). **The Arc half of that acquisition HAPPENED: see +§4 and §6.2 for the board (`garlic-clove`, Intel Arc Pro B60) and for the +measurement that answers `VK-I`'s staging-path half without code. The RDNA arm +is still unacquired, and the gate re-run — the half that actually matters — +is still owed.** --- @@ -248,7 +252,8 @@ our own CUDA paged kernel**, recorded as a partial-from-scratch entry in |---|---|---| | **GB10 on `dgx.casa`** | YES — `NVIDIA GB10`, `INTEGRATED_GPU`, Vulkan API 1.4.312, vendor `0x10de`, 249 device extensions incl. **`VK_KHR_cooperative_matrix` v2** and **`VK_NV_cooperative_matrix2`**, `VK_KHR_shader_float16_int8`, `VK_KHR_{8,16}bit_storage`, `VK_KHR_shader_integer_dot_product`, `VK_KHR_timeline_semaphore`, `VK_EXT_memory_budget`, `VK_KHR_buffer_device_address`. One 89.72 GiB `DEVICE_LOCAL` heap with a `DEVICE_LOCAL|HOST_VISIBLE` type — unified | **PRIMARY. Correctness oracle box AND the optimization target** (user decision 2026-08-06). Both llama.cpp coopmat tiers are reachable | | **`llvmpipe` (dev box)** | YES — Vulkan 1.4.318, CPU, `mesa-vulkan-drivers` | GPU-free CI correctness. **Never a speed venue** | -| **AMD RDNA / Intel Arc** | **NO — none on any box** | `VK-I`. Deferred by user decision; the only venue where a Vulkan win means something to a user, and the only thing that exercises the missing staging path | +| **Intel Arc Pro B60 on `garlic-clove`** | **YES — acquired, measured 2026-09-27.** `Intel(R) Arc(tm) Pro B60 Graphics (BMG G21)`, `8086:e211`, ASUS subsys `1849:6023`, `xe` driver, `PHYSICAL_DEVICE_TYPE_DISCRETE_GPU`, device API **1.4.354** (conformance 1.4.0.0), Mesa 26.2.3 / LLVM 21.1.8, `VK_KHR_cooperative_matrix` = true. Heaps 20.91 GiB `DEVICE_LOCAL` + 23.44 GiB host. **NOT an `rc` fleet device** — file mutex, not a lease | `VK-I`, PARTIALLY ANSWERED — see §6.2. The "acquire later" half of the 2026-08-06 decision is DONE | +| ~~**AMD RDNA**~~ | **NO — still none on any box** | `VK-I`'s RDNA arm. A discrete AMD board would need ReBAR for the same reason, and is still unacquired | **Premise update — the 2026-07-22 toolchain constraint is STALE.** That spec determined the committed-SPIR-V route partly because *"neither box grants sudo"* @@ -327,6 +332,71 @@ the umbrella, not a substitute for them. | **VK-H** | **Attention variants + samplers** (16 ops) | B (samplers), G (attn variants) | **83/83 — closes the op surface** | | **VK-I** | **AMD/RDNA (or Arc) bring-up** | hardware acquisition | The staging path for non-host-visible memory, and the gate re-run where Vulkan actually matters | +### 6.2 `VK-I` PARTIALLY ANSWERED — ReBAR retires the staging path — 2026-09-27 + +`VK-I` was the one sub-project blocked purely on acquisition, and §4 recorded +the decision as **"GB10 first, acquire later"** (2026-08-06). The board is now +on the estate: **`garlic-clove`, an Intel Arc Pro B60.** Its deliverable splits, +and the split is not the one the row was written against. + +**THE STAGING-PATH HALF IS ANSWERED BY HARDWARE, NOT BY CODE, AND NO CODE +SHOULD BE WRITTEN FOR IT.** The deliverable was "the staging path for +**non-host-visible** memory". MEASURED on `garlic-clove` 2026-09-27 with +`vulkaninfo`: the device is `PHYSICAL_DEVICE_TYPE_DISCRETE_GPU` with two heaps +(20.91 GiB `DEVICE_LOCAL`, 23.44 GiB host), **but** `memoryTypes[3]` and +`memoryTypes[6]` expose `DEVICE_LOCAL | HOST_VISIBLE | HOST_COHERENT` +(propertyFlags `0x0007`) on the device-local heap. Resizable BAR maps the VRAM +into the host address space, so the condition `VK-I` was built to survive — +device-local memory the host may not dereference — **does not arise on this +card**. `VulkanContext`'s existing preference for a device-local host-visible +type (`vulkan_context.cpp:873`) already selects it, so +`VulkanBackend::DeviceMemoryIsHostAddressable()` is sound here without a +staging copy. + +**THIS RETIRES A RISK; IT DOES NOT DELETE A REQUIREMENT.** The property is a +property of **ReBAR**, not of Arc and not of this backend — and it is a +**performance** property, not a correctness one. `VulkanContext` +(`vulkan_context.cpp:872-875`) does not require `DEVICE_LOCAL` at all: it +PREFERS `DEVICE_LOCAL | HOST_VISIBLE | HOST_COHERENT`, and when that lookup +fails it FALLS BACK to `HOST_VISIBLE | HOST_COHERENT` without `DEVICE_LOCAL`, +initializing only if that fails too. `AllocBuffer` allocates and +`vkMapMemory`-maps whichever type was selected, and +`DeviceMemoryIsHostAddressable()` returns true unconditionally +(`vulkan_backend.cpp:130-135`) because every allocation is mapped. + +So a discrete card without ReBAR does **not** turn the GGUF keep-quant CPU +fall-through or the portable reference tier into corruption. Host visibility +is what correctness needs, the fallback supplies it, and the host vec_dot +kernel keeps reading memory it can address. What the card without ReBAR costs +is that the GPU reaches that memory across the bus rather than hitting VRAM +locally — **slowness, not unsafety**, and the staging path `VK-I` exists to +buy that back. So the standing requirement is the one already in the code: +keep *preferring* a device-local host-visible type and let the ordered +fallback carry correctness where none exists. `unified_memory_` records which +of the two happened, and it is the performance branch, not a safety one. That +is now written down at the `BACKEND-VULKAN-KEEPQUANT` block in +`src/vt/vulkan/vulkan_ops.cpp`, which until 2026-09-27 carried the false claim +**"The B60 is integrated"** — right conclusion, wrong mechanism, and dangerous +as a generalisation. An earlier revision of this paragraph went further and +called the no-ReBAR case *corruption*; that was wrong, and the correction is +the host-visibility/device-locality distinction above. + +**THE SECOND HALF IS THE REAL ONE AND IT IS STILL OWED.** "The gate re-run +where Vulkan actually matters" is now reachable and remains the entire value of +`VK-I`. §0's framing is why: on GB10 *"llama.cpp's own CUDA backend will beat +both of us there. Vulkan on an NVIDIA part is nobody's fastest path; it is the +portability path."* Every Vulkan speed number in this spec is therefore either +llvmpipe (a software rasteriser) or the wrong chip. `garlic-clove` is the first +venue where **Vulkan is the only accelerator path on the box**, so a Vulkan win +here is a win a user would actually feel. Concretely still owed, unchanged: +the 27B prefill/decode re-run and its reference-tier count (named in §6.0b as +"the only thing that can turn the structure above into a result"), `VK-C`'s +coopmat tactic selection on a non-NVIDIA part, and re-taking the +`BENCH-VK-LLAMA` verdict the records call the most fragile in the enumeration +(0.23% margin inside a 0.69% spread, #1003). **20.91 GiB of device-local memory +is the binding constraint** and 27B does not fit; the reachability question is +which model arm does, and the box holds no weights at all. + ### 6.0b The DECODE GEMV lever, measured to its floor — 2026-08-09 `row/BACKEND-VULKAN-GEMVROWS`. `vt_matmul_vec` is 89-90% of 27B decode GPU time @@ -517,9 +587,16 @@ gated at — and IDENTICAL to what the 128-wide module scores on the same inputs The residual STREAM is held to the bit-exact tier by `memcmp`, which is what proves `vt_round_through` is the memory round trip rather than an approximation of it. `test_vulkan_backend` 30/30 (2371 assertions) on GB10 and 30/30 (1828) on -llvmpipe; `test_opt_paged_engine` with `VLLM_CPP_DEVICE=vulkan` still 6/6 +llvmpipe; `test_opt_paged_engine` still 6/6 token-exact (96/96) with 0 declines on BOTH arms; `test_backend_cross_device` -11/11. +11/11. **The device is asserted from the test's own printed `BACKEND PROOF` line +(`the engine selected device type 3`), NOT from an env var — `VLLM_CPP_DEVICE` +is read NOWHERE in the tree.** This line used to read +``test_opt_paged_engine` with `VLLM_CPP_DEVICE=vulkan``, which is a placebo +that happened to select Vulkan only because those builds had no CUDA compiler; +see [benchmark-record.md](../benchmark-record.md) §0 of the +`BACKEND-VULKAN-LOADMEM` entry. The corrected invocation is +[`load-direct-upload.md`](load-direct-upload.md)'s. **AN HONEST LIMIT ON THAT e2e GATE.** `opt-125m` is a LayerNorm model — two `vt::LayerNorm` calls per layer and ZERO `vt::RmsNorm` — so the standing @@ -714,8 +791,10 @@ and it stays sequential inside the workgroup. NMSE vs the CPU oracle in the same binary: prefill out `1.47e-14`, prefill carried state `6.43e-15`, bf16 arm `0`; decode (indexed cache) out `1.64e-14` and cache `3.31e-15`, decode (compact state) out `1.75e-14`. `test_vulkan_backend` -25/25 cases, 1020/1020 assertions. `test_opt_paged_engine` with -`VLLM_CPP_DEVICE=vulkan` still 6/6 token-exact (96/96), 0 declines. +25/25 cases, 1020/1020 assertions. `test_opt_paged_engine` still 6/6 +token-exact (96/96), 0 declines, device asserted from the printed +`BACKEND PROOF` line — `VLLM_CPP_DEVICE` is read nowhere and does not select +it. **NOT measured: any speed number.** Local Vulkan is llvmpipe. The 27B prefill/decode re-run on GB10, and the reference-tier count that goes with it, diff --git a/src/vt/vulkan/vulkan_ops.cpp b/src/vt/vulkan/vulkan_ops.cpp index b33d1bec8a..b2e613e765 100644 --- a/src/vt/vulkan/vulkan_ops.cpp +++ b/src/vt/vulkan/vulkan_ops.cpp @@ -2108,12 +2108,12 @@ void AttnQkNormRopeGateKernel(Queue&, Tensor& q_out, Tensor& k_out, Tensor& gate // --------------------------------------------------------------------------- // BACKEND-VULKAN-KEEPQUANT — the GGUF keep-quant tier, by CPU FALL-THROUGH. // -// WHY A HOST KERNEL IS THE CORRECT VULKAN REGISTRATION (for now). The B60 is -// integrated: VulkanContext allocates every tensor from HOST_VISIBLE | -// HOST_COHERENT memory and VulkanBackend::DeviceMemoryIsHostAddressable() -// answers true unconditionally, so the host vec_dot kernels dereference the -// SAME bytes the rest of the graph dispatches shaders over. That is precisely -// the property the portable reference tier already relies on to run the CPU +// WHY A HOST KERNEL IS THE CORRECT VULKAN REGISTRATION (for now). VulkanContext +// allocates every tensor from HOST_VISIBLE | HOST_COHERENT memory and +// VulkanBackend::DeviceMemoryIsHostAddressable() answers true unconditionally, +// so the host vec_dot kernels dereference the SAME bytes the rest of the graph +// dispatches shaders over. That is precisely the property the portable +// reference tier already relies on to run the CPU // kMatmulBTQuant lazily at GetOp time (op_provider.cpp // MaybeInstallReferenceTier) — this registration makes that arrangement // EAGER, so it also flips vt::OpRegistered(kMatmulBTQuant, kVULKAN), which @@ -2131,6 +2131,35 @@ void AttnQkNormRopeGateKernel(Queue&, Tensor& q_out, Tensor& k_out, Tensor& gate // discharged here on the NATIVE path too because the native kernel here IS a // host kernel. // +// WHY `DeviceMemoryIsHostAddressable()` IS TRUE HERE, AND WHY IT IS NOT FREE. +// It is NOT because the board is integrated. MEASURED on `garlic-clove` +// (Intel Arc Pro B60, `8086:e211`, `xe` driver) 2026-09-27 with `vulkaninfo`: +// the device reports `PHYSICAL_DEVICE_TYPE_DISCRETE_GPU` with two heaps — +// 20.91 GiB `DEVICE_LOCAL` and 23.44 GiB host — so "integrated" was the wrong +// reason and this comment used to say so. The property holds anyway, and by a +// different mechanism: Resizable BAR maps VRAM into the host address space, so +// `memoryTypes[3]` and `memoryTypes[6]` expose `DEVICE_LOCAL | HOST_VISIBLE | +// HOST_COHERENT` (propertyFlags 0x0007), which is the type `vulkan_context.cpp:873` +// prefers. So the allocation IS ordinary host memory the GPU also reads. +// +// THIS IS A PROPERTY OF ReBAR, NOT OF ARCS, AND IT IS A PERFORMANCE +// PROPERTY, NOT A CORRECTNESS ONE. `VulkanContext` (vulkan_context.cpp:872-875) +// PREFERS `DEVICE_LOCAL | HOST_VISIBLE | HOST_COHERENT` and, when that lookup +// fails, FALLS BACK to `HOST_VISIBLE | HOST_COHERENT` without `DEVICE_LOCAL`, +// refusing to initialize only if that fails too. `AllocBuffer` then allocates +// and `vkMapMemory`-maps whichever type was selected, and +// `DeviceMemoryIsHostAddressable()` reports true unconditionally +// (vulkan_backend.cpp:130-135). So a discrete card WITHOUT ReBAR does NOT +// make this registration memory corruption: the fallback allocation is still +// host-visible, and the host vec_dot kernel still reads memory it can +// address. What it costs is that the GPU reaches that memory across the bus +// instead of hitting VRAM locally, which is exactly the staging path `VK-I` +// was scoped to build (vulkan-full-support.md §4/§6). The requirement is +// therefore the one already in the code: keep PREFERRING a device-local +// host-visible type, and let the ordered fallback carry correctness on a +// board that has none. `unified_memory_` is the flag that records which of +// the two happened, and it is the performance branch, not a safety one. +// // Mirrors cuda_quant_dot.cu:1835 (CUDA falls through to GetOp(...,kCPU) for // dtypes its GPU kernel lacks; on unified memory that fall-through is free). // Local activations-reader (mirrors cpu_quant_gemm.cpp LoadActF32): the keep-quant