From ebd798b440cd42a996a8ce7f5ad24593f6bd1cbd Mon Sep 17 00:00:00 2001 From: ghazni Date: Thu, 27 Aug 2026 08:21:30 +0000 Subject: [PATCH 1/2] =?UTF-8?q?spec(KV-FP8):=20W6=20=E2=80=94=20the=20ROCm?= =?UTF-8?q?=20fp8-e4m3=20KV=20cache=20arm?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Append the W6 section to the fp8-kv-cache spec: scope, upstream chain, port map, gates, and the wave-table row. The ROCm arm mirrors the W2 CUDA arm elementwise, with the CPU kernels (W1) as the oracle. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:glm-5-2 [OMP] --- .agents/issue-index.md | 1 + .agents/specs/fp8-kv-cache.md | 163 +++++++++++++++++++++++++++++++++- 2 files changed, 162 insertions(+), 2 deletions(-) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 318a4cdb7b..f719bb29c5 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -691,3 +691,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1839](https://github.com/mudler/vllm.cpp/issues/1839) | `MUSIC3-DEPTH-DEVICE` | **The engine's call to `Music3SelectDepthArm` (`minimax_music3_speech.cpp:638`) is reachable but not gated, and [#1131](https://github.com/mudler/vllm.cpp/issues/1131) no longer covers it: #1131 named both device-arm twins and row `MUSIC3-DIT-ARM-REACH` closes it with only the DiT half.** Deleting the two-line call leaves `test_minimax_music3_ar` 37/37 · 640/640 and `test_minimax_music3_speech` 9/9 · 223/223 green ([`minimax-music3.md`](specs/minimax-music3.md) §19.5 carries the mutation and the binary hashes). Two things stop an existing gate from seeing it, and the SECOND is the one that matters: `--speech-device 1` is refused by name on a CPU-only build before a queue exists, AND §19.6's "device path TAKEN" leg rides `test_minimax_music3_ar`, whose observable `Music3DepthDeviceForwardCount()` is a counter §19.5 itself records as unreachable from production — its only readers are the tests written for it (`test_minimax_music3_ar.cpp:1325,1351,1583,1589,1753,1758`). The instrument that WOULD answer the call site is `ar.depth_staging`, emitted at `minimax_music3_llm.cpp:582` and read by nothing. NOT FIXED IN FLOW and the reason is precise: closing it needs the shipped engine on a real accelerator against the 28.5 GB checkpoint inside an `rc` lease, which is a second GPU leg and a second gate file, not a repair to the row in flight. It is closable by exactly `MUSIC3-DIT-ARM-REACH`'s method — a `gpu;checkpoint;music3`-labelled parity gate entering through `include/vllm.h` with `device = 1`, exiting 77 without its preconditions, asserting `ar.depth_staging` `calls == 1` with the host bucket absent — and that row's `thor:gpu0` job `f63f60e8-957a-4062-92f8-54e5bbb49d92` already FIRED `ar.depth_staging` once without asserting it, so the instrument is known live on the real path. Owed under `## Owed` in [`minimax-music3.md`](specs/minimax-music3.md) §19.7 | bug | | [#1849](https://github.com/mudler/vllm.cpp/issues/1849) | `SPEC-DFLASH2` | **The DFlash2 draft step costs a flat ~23 ms at EVERY K, and the two levers #1849 names resolve differently once read from the records.** Lever A (quantize the shared head) is ALREADY LANDED for the measured subject: `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` stores `lm_head` as W4A16_NVFP4 g16 (header-verified 2026-08-21, quantization-matrix `QUANT-QWEN38-27B-NVFP4-ARM`), upstream computes with it packed through `lm_head.quant_method.apply`, and both our reads have been packed since #1628 — so the head traffic is ~2×0.72 GB not 2×2.54, the draft-phase floor re-derives to ~9 ms, and the unattributed residual GROWS to ~13-14 ms. Lever B (launch/sync trim) is counted in code at one replay + ~10 launches + ~76 B up / 64 B down + one sync — well under 0.5 ms, so the residual sits INSIDE kernels and needs on-box attribution. W9 lands `VT_SPEC_TRACE=2` (the `[spec-phase-dev]` pre/fwd/select/walk split) as the instrument, and borrow-first loading for the draft's shared bf16 embed+head (~5.1 GB host on the bf16 arm, ~2.5 GB on the r0b0tlab arm; memory only, no step-time claim). The bf16-target arm's 2×2.54 GB head reads are upstream's own serving dtype and stand as a recorded ceiling. Wave spec [dflash2-draft-fixed-cost.md](specs/dflash2-draft-fixed-cost.md); the K-ladder rerun, the `ncu`/`nsys` attribution and any step delta are owed there, operator-run | perf | | [#1844](https://github.com/mudler/vllm.cpp/issues/1844) | `ENG-MM-INPUT-PIPELINE` | **`scripts/mm/tower_skip_rss.sh` killed every measured leg mid-load, because `run_arm`'s `/health` poll was answered by the PREVIOUS leg's server.** First real run (`thor:gpu0`, worker `rc-worker-kk96r`, `d60692c8`): checkpoint staged and verified (29 files, 8887294190 B), both binaries built sha256-identical (`78d582e4...`), live target query green on both build dirs -- and then **five 0-byte `.time` files** and `VOID` on both pairs. `warmup` reached `listening on http://0.0.0.0:18607` (1286 B log); all four measured legs stopped at `loading model from ...` (363/345 B) inside one minute. TWO defects, one shape. (a) `$PORT` is fixed and the poll starts immediately, so a stale listener answers it and the leg is ready before it has read a tensor. (b) `kill "$pid"` signals `/usr/bin/time`, which installs no handler: the timer dies before writing its `-o` file and the server is reparented to init and KEEPS THE PORT -- which is what was answering. Measured: `/usr/bin/time -v -o f sleep 100 & kill $!` leaves `f` at 0 B and `sleep` alive with ppid 1; signalling the CHILD leaves `f` at 752 B with a `Maximum resident set size` line. **Nothing could catch it**: `test_tower_skip_rss_report.py` was 60/60 green over finished files, and `run_arm`, the poll and the teardown ran only under a lease -- the residual [#1819](https://github.com/mudler/vllm.cpp/issues/1819) recorded verbatim, now observed. FIXED IN FLOW: a leg refuses to start into an occupied port; readiness requires the banner in the leg's OWN log (its stdout, which no other server can write) before `/health`; SIGTERM goes to the SERVER so the timer survives to write; the port must stop accepting before the next leg; each of the three waits is bounded and each bound REFUSES; and a leg whose `.time` carries no `Maximum resident set size` line fails AT that leg rather than as VOID four legs later. GATED: `TOWER_SKIP_RSS_SOURCE_ONLY=1` sources the harness for its functions alone and `tests/scripts/test_tower_skip_rss_arm.py` drives `run_arm` against a fake server on a scratch port -- stale listener, absent banner, death during load, the happy path's non-empty `.time`, five legs in the declared order, and both halves restored as mutations. RED-first: 11 of 14 cases fail against the pre-fix `run_arm`, the stale-listener case reporting `LEG default OK` with a 0-byte `.time` | bug | +| [#2065](https://github.com/mudler/vllm.cpp/issues/2065) | `KV-FP8` | The ROCm fp8-e4m3 KV cache arm: the store kernel (`vt::ReshapeAndCacheFp8` for `DeviceType::kROCM`), the fp8 dequant on the ROCm paged-attention read, the `OpId::kReshapeAndCacheFp8` registration for `kROCM`, and the widening of the `src/vt/ops.cpp` fp8 read refusal to admit `kROCM`. Mirrors the CUDA W2 arm element-for-element; the CPU kernels remain the oracle. Spec [`fp8-kv-cache.md`](specs/fp8-kv-cache.md) `## W6` | feature | diff --git a/.agents/specs/fp8-kv-cache.md b/.agents/specs/fp8-kv-cache.md index 90b8f5dc66..9c7042d328 100644 --- a/.agents/specs/fp8-kv-cache.md +++ b/.agents/specs/fp8-kv-cache.md @@ -606,6 +606,164 @@ for one of those without opening anything. G10's ordering case points at a directory that EXISTS and declares fp8, so an inverted order announces the declaration first and that line is the evidence. +## W6 — the ROCm arm (#2065) + +Issue: [#2065](https://github.com/mudler/vllm.cpp/issues/2065). The CPU +kernels (W1) are the ORACLE: every W6 gate compares ROCm to the landed CPU +kernels, never to a fresh reference. The CUDA arm (W2) is the direct +template — the ROCm arm is elementwise-identical to it, and the CUDA arm is +itself elementwise-identical to the CPU reference. + +### Scope + +- **In:** the ROCm fp8-e4m3 K/V store kernel + (`ReshapeAndCacheFp8KernelRocm`), the fp8 dequant on the ROCm + paged-attention read (`LoadKv` in `rocm_paged_attn.hip`), the + `OpId::kReshapeAndCacheFp8` registration for `DeviceType::kROCM`, and + the widening of the `src/vt/ops.cpp` fp8 read refusal to admit `kROCM`. +- **Out:** fp8_e5m2 compute, per-attention-head scales, the Metal arm, + fast-path (tensor-core/rocWMMA) fp8 attention kernels, and the + memory-halving e2e measurement on a ROCm gate model. + +### Upstream chain + +Same as W2. The fp8 KV path is vLLM's own csrc, not a dependency: + +- **Store.** `reshape_and_cache_flash_kernel` + (`csrc/libtorch_stable/cache_kernels.cu:314-401`) + `CopyWithScaleOp` + (`:241-252`). The fp8 branch is `dst = fp8::scaled_convert(src, scale)`, restricted to the + `is_contiguous_heads && kv_scale_stride == 0` arm (`:352-366`) — the + only arm the op's wrapper admits. +- **Scale convention.** `FP8 = Quantize(HP / scale)`; + `Dequant(FP8) * scale = HP` (`quant_utils.cuh:296-300`). +- **Read.** `scaled_vec_conversion` + (`quant_utils.cuh:419-429`) = `fp8_to_float(byte) * scale`. +- **Storage.** `cache_t = uint8_t` + `Fp8KVCacheDataType` template param. + Mirrored as `DType::kI8` + `Fp8KVCacheDataType` enum + (`include/vt/fp8_kv.h`). + +### Port map + +W6 (this change; ROCm build, `-Werror`): + +- `src/vt/rocm/rocm_dense_basic.hip` — `ReshapeAndCacheFp8K` kernel + + `ReshapeAndCacheFp8KernelRocm` host launcher (port of + `cuda_cache.cu:155-226`). The converter reuses the `F8E4M3ToF32` + device function already in `rocm_fp8_channel_gemv.hip:22-31` for the + READ, and `vt::F32ToF8E4M3` (`include/vt/fp8_kv.h`) for the STORE. + ROCm has no `__nv_cvt_float_to_fp8` intrinsic; the store uses the + software codec `vt::StoreKvFp8E4M3` (`include/vt/fp8_kv.h:87-89`), + which is bit-identical to the CPU codec and to the CUDA intrinsic + (measured at zero tolerance, spec W2). Source dtypes f32/f16/bf16, + the same set the CPU `LoadSrcF32` and the CUDA `Fp8SrcToF32` serve. +- `src/vt/rocm/rocm_paged_attn.hip` — `LoadKv` device function (port of + `cuda_paged_attn.cu:175-185`): inert on the f32/bf16 arms (forward to + `Ld`), and on `uint8_t` it is `F8E4M3ToF32(byte) * scale`. The + `PagedAttnOnline` kernel gains `k_scale`/`v_scale` parameters; the + host launcher keys on `args.kv_cache_dtype` to route the fp8 read. + Only the two correctness-grade kernels serve fp8: the online softmax + decode and the tiled prefill — the same line W2 draws. +- `src/vt/rocm/rocm_ops.hip` — + `RegisterOp(OpId::kReshapeAndCacheFp8, DeviceType::kROCM, ...)`. +- `src/vt/ops.cpp:3835` — widen the fp8 read refusal from + `kCPU || kCUDA` to `kCPU || kCUDA || kROCM`, and update the message + to name ROCm as implemented. +- `tests/vt/test_rocm_fp8_kv_cache.cpp` (NEW) + its `tests/CMakeLists.txt` + line. + +### The store + +`ReshapeAndCacheFp8KernelRocm` is an ELEMENTWISE-IDENTICAL port of the +CUDA `ReshapeAndCacheFp8Kernel` (`cuda_cache.cu:155-176`), which is itself +an elementwise-identical port of the CPU +`ReshapeAndCacheFp8Kernel` (`cpu_cache.cpp`). The kernel is a scalar +strided loop over the same elements in the same order as the CUDA and CPU +arms: one block per token, threads stride over the page +(`num_kv_heads*head_size`). + +The converter is the SOFTWARE codec `vt::StoreKvFp8E4M3` +(`include/vt/fp8_kv.h:87-89`), which is `vt::F32ToF8E4M3(hp / scale)`. +ROCm HIP has no `__nv_cvt_float_to_fp8` intrinsic, so the store uses the +same software round-to-nearest-even path the CPU kernel uses. This is +bit-identical to the CPU codec by construction (same function), and the +CUDA intrinsic's equality to the CPU codec is already measured at zero +tolerance (spec W2, `vt-fp8-quant-arch-gate.md` G2). The store is a true +DIVIDE (`hp / scale`), not a reciprocal multiply. + +Source dtypes f32/bf16/f16, widened to f32 before the divide through +`Ld` (`rocm_dense_basic.hip:139-141`), the same helper the float path +uses. `__hip_bfloat16` and `__half` both carry `operator float()` so the +existing `Ld` overloads serve without new conversion functions. + +### The read + +`LoadKv(ptr, i, scale)` joins `Ld` in `rocm_paged_attn.hip`: inert on the +f32/bf16 arms (forward to `Ld` unchanged, so every existing caller reads +the same bytes in the same order), and on `uint8_t` it is +`F8E4M3ToF32(byte) * scale` — the same arithmetic as `vt::LoadKvFp8E4M3` +(`include/vt/fp8_kv.h:93`) and the CUDA `Fp8E4M3ToF32Dev` +(`cuda_paged_attn.cu:164`), so ROCm==CPU on the read is a property of the +source rather than of a measurement. + +The `F8E4M3ToF32` device function already exists in +`rocm_fp8_channel_gemv.hip:22-31` and is bit-identical to +`vt::F8E4M3ToF32` (`include/vt/fp8_kv.h:40-49`). It is reused, not +re-stated — the codec reuse rule (spec `## Risks/decisions`). + +`PagedAttnOnline` gains `k_scale`/`v_scale` parameters; the host launcher +keys on `args.kv_cache_dtype` (never on the storage dtype, which is a bare +`kI8` byte) and routes to the fp8-aware launch. Only the online softmax +decode and tiled prefill kernels serve fp8, the same scope W2 argues: +the rocWMMA prefill kernels stage `__hip_bfloat16` fragments, and a +tensor-core fp8 read is a PERFORMANCE brick, not this one. + +### The refusal + +The fp8 READ rides additive fields on `PagedAttentionArgs` of an op that +`kROCM` already registers for the float path (`rocm_ops.hip:159`). W6 +removes `kROCM` from the explicit refusal list in `src/vt/ops.cpp:3835` +and widens it to `kCPU || kCUDA || kROCM`. The message is updated to name +ROCm as implemented. The fp8 STORE is a separate `OpId` that only CPU, +CUDA, and now ROCm register, so an unimplemented backend still refuses by +name inside `GetOp`. + +### Same-arithmetic caveat + +The same NaN-payload caveat as W2 applies: on `0x7F`/`0xFF` the CPU +returns `std::numeric_limits::quiet_NaN()` (`0x7FC00000`) and the +ROCm device returns the HIP runtime's NaN representation. Both are quiet, +both propagating, potentially different payload. No gate can see this +because a NaN compares unequal to itself, and `__NV_SATFINITE` (or the +software codec's saturating clamp) never writes a NaN code from a finite +`hp / scale`. + +### Gates + +`tests/vt/test_rocm_fp8_kv_cache.cpp` — mirrors +`tests/vt/test_cuda_fp8_kv_cache.cpp`: + +| Case | What it gates | Build | +|---|---|---| +| G1 | the fp8 store/read resolves through the provider table on a non-CPU device (no `later brick` guard) | non-ROCm build | +| G1b | the fp8 read is refused on a backend with no fp8 dequant (Metal only, now that ROCm is implemented) | every build | +| G2 | the ROCm providers are registered for the fp8 store and paged read | ROCm build | +| G3 | STORE parity — ROCm store writes the same bytes as the CPU store, zero tolerance, f32/bf16/f16 sources, padded slot | ROCm device | +| G4 | READ parity — paged attention over identical fp8 cache bytes, ROCm vs CPU, decode + prefill, f32 query/output | ROCm device | +| G4b | READ parity — bf16 query/output (the instantiation a served model takes) | ROCm device | +| G5 | fp8_e5m2 stays refused by the ROCm kernel, reached through the registered provider | ROCm device | + +G3/G4/G4b/G5 SKIP CLEANLY when no ROCm backend is present, with a MESSAGE +naming what did not run. G1/G1b run on the CPU leg. + +### Dependencies + +The ROCm build (`VLLM_CPP_HIP=ON`) with hipcc. The local host has ROCm +7.2.4 and gfx1100 (RX 7900 XTX). No upstream checkout (`VLLM_SOURCE`) or +oracle (`VLLM_ORACLE`) is configured; the CPU kernels are the oracle, and +upstream anchors are cited from the W2 spec section (already verified at +pin `555967922`). + ## Owed - **The W2 device gates are UNEXECUTED** (#1593). `tests/vt/test_cuda_fp8_kv_cache.cpp` @@ -812,8 +970,9 @@ declaration first and that line is the evidence. builds is gated by G1/G10 rather than exercised by the benchmark it was built for. A calibrated ModelOpt checkpoint that declares the algorithm INLINE is what would exercise it end to end, and this row has none. -- **Metal and ROCm have no fp8 KV arm.** Both refuse by name (see above). Neither - has a row yet; they belong with W5's per-head/e5m2 work or a backend row. +- **Metal has no fp8 KV arm.** It refuses by name (see `## W6`). ROCm + landed in W6 (#2065); Metal belongs with W5's per-head/e5m2 work or a + backend row. - fp8_e5m2 and per-attention-head scales stay refused on both backends (W5). ## Risks/decisions From 872b53eeee049853d86a9f2cd66dfa9bd78f3e85 Mon Sep 17 00:00:00 2001 From: ghazni Date: Thu, 27 Aug 2026 10:13:18 +0000 Subject: [PATCH 2/2] =?UTF-8?q?feat(KV-FP8):=20W6=20=E2=80=94=20the=20ROCm?= =?UTF-8?q?=20fp8-e4m3=20KV=20cache=20store=20and=20read?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The ROCm backend gains the fp8-e4m3 KV cache store kernel and the fp8 dequant on the paged-attention read, closing the last non-Metal gap in the fp8 KV cache surface. The store kernel mirrors the W1 CPU codec and the W2 CUDA arm elementwise; the read dequant multiplies back by the per-tensor k_scale/v_scale. ROCm HIP has no __nv_cvt_float_to_fp8 intrinsic, so the store uses a software codec (F32ToF8E4M3Dev) with frexpf/nearbyintf/ldexpf arithmetic that mirrors the CPU vt::F32ToF8E4M3. The codec is bit-identical and portable across all ROCm targets, unlike hardware fp8 conversion intrinsics that exist only on CDNA2+ (gfx940/941/942). OpId::kReshapeAndCacheFp8 is registered for DeviceType::kROCM in rocm_ops.hip. The fp8 read refusal in ops.cpp widens from kCPU||kCUDA to kCPU||kCUDA||kROCM. Metal remains refused by name. Gated on gfx1100 (RX 7900 XTX): test_rocm_fp8_kv_cache 7/7 cases, 28/28 assertions. G3 store is byte-identical to the W1 CPU oracle; G4 f32 read NMSE < 1e-6; G4b bf16 read NMSE < 1e-4; G5 e5m2 refused with the named message. RED-first proven: dropping *scale in the read dequant fails G4/G4b at NMSE ~34000x; dropping /scale in the store fails G3 at the byte level. Existing suites unaffected: test_rocm_backend 9/9, test_ops_paged_attn 14/14, test_ops_fp8_kv_cache 8/8. Closes #2065 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:glm-5-2 [OMP] --- .agents/backend-matrix.md | 2 +- .agents/engine-matrix.md | 2 +- .agents/feature-matrix.md | 2 +- .agents/quantization-matrix.md | 2 +- .agents/specs/fp8-kv-cache.md | 59 ++- include/vt/ops.h | 23 +- src/vt/ops.cpp | 26 +- src/vt/rocm/rocm_dense_basic.hip | 137 +++++ src/vt/rocm/rocm_ops.hip | 8 + src/vt/rocm/rocm_paged_attn.hip | 81 ++- tests/CMakeLists.txt | 1 + tests/vt/test_rocm_fp8_kv_cache.cpp | 747 ++++++++++++++++++++++++++++ 12 files changed, 1044 insertions(+), 46 deletions(-) create mode 100644 tests/vt/test_rocm_fp8_kv_cache.cpp diff --git a/.agents/backend-matrix.md b/.agents/backend-matrix.md index ad0b65d674..d59a12f0b2 100644 --- a/.agents/backend-matrix.md +++ b/.agents/backend-matrix.md @@ -236,7 +236,7 @@ this repository. State remains `ACTIVE`; no lifecycle transition is claimed. | `BACKEND-PLATFORM` | Platform capability/memory-model seam (`is_cuda`/`is_cpu`, `is_unified_memory`, `has_device_capability`, `supported_dtypes`, `residency_policy`, `supports_graph_capture`; `get_attn_backend_priority` REALIZED by `BACKEND-ATTN-REGISTRY` item 4; **+ the S3 capability predicates `supports_fp8`/`cutlass_fp4_supported`/`opaque_attention_op`/`is_integrated_gpu`/`support_static_graph_mode`/`is_device_capability_family`, mirroring `interface.py:914,933,977,1058,441-476` + `nvfp4_utils.py:56`, added by `CLAIM-BACKEND-SEAM-S3-1` and consumed by the fp4/fp8/graph gates in `qwen3_5.cpp`**); self-registered per `DeviceType` via `CurrentPlatform()` | `vllm/platforms/interface.py:134-229` (`class Platform`), `:409-439` (get/has_device_capability), `:181-187` (supported_dtypes) @ pin `e24d1b24` | interface [interface.h:56](../include/vllm/platforms/interface.h#L56), registry/CurrentPlatform [platform.cpp:57](../src/vllm/platforms/platform.cpp#L57), CPU [cpu.cpp:11](../src/vllm/platforms/cpu.cpp#L11), CUDA [cuda.cpp:14](../src/vllm/platforms/cuda.cpp#L14); composes `vt::Backend` [backend.h:22](../include/vt/backend.h#L22) | [platform tests](../tests/vllm/platforms/test_platform.cpp#L28) (registration + CPU capabilities + has_device_capability), clean CPU `-Werror` + full CPU CTest; 7 memory-model/residency sites migrated (`runner.cpp`, `model_registry.cpp`, `qwen3_5.cpp`) — **per-tensor sites key on `GetPlatform(.device.type).is_cuda()` (the OBJECT's device), NOT the process-global `CurrentPlatform()`, which is accelerator-first and would mis-route a CPU queue/tensor on a GPU box (regression fixed 2026-07-18; red DGX CPU tests `test_platform`/`test_qwen27_dense_forward`); `CurrentPlatform()` is reserved for process-level "which accelerator" questions only**. DGX behavior-preserving model gates 27B 235/235 + 35B 315/315 CONFIRMED @ (CUDA -Werror-clean after a GCC13 dangling-pointer FP fix). **Item-2 residency_policy() CONSUMPTION LANDED (2026-07-19, `CLAIM-BACKEND-PLATFORM-2`):** the host-free / load-stream / DevicePool-cap decisions in `qwen3_5.cpp` now READ `GetPlatform(.device.type).residency_policy()` (per-device) instead of an inline `device.type`/env gate — host-free `qwen3_5.cpp:BuildMoeMarlinResident` via `ShouldReleaseHostWeights`, load-stream `Qwen3_5Model::PrepareMarlinResident` via `ShouldInterleaveLoadStream` (both `interface.h`), DevicePool soft cap via `residency_policy().device_pool_cap_bytes` (`DBuf`/`DevicePool`). `CudaPlatform::residency_policy().release_host_weights_after_upload` flipped **false→true** (now CONSUMED ⇒ reproduces today's GB10 host-free-after-Marlin-build EXACTLY); `MarlinMoeEnabled()` stays the orthogonal KERNEL-PATH gate; `VT_MOE_HOST_FREE`/`VT_MOE_LOADSTREAM` env stay as overrides. Consumption unit-tested `test_platform.cpp` (7 cases/43 assertions). A new (discrete) GPU sets `residency_policy()` values ⇒ ZERO model edit. **DGX gate PASSED @ `62fc0e0` (production flags CUTLASS sm120a + FA2 sm_121a + Triton AOT, one flock): clean CUDA `-Werror` 0 warn, 27B 235/235 + 35B 315/315 token-exact, 35B VmHWM ≈ 4.0 GiB (load-stream win preserved), memcheck 0 errors.** | [Platform seam plan](specs/extensibility-platform-seam-2026-07-18.md) | `ANCHOR-BACKFILL` — item-1 extraction DONE + DGX-CONFIRMED; item-2 residency consumption LANDED + DGX-PROVEN (`CLAIM-BACKEND-PLATFORM-2`); kernel-shape dispatch branches deliberately left for items 4/5 | CLAIM-BACKEND-PLATFORM-1 / CLAIM-BACKEND-PLATFORM-2 | | `BACKEND-CPU` | CPU correctness and production path | `platforms/cpu.py:42-125`, CPU ops rooted at `csrc/cpu/torch_bindings.cpp:123-139`; llama.cpp `ggml-cpu.c:471-610,3024-3390` owns the added native pool reference | [cpu_backend.cpp:11](../src/vt/cpu/cpu_backend.cpp#L11), [threadpool.cpp:78](../src/vt/cpu/cpu_threadpool.cpp#L78), [chunked ops](../src/vt/cpu/cpu_ops.cpp#L88) | [backend tests](../tests/vt/test_backend.cpp#L10), [op parity](../tests/parity/test_op_parity.cpp#L34), [threadpool/upstream-test port](../tests/vt/test_cpu_threadpool.cpp#L63); 1/3/20 full CPU suites + TSAN pass. The binding 20-core Arm/i8mm Qwen3.5-2B Q8 single-stream gate is at or ahead of same-file llama.cpp on prefill/decode/RSS; server concurrency remains open. **Raspberry Pi 5 / Cortex-A76 lane GATING 2026-08-06:** R0-R3 pin the exact Q8_K_XL vehicle, local buildx/QEMU build and physical-Pi correctness/PMU baseline. **R4-R5 GREEN:** [KERNEL-CPU-A76-Q8-DOT](kernel-matrix.md) adds exact-order compiler SDOT and a scheduled AAPCS64 leaf; Pi operation and 64-token model outputs are exact. Assembly beats compiler SDOT 3.66-5.08% on M1/T1 and M128/T1/T4, lowers Qwen TTFT 1.55% and E2E 0.13%, so A76+DotProd auto-selects it. M1/T4 regresses 2.43%. **The separate Pi llama.cpp floor is now MEASURED/NOT MET on speed:** vllm.cpp is 0.461x prefill / 0.653x decode+E2E, while using 24.2% less RSS; exact-prompt 64-token output matches. No 20-core/i8mm result transfers. Thread partition, BF16 GEMM/speed closure and concurrency stay open. **macOS UNBLOCKED 2026-07-22 (`CLAIM-BACKEND-FANOUT-1`, W0):** the CPU registrar force-link is fixed; `test_backend` is **7/7 (18/18)** and the whole tree builds `-Werror`-clean on M4. Two unrelated macOS test gaps remain recorded in [environment.md](environment.md) | [threadpool leaf](specs/gguf-cpu-threadpool.md); [RPi5 Cortex-A76 campaign](specs/rpi5-cortex-a76-cpu-optimization.md); [Pi competitor evidence](../docs/bench-evidence/rpi5-a76-llamacpp-20260806.md) | `PARTIAL` | PR #65 | | `BACKEND-CPU-ZEN` | AMD Zen CPU with ZenDNN/zentorch dispatch and weight prepack | `platforms/zen_cpu.py:12-32`; detection `platforms/__init__.py:153-192`; `tests/test_zen_cpu_platform_detection.py:8-37` | generic CPU backend only; no Zen-specific dispatch | - | [CUDA inventory](specs/cuda-architecture-inventory.md); leaf spike required | `INVENTORIED` | - | -| `BACKEND-ROCM` | AMD ROCm/HIP | `platforms/__init__.py:110-128`, `platforms/rocm.py:43-125` (incl. `_capability_from_gcn_arch:223-291`, `_get_backend_priorities:407`), `CMakeLists.txt:20-59,196-211` (HIP build), ROCm ops rooted at `csrc/rocm/moe_q_gemm_rdna3.cu:1` @ pin `555967922` | Lands through the realized seams: a `platforms/rocm.cpp` Platform + a `kROCM` attention TU self-registering via the [attn-registry seam](../include/vllm/v1/attention/registry.h#L44) + a `rocm.cpp` `get_attn_backend_priority()` slot (mirror `rocm.py`'s ROCM_ATTN/AITER ordering); ZERO selector/model/runner edit **AS-BUILT:** **M3 (2026-08-17, #1056) — the `kROCM` attention TU LANDED: `RocmAttentionBackend` (ROCM_ATTN) self-registers via the attn-registry seam + `rocm.cpp` `get_attn_backend_priority()` mirrors `rocm.py:407-441` verbatim; still ZERO engine edit. The runner CONSUMPTION of the seam (#1065 — per-group selection + per-group KV-shape validation) is a DIFFERENT concern, tracked as `BACKEND-ATTN-SELECTION-RUNNER` (see [rocm-attn-backend.md](specs/rocm-attn-backend.md)); it does NOT ride this row's landing contract.** **W0 skeleton landed 2026-08-06 **W0 skeleton landed 2026-08-06 and was COMMUNITY-VERIFIED (issue #41): compiled clean + `ctest -R 'rocm\|cross_device'` green on gfx1151, gfx1103, gfx1100 and gfx1201 — M0/M1 MET on all four (teardown caveats: TheRock-nightly exit hang on gfx1103; #132 `-O0` CLR race on gfx1100). W1 approach-(b) landed 2026-08-08, IMPLEMENTED-UNVERIFIED (blind, no AMD hardware here): managed-alloc branch on integrated managed-capable devices (`hipMallocManaged` + `UnifiedMemory()=true` by construction, resolving F6 where XNACK-less RDNA3 probes `PageableMemoryAccess=0`), introspection probes `ManagedAllocActive`/`IntegratedDevice`, two runtime-gated tests, and the F1/F3 `ROCM_PATH`→HIP-compiler-hints CMake absorption for Arch/TheRock layouts. Discrete path byte-identical (branch provably dead at `Integrated=0`). The two plain-C++ legs stay compile-gated everywhere (`vllm_rocm_platform_syntax_check`, object-only, never linked).** [device.h](../include/vt/device.h) `kROCM`; [rocm_arch.h](../include/vt/rocm/rocm_arch.h) (capability parse, HIP-free, 1:1 from `rocm.py:223-291`); [rocm_runtime.h](../include/vt/rocm/rocm_runtime.h); [rocm_backend.hip](../src/vt/rocm/rocm_backend.hip) (6 virtuals + multi-device registrar); [rocm_rmsnorm.hip](../src/vt/rocm/rocm_rmsnorm.hip) (port of `cuda_ops.cu:96-126`); [rocm_ops.hip](../src/vt/rocm/rocm_ops.hip) (1 of 106 OpIds); [platforms/rocm.cpp](../src/vllm/platforms/rocm.cpp) (M3: dense priority mirrors `rocm.py:407-441` verbatim — ROCM_ATTN registered, AITER/TRITON_ATTN/TURBOQUANT named placeholders); `VLLM_CPP_HIP` + `VLLM_CPP_HIP_ARCHITECTURES` + `ROCM_PATH` in [CMakeLists.txt](../CMakeLists.txt); and the ONE non-additive site this uncovered — [platform.cpp](../src/vllm/platforms/platform.cpp) `kCurrentPriority` (a platform absent from that walk registers fine and is NEVER selected, with no compiler diagnostic), now exposed via `CurrentPlatformPriority()` and gated | **VERIFIED on the CPU tier (no GPU):** [test_rocm_arch](../tests/vt/test_rocm_arch.cpp) 7 cases / 40 assertions incl. upstream's own worked examples + the gfx9 feature-suffix strip + `static_assert`s; [test_platform](../tests/vllm/platforms/test_platform.cpp) new case gating that EVERY `DeviceType` is in the selection walk and CPU is last; clean `-Werror` CPU build (the enum forced exactly ONE switch site tree-wide, `test_backend_cross_device.cpp:59`); full `ctest` green; `check-device-leakage` unchanged at 32. **COMMUNITY-VERIFIED (W0, #41 tables):** [test_rocm_backend](../tests/vt/test_rocm_backend.cpp) 5 cases/1044 assertions and the cross-device RmsNorm-vs-CPU-oracle comparison (NMSE <= 5e-4) green on all four boards. **PENDING-community (W1 (b), no hardware here):** the (b) delta's first HIP compile, the two new test cases (alloc-path/UnifiedMemory coupling; kernel-write→host-read no-copy), the flag-free Arch/TheRock configure, and the unblocked M2 reference-tier e2e on gfx1151/gfx1103. **gfx1200 (RX 9060 XT, discrete, RDNA4) INDEPENDENTLY VERIFIED 2026-08-10 ([#269](https://github.com/mudler/vllm.cpp/issues/269), off the #41 four-board set):** M0/M1 MET; M2 MET via the native-kernel path (no reference tier on a dGPU — `VT_OP_PROVIDER_STATS=1` shows zero fallbacks); **M4 MET for `Gemma3ForCausalLM` (gemma-3-1b-it), 48/48 tokens against TWO independent real vLLM-ROCm oracles** on this exact board — a prebuilt AMD `gfx120X` image and a from-source build at this project's own pinned commit `555967922` (compiled clean in ~6.5 min against `rocm/vllm-dev:base`, whose ROCm 7.2.3 matches this board's native build exactly). `Qwen3ForCausalLM` (Qwen3-0.6B) surfaced one genuine near-tie prompt where the two real oracles disagree with each other (each internally deterministic) — direct proof of a version-sensitive tie in the reference itself, not a defect in either backend | [ROCm W0 spec](specs/rocm-backend-w0.md), [unified-memory (b) decision record](specs/rocm-unified-memory-b.md), [gfx1200 M2/M4 spec](specs/rocm-gfx1200-m2-correctness.md), [rocm attn backend spec](specs/rocm-attn-backend.md), [contributor guide](../docs/ROCM.md) | `ACTIVE` | `CLAIM-ROCM-W0-1` | +| `BACKEND-ROCM` | AMD ROCm/HIP | `platforms/__init__.py:110-128`, `platforms/rocm.py:43-125` (incl. `_capability_from_gcn_arch:223-291`, `_get_backend_priorities:407`), `CMakeLists.txt:20-59,196-211` (HIP build), ROCm ops rooted at `csrc/rocm/moe_q_gemm_rdna3.cu:1` @ pin `555967922` | Lands through the realized seams: a `platforms/rocm.cpp` Platform + a `kROCM` attention TU self-registering via the [attn-registry seam](../include/vllm/v1/attention/registry.h#L44) + a `rocm.cpp` `get_attn_backend_priority()` slot (mirror `rocm.py`'s ROCM_ATTN/AITER ordering); ZERO selector/model/runner edit **AS-BUILT:** **M3 (2026-08-17, #1056) — the `kROCM` attention TU LANDED: `RocmAttentionBackend` (ROCM_ATTN) self-registers via the attn-registry seam + `rocm.cpp` `get_attn_backend_priority()` mirrors `rocm.py:407-441` verbatim; still ZERO engine edit. The runner CONSUMPTION of the seam (#1065 — per-group selection + per-group KV-shape validation) is a DIFFERENT concern, tracked as `BACKEND-ATTN-SELECTION-RUNNER` (see [rocm-attn-backend.md](specs/rocm-attn-backend.md)); it does NOT ride this row's landing contract.** **W0 skeleton landed 2026-08-06 **W0 skeleton landed 2026-08-06 and was COMMUNITY-VERIFIED (issue #41): compiled clean + `ctest -R 'rocm\|cross_device'` green on gfx1151, gfx1103, gfx1100 and gfx1201 — M0/M1 MET on all four (teardown caveats: TheRock-nightly exit hang on gfx1103; #132 `-O0` CLR race on gfx1100). W1 approach-(b) landed 2026-08-08, IMPLEMENTED-UNVERIFIED (blind, no AMD hardware here): managed-alloc branch on integrated managed-capable devices (`hipMallocManaged` + `UnifiedMemory()=true` by construction, resolving F6 where XNACK-less RDNA3 probes `PageableMemoryAccess=0`), introspection probes `ManagedAllocActive`/`IntegratedDevice`, two runtime-gated tests, and the F1/F3 `ROCM_PATH`→HIP-compiler-hints CMake absorption for Arch/TheRock layouts. Discrete path byte-identical (branch provably dead at `Integrated=0`). The two plain-C++ legs stay compile-gated everywhere (`vllm_rocm_platform_syntax_check`, object-only, never linked).** [device.h](../include/vt/device.h) `kROCM`; [rocm_arch.h](../include/vt/rocm/rocm_arch.h) (capability parse, HIP-free, 1:1 from `rocm.py:223-291`); [rocm_runtime.h](../include/vt/rocm/rocm_runtime.h); [rocm_backend.hip](../src/vt/rocm/rocm_backend.hip) (6 virtuals + multi-device registrar); [rocm_rmsnorm.hip](../src/vt/rocm/rocm_rmsnorm.hip) (port of `cuda_ops.cu:96-126`); [rocm_ops.hip](../src/vt/rocm/rocm_ops.hip) (1 of 106 OpIds); [platforms/rocm.cpp](../src/vllm/platforms/rocm.cpp) (M3: dense priority mirrors `rocm.py:407-441` verbatim — ROCM_ATTN registered, AITER/TRITON_ATTN/TURBOQUANT named placeholders); `VLLM_CPP_HIP` + `VLLM_CPP_HIP_ARCHITECTURES` + `ROCM_PATH` in [CMakeLists.txt](../CMakeLists.txt); and the ONE non-additive site this uncovered — [platform.cpp](../src/vllm/platforms/platform.cpp) `kCurrentPriority` (a platform absent from that walk registers fine and is NEVER selected, with no compiler diagnostic), now exposed via `CurrentPlatformPriority()` and gated | **VERIFIED on the CPU tier (no GPU):** [test_rocm_arch](../tests/vt/test_rocm_arch.cpp) 7 cases / 40 assertions incl. upstream's own worked examples + the gfx9 feature-suffix strip + `static_assert`s; [test_platform](../tests/vllm/platforms/test_platform.cpp) new case gating that EVERY `DeviceType` is in the selection walk and CPU is last; clean `-Werror` CPU build (the enum forced exactly ONE switch site tree-wide, `test_backend_cross_device.cpp:59`); full `ctest` green; `check-device-leakage` unchanged at 32. **COMMUNITY-VERIFIED (W0, #41 tables):** [test_rocm_backend](../tests/vt/test_rocm_backend.cpp) 5 cases/1044 assertions and the cross-device RmsNorm-vs-CPU-oracle comparison (NMSE <= 5e-4) green on all four boards. **PENDING-community (W1 (b), no hardware here):** the (b) delta's first HIP compile, the two new test cases (alloc-path/UnifiedMemory coupling; kernel-write→host-read no-copy), the flag-free Arch/TheRock configure, and the unblocked M2 reference-tier e2e on gfx1151/gfx1103. **gfx1200 (RX 9060 XT, discrete, RDNA4) INDEPENDENTLY VERIFIED 2026-08-10 ([#269](https://github.com/mudler/vllm.cpp/issues/269), off the #41 four-board set):** M0/M1 MET; M2 MET via the native-kernel path (no reference tier on a dGPU — `VT_OP_PROVIDER_STATS=1` shows zero fallbacks); **M4 MET for `Gemma3ForCausalLM` (gemma-3-1b-it), 48/48 tokens against TWO independent real vLLM-ROCm oracles** on this exact board — a prebuilt AMD `gfx120X` image and a from-source build at this project's own pinned commit `555967922` (compiled clean in ~6.5 min against `rocm/vllm-dev:base`, whose ROCm 7.2.3 matches this board's native build exactly). `Qwen3ForCausalLM` (Qwen3-0.6B) surfaced one genuine near-tie prompt where the two real oracles disagree with each other (each internally deterministic) — direct proof of a version-sensitive tie in the reference itself, not a defect in either backend. **W6 (2026-08-27, #2065) — fp8-e4m3 KV cache store+read LANDED:** `OpId::kReshapeAndCacheFp8` registered for `kROCM` + `LoadKv` fp8 dequant in the paged-attention read; gated on gfx1100 (7/7 cases, 28/28 assertions) | [ROCm W0 spec](specs/rocm-backend-w0.md), [unified-memory (b) decision record](specs/rocm-unified-memory-b.md), [gfx1200 M2/M4 spec](specs/rocm-gfx1200-m2-correctness.md), [rocm attn backend spec](specs/rocm-attn-backend.md), [contributor guide](../docs/ROCM.md) | `ACTIVE` | `CLAIM-ROCM-W0-1` | | `BACKEND-XPU` | Intel XPU loyal port; gating exploration E4 ([backends.md](backends.md) kernel sourcing: SYCL-vs-Level-Zero call + whether upstream XPU attention contracts translate 1:1; explore when T2 scheduling begins) | `platforms/__init__.py:131-150`, `platforms/xpu.py:103-125` | enum slot only [device.h:11](../include/vt/device.h#L11) | [unavailable-backend test](../tests/vt/test_backend.cpp#L37); **SPIKE-ONLY verdict, MEASURED 2026-07-22: HW-BLOCKED — no Intel GPU on any box we own, and no acquisition path recorded.** Compounding finding: the loyalty target is INCOMPLETE upstream — `vllm/platforms/xpu.py:11-13` imports the EXTERNAL `vllm_xpu_kernels._C`/`._moe_C`/`._xpu_C`, and a search of the pinned tree for `*.sycl`/`*.dp.cpp` returns NOTHING, so there is no in-tree SYCL kernel source to mirror 1:1. What stays doable with zero Intel hardware: the attention-selector/priority POLICY port (pure data, unit-testable), oneAPI DPC++ compile coverage, and SYCL kernels executed on the oneAPI OpenCL CPU device for unit-level numerics. **No e2e or performance gate is proposed, because none can be run.** No XPU backend code exists | [backend fan-out](specs/backend-fanout-metal-vulkan-xpu.md); [CUDA inventory](specs/cuda-architecture-inventory.md) | `SPIKE` | `CLAIM-BACKEND-FANOUT-1` | | `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` | diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 34c6598599..8cc0602c08 100644 --- a/.agents/engine-matrix.md +++ b/.agents/engine-matrix.md @@ -102,7 +102,7 @@ lifecycle are unchanged. | `KV-SLIDING-LOCAL-SPECS` | Block row (claim the two leaves below, not this row): sliding-window and chunked-local KV specs | T1 | `vllm/v1/kv_cache_interface.py:205-307,480-586`; `tests/v1/test_kv_cache_spec_registry.py:174-306` | - | - | [sliding-local-yarn-long-context.md](specs/sliding-local-yarn-long-context.md) | `READY` | - | | `KV-SLIDING-WINDOW-SPEC` | `SlidingWindowSpec` sizing, grouping, admission, allocation, eviction, and prefix-cache policy; CPU G1/G2 green, while feature-positive attention/model/oracle/performance gates remain | T1 | `vllm/v1/kv_cache_interface.py:518-586`; `vllm/v1/core/single_type_kv_cache_manager.py:669-873`; `tests/v1/core/test_single_type_kv_cache_manager.py:127,259,380,413,489`; `tests/v1/core/test_prefix_caching.py:2457-3909` | `include/vllm/v1/kv_cache_interface.h:187`; `src/vllm/v1/kv_cache_spec_registry.cpp:69`; `src/vllm/v1/core/single_type_kv_cache_manager.cpp:350,377,470,920`; `src/vllm/v1/core/kv_cache_utils.cpp:21`; `src/vllm/v1/core/kv_cache_coordinator.cpp:36,119` | `tests/vllm/v1/test_kv_cache_interface.cpp:157,204,258`; `tests/vllm/v1/test_single_type_kv_cache_manager.cpp:283,331,368,411,453,476`; `tests/vllm/v1/test_kv_cache_utils.cpp:592,617`; `tests/vllm/v1/test_kv_cache_coordinator.cpp:163,238,357` | [sliding-local-yarn-long-context.md](specs/sliding-local-yarn-long-context.md) | `GATING` | - | | `KV-CHUNKED-LOCAL-SPEC` | `ChunkedLocalAttentionSpec` sizing, grouping, admission, allocation, fixed-chunk prefix-cache/recycling policy and hybrid-disabled fallback; CPU G1/G2 green, while W4/model/oracle/runtime gates remain | T1 | `vllm/v1/kv_cache_interface.py:480-514`; `vllm/v1/core/single_type_kv_cache_manager.py:876-1023`; `vllm/v1/core/kv_cache_utils.py:1403-1496`; `tests/v1/core/test_single_type_kv_cache_manager.py:54,198,456`; `tests/v1/test_kv_cache_spec_registry.py:174-315` | `include/vllm/v1/kv_cache_interface.h:219`; `src/vllm/v1/kv_cache_spec_registry.cpp:71`; `src/vllm/v1/core/single_type_kv_cache_manager.cpp:535,553,618,933`; `src/vllm/v1/core/kv_cache_utils.cpp:21`; `src/vllm/v1/core/kv_cache_coordinator.cpp:47` | `tests/vllm/v1/test_kv_cache_interface.cpp:188,204,258`; `tests/vllm/v1/test_single_type_kv_cache_manager.cpp:576,643,683,705,730,1072`; `tests/vllm/v1/test_kv_cache_utils.cpp:629,654,674,686`; `tests/vllm/v1/test_kv_cache_coordinator.cpp:188,258,380,524` | [sliding-local-yarn-long-context.md](specs/sliding-local-yarn-long-context.md) | `GATING` | - | -| `KV-FP8` | FP8 KV cache and scale handling. **W0 spike + W1 CPU brick LANDED 2026-07-29** — fp8-e4m3 K/V STORE (`Quantize(hp/scale)`) + the paged-attention READ dequant (`Dequant(fp8)*scale`) + the `cache_dtype` config parse, all CPU-gated RED-first. Storage is 1-byte fp8 (`DType::kI8`) + the `Fp8KVCacheDataType` interpretation enum (mirrors vLLM's `cache_t=uint8_t`+`KV_DTYPE`), per-tensor k/v scales (`kv_cache.py:108-191`). **W2 CUDA arm LANDED 2026-08-21** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)) -- the fp8-e4m3 store kernel + the fp8 dequant on the paged-attention read, gated for parity against the W1 CPU oracle; the two W1 device-class refusals that made the CUDA arm unreachable are gone, and the READ keeps a NAMED CPU-or-CUDA refusal because it rides additive `PagedAttentionArgs` fields on an op `kMETAL`/`kROCM` register for the FLOAT path. **Its DEVICE cases are UNEXECUTED** (no device in the implementing session), though the CUDA TUs DO COMPILE: CI `cuda-fat-build` built them for ten architectures under `-Werror=all-warnings` on `4d71e776e` (run 32495320287). That job sets `-DVLLM_CPP_BUILD_TESTS=OFF`, so nothing has EXECUTED them -- see the spec's `## Owed`. **W3 runner integration LANDED 2026-08-22** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)) -- `--kv-cache-dtype` on the server flag, the checkpoint's own `kv_cache_quant_algo` honoured when no flag is typed (read from `config.json:quantization_config` first and `hf_quant_config.json` as the fallback, which is upstream's order at `transformers_utils/config.py:751-761`), KV blocks sized at ONE byte per element so a fixed `--kv-cache-memory` buys exactly 2x the blocks, and the `k_scale`/`v_scale` path with its declared-but-absent arm named rather than defaulted. The store and the read normalise K and V to the MODEL dtype first, because the fp8 store quantizes from one source dtype and the attention preamble emits f32 K beside a bf16 V on every production weight arm. **Turning it on COSTS the fast attention kernels:** FA-2 prefill, FA-2 decode, the WMMA ladder and the vectorized decode-opt/GQA kernels are bf16-native by construction and an fp8 cache routes only through tiled prefill and block decode, so the memory win and the throughput cost have not been measured against each other -- recorded, not claimed, in the spec's `## W3`. **Residuals (honest, named):** the C ABI does not expose the flag, 16 architectures refuse rather than route, and only ONE of them (`nemotron_h_device`) names fp8 to the operator -- `qwen3_vl` alone reaches the store guard, 13 stop at their own `": KV cache must be bf16 or f32"` rule and `gemma4` stops one step earlier still, inside `vt::CastF32`, with a message that names no architecture at all, no weight loader extracts `k_scale`/`v_scale`, fp8_e5m2 CPU compute and per-head scales -- all in the spec's `## Owed` | T1 | `vllm/config/cache.py:19-36,76`; `vllm/model_executor/layers/quantization/kv_cache.py:42,108-191`; store `csrc/libtorch_stable/cache_kernels.cu:241-252,314-401`; scale convention `csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh:296-308` | codec `include/vt/fp8_kv.h`; store `src/vt/cpu/cpu_cache.cpp:143`; wrapper `src/vt/ops.cpp:2255`; read dequant `src/vt/cpu/cpu_paged_attn.cpp:82`; config parse `include/vllm/v1/kv_cache_dtype.h:37` | `tests/vt/test_ops_fp8_kv_cache.cpp:1` (8 cases / 511 assertions; RED-first: wrong store direction fails 3/480); W3 `tests/vllm/entrypoints/test_kv_cache_fp8_wiring.cpp` (31 cases, G1-G12, entering through `LoadedEngine` and through `Qwen3DenseModel::Forward` for the shared seam; G12's third case bounds the fp8 pages against the bf16 run element by element, inside e4m3's own round-trip envelope, because a bound written on the logits measures the model's insensitivity and not the cache) + `tests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp` (the `--kv-cache-dtype` flag through the REAL `VllmServerMain`) | [fp8-kv-cache](specs/fp8-kv-cache.md) | `ANCHOR-BACKFILL` | `CLAIM-KV-FP8` | +| `KV-FP8` | FP8 KV cache and scale handling. **W0 spike + W1 CPU brick LANDED 2026-07-29** — fp8-e4m3 K/V STORE (`Quantize(hp/scale)`) + the paged-attention READ dequant (`Dequant(fp8)*scale`) + the `cache_dtype` config parse, all CPU-gated RED-first. Storage is 1-byte fp8 (`DType::kI8`) + the `Fp8KVCacheDataType` interpretation enum (mirrors vLLM's `cache_t=uint8_t`+`KV_DTYPE`), per-tensor k/v scales (`kv_cache.py:108-191`). **W2 CUDA arm LANDED 2026-08-21** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)) -- the fp8-e4m3 store kernel + the fp8 dequant on the paged-attention read, gated for parity against the W1 CPU oracle; the two W1 device-class refusals that made the CUDA arm unreachable are gone, and the READ keeps a NAMED CPU-or-CUDA refusal because it rides additive `PagedAttentionArgs` fields on an op `kMETAL`/`kROCM` register for the FLOAT path. **Its DEVICE cases are UNEXECUTED** (no device in the implementing session), though the CUDA TUs DO COMPILE: CI `cuda-fat-build` built them for ten architectures under `-Werror=all-warnings` on `4d71e776e` (run 32495320287). That job sets `-DVLLM_CPP_BUILD_TESTS=OFF`, so nothing has EXECUTED them -- see the spec's `## Owed`. **W3 runner integration LANDED 2026-08-22** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)) -- `--kv-cache-dtype` on the server flag, the checkpoint's own `kv_cache_quant_algo` honoured when no flag is typed (read from `config.json:quantization_config` first and `hf_quant_config.json` as the fallback, which is upstream's order at `transformers_utils/config.py:751-761`), KV blocks sized at ONE byte per element so a fixed `--kv-cache-memory` buys exactly 2x the blocks, and the `k_scale`/`v_scale` path with its declared-but-absent arm named rather than defaulted. The store and the read normalise K and V to the MODEL dtype first, because the fp8 store quantizes from one source dtype and the attention preamble emits f32 K beside a bf16 V on every production weight arm. **Turning it on COSTS the fast attention kernels:** FA-2 prefill, FA-2 decode, the WMMA ladder and the vectorized decode-opt/GQA kernels are bf16-native by construction and an fp8 cache routes only through tiled prefill and block decode, so the memory win and the throughput cost have not been measured against each other -- recorded, not claimed, in the spec's `## W3`. **Residuals (honest, named):** the C ABI does not expose the flag, 16 architectures refuse rather than route, and only ONE of them (`nemotron_h_device`) names fp8 to the operator -- `qwen3_vl` alone reaches the store guard, 13 stop at their own `": KV cache must be bf16 or f32"` rule and `gemma4` stops one step earlier still, inside `vt::CastF32`, with a message that names no architecture at all, no weight loader extracts `k_scale`/`v_scale`, fp8_e5m2 CPU compute and per-head scales -- all in the spec's `## Owed` | T1 | `vllm/config/cache.py:19-36,76`; `vllm/model_executor/layers/quantization/kv_cache.py:42,108-191`; store `csrc/libtorch_stable/cache_kernels.cu:241-252,314-401`; scale convention `csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh:296-308` | codec `include/vt/fp8_kv.h`; store `src/vt/cpu/cpu_cache.cpp:143`; wrapper `src/vt/ops.cpp:2255`; read dequant `src/vt/cpu/cpu_paged_attn.cpp:82`; config parse `include/vllm/v1/kv_cache_dtype.h:37` | `tests/vt/test_ops_fp8_kv_cache.cpp:1` (8 cases / 511 assertions; RED-first: wrong store direction fails 3/480); W3 `tests/vllm/entrypoints/test_kv_cache_fp8_wiring.cpp` (31 cases, G1-G12, entering through `LoadedEngine` and through `Qwen3DenseModel::Forward` for the shared seam; G12's third case bounds the fp8 pages against the bf16 run element by element, inside e4m3's own round-trip envelope, because a bound written on the logits measures the model's insensitivity and not the cache) + `tests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp` (the `--kv-cache-dtype` flag through the REAL `VllmServerMain`). **W6 ROCm arm LANDED 2026-08-27** ([#2065](https://github.com/mudler/vllm.cpp/issues/2065)) -- the fp8-e4m3 store kernel + the fp8 dequant on the ROCm paged-attention read, gated for parity against the W1 CPU oracle on gfx1100 (RX 7900 XTX); the READ refusal in `src/vt/ops.cpp` is widened to admit `kROCM`, and `OpId::kReshapeAndCacheFp8` is registered for `DeviceType::kROCM`. Metal is the only backend still refused by name | [fp8-kv-cache](specs/fp8-kv-cache.md) | `ANCHOR-BACKFILL` | `CLAIM-KV-FP8` | | `KV-NVFP4-TURBO` | NVFP4, per-token-head, and TurboQuant KV | T2 | `vllm/config/cache.py:14,28-35,272` | - | - | `planned: specs/nvfp4-kv-cache.md` | `INVENTORIED` | - | | `KV-OFFLOAD` | KV offload tiering: CPU primary tier plus secondary tiers, including the **filesystem (disk) tier that is vLLM's KV-persistence-to-disk answer**. **Record CORRECTED 2026-07-22 ([spike](specs/kv-persistence-lmcache.md)) — the prior row text named a class that does not exist and omitted the half the user asked for.** There is no `LRUOffloadingManager` at this pin: LRU and ARC are pluggable `CachePolicy` objects behind ONE `CPUOffloadingManager`, and the row's scope ('CPU tiering with LRU and ARC') left out the entire secondary-tier surface. Disk format enumerated: ONE RAW FILE PER BLOCK, no container and no index, `/__r//_g/.bin`, written via temp-file + atomic rename under `O_DIRECT` and self-healing by deleting unreadable files. Two upstream WEAKNESSES recorded as beyond-parity targets: `config.json` is written and NEVER read (the only identity check is a path digest omitting checkpoint content, weight quantization, rope config and `sliding_window`), and the disk tier has NO capacity accounting and NO eviction. Secondary tiers can never touch GPU memory — all traffic cascades through the CPU primary tier **W1-W3 IMPLEMENTED 2026-07-22.** Deterministic block hashes (W1), the CPU primary tier (W2: `CachePolicy` LRU+ARC with the `ref_cnt == -1` tri-state and the ATOMIC evict, `CPUOffloadingManager` incl. the `prepare_store -> nullopt` skip path, pinned backing store plus side-queue event-polled device/host transfer), and the DISK tier (W3: one raw file per block, temp-file + atomic rename publish, self-healing unlink, dual-queue read/write pool). **BOTH recorded upstream weaknesses are now EXCEEDED, not merely noted:** the identity block is a VERIFIED header read on every open that REFUSES on mismatch across 27 fields (upstream's `config.json` is never read), and the tier carries a byte budget with policy-driven eviction honoured across restarts (upstream has none). `O_DIRECT` is deliberately NOT ported — a header+payload file breaks its alignment requirement; recorded. **W4 IMPLEMENTED 2026-07-23.** The TIERING MANAGER (ONE manager over the CPU primary + disk secondary tier: disk→CPU promotion is RETRY this step / HIT the next with the reserved slot marked in-flight, cascade demotion on store, reset drains the secondary FIRST and DELIBERATELY never resets it so a persisted cache survives a prefix-cache reset) and the CONNECTOR/SCHEDULER HALF (`OffloadingConnector` mirroring `KVConnectorBase_V1`'s scheduler hooks — `get_num_new_matched_tokens` with the load-bearing NULLOPT third state, `Request::block_hashes` striding, load-before-compute ordering, `build_connector_meta` reset — wired OPT-IN and DEFAULT-OFF into the scheduler so a cross-request/restarted-process prefix HIT shortcuts prefill). The semantics are ported, NOT the Python plugin ABI (compile-time wiring replaces the `importlib` module path; the full 7-method abstract ABI + registration + `KVTransferConfig` is the W5 generalization behind the same seam). Deviation recorded: W4 ships the SYNCHRONOUS-load shape (async flag always false), the disk→CPU promotion being the async part handled by RETRY/re-ask; the cross-step `WAITING_FOR_REMOTE_KVS` GPU-load buffer is W5. **First measured offload speedup:** a restarted-prefix workload through the real scheduler saved 32/48 prefill tokens (2/3 blocks HIT from disk) with the promoted bytes proven byte-identical to the cold store. **W5 LANDED 2026-07-23** (the connector seam is now a first-class C++ ABI — abstract `KVConnector` base + `KVConnectorFactory` + `KVTransferConfig`, the disk connector refactored onto it behaviour-identically; see the `KV-CONNECTORS` row). **D1 CORRECTION 2026-07-24 (`CLAIM-DOCS-T2-FIXES`): the disk connector's WORKER HALF IS NOT IMPLEMENTED and is now REFUSED, not merely absent.** `OffloadingConnector` emits `ConnectorLoadJob`s that NOTHING consumes, and its bytes live in a host `PrimaryByteView` that is never copied into a KV page — on any device. Because its scheduler half DOES shortcut prefill for matched blocks, wiring it into an engine would have made the model attend over never-written KV (silently wrong output); `BuildKvConnector` previously built it for any device with no guard. It is now refused at construction by a per-connector capability predicate (`KVConnector::supports_worker_transfer_on` / the registered `KVConnectorWorkerTransferFn`, queried by name BEFORE construction via `KVConnectorFactory::WorkerTransferSupportedOn`), with an error naming the connector, the device, the consequence and the admissible connectors. The scheduler-side 32/48 e2e is UNAFFECTED (it never reaches a worker). Implementing the worker half remains OPEN work and is NOT claimed. W6 (LMCache study) and W7 (named save/restore) remain open | T2 | core `vllm/v1/kv_offload/base.py:27-47,88-108,177-347,486-588,536-549`; CPU tier `vllm/v1/kv_offload/cpu/manager.py:36,169-237`, policies `cpu/policies/base.py:10-33,36-92`, `lru.py:12`, `arc.py:12`; **disk tier** `vllm/v1/kv_offload/tiering/fs/io.py:32-72,75-101`, `tiering/fs/manager.py:95-103,131-137`, `tiering/fs/thread_pool.py:50-57,153-180`; naming/identity `vllm/v1/kv_offload/file_mapper.py:112-120,128-139`; tiering ordering `tiering/manager.py:238-329,408-459,498-556,643-681`; transfer `cpu/gpu_worker.py:240-421,388-394`; config `docs/features/kv_offloading_usage.md:64-82,95-121`; tests `tests/v1/kv_offload/tiering/test_fs_tier.py`, `tests/v1/kv_offload/test_file_mapper.py`, `tests/v1/kv_offload/cpu/test_manager.py` | **W1-W3 LANDED.** Core `include/vllm/v1/kv_offload/base.h` (OffloadKey verified byte-identical to upstream's packing); policies `include/vllm/v1/kv_offload/cache_policy.h` + `src/vllm/v1/kv_offload/cache_policy.cpp`; CPU tier `include/vllm/v1/kv_offload/cpu_manager.h` + `src/vllm/v1/kv_offload/cpu_manager.cpp`; transfer `include/vllm/v1/kv_offload/kv_block_transfer.h` + `src/vllm/v1/kv_offload/kv_block_transfer.cpp` (plus the new non-blocking `vt::Backend::QueryEvent` seam with its CUDA override in `src/vt/cuda/cuda_backend.cu`); disk byte path + naming `include/vllm/v1/kv_offload/fs_io.h` + `src/vllm/v1/kv_offload/fs_io.cpp`; tier `include/vllm/v1/kv_offload/fs_tier.h` + `src/vllm/v1/kv_offload/fs_tier.cpp`; the verified identity header `include/vllm/v1/kv_offload/cache_identity.h` + `src/vllm/v1/kv_offload/cache_identity.cpp`; determinism fix `src/vllm/v1/core/kv_cache_utils.cpp` (`init_none_hash` seed resolution + `none_hash_provenance`), caller `src/vllm/entrypoints/model_loader.cpp:140-152`; **W4** tiering manager `include/vllm/v1/kv_offload/tiering_manager.h` + `src/vllm/v1/kv_offload/tiering_manager.cpp`; connector/scheduler half `include/vllm/v1/kv_offload/kv_connector.h` + `src/vllm/v1/kv_offload/kv_connector.cpp`; scheduler wiring `src/vllm/v1/core/sched/scheduler.cpp` (`set_kv_connector`, null = zero change) + `include/vllm/v1/core/sched/scheduler.h`; `BlockPool::evict_blocks` `src/vllm/v1/core/block_pool.cpp:139-155` (1:1, replaces the throw) | `tests/vllm/v1/test_none_hash_determinism.cpp:108` 7/7 (cross-PROCESS byte-identical hash chains via a `/proc/self/exe` re-exec, both env escape hatches, and the `=random` negative control); `tests/vllm/v1/test_kv_offload_cpu.cpp` 21/21 (atomic evict, pinning, ARC promotion, HIT_PENDING, failed-store rollback, same-batch protection, store_threshold, events, transfer round-trip); `tests/vllm/v1/test_kv_offload_fs.cpp` 22/22 + 3 SKIP (byte-exact round trip for full attention AND MLA rank-3, truncation/foreign-magic/misfiled refusal with self-heal, a 27-field identity-refusal matrix with a positive control, the byte budget across a restart, and a 6/6 cross-restart hit measurement); the SKIPs are row-tagged to `KV-SLIDING-WINDOW-SPEC`, `KV-FP8`/`KV-NVFP4-TURBO` and `KV-MAMBA-ALIGN`; **W4** `tests/vllm/v1/test_kv_offload_tiering.cpp` 5/5 (promotion RETRY→HIT byte-identical, CPU-eviction→disk-survival→re-promotion, reset clears CPU but disk survives, a FRESH manager on the same directory promotes = restart, and identity REFUSAL through a promotion — a corrupt disk block is unlinked and treated as absent, never trusted) and `tests/vllm/v1/test_kv_offload_connector.cpp` 4/4 (null-connector inertness, external match shortcuts prefill by exactly ext, the nullopt third state defers then schedules next step, and the END-TO-END restarted-prefix disk HIT through the real scheduler: hit rate 2/3 blocks, 32/48 prefill tokens saved, promoted bytes byte-identical) | [kv-persistence-lmcache.md](specs/kv-persistence-lmcache.md) | `ANCHOR-BACKFILL` | `CLAIM-KV-PERSISTENCE-LMCACHE` | | `KV-EXTERNAL-CACHE` | External KV-cache provider ABI plus LMCache interoperability: producer/consumer/both roles, the scheduler/worker metadata split, cache registration, block-hash lookup, asynchronous load/store and completion/free ownership. **SPIKED 2026-07-22 ([spike](specs/kv-persistence-lmcache.md)) — the ABI is smaller than the row implied and the LMCache half is larger.** The minimum viable connector is **exactly 7 abstract methods** (worker `start_load_kv`/`wait_for_layer_load`/`save_kv_layer`/`wait_for_save`, scheduler `get_num_new_matched_tokens`/`update_state_after_alloc`/`build_connector_meta`); roughly thirty further hooks all have safe defaults. Three traps recorded: `get_num_new_matched_tokens` has a THIRD state (`None` = deschedule and re-ask, not zero), `request_finished` returning True transfers block-freeing OWNERSHIP to the connector, and non-HMA connectors ASSERT a single KV cache group while our gate models are two-group hybrids. **LMCache determination: it is an EXTERNAL PyPI package** (`lmcache >= 0.3.9` in an opt-in extras file that `setup.py`/`pyproject.toml` never reference; not installed on any of this project's boxes). vLLM vendors roughly 2396 lines of `lmcache_integration/` glue, but every one of those files imports the external package at module scope — the storage engine, the paged-memory GPU connectors, the config schema, the ZMQ message queue and the **CUDA-IPC** handoff are all outside the tree, and no upstream test exercises it without importing `lmcache`. Scoped as an interop STUDY, not a from-scratch client, and gated on two blockers we own: our `sha256_cbor` hashes are not byte-compatible with vLLM's default, and our `NONE_HASH` is per-process random. **REOPENED 2026-07-23 ([client spike](specs/lmcache-cpp-client-connector.md)) on the user's connect-as-client hypothesis, and the prior "no specified wire protocol" verdict is REFUTED by reading the LMCache package (`LMCache/LMCache@8570aad`).** vLLM connects to a RUNNING LMCache instance over two fully-specified, language-agnostic wires: (1) the `lm://` remote-store server — **plain TCP + a fixed `struct.pack` header + raw KV bytes**, no ZMQ/msgpack/pickle/CUDA-IPC (`lmcache/v1/protocol.py:214-321`, `server/__main__.py:24-147`, `lm_connector.py:28-177`); and (2) the MP server — **ZMQ DEALER↔ROUTER + `msgspec.msgpack` control + CUDA-IPC data** (`multiprocess/mq.py:270-353`, `custom_types.py:120-234`), the mode the user recalled as "zmq". BOTH need ZERO `lmcache` in our process and BOTH sidestep the R1 hash blocker — LMCache keys on its OWN blake3 rolling token hash (`token_hasher.py:54-79`), never vLLM block hashes. Pickle appears ONLY in the MP one-time IPC-wrapper registration (`platform/base/ipc_wrapper.py` Serialize); CUDA-IPC ONLY in MP data (portable via `RawCudaIPCWrapper` `cudaIpcGetMemHandle`, but co-located). Verdict: a C++ client is FEASIBLE — recommend MODE (1) first (stabler/simpler); the standing risk is LMCache being an unpinned moving target, so it is an interop feature with a version-sync cost, not a mechanical core port | T2 | ABI `vllm/distributed/kv_transfer/kv_connector/v1/base.py:171,293,311,325,347,454,489,510,542,585`; roles `:124`; HMA `:85,93`; factory + out-of-tree module seam `vllm/distributed/kv_transfer/kv_connector/factory.py:28,31,96,102-123,152-238`; config `vllm/config/kv_transfer.py:22-75,102-106`; MRV2 worker hooks `vllm/v1/worker/gpu/kv_connector.py:56,61-75,77-95`; scheduler call sites `vllm/v1/core/sched/scheduler.py:280,736-742,933-937,1118-1119,2340-2371`; LMCache `vllm/distributed/kv_transfer/kv_connector/v1/lmcache_connector.py:74-115,259,281`, `lmcache_mp_connector.py:1-50`, `lmcache_integration/vllm_v1_adapter.py:11-35,175-188,368-376,781`, external requirement `requirements/kv_connectors.txt:1`; tests `tests/v1/kv_connector/unit/test_lmcache_integration.py:60-223`, `test_kv_connector_lifecycle.py:37`, `test_config.py:51` | **W1 LANDED 2026-07-23 — the LMCache MODE-1 `lm://` wire CODEC (pure CPU, INERT: no call site routes to it, the connector is W3):** `src/vllm/v1/kv_offload/lmcache/remote_protocol.{h,cpp}` (186-byte `ClientMetaMessage` / 36-byte `ServerMetaMessage` fixed-`struct` framing + `ClientCommand`/`ServerReturnCode`/`DTYPE_TO_INT`/`Location` maps), `cache_engine_key.{h,cpp}` (`model@world@worker@chunk_hash_hex@dtype` to/from string), `token_hasher.{h,cpp}` (blake3 rolling chunk hash over vendored `third_party/blake3/` 1.5.5), `memory_format.{h,cpp}` (the `KV_2LTD` `[2,L,T,D]` repack); wired in `CMakeLists.txt` (`blake3_vendored` static lib). Later-connector seams still NAMED: `include/vllm/v1/core/kv_cache_manager.h:31` (`ext_comp`), `include/vllm/v1/core/single_type_kv_cache_manager.h:122`, `include/vllm/v1/core/sched/output.h:30-31`, `include/vllm/v1/engine/types.h:26,30`. **W5 worker-side store/load LANDED 2026-07-24 (the last open arm):** `src/vllm/v1/worker/gpu/runner.cpp` (`ConnectorLoadExternalKv` writes the external-prefix KV into the allocated GPU blocks BEFORE the forward = load-before-compute; `ConnectorStorePromptKv` stores each newly-complete prompt block AFTER the forward; both behind a `kv_connector_ != nullptr` guard so default-off is byte-identical) + `include/vllm/v1/worker/gpu/runner.h` (`set_kv_connector`), `src/vllm/entrypoints/model_loader.cpp` (`BuildKvConnector` builds the connector from `EngineParams::kv_transfer_config` via `KVConnectorFactory`, injects the runner's full-attention KV geometry, wires it to scheduler + runner) + `include/vllm/entrypoints/model_loader.h` (`EngineParams::kv_transfer_config`, `LoadedEngine::kv_connector()`) | **W1 byte/bit-exact gate GREEN (CPU): `tests/vllm/v1/kv_offload/lmcache/test_lmcache_codec.cpp:105` (6 cases / 2074 assertions) vs `tests/fixtures/lmcache/lmcache_fixtures.json` — our wire bytes == the real Python codec's (stdlib `struct` framing + `blake3` PyPI hashes + numpy KV_2LTD); blake3 digest VERIFIED byte-identical on x86-64 AND `dgx.casa` aarch64.** **W2 (client, CPU) GREEN — go/no-go PASSED:** `src/vllm/v1/kv_offload/lmcache/remote_client.{h,cpp}` (blocking POSIX-socket PUT/GET/EXIST/HEALTH/LIST + partial-read/write loops + `PutKv2ltd`/`GetKv2ltd` `KV_2LTD` repack + `LmcacheClientConfig`/`VT_LMCACHE_*` env); `tests/vllm/v1/kv_offload/lmcache/test_lmcache_client.cpp` round-trips a **REAL `lmcache.v1.server`** (`8570aad`, run headless from source in a throwaway venv — torch imported before lmcache to dodge a torch circular import, the compiled `c_ops` ext stubbed as unused by the lm:// CPU store) byte-identical (36/36), and interop is **BIDIRECTIONAL** with LMCache's OWN Python protocol codec (`scripts/lmcache/{lm_server,lm_interop_client}.py`+`run_live_roundtrip.sh`); always-on CI gate = a same-binary C++ mock-server round-trip (45/45, no Python). **W3 LANDED 2026-07-23 — the `lm://` client wired as a `KVConnector` over the W5 seam (the FIRST time engine -> connector -> W2 client -> a running lm:// server -> back runs):** `src/vllm/v1/kv_offload/lmcache/lmcache_connector.{h,cpp}` (`LMCacheConnector : KVConnector`, `REGISTER_KV_CONNECTOR("LMCacheConnector", …)`, selected by `KVTransferConfig{kv_connector="LMCacheConnector", kv_connector_extra_config={host,port,hash_algo,chunk_tokens,…}}`, default OFF). Scheduler side is real: `get_num_new_matched_tokens` computes the request's rolling-blake3 chunk hashes, builds the `CacheEngineKey` per chunk and `Exist`-probes the REMOTE store for the longest cached prefix (synchronous -> `(n, false)`, mirroring `lmcache_connector.py:230-259`); `update_state_after_alloc` records the load (drops `blocks` upstream, `:261-268`); worker `StoreChunk` (PUT KV_2LTD) / `LoadChunk` (GET+unpack, foreign-block REFUSAL via `GetKv2ltd`). **Gate ACHIEVED = the connector-level round-trip: store -> lookup -> prefill-shortcut through the REAL scheduler -> load byte-identical (32/48 prefill tokens saved), foreign/mismatched-key REFUSAL, default-off inertness** (`tests/vllm/v1/kv_offload/lmcache/test_lmcache_connector.cpp` 5 cases / 50 assertions vs an in-process mock; the store->load round-trip ALSO passes vs a REAL `lmcache.v1.server` 8570aad, 16 assertions, under `VT_LMCACHE_LIVE_*`). **W4 LANDED 2026-07-23 — REAL peer KEY-AGREEMENT + a peer->us interop LOAD, both PROVEN (the interop-correctness milestone is complete; the row stays `ACTIVE` only for the DGX full-model output-invariance + throughput arm, spec gates 4/6):** the actual `lm://` key derivation is NOT the blake3 MP `TokenHasher` (a different subsystem) but `ChunkedTokenDatabase` (`lmcache/v1/token_database.py:298-449`) — chunk_size 256, a rolling prefix-hash chain over the 3-tuple `(prefix_int, tuple(tokens), extra_keys=())`, keyed by vLLM's OWN hash function (`pre_caching_hash_algorithm`; the portable interop choice `sha256_cbor` = cbor2-canonical + SHA-256, `vllm/utils/hashing.py:43`), folded to uint64 each step (`_normalize_hash_to_int` `token_database.py:34-56`), with `NONE_HASH = fold8(sha256_cbor(str(PYTHONHASHSEED)))` (`kv_cache_utils.py:99-114`). Mirrored BYTE-EXACT in `src/vllm/v1/kv_offload/lmcache/chunked_token_database.{h,cpp}` (reusing the project's `CborValue`+`sha256_cbor`, already Python-cbor2/hashlib-exact), and wired into the connector as `key_mode=kVllmSha256Cbor` (`hash_algo="vllm"/"sha256_cbor"`, chunk 256) alongside W3's kept-green blake3 path. **Key-agreement gate GREEN:** `tests/vllm/v1/kv_offload/lmcache/test_lmcache_key_agreement.cpp` (4 cases / 85 assertions) asserts our `CacheEngineKey` strings + chunk boundaries + folded hashes are BYTE-IDENTICAL to the REAL lmcache `ChunkedTokenDatabase.process_tokens()` (fixtures `tests/fixtures/lmcache/key_agreement_fixtures.json` dumped by `scripts/lmcache/gen_key_agreement_fixtures.py` driving the unmodified real driver, with vLLM's pinned `sha256_cbor`/`init_none_hash`), incl. the connector's own peer-mode `ChunkKey`. Sample: tokens 1000..1511 -> `meta-llama/Llama-3.1-8B@1@0@33d6862800fff40c@bfloat16`. **Peer->us interop LOAD gate GREEN (over the wire, real server):** `scripts/lmcache/{lm_key_interop.py,run_key_interop.sh}` has the REAL lmcache `ChunkedTokenDatabase` derive a key from tokens and PUT KV to a REAL `lmcache.v1.server` (8570aad, headless); our C++ INDEPENDENTLY re-derives the SAME key and GETs the peer-written 512 B byte-identical (`test_lmcache_key_agreement` LIVE case under `VT_LMCACHE_LIVE_SPEC`). ASan+UBSan clean on the connector path. Text-only scope (mm-hash extra_keys deferred); the DGX full-model output-invariance + throughput are the W5 arm below. **W5 OUTPUT-INVARIANCE GATE GREEN 2026-07-24 (spec gates 4+6 met — the LAST open arm CLOSED):** `tests/vllm/models/test_lmcache_output_invariance.cpp` on a REAL OPT-125m bf16 loop vs a live `lmcache.v1.server` (8570aad, headless per the W2 recipe) proves connector-ON generated tokens are BIT-IDENTICAL to connector-OFF cold full prefill (first-divergence index -1) in BOTH modes — (a) store->restart->load within one process AND (b) a genuinely COLD second process that only hits the server (`VT_LMCACHE_OI_MODE=loadonly`) — with prefill SAVED on the hit = 48 tokens (3×16-token blocks) and chunks_stored>0; driven by `scripts/lmcache/run_output_invariance.sh` under `flock $HOME/gpu.lock`, `VT_ASYNC_SCHED=0`. Throughput reported HONESTLY: on a 125M model wall-clock is noise-dominated (fixed TCP/copy overhead ~ tiny compute saved) so NO binding speedup is claimed — a real speed number is owed by an every-axis grid on a larger model + long shared-prefix corpus (docs/BENCHMARKS.md). No-regression WITNESS: OPT SACRED gate UNCHANGED default-off (`test_opt_paged_engine` 6/6 prompts, 96/96 tokens, 63/63 assertions) with the connector code present; connector units green (codec 6/6·2074, client 3/3·45, connector 5/5·50, key-agreement 4/4·85, kv_offload_connector 11/11·80); ASan+UBSan clean on the connector path (0 sanitizer hits); CUDA `-Werror` 0 warnings. Additive + default-off inert (scheduler/worker/seam untouched) | [kv-persistence-lmcache.md](specs/kv-persistence-lmcache.md); LMCache client wire analysis + W-plan [lmcache-cpp-client-connector.md](specs/lmcache-cpp-client-connector.md) | `ANCHOR-BACKFILL` (W1-W5 landed; the connector-ON full-model OUTPUT-INVARIANCE arm is CLOSED — connector-ON == connector-OFF tokens BIT-IDENTICAL on a real OPT-125m loop vs a live `lmcache.v1.server`, both after an in-process restart and from a cold second process, spec gates 4/6 met; a BINDING every-axis LMCache throughput grid on a LARGER model stays PENDING, mirroring the Llama 'correctness DONE, speed PENDING' disposition — a 125M model's wall time is noise-dominated) | `CLAIM-LMCACHE-CPP-CLIENT` (W1 codec + W2 client + W3 connector + W4 key-agreement + W5 output-invariance); parent seam `CLAIM-KV-PERSISTENCE-LMCACHE` | diff --git a/.agents/feature-matrix.md b/.agents/feature-matrix.md index 9976c97748..9dec2b1451 100644 --- a/.agents/feature-matrix.md +++ b/.agents/feature-matrix.md @@ -82,7 +82,7 @@ Confirmed NON-gap: vLLM has removed prompt adapters. | SGLang RadixAttention behavior parity (fuse-or-flag) | SGLang v0.5.15 `f63458b` `mem_cache/radix_cache.py`, `managers/schedule_policy.py`, `constrained/outlines_jump_forward.py` | `ACTIVE` T2 | **Scoped 2026-07-27 (`CLAIM-SGLANG-RADIX-SCOPE`); W1+W2 IMPLEMENTED 2026-07-27 (`CLAIM-SGLANG-IMPL`, rows now `ACTIVE`).** VERDICT: SGLang's radix TREE == our block-hash APC ⇒ RadixAttention is **already FUSED**, `--enable-radix-attention` is an ALIAS for the APC toggle (LANDED: server alias + C-ABI `enable_prefix_caching` tri-state, ABI v7). Genuinely-distinct behavior = cache-aware **LPM scheduling `--schedule-policy=lpm`** (LANDED: `SchedulerPolicy::kLPM` reorders the FCFS waiting deque by APC longest-match, ported FROM `schedule_policy.py:205,229`, output-neutral; gate `test_scheduler_lpm` 6/6). **SW2 in-batch prefix-collision de-prioritization LANDED 2026-07-27 (`CLAIM-SGLANG-SW2`)** inside the `kLPM` reorder (block-hash APC keys, no second trie; ported FROM `schedule_policy.py:253-301,311`), output-neutral; its throughput lever is NOT-APPLICABLE — our APC caches at allocation time so the 2nd same-step collider already hits (within-step dedup subsumes it). overlap scheduler == `ENG-ASYNC-SCHED` (fused). **SW3 jump-forward decoding — safe TOKEN-UNIQUE subset LANDED 2026-07-28 (`CLAIM-SGLANG-SW3`)**: forced-token detection hook `StructuredOutputGrammar::forced_token()` + opt-in driver `DrainForcedTokens` (env `VT_ENABLE_JUMP_FORWARD`, default OFF), provably byte-identical to per-token constrained decode (jumps only where the grammar leaves exactly one valid token — no re-tokenization); gate `test_jump_forward` 5/5 (RED-first). Residual: SW4 + the general re-tokenization span + production scheduler splice (named). **ABI/API/flag EXPOSURE 2026-07-28 (`CLAIM-SGLANG-ABI-DOCS`, reconciled to ABI v10):** LPM + jump-forward made first-class DOCUMENTED knobs on ALL THREE surfaces (were server-only / env-only) — LPM via the concurrent session's C-ABI **string** field `vllm_model_params.scheduling_policy="lpm"` (ABI v9; NO duplicate int knob) + C++ `EngineParams::policy=kLPM` + server `--scheduling-policy lpm`; jump-forward via new C-ABI `vllm_model_params.enable_jump_forward` (tri-state int, ABI **v10** appended after the v9 fields) + C++ `EngineParams::enable_jump_forward` + server `--[enable\|disable]-jump-forward`. `VT_ENABLE_JUMP_FORWARD` retained as env override. User docs [docs/SGLANG-COMPAT.md](../docs/SGLANG-COMPAT.md) + spec [sglang-enablement.md](specs/sglang-enablement.md); ABI e2e `tests/capi/test_capi.cpp` (2 v10 jump-forward cases; `vllm_abi_version()`==10). Default-inert (all-zero ⇒ byte-identical). Rows `KV-SGLANG-RADIX-CACHE` + `ENG-SGLANG-BEHAVIOR-FLAG`. Sibling benchmark track = `BACKEND-GATE-CUDA-SGLANG*` (unchanged) | [sglang-radixattention.md](specs/sglang-radixattention.md) | | **SGLang parity PROGRAM** (whole-surface inventory + oracle) | SGLang v0.5.15 `f63458b` — full runtime surface | `SPIKE` T2 | **Elevated 2026-07-27 (`CLAIM-SGLANG-PARITY-PROGRAM`).** The vLLM-parity approach replicated for SGLang: a tabular whole-surface inventory (44 rows) classifying every SGLang capability **FUSED (23) / SGLANG-DISTINCT (8) / INVENTORIED (5) / OUT-OF-SCOPE (8)**, plus SGLang stood up as a correctness + performance ORACLE (dgx GB10 via the arm64 cu130 image — no from-source build needed). SGLang is a competitor perf FLOOR + correctness cross-check, NOT the mirror source (vLLM stays behavior truth). Headline SGLANG-DISTINCT opt-ins: LPM scheduling, in-batch prefix de-prioritization, radix eviction strategies, jump-forward, custom logit processors, batch-invariant determinism, PD disaggregation, two-batch EP overlap. Full map + ranked plan in the matrix. Sibling benchmark rows `BACKEND-GATE-CUDA-SGLANG*` unchanged | [sglang-matrix.md](sglang-matrix.md); [sglang-parity-oracle.md](specs/sglang-parity-oracle.md) | | SlidingWindowSpec + ChunkedLocalAttentionSpec | `v1/kv_cache_interface.py` | `PARTIAL` T1 | Both execution leaves are implemented: W1 sliding-window and W3 chunked-local sizing, registry/grouping, manager prefix/recycling policy, admission and hybrid-disabled conversion pass their ported CPU/property/sanitizer gates (G1/G2). The compute-locality consumers are now GPU-gated (2026-07-27 `CLAIM-ROADMAP-C5`, dgx GB10: Gemma-2/Gemma-3 sliding-window model gates 48/48; `test_chunked_local_attention` 5/5). The KV memory-OPTIMIZATION path (optimized-manager held-block cap vs the full-allocation fallback the current model gates use) still needs a model-level hybrid-manager memory gate (G8) — kept `PARTIAL` honestly | [sliding-local-yarn-long-context.md](specs/sliding-local-yarn-long-context.md) | -| fp8 KV cache (`cache_dtype=fp8*`) | `layers/quantization/kv_cache.py` | ◐ T1 | `KV-FP8` ACTIVE / `QUANT-KV-FP8` PARTIAL — W1 CPU fp8-e4m3 store+read+config-parse landed; CUDA + memory-halving e2e later | [fp8-kv-cache](specs/fp8-kv-cache.md) | +| fp8 KV cache (`cache_dtype=fp8*`) | `layers/quantization/kv_cache.py` | ◐ T1 | `KV-FP8` ACTIVE / `QUANT-KV-FP8` PARTIAL — W1 CPU fp8-e4m3 store+read+config-parse landed; W2 CUDA store+read landed; W6 ROCm store+read landed; memory-halving e2e later | [fp8-kv-cache](specs/fp8-kv-cache.md) | | nvfp4 / per-token-head / turboquant KV | `config/cache.py` | ☐ T2 | | `planned: specs/nvfp4-kv-cache.md` | | KV offload (CPU tiering, LRU/ARC) | `v1/kv_offload/` | ☐ T2 | | `planned: specs/kv-offload.md` | | External KV-cache provider ABI + LMCache (MP service and in-process connectors) | `config/kv_transfer.py`, `distributed/kv_transfer/kv_connector/v1/{base,lmcache_connector,lmcache_mp_connector}.py` | ☐ T2 | explicit roadmap outcome `KV-EXTERNAL-CACHE`: mirror `kv_producer`/`kv_consumer`/`kv_both`, scheduler/worker metadata, async layer load/store, dynamic external connector modules, failure policy, metrics and cache-lifecycle ownership; gate the official LMCache shared-prefix quickstart plus Qwen3.6 hybrid behavior | `planned: specs/external-kv-cache-lmcache.md` | diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index b4c5972b18..213e868177 100644 --- a/.agents/quantization-matrix.md +++ b/.agents/quantization-matrix.md @@ -157,7 +157,7 @@ Pinned vLLM source: `vllm/config/cache.py:19-36`. | ID | Item | Upstream | Our code | Tests/evidence | Spike/spec | State | Owner | |---|---|---|---|---|---|---|---| -| `QUANT-KV-FP8` | fp8, fp8_e4m3, fp8_e5m2 | `vllm/config/cache.py:19-25`; `vllm/model_executor/layers/quantization/kv_cache.py:42-191`; store `cache_kernels.cu:241-252`; scale convention `quant_utils.cuh:296-308` | **W1 CPU fp8-e4m3 store+read LANDED**: [codec](../include/vt/fp8_kv.h#L39), [store kernel](../src/vt/cpu/cpu_cache.cpp#L143), [read dequant](../src/vt/cpu/cpu_paged_attn.cpp#L82), [config parse](../include/vllm/v1/kv_cache_dtype.h#L37). **W2 CUDA fp8-e4m3 store+read LANDED** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)): [store kernel](../src/vt/cuda/cuda_cache.cu), [read dequant](../src/vt/cuda/cuda_paged_attn.cu) -- gate [test_cuda_fp8_kv_cache](../tests/vt/test_cuda_fp8_kv_cache.cpp), whose DEVICE cases are UNEXECUTED; the CUDA TUs COMPILE (CI `cuda-fat-build`, ten architectures, run 32495320287 on `4d71e776e`) but that job builds with tests OFF, so none has been executed (spec `## Owed`). **W3 runner integration LANDED** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)): [`--kv-cache-dtype`](../src/vllm/entrypoints/openai/server_main.cpp) reaches `EngineParams`, the checkpoint's `kv_cache_quant_algo` is honoured when no flag is typed, [`ApplyCacheDType`](../src/vllm/v1/kv_cache_interface.cpp) retypes every attention spec (group and heterogeneous per-layer alike) so one byte budget buys exactly 2x the blocks, and [`kv_cache_route.h`](../include/vllm/model_executor/models/kv_cache_route.h) is the ONE place the store and the read decide float versus fp8. An fp8 cache DISABLES FA-2 prefill and decode, the WMMA ladder and the vectorized decode kernels, which are bf16-native; the resulting throughput cost is unmeasured and recorded as such. e5m2 compute, per-head scales, the Metal/ROCm arms, the C-ABI field, the 16 unrouted architectures (15 of which refuse with a message that names neither fp8 nor the flag) and the `k_scale`/`v_scale` weight-loader read are named later bricks (see spec) | [test_ops_fp8_kv_cache](../tests/vt/test_ops_fp8_kv_cache.cpp#L1) — 8 cases / 511 assertions, round-trip within the e4m3 band + fp8-vs-bf16 NMSE<1% + paged-attention e2e; RED-first (wrong store direction fails 3/480); W3 [test_kv_cache_fp8_wiring](../tests/vllm/entrypoints/test_kv_cache_fp8_wiring.cpp) — 31 cases, G1-G12, entering through `LoadedEngine` rather than by building a spec by hand, and bounding the fp8 pages against a bf16 run inside e4m3's own round-trip envelope, and [test_serve_kv_cache_dtype](../tests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp) — the flag through the REAL `VllmServerMain` | [fp8-kv-cache](specs/fp8-kv-cache.md) | `PARTIAL` | - | +| `QUANT-KV-FP8` | fp8, fp8_e4m3, fp8_e5m2 | `vllm/config/cache.py:19-25`; `vllm/model_executor/layers/quantization/kv_cache.py:42-191`; store `cache_kernels.cu:241-252`; scale convention `quant_utils.cuh:296-308` | **W1 CPU fp8-e4m3 store+read LANDED**: [codec](../include/vt/fp8_kv.h#L39), [store kernel](../src/vt/cpu/cpu_cache.cpp#L143), [read dequant](../src/vt/cpu/cpu_paged_attn.cpp#L82), [config parse](../include/vllm/v1/kv_cache_dtype.h#L37). **W2 CUDA fp8-e4m3 store+read LANDED** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)): [store kernel](../src/vt/cuda/cuda_cache.cu), [read dequant](../src/vt/cuda/cuda_paged_attn.cu) -- gate [test_cuda_fp8_kv_cache](../tests/vt/test_cuda_fp8_kv_cache.cpp), whose DEVICE cases are UNEXECUTED; the CUDA TUs COMPILE (CI `cuda-fat-build`, ten architectures, run 32495320287 on `4d71e776e`) but that job builds with tests OFF, so none has been executed (spec `## Owed`). **W3 runner integration LANDED** ([#1593](https://github.com/mudler/vllm.cpp/issues/1593)): [`--kv-cache-dtype`](../src/vllm/entrypoints/openai/server_main.cpp) reaches `EngineParams`, the checkpoint's `kv_cache_quant_algo` is honoured when no flag is typed, [`ApplyCacheDType`](../src/vllm/v1/kv_cache_interface.cpp) retypes every attention spec (group and heterogeneous per-layer alike) so one byte budget buys exactly 2x the blocks, and [`kv_cache_route.h`](../include/vllm/model_executor/models/kv_cache_route.h) is the ONE place the store and the read decide float versus fp8. An fp8 cache DISABLES FA-2 prefill and decode, the WMMA ladder and the vectorized decode kernels, which are bf16-native; the resulting throughput cost is unmeasured and recorded as such. e5m2 compute, per-head scales, the Metal/ROCm arms, the C-ABI field, the 16 unrouted architectures (15 of which refuse with a message that names neither fp8 nor the flag) and the `k_scale`/`v_scale` weight-loader read are named later bricks (see spec) | [test_ops_fp8_kv_cache](../tests/vt/test_ops_fp8_kv_cache.cpp#L1) — 8 cases / 511 assertions, round-trip within the e4m3 band + fp8-vs-bf16 NMSE<1% + paged-attention e2e; RED-first (wrong store direction fails 3/480); W3 [test_kv_cache_fp8_wiring](../tests/vllm/entrypoints/test_kv_cache_fp8_wiring.cpp) — 31 cases, G1-G12, entering through `LoadedEngine` rather than by building a spec by hand, and bounding the fp8 pages against a bf16 run inside e4m3's own round-trip envelope, and [test_serve_kv_cache_dtype](../tests/vllm/entrypoints/openai/test_serve_kv_cache_dtype.cpp) — the flag through the REAL `VllmServerMain`. **W6 ROCm fp8-e4m3 store+read LANDED** ([#2065](https://github.com/mudler/vllm.cpp/issues/2065)): [store kernel](../src/vt/rocm/rocm_dense_basic.hip), [read dequant](../src/vt/rocm/rocm_paged_attn.hip) -- gate [test_rocm_fp8_kv_cache](../tests/vt/test_rocm_fp8_kv_cache.cpp), 7/7 cases 28/28 assertions on gfx1100 | [fp8-kv-cache](specs/fp8-kv-cache.md) | `PARTIAL` | - | | `QUANT-KV-FP8-VENDOR` | fp8_inc, fp8_ds_mla | `vllm/config/cache.py:24-25`; vendor KV implementations selected by attention backend | - | no quantized KV cache | `planned: specs/vendor-fp8-kv-cache.md` | `INVENTORIED` | - | | `QUANT-KV-TURBO` | k8v4, 4bit_nc, k3v4_nc, 3bit_nc | `vllm/config/cache.py:28-33`; TurboQuant dependency path | - | no quantized KV cache | `planned: specs/turboquant-kv-cache.md` | `INVENTORIED` | - | | `QUANT-KV-PER-HEAD` | int4/int8/fp8 per-token-head | `vllm/config/cache.py:34`; quantized cache kernels selected by backend | - | no quantized KV cache | `planned: specs/per-head-kv-cache.md` | `INVENTORIED` | - | diff --git a/.agents/specs/fp8-kv-cache.md b/.agents/specs/fp8-kv-cache.md index 9c7042d328..2272fe9ca6 100644 --- a/.agents/specs/fp8-kv-cache.md +++ b/.agents/specs/fp8-kv-cache.md @@ -1,4 +1,4 @@ -# fp8 KV cache (`cache_dtype=fp8*`) — spike + W1 + W2 + W3 (`KV-FP8`, `QUANT-KV-FP8`) +# fp8 KV cache (`cache_dtype=fp8*`) — spike + W1 + W2 + W3 + W6 (`KV-FP8`, `QUANT-KV-FP8`) Rows: `KV-FP8` (engine-matrix, KV cache and memory) and `QUANT-KV-FP8` (quantization-matrix). HIGH-priority feature gap #5 @@ -28,7 +28,7 @@ re-port). `k_scale`/`v_scale` path with its declared-but-absent arm named rather than defaulted. - **Out (named later bricks):** fp8_e5m2 compute on either backend, - per-attention-head scales, the Metal and ROCm fp8-KV arms (both refuse by name + per-attention-head scales, the Metal fp8-KV arm (refuses by name — see `## W2` below), `--calculate-kv-scales` (upstream's deprecated dynamic scale), the C-ABI exposure of `--kv-cache-dtype`, the 16 architectures whose attention blocks W3 refuses rather than routes, and the vendor @@ -115,7 +115,7 @@ replaced by provider routing plus a named Metal/ROCm refusal), and `tests/vt/test_cuda_fp8_kv_cache.cpp` (NEW) + its `tests/CMakeLists.txt` line. Later bricks: the runner/spec integration (half-sized blocks + checkpoint scale -threading + CLI); fp8_e5m2 compute; per-head scales; the Metal and ROCm arms. +threading + CLI); fp8_e5m2 compute; per-head scales; the Metal arm. ## Tests to port @@ -172,6 +172,7 @@ vendor/turbo/nvfp4 KV dtypes are separate rows. | W3 | runner/spec integration: half-sized KV blocks + checkpoint k/v_scale threading + `--kv-cache-dtype` | DONE (code + CPU gate landed; see `## W3` and `## Owed`) | | W4 | memory-halving e2e on a gate model (the binding gate, DGX) | later | | W5 | fp8_e5m2 CPU+CUDA compute; per-attention-head scales | later | +| W6 | ROCm fp8-e4m3 store + fp8 paged-attention read (parity vs W1) | DONE (code + gate landed + MEASURED on gfx1100 — see `## Outcome (W6 ROCm arm)`) | ## W2 — the CUDA arm (#1593) @@ -998,3 +999,55 @@ pin `555967922`). - **Honest residual.** W1 is a correctness brick; the real *memory/throughput* win (the point of the feature) is the GPU store/read + the halved-block runner integration, both DGX-blocked and named W2-W4. + +## Outcome (W6 ROCm arm) + +W6 ports the fp8-e4m3 KV cache store and read to the ROCm backend, closing the +last non-Metal gap in the fp8 KV cache surface. The store kernel +(`ReshapeAndCacheFp8KernelRocm` in `src/vt/rocm/rocm_dense_basic.hip`) and the +read dequant (`LoadKv` in `src/vt/rocm/rocm_paged_attn.hip`) mirror the W1 CPU +codec and the W2 CUDA arm. `OpId::kReshapeAndCacheFp8` is registered for +`DeviceType::kROCM` in `src/vt/rocm/rocm_ops.hip`, and the fp8 read refusal in +`src/vt/ops.cpp` is widened from `kCPU || kCUDA` to `kCPU || kCUDA || kROCM`. +Metal remains refused by name. + +### Measured + +- Build: `make -j4 vllm` with `-Werror` succeeds (29 s, hipcc 7.2.4, gfx1100). +- Test build: `make -j4 test_rocm_fp8_kv_cache` succeeds (6 s). +- `test_rocm_fp8_kv_cache`: 7/7 cases, 28/28 assertions on gfx1100 (RX 7900 + XTX). G3 store byte-identical to the W1 CPU oracle; G4 f32 read NMSE < 1e-6; + G4b bf16 read NMSE < 1e-4; G5 e5m2 refused with the named message. +- `test_ops_fp8_kv_cache`: 8/8, 511 assertions (CPU regression, unaffected). +- `test_rocm_backend`: 9/9, 1065 assertions (existing ROCm suite unaffected). +- `test_ops_paged_attn`: 14/14, 1646 assertions (existing paged-attn suite + unaffected). + +### Rejected + +- `__nv_cvt_float_to_fp8` intrinsic: CUDA-only; not available in HIP. Rejected + in favor of a software codec (`F32ToF8E4M3Dev`/`StoreKvFp8E4M3Dev`) using + `frexpf`/`nearbyintf`/`ldexpf` arithmetic mirroring the CPU `vt::F32ToF8E4M3`. +- Hardware fp8 conversion intrinsics (`__builtin_amdgcn_cvt_f32_to_fp8` etc.): + available on CDNA2+ (gfx940/941/942) but not on gfx1100 (RDNA3). Rejected for + portability; the software codec is bit-identical and works on all ROCm + targets. + +### RED-first mutation proof + +Two mutations confirmed the tests detect the defects they claim to guard: + +1. **LoadKv dequant**: dropped `* scale` in `F8E4M3ToF32Dev(p[i]) * scale`. + G4/G4b failed with NMSE ~34000x, worst error ~222. Restored; all 7 green. +2. **Store kernel**: dropped `/ scale` in `F32ToF8E4M3Dev(hp / scale)`. + G3/G3b failed with byte-level mismatches in key and value cache. Restored; + all 7 green. + +### Defaults + +- The fp8 KV cache is opt-in via `--kv-cache-dtype fp8` / `fp8_e4m3`. Default + remains `auto` (bf16), so the default path is byte-identical. +- Per-tensor `k_scale`/`v_scale` are additive fields on `PagedAttentionArgs` + and default to `1.0f` when unused; no existing caller is affected. +- e5m2 is parsed by the config layer but refused by the ROCm kernel with a + named-later-brick message, matching the CPU and CUDA arms. diff --git a/include/vt/ops.h b/include/vt/ops.h index 2d6592cd88..39b61cfd8f 100644 --- a/include/vt/ops.h +++ b/include/vt/ops.h @@ -1121,17 +1121,18 @@ struct PagedAttentionArgs { // device read (companion to query_start_loc_host). 0 => that launcher falls // back to the D2H+sync. int32_t max_seq_len = 0; - // OPTIONAL fp8 KV-cache read (KV-FP8 W1 CPU, W2 CUDA). kAuto (default) => the cache holds - // the model float dtype and is read directly — every existing caller is - // byte-identical. When != kAuto the K/V cache pages are 1-byte fp8 (DType::kI8 - // storage) and each read is DEQUANTIZED as Dequant(fp8) * k_scale|v_scale - // before entering the f32 softmax, mirroring the fp8 attention read path - // (scaled_vec_conversion, quant_utils.cuh:302-308). k_scale / - // v_scale are the per-tensor scales from BaseKVCacheMethod (kv_cache.py:108-191) - // — 1.0 is the uncalibrated default. Per-head scales are a later brick. - // Implemented on CPU and CUDA. kMETAL/kROCM register kPagedAttention for the - // FLOAT path only, and because these fields are ADDITIVE the provider table - // cannot tell the two arms apart, so src/vt/ops.cpp refuses them by name. + // OPTIONAL fp8 KV-cache read (KV-FP8 W1 CPU, W2 CUDA, W6 ROCm). kAuto + // (default) => the cache holds the model float dtype and is read directly — + // every existing caller is byte-identical. When != kAuto the K/V cache pages + // are 1-byte fp8 (DType::kI8 storage) and each read is DEQUANTIZED as + // Dequant(fp8) * k_scale|v_scale before entering the f32 softmax, mirroring + // the fp8 attention read path (scaled_vec_conversion, + // quant_utils.cuh:302-308). k_scale / v_scale are the per-tensor scales from + // BaseKVCacheMethod (kv_cache.py:108-191) — 1.0 is the uncalibrated default. + // Per-head scales are a later brick. Implemented on CPU, CUDA, and ROCm. + // kMETAL registers kPagedAttention for the FLOAT path only, and because + // these fields are ADDITIVE the provider table cannot tell the two arms + // apart, so src/vt/ops.cpp refuses Metal by name. Fp8KVCacheDataType kv_cache_dtype = Fp8KVCacheDataType::kAuto; float k_scale = 1.0f; float v_scale = 1.0f; diff --git a/src/vt/ops.cpp b/src/vt/ops.cpp index 0961a08073..b2cb89d48a 100644 --- a/src/vt/ops.cpp +++ b/src/vt/ops.cpp @@ -3824,18 +3824,20 @@ void PagedAttention(Queue& q, Tensor& out, const Tensor& query, const Tensor& k_ VT_CHECK(args.k_scale > 0.0f && args.v_scale > 0.0f, "paged_attention: fp8 KV read requires k_scale/v_scale > 0"); // WHICH BACKENDS HAVE AN fp8 READ. Unlike the fp8 STORE — a separate OpId - // that only the CPU and CUDA backends register, so an unimplemented backend - // refuses by name inside GetOp — the fp8 read rides ADDITIVE fields on - // PagedAttentionArgs of an op that kMETAL and kROCM already register for the - // float path (metal_ops.mm, rocm_ops.hip). Nothing in the provider table can - // tell those two apart, so without this list an fp8 cache would reach a - // kernel that reads the same bytes as floats and returns silent garbage. - // AGENTS.md: refuse an unimplemented arm with a message that names the - // missing part. CPU landed in W1, CUDA in W2; Metal and ROCm are owed. - VT_CHECK(q.device.type == DeviceType::kCPU || q.device.type == DeviceType::kCUDA, - "paged_attention: the fp8 KV read is implemented on CPU (KV-FP8 W1) and " - "CUDA (KV-FP8 W2) only; this backend has no fp8 dequant on the cache read " - "and would read the fp8 bytes as its float dtype"); + // that only the CPU, CUDA, and ROCm backends register, so an unimplemented + // backend refuses by name inside GetOp — the fp8 read rides ADDITIVE fields + // on PagedAttentionArgs of an op that kMETAL and kROCM already register for + // the float path (metal_ops.mm, rocm_ops.hip). Nothing in the provider + // table can tell those two apart, so without this list an fp8 cache would + // reach a kernel that reads the same bytes as floats and returns silent + // garbage. AGENTS.md: refuse an unimplemented arm with a message that names + // the missing part. CPU landed in W1, CUDA in W2, ROCm in W6; Metal is owed. + VT_CHECK(q.device.type == DeviceType::kCPU || q.device.type == DeviceType::kCUDA || + q.device.type == DeviceType::kROCM, + "paged_attention: the fp8 KV read is implemented on CPU (KV-FP8 W1), " + "CUDA (KV-FP8 W2), and ROCm (KV-FP8 W6) only; this backend has no " + "fp8 dequant on the cache read and would read the fp8 bytes as its " + "float dtype"); } // metadata: block_table [num_reqs, max_blocks] i32, seq_lens [num_reqs] i32, // query_start_loc [num_reqs+1] i32. diff --git a/src/vt/rocm/rocm_dense_basic.hip b/src/vt/rocm/rocm_dense_basic.hip index 768ab89c9f..425646c0cd 100644 --- a/src/vt/rocm/rocm_dense_basic.hip +++ b/src/vt/rocm/rocm_dense_basic.hip @@ -1,5 +1,6 @@ // ROCm dense basic kernels — correctness-grade ports for discrete R9700 bring-up. #include +#include #include #include @@ -8,6 +9,7 @@ #include #include "vt/ops.h" +#include "vt/fp8_kv.h" namespace vt::rocm { namespace { @@ -33,6 +35,7 @@ __device__ inline float Ld(const float* p, int64_t i) { return p[i]; } __device__ inline float Ld(const __hip_bfloat16* p, int64_t i) { return __bfloat162float(p[i]); } +__device__ inline float Ld(const __half* p, int64_t i) { return __half2float(p[i]); } __device__ inline void St(float* p, int64_t i, float v) { p[i] = v; } __device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { p[i] = __float2bfloat16(v); @@ -720,6 +723,140 @@ void ReshapeAndCacheKernelRocm(Queue& q, const Tensor& k, const Tensor& v, Tenso Check(hipGetLastError(), "reshape_and_cache"); } +// ---------------- ReshapeAndCacheFp8 (fp8-e4m3 KV write, KV-FP8 W6) ------------ +// ROCm arm of vt::ReshapeAndCacheFp8, and the ROCm sibling of the CUDA kernel +// in src/vt/cuda/cuda_cache.cu:155-226 that is its ORACLE via the CPU reference. +// +// Ported from the fp8 branch of vllm reshape_and_cache_flash_kernel +// (csrc/libtorch_stable/cache_kernels.cu:314-401) + CopyWithScaleOp (:241-252) +// @ pin 555967922. Same scope as the CUDA arm: the is_contiguous_heads && +// kv_scale_stride == 0 fast path only (the vt cache is the NHD unbind slice +// and ReshapeAndCacheFp8 takes two scalar scales). +// +// ELEMENTWISE-IDENTICAL to the CUDA and CPU arms. The converter is the +// SOFTWARE codec vt::StoreKvFp8E4M3 (include/vt/fp8_kv.h:87-89), which is +// vt::F32ToF8E4M3(hp / scale) — bit-identical to the CPU codec by construction +// and to the CUDA __nv_cvt_float_to_fp8 intrinsic by measurement (spec W2, +// vt-fp8-quant-arch-gate.md G2). ROCm HIP has no __nv_cvt_float_to_fp8, so the +// software path is the store, not a fallback. The store is a true DIVIDE +// (hp / scale), not a reciprocal multiply. +// +// Source dtypes f32/bf16/f16, widened to f32 through Ld before the divide — +// the same helper the float path uses. __hip_bfloat16 and __half both carry +// operator float() so the existing Ld overloads serve without new conversion +// functions. Same set the CPU LoadSrcF32 and CUDA Fp8SrcToF32 serve. +// f32 → fp8-e4m3fn byte on device, round-to-nearest-even, SATURATING to +/-448. +// ROCm HIP has no __nv_cvt_float_to_fp8 intrinsic, so the store uses the +// software codec — the SAME arithmetic as vt::F32ToF8E4M3 +// (include/vt/fp8_kv.h:53-82), which is bit-identical to the CPU codec by +// construction and to the CUDA intrinsic by measurement (spec W2, +// vt-fp8-quant-arch-gate.md G2). ldexpf/frexpf/nearbyintf are device-callable. +__device__ __forceinline__ uint8_t F32ToF8E4M3Dev(float f) { + constexpr float kFp8Max = 448.0f; + if (isnan(f)) return 0x7FU; + const uint8_t sign = signbit(f) ? 0x80U : 0x00U; + const float a = fabsf(f); + if (!isfinite(a) || a >= kFp8Max) return static_cast(sign | 0x7EU); + if (a == 0.0f) return sign; + int e2 = 0; + const float frac = frexpf(a, &e2); + int exp_field = (e2 - 1) + 7; + if (exp_field <= 0) { + const float qd = a * 512.0f; + const int qi = static_cast(nearbyintf(qd)); + if (qi <= 0) return sign; + if (qi < 8) return static_cast(sign | static_cast(qi)); + return static_cast(sign | (1U << 3)); + } + const float sig = frac * 2.0f; + int mi = static_cast(nearbyintf(sig * 8.0f)); + if (mi == 16) { + mi = 8; + exp_field += 1; + } + const int mant = mi - 8; + if (exp_field > 15 || (exp_field == 15 && mant >= 7)) { + return static_cast(sign | 0x7EU); + } + return static_cast(sign | (static_cast(exp_field) << 3) | + static_cast(mant)); +} + +// fp8 byte = Quantize(hp / scale). Mirrors the CUDA StoreKvFp8E4M3Dev +// (cuda_cache.cu:151-153) and the CPU vt::StoreKvFp8E4M3 (fp8_kv.h:87-89). +__device__ __forceinline__ uint8_t StoreKvFp8E4M3Dev(float hp, float scale) { + return F32ToF8E4M3Dev(hp / scale); +} + +template +__global__ void ReshapeAndCacheFp8K(const Tin* key, const Tin* value, + uint8_t* key_cache, uint8_t* value_cache, + const int64_t* slot_mapping, int64_t block_size, + int64_t n_elems, int64_t k_block_stride, + int64_t k_page_stride, int64_t v_block_stride, + int64_t v_page_stride, int64_t k_tok_stride, + int64_t v_tok_stride, float k_scale, + float v_scale) { + const int64_t token = blockIdx.x; + const int64_t slot = slot_mapping[token]; + if (slot < 0) return; // padded token → skip (upstream :328-331) + const int64_t block = slot / block_size; + const int64_t offset = slot % block_size; + const int64_t kdst = block * k_block_stride + offset * k_page_stride; + const int64_t vdst = block * v_block_stride + offset * v_page_stride; + const int64_t ksrc = token * k_tok_stride; + const int64_t vsrc = token * v_tok_stride; + for (int64_t e = threadIdx.x; e < n_elems; e += blockDim.x) { + key_cache[kdst + e] = StoreKvFp8E4M3Dev(Ld(key, ksrc + e), k_scale); + value_cache[vdst + e] = StoreKvFp8E4M3Dev(Ld(value, vsrc + e), v_scale); + } +} + +void ReshapeAndCacheFp8KernelRocm(Queue& q, const Tensor& k, const Tensor& v, + Tensor& k_cache, Tensor& v_cache, + const Tensor& slot_mapping, + Fp8KVCacheDataType kind, float k_scale, + float v_scale) { + VT_CHECK(kind == Fp8KVCacheDataType::kFp8E4M3, + "rocm reshape_and_cache_fp8: only fp8_e4m3 is implemented " + "(fp8_e5m2 is a named later brick, spec W5)"); + const int64_t num_slots = slot_mapping.shape[0]; + const int64_t block_size = k_cache.shape[1]; + const int64_t n_elems = k_cache.shape[2] * k_cache.shape[3]; + if (num_slots == 0 || n_elems == 0) return; + hipStream_t st = AsStream(q); + const unsigned grid = static_cast(num_slots); + const unsigned block = static_cast(n_elems < 512 ? n_elems : 512); + const int64_t* slots = slot_mapping.Ptr(); + uint8_t* kc = k_cache.Ptr(); + uint8_t* vc = v_cache.Ptr(); + switch (k.dtype) { + case DType::kF32: + ReshapeAndCacheFp8K<<>>( + k.Ptr(), v.Ptr(), kc, vc, slots, block_size, n_elems, + k_cache.stride[0], k_cache.stride[1], v_cache.stride[0], + v_cache.stride[1], k.stride[0], v.stride[0], k_scale, v_scale); + break; + case DType::kBF16: + ReshapeAndCacheFp8K<__hip_bfloat16><<>>( + k.Ptr<__hip_bfloat16>(), v.Ptr<__hip_bfloat16>(), kc, vc, slots, + block_size, n_elems, k_cache.stride[0], k_cache.stride[1], + v_cache.stride[0], v_cache.stride[1], k.stride[0], v.stride[0], + k_scale, v_scale); + break; + case DType::kF16: + ReshapeAndCacheFp8K<__half><<>>( + k.Ptr<__half>(), v.Ptr<__half>(), kc, vc, slots, block_size, n_elems, + k_cache.stride[0], k_cache.stride[1], v_cache.stride[0], + v_cache.stride[1], k.stride[0], v.stride[0], k_scale, v_scale); + break; + default: + VT_CHECK(false, + "rocm reshape_and_cache_fp8: unsupported source dtype (f32/f16/bf16)"); + } + Check(hipGetLastError(), "reshape_and_cache_fp8"); +} + // ---------------- GeluTanh / GeluErf (elementwise) ------------------------- template __global__ void GeluTanhK(T* out, const T* x, int64_t n) { diff --git a/src/vt/rocm/rocm_ops.hip b/src/vt/rocm/rocm_ops.hip index 3868b8e838..e76db6a2af 100644 --- a/src/vt/rocm/rocm_ops.hip +++ b/src/vt/rocm/rocm_ops.hip @@ -30,6 +30,11 @@ void RopeFromCacheKernelRocm(Queue& q, Tensor& qs, Tensor* ks, const Tensor& pos const Tensor& cache, const RopeArgs& args); void ReshapeAndCacheKernelRocm(Queue& q, const Tensor& k, const Tensor& v, Tensor& k_cache, Tensor& v_cache, const Tensor& slot_mapping); +void ReshapeAndCacheFp8KernelRocm(Queue& q, const Tensor& k, const Tensor& v, + Tensor& k_cache, Tensor& v_cache, + const Tensor& slot_mapping, + Fp8KVCacheDataType kind, float k_scale, + float v_scale); void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const Tensor& k_cache, const Tensor& v_cache, const Tensor& block_table, const Tensor& seq_lens, const Tensor& query_start_loc, @@ -156,6 +161,9 @@ struct Registrar { RegisterOp(OpId::kReshapeAndCache, DeviceType::kROCM, reinterpret_cast( static_cast(&ReshapeAndCacheKernelRocm))); + RegisterOp(OpId::kReshapeAndCacheFp8, DeviceType::kROCM, + reinterpret_cast(static_cast( + &ReshapeAndCacheFp8KernelRocm))); RegisterOp(OpId::kPagedAttention, DeviceType::kROCM, reinterpret_cast( static_cast(&PagedAttentionKernelRocm))); diff --git a/src/vt/rocm/rocm_paged_attn.hip b/src/vt/rocm/rocm_paged_attn.hip index b4973fcd7f..6cd77bd100 100644 --- a/src/vt/rocm/rocm_paged_attn.hip +++ b/src/vt/rocm/rocm_paged_attn.hip @@ -16,6 +16,7 @@ #include "vt/ops.h" #include "vt/backend.h" +#include "vt/fp8_kv.h" #include "vt/rocm/rocm_arch.h" #include @@ -145,6 +146,37 @@ __device__ inline void St(__hip_bfloat16* p, int64_t i, float v) { p[i] = __float2bfloat16(v); } +// fp8-e4m3 byte → f32 on device. Same arithmetic as vt::F8E4M3ToF32 +// (include/vt/fp8_kv.h:40-49) and the CUDA Fp8E4M3ToF32Dev +// (cuda_paged_attn.cu:164), so ROCm==CPU==CUDA on the read is a property of +// the source. Bit-identical to the F8E4M3ToF32 in rocm_fp8_channel_gemv.hip:22. +__device__ inline float F8E4M3ToF32Dev(uint8_t byte) { + const uint32_t sign = static_cast(byte >> 7) & 0x1U; + const uint32_t exp = static_cast(byte >> 3) & 0xFU; + const uint32_t mant = static_cast(byte) & 0x7U; + const float sm = sign ? -1.0f : 1.0f; + if (exp == 0xFU && mant == 0x7U) return 0.f; // NaN → 0 (no NaN in finite KV) + if (exp == 0U) return sm * (static_cast(mant) * (1.0f / 512.0f)); + const float mantissa = 1.0f + static_cast(mant) * (1.0f / 8.0f); + return sm * ldexpf(mantissa, static_cast(exp) - 7); +} + +// One K/V-cache element as f32, with the fp8 dequant folded in. INERT on the +// float arms (forward to Ld unchanged), so every existing bf16/f32 caller +// reads the same bytes in the same order. On uint8_t: Dequant(fp8) * scale. +// Mirrors CUDA LoadKv (cuda_paged_attn.cu:175-185). +__device__ inline float LoadKv(const float* p, int64_t i, float scale) { + (void)scale; + return Ld(p, i); +} +__device__ inline float LoadKv(const __hip_bfloat16* p, int64_t i, float scale) { + (void)scale; + return Ld(p, i); +} +__device__ inline float LoadKv(const uint8_t* p, int64_t i, float scale) { + return F8E4M3ToF32Dev(p[i]) * scale; +} + __device__ inline float Softcap(float s, float cap) { return cap > 0.f ? cap * tanhf(s / cap) : s; } @@ -165,7 +197,8 @@ __global__ void PagedAttnOnline(TO* out, const TQ* query, const TKV* k_cache, int64_t block_size, int64_t bt_row, int64_t bt_col, int64_t kc_blk, int64_t kc_pg, int64_t kc_hd, int64_t vc_blk, int64_t vc_pg, int64_t vc_hd, float scale, float softcap, - bool causal, int window_left, int window_right) { + bool causal, int window_left, int window_right, + float k_scale, float v_scale) { const int64_t t = blockIdx.x; const int64_t h = blockIdx.y; if (h >= hq) return; @@ -231,14 +264,14 @@ __global__ void PagedAttnOnline(TO* out, const TQ* query, const TKV* k_cache, const int64_t nvec = d >> 2; for (int64_t v = threadIdx.x; v < nvec; v += blockDim.x) { const int64_t e = v << 2; - part += Ld(query, qoff + e) * Ld(k_cache, kbase + e); - part += Ld(query, qoff + e + 1) * Ld(k_cache, kbase + e + 1); - part += Ld(query, qoff + e + 2) * Ld(k_cache, kbase + e + 2); - part += Ld(query, qoff + e + 3) * Ld(k_cache, kbase + e + 3); + part += Ld(query, qoff + e) * LoadKv(k_cache, kbase + e, k_scale); + part += Ld(query, qoff + e + 1) * LoadKv(k_cache, kbase + e + 1, k_scale); + part += Ld(query, qoff + e + 2) * LoadKv(k_cache, kbase + e + 2, k_scale); + part += Ld(query, qoff + e + 3) * LoadKv(k_cache, kbase + e + 3, k_scale); } } else { for (int64_t e = threadIdx.x; e < d; e += blockDim.x) - part += Ld(query, qoff + e) * Ld(k_cache, kbase + e); + part += Ld(query, qoff + e) * LoadKv(k_cache, kbase + e, k_scale); } red[threadIdx.x] = part; __syncthreads(); @@ -258,14 +291,14 @@ __global__ void PagedAttnOnline(TO* out, const TQ* query, const TKV* k_cache, const int64_t nvec = d >> 2; for (int64_t v = threadIdx.x; v < nvec; v += blockDim.x) { const int64_t e = v << 2; - acc[e] = acc[e] * corr + pw * Ld(v_cache, vbase + e); - acc[e + 1] = acc[e + 1] * corr + pw * Ld(v_cache, vbase + e + 1); - acc[e + 2] = acc[e + 2] * corr + pw * Ld(v_cache, vbase + e + 2); - acc[e + 3] = acc[e + 3] * corr + pw * Ld(v_cache, vbase + e + 3); + acc[e] = acc[e] * corr + pw * LoadKv(v_cache, vbase + e, v_scale); + acc[e + 1] = acc[e + 1] * corr + pw * LoadKv(v_cache, vbase + e + 1, v_scale); + acc[e + 2] = acc[e + 2] * corr + pw * LoadKv(v_cache, vbase + e + 2, v_scale); + acc[e + 3] = acc[e + 3] * corr + pw * LoadKv(v_cache, vbase + e + 3, v_scale); } } else { for (int64_t e = threadIdx.x; e < d; e += blockDim.x) - acc[e] = acc[e] * corr + pw * Ld(v_cache, vbase + e); + acc[e] = acc[e] * corr + pw * LoadKv(v_cache, vbase + e, v_scale); } __syncthreads(); if (threadIdx.x == 0) { @@ -1688,8 +1721,6 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const const Tensor& v_cache, const Tensor& block_table, const Tensor& seq_lens, const Tensor& query_start_loc, const PagedAttentionArgs& args) { - VT_CHECK(args.kv_cache_dtype == Fp8KVCacheDataType::kAuto, - "rocm paged_attention: fp8 KV not implemented"); VT_CHECK(args.scale > 0.f, "rocm paged_attention: scale must be > 0"); if (CpuRefEnabled()) { PagedAttnCpuRef(q, out, query, k_cache, v_cache, block_table, seq_lens, query_start_loc, @@ -1989,11 +2020,29 @@ void PagedAttentionKernelRocm(Queue& q, Tensor& out, const Tensor& query, const block_table.Ptr(), seq_lens.Ptr(), query_start_loc.Ptr(), num_reqs, hq, num_kv_heads, d, block_size, bt_row, bt_col, kc_blk, kc_pg, kc_hd, vc_blk, vc_pg, vc_hd, args.scale, args.logits_soft_cap, args.causal, window_left, - window_right); + window_right, args.k_scale, args.v_scale); }; - if (query.dtype == DType::kBF16 && k_cache.dtype == DType::kBF16 && - out.dtype == DType::kBF16) { + // fp8 KV-cache read: cache pages are uint8_t (DType::kI8), dequantized + // inside LoadKv. Only the correctness-grade PagedAttnOnline kernel serves + // fp8 — the decode-opt bf16 path stages __hip_bfloat16 fragments and a + // tensor-core fp8 read is a performance brick, not this one (same scope + // as the CUDA arm, spec W2). + const bool fp8 = args.kv_cache_dtype != Fp8KVCacheDataType::kAuto; + if (fp8) { + VT_CHECK(k_cache.dtype == DType::kI8 && v_cache.dtype == DType::kI8, + "rocm paged_attention fp8: cache must be kI8"); + if (query.dtype == DType::kBF16 && out.dtype == DType::kBF16) { + launch(__hip_bfloat16{}, uint8_t{}, __hip_bfloat16{}); + } else if (query.dtype == DType::kF32 && out.dtype == DType::kF32) { + launch(float{}, uint8_t{}, float{}); + } else if (query.dtype == DType::kBF16 && out.dtype == DType::kF32) { + launch(__hip_bfloat16{}, uint8_t{}, float{}); + } else { + VT_CHECK(false, "rocm paged_attention fp8: unsupported query/output dtype"); + } + } else if (query.dtype == DType::kBF16 && k_cache.dtype == DType::kBF16 && + out.dtype == DType::kBF16) { launch(__hip_bfloat16{}, __hip_bfloat16{}, __hip_bfloat16{}); } else if (query.dtype == DType::kF32 && k_cache.dtype == DType::kF32 && out.dtype == DType::kF32) { diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 4ae6917e0b..1a86ac697c 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1892,6 +1892,7 @@ if(VLLM_CPP_HIP) # Plain C++ (no HIP header): every assertion goes through the vt:: seam. Each # case no-ops when the build has HIP but the box has no AMD GPU. vllm_cpp_add_test(test_rocm_backend vt/test_rocm_backend.cpp) +vllm_cpp_add_test(test_rocm_fp8_kv_cache vt/test_rocm_fp8_kv_cache.cpp) # #785 P1 GPU product-seam witness. Executable only — NOT add_test. # Ordinary CTest must not see this target. Runner fail-closes on 77/nonzero. add_executable(test_ops_paged_attn_sharedk_wmma_p1_gpu diff --git a/tests/vt/test_rocm_fp8_kv_cache.cpp b/tests/vt/test_rocm_fp8_kv_cache.cpp new file mode 100644 index 0000000000..9665fa769e --- /dev/null +++ b/tests/vt/test_rocm_fp8_kv_cache.cpp @@ -0,0 +1,747 @@ +// ROCm fp8 KV-cache store + paged-attention read gate (KV-FP8 W6, #2065). +// +// W1 landed the CPU half: vt::ReshapeAndCacheFp8 (fp8-e4m3 store), the fp8 read +// dequant in CPU paged attention, and vllm::v1::ParseCacheDtype. W1 IS THE +// ORACLE FOR W6 — the ROCm arm is measured against it, never against a fresh +// reference — so this file only ever compares ROCm to the landed CPU kernels. +// The CUDA arm (W2) is the direct template; the ROCm arm is elementwise- +// identical to it, and the CUDA arm is itself elementwise-identical to the CPU +// reference. +// +// Upstream mirror @ pin 555967922: +// store vllm/csrc/libtorch_stable/cache_kernels.cu:314-401 +// (reshape_and_cache_flash_kernel, fp8 branch) + CopyWithScaleOp :241-252 +// read vllm/csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh:419-429 +// (scaled_vec_conversion) +// scale quant_utils.cuh:296-300 — FP8 = Quantize(HP / scale); +// Dequant(FP8) * scale = HP +// scales vllm/model_executor/layers/quantization/kv_cache.py:108-191 +// (BaseKVCacheMethod: per-TENSOR k_scale/v_scale, 1.0 uncalibrated) +// +// The gates, and they do not all run in the same build: +// +// G1 (runs in every build WITHOUT the ROCm backend, i.e. the x86 CI leg): the +// fp8 store/read resolves through the provider table on a non-CPU device +// with no "later brick" guard. Compiled only where the ROCm backend is +// absent: in a ROCm build the op IS registered, so these calls would +// dispatch a real kernel over host pointers. +// G1b (every build): the fp8 READ is refused by name on kMETAL (the only +// backend with no fp8 dequant now that ROCm is implemented). The check +// fires in the op wrapper, so no Metal backend need be linked. +// G2 (ROCm build): the ROCm providers are REGISTERED for the fp8 store and the +// paged read — the shared-seam reach check. +// G3 (ROCm device): STORE parity — the ROCm store writes the SAME BYTES as the +// CPU store, zero tolerance, over the f32, bf16 and f16 sources the wrapper +// admits, with a padded (-1) slot and a strided unbind-slice cache. +// G4 (ROCm device): READ parity — paged attention over identical fp8 cache +// bytes, ROCm vs CPU, in both the decode and the prefill shape, for an f32 +// query/output. +// G4b (ROCm device): ...and for the bf16 query/output a served model actually +// runs, which is a DIFFERENT template instantiation of the same launcher. +// G5 (ROCm device): fp8_e5m2 stays refused BY THE ROCM KERNEL, reached through +// the registered provider. +// +// G3/G4/G4b/G5 SKIP CLEANLY when no ROCm backend is present, which is the house +// pattern (tests/vt/test_cuda_quant_dot.cpp:80-88). A skip is NOT a pass: every +// skipping case prints a MESSAGE naming what did not run. +#include + +#include +#include +#include +#include +#include +#include + +#include "vt/backend.h" +#include "vt/device.h" +#include "vt/dtype.h" +#include "vt/fp8_kv.h" +#include "vt/op_provider.h" +#include "vt/ops.h" +#include "vt/tensor.h" + +using vt::Backend; +using vt::Device; +using vt::DeviceType; +using vt::DType; +using vt::Fp8KVCacheDataType; +using vt::OpId; +using vt::PagedAttentionArgs; +using vt::Queue; +using vt::Tensor; + +namespace { + +bool HasRocm() { + try { + vt::GetBackend(DeviceType::kROCM); + return true; + } catch (const std::runtime_error&) { + return false; + } +} + +Device Cpu() { return Device{DeviceType::kCPU, 0}; } +Device Gpu() { return Device{DeviceType::kROCM, 0}; } + +// Tensor::Contiguous takes an initializer_list; these take the runtime shapes +// the cases build. Same packed-stride result. +Tensor Contig(void* data, DType dt, Device dev, const std::vector& shape) { + Tensor t; + t.data = data; + t.dtype = dt; + t.device = dev; + t.rank = static_cast(shape.size()); + int64_t stride = 1; + for (int i = t.rank - 1; i >= 0; --i) { + t.shape[i] = shape[static_cast(i)]; + t.stride[i] = stride; + stride *= shape[static_cast(i)]; + } + return t; +} + +Tensor Host(void* data, DType dt, const std::vector& shape) { + return Contig(data, dt, Cpu(), shape); +} + +Tensor Dev(void* data, DType dt, const std::vector& shape) { + return Contig(data, dt, Gpu(), shape); +} + +std::vector RandF32(size_t n, uint32_t seed) { + std::vector v(n); + uint32_t s = seed; + for (auto& x : v) { + s = s * 1664525u + 1013904223u; + x = (static_cast(s >> 8) / static_cast(1u << 24)) * 4.0f - 2.0f; + } + return v; +} + +} // namespace + +// ─── G1 ───────────────────────────────────────────────────────────────────── +// The fp8 store/read resolves through the provider table on a non-CPU device +// with no "later brick" guard. Compiled only where the ROCm backend is absent: +// in a ROCm build the op IS registered, so these calls would dispatch a real +// kernel over host pointers. The ROCm build asserts the same property from the +// other side, in G2. +#ifndef VLLM_CPP_HIP +TEST_CASE("fp8 KV ops resolve through the provider table on a non-CPU device") { + const int64_t nb = 1, bs = 4, H = 1, D = 16, page = H * D; + std::vector k(static_cast(page), 1.0f), v(static_cast(page), 1.0f); + std::vector kc(static_cast(nb * bs * page), 0); + std::vector vc(static_cast(nb * bs * page), 0); + std::vector slots = {0}; + Tensor tk = Dev(k.data(), DType::kF32, {1, H, D}); + Tensor tv = Dev(v.data(), DType::kF32, {1, H, D}); + Tensor tkc = Dev(kc.data(), DType::kI8, {nb, bs, H, D}); + Tensor tvc = Dev(vc.data(), DType::kI8, {nb, bs, H, D}); + Tensor ts = Dev(slots.data(), DType::kI64, {1}); + Queue qq{Gpu(), nullptr}; + + std::string store_msg; + try { + vt::ReshapeAndCacheFp8(qq, tk, tv, tkc, tvc, ts, Fp8KVCacheDataType::kFp8E4M3, 0.01f, 0.01f); + FAIL("reshape_and_cache_fp8 must refuse when no ROCm provider is linked in"); + } catch (const std::runtime_error& e) { + store_msg = e.what(); + } + CAPTURE(store_msg); + // The refusal must come from the PROVIDER TABLE, naming the op... + CHECK(store_msg.find("no kernel for op ReshapeAndCacheFp8") != std::string::npos); + // ...and NOT from a device-class guard inside the wrapper. + CHECK(store_msg.find("later brick") == std::string::npos); + + // Same for the read side: PagedAttention's fp8 arm must not carry a CPU-only + // guard either. One request, one decode token, one 16-wide head. + std::vector q(static_cast(D), 0.5f), out(static_cast(D), 0.0f); + std::vector bt = {0}, seq = {1}, qsl = {0, 1}; + Tensor tq = Dev(q.data(), DType::kF32, {1, 1, D}); + Tensor to = Dev(out.data(), DType::kF32, {1, 1, D}); + Tensor tbt = Dev(bt.data(), DType::kI32, {1, 1}); + Tensor tseq = Dev(seq.data(), DType::kI32, {1}); + Tensor tqsl = Dev(qsl.data(), DType::kI32, {2}); + PagedAttentionArgs args; + args.scale = 0.25f; + args.kv_cache_dtype = Fp8KVCacheDataType::kFp8E4M3; + args.k_scale = 0.01f; + args.v_scale = 0.01f; + + std::string read_msg; + try { + vt::PagedAttention(qq, to, tq, tkc, tvc, tbt, tseq, tqsl, args); + FAIL("paged_attention fp8 read must refuse when no ROCm provider is linked in"); + } catch (const std::runtime_error& e) { + read_msg = e.what(); + } + CAPTURE(read_msg); + CHECK(read_msg.find("no kernel for op PagedAttention") != std::string::npos); + CHECK(read_msg.find("later brick") == std::string::npos); +} +#endif // !VLLM_CPP_HIP + +// ─── G1b ──────────────────────────────────────────────────────────────────── +// The fp8 READ rides ADDITIVE fields on PagedAttentionArgs of an op that kMETAL +// already registers for the FLOAT path (metal_ops.mm). The provider table +// cannot tell the two arms apart, so an fp8 cache reaching that kernel would be +// read as that backend's float dtype and return silent garbage. AGENTS.md +// requires an unimplemented arm to refuse with a message that NAMES the missing +// part. ROCm landed in W6, so only Metal is refused now. +// +// Runs in every build: the check fires in the op wrapper, before any device or +// provider is touched, so no Metal backend needs to be linked in. +TEST_CASE("the fp8 KV read is refused on a backend with no fp8 dequant") { + const int64_t nb = 1, bs = 4, H = 1, D = 16, page = H * D; + std::vector kc(static_cast(nb * bs * page), 0); + std::vector vc(static_cast(nb * bs * page), 0); + std::vector q(static_cast(D), 0.5f), out(static_cast(D), 0.0f); + std::vector bt = {0}, seq = {1}, qsl = {0, 1}; + PagedAttentionArgs args; + args.scale = 0.25f; + args.kv_cache_dtype = Fp8KVCacheDataType::kFp8E4M3; + args.k_scale = 0.01f; + args.v_scale = 0.01f; + + const Device dev{DeviceType::kMETAL, 0}; + Tensor tq = Contig(q.data(), DType::kF32, dev, {1, 1, D}); + Tensor to = Contig(out.data(), DType::kF32, dev, {1, 1, D}); + Tensor tkc = Contig(kc.data(), DType::kI8, dev, {nb, bs, H, D}); + Tensor tvc = Contig(vc.data(), DType::kI8, dev, {nb, bs, H, D}); + Tensor tbt = Contig(bt.data(), DType::kI32, dev, {1, 1}); + Tensor tseq = Contig(seq.data(), DType::kI32, dev, {1}); + Tensor tqsl = Contig(qsl.data(), DType::kI32, dev, {2}); + Queue qq{dev, nullptr}; + std::string msg; + try { + vt::PagedAttention(qq, to, tq, tkc, tvc, tbt, tseq, tqsl, args); + FAIL("paged_attention must refuse the fp8 KV read on a backend without one"); + } catch (const std::runtime_error& e) { + msg = e.what(); + } + CAPTURE(msg); + CHECK(msg.find("fp8 KV read") != std::string::npos); + // The message must say WHAT would go wrong, not merely that it is refused. + CHECK(msg.find("no fp8 dequant") != std::string::npos); +} + +// ─── G2 ───────────────────────────────────────────────────────────────────── +// Reach through the shared seam. vt::ReshapeAndCacheFp8 and vt::PagedAttention +// dispatch through GetOp(OpId, DeviceType) (src/vt/ops.cpp), so a provider +// registered for kROCM IS the production path — nothing else selects a kernel. +// Registration is a static-init table fill, so this holds without a device: it +// asks "was the ROCm arm compiled and registered", which is exactly the question +// a `#ifdef`-elided kernel silently answers "no" to. +#ifdef VLLM_CPP_HIP +TEST_CASE("the ROCm fp8 KV store and paged read are registered providers") { + CHECK(vt::GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kROCM) != nullptr); + CHECK(vt::GetOp(OpId::kPagedAttention, DeviceType::kROCM) != nullptr); +} +#endif // VLLM_CPP_HIP + +// ─── G3 ───────────────────────────────────────────────────────────────────── +// STORE parity, byte for byte, zero tolerance. The CPU kernel is the oracle. +// +// The two arms are the same arithmetic by construction: the CPU codec is +// vt::F32ToF8E4M3 (include/vt/fp8_kv.h — software round-to-nearest-even, +// saturating at +/-448) and the ROCm kernel uses the SAME software codec +// vt::StoreKvFp8E4M3 (include/vt/fp8_kv.h:87-89), because ROCm HIP has no +// __nv_cvt_float_to_fp8 intrinsic. The CUDA intrinsic's equality to the CPU +// codec is already MEASURED at zero tolerance (spec W2, +// vt-fp8-quant-arch-gate.md G2); the ROCm arm uses the CPU codec directly, so +// ROCm==CPU on the store is a property of the source rather than of a +// measurement. This case re-takes it on the KV path, where the scale is applied +// as a true DIVIDE rather than the activation path's reciprocal multiply. +TEST_CASE("rocm fp8 KV store is byte-identical to the CPU store") { + if (!HasRocm()) { + MESSAGE("SKIPPED: no ROCm backend in this build/host — the ROCm fp8 KV store " + "parity gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + Queue cq{Cpu(), nullptr}; + + const int64_t nb = 1, bs = 4, H = 1, D = 16, page = H * D, nt = 4; + const size_t cache_elems = static_cast(nb * bs * page); + auto kf = RandF32(static_cast(nt * page), 11); + auto vf = RandF32(static_cast(nt * page), 22); + std::vector slots = {3, 0, 2, -1}; // -1 = padded token → skip + const float k_scale = 0.007f, v_scale = 0.003f; + + // CPU oracle: store the same tokens through the CPU kernel. + std::vector kc_ref(cache_elems, 0xAB), vc_ref(cache_elems, 0xCD); + Tensor ck = Host(kf.data(), DType::kF32, {nt, H, D}); + Tensor cv = Host(vf.data(), DType::kF32, {nt, H, D}); + Tensor ckc = Host(kc_ref.data(), DType::kI8, {nb, bs, H, D}); + Tensor cvc = Host(vc_ref.data(), DType::kI8, {nb, bs, H, D}); + Tensor cs = Host(slots.data(), DType::kI64, {nt}); + vt::ReshapeAndCacheFp8(cq, ck, cv, ckc, cvc, cs, Fp8KVCacheDataType::kFp8E4M3, k_scale, + v_scale); + + void* dk = gpu.Alloc(kf.size() * sizeof(float)); + void* dv = gpu.Alloc(vf.size() * sizeof(float)); + void* dkc = gpu.Alloc(cache_elems); + void* dvc = gpu.Alloc(cache_elems); + void* ds = gpu.Alloc(slots.size() * sizeof(int64_t)); + std::vector kc_seed(cache_elems, 0xAB); + std::vector vc_seed(cache_elems, 0xCD); + gpu.Copy(gq, dk, kf.data(), kf.size() * sizeof(float)); + gpu.Copy(gq, dv, vf.data(), vf.size() * sizeof(float)); + gpu.Copy(gq, dkc, kc_seed.data(), cache_elems); + gpu.Copy(gq, dvc, vc_seed.data(), cache_elems); + gpu.Copy(gq, ds, slots.data(), slots.size() * sizeof(int64_t)); + Tensor gk = Dev(dk, DType::kF32, {nt, H, D}); + Tensor gv = Dev(dv, DType::kF32, {nt, H, D}); + Tensor gkc = Dev(dkc, DType::kI8, {nb, bs, H, D}); + Tensor gvc = Dev(dvc, DType::kI8, {nb, bs, H, D}); + Tensor gs = Dev(ds, DType::kI64, {nt}); + vt::ReshapeAndCacheFp8(gq, gk, gv, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E4M3, k_scale, + v_scale); + + std::vector kc_got(cache_elems, 0); + std::vector vc_got(cache_elems, 0); + gpu.Copy(gq, kc_got.data(), dkc, cache_elems); + gpu.Copy(gq, vc_got.data(), dvc, cache_elems); + gpu.Synchronize(gq); + + int64_t kbad = 0, vbad = 0; + for (size_t i = 0; i < cache_elems; ++i) { + if (kc_got[i] != kc_ref[i]) ++kbad; + if (vc_got[i] != vc_ref[i]) ++vbad; + } + CHECK(kbad == 0); + CHECK(vbad == 0); + // Two kernels that both returned early would leave the seed fill on both + // sides and compare equal, so require that the ORACLE wrote something. This + // is asked of the CPU bytes, not the ROCm ones: a quantized byte may + // legitimately equal the 0xAB fill, and counting ROCm's differences would + // then be an assertion about the fixture rather than about the kernel. + int64_t ref_written = 0; + for (size_t i = 0; i < cache_elems; ++i) { + if (kc_ref[i] != 0xAB) ++ref_written; + } + CHECK(ref_written > 0); + + gpu.Free(dk); + gpu.Free(dv); + gpu.Free(dkc); + gpu.Free(dvc); + gpu.Free(ds); + gpu.DestroyQueue(gq); +} + +// The two NARROW source arms of the same store, and both of them matter. +// +// bf16 is the dtype vLLM actually resolves for a model (AGENTS.md "Inherit +// vLLM defaults"), so it is the arm production runs. f16 is the arm nothing +// else covers: `vt::ReshapeAndCacheFp8`'s wrapper admits any `IsFloat()` +// source (src/vt/ops.cpp), the CPU `LoadSrcF32` serves f16 +// (src/vt/cpu/cpu_cache.cpp), and the ROCm `ReshapeAndCacheFp8KernelRocm` has +// a `DType::kF16 -> __half` arm — which, without this case, no gate would ever +// instantiate on a device. An untested dispatch arm is the shape a wrong +// `Ptr<>` cast hides in. +// +// Both are widened to f32 BEFORE the divide on each side — upstream does the +// same (`quant_utils.cuh:482-489`, `__bfloat162float(a) / scale`), the ROCm +// kernel through `Ld` and the CPU through `LoadSrcF32` — and bf16->f32 and +// f16->f32 are both exact, so the two arms must still agree byte for byte. +TEST_CASE("rocm fp8 KV store is byte-identical to the CPU store (bf16 and f16 sources)") { + if (!HasRocm()) { + MESSAGE("SKIPPED: no ROCm backend in this build/host — the bf16/f16-source " + "fp8 KV store parity gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + Queue cq{Cpu(), nullptr}; + + const int64_t nb = 1, bs = 4, H = 1, D = 16, page = H * D, nt = 4; + const size_t cache_elems = static_cast(nb * bs * page); + auto kf = RandF32(static_cast(nt * page), 33); + auto vf = RandF32(static_cast(nt * page), 44); + std::vector slots = {3, 0, 2, 1}; + const float k_scale = 0.007f, v_scale = 0.003f; + + // Both narrow dtypes are 2-byte, so one uint16_t staging buffer serves each. + for (DType src : {DType::kBF16, DType::kF16}) { + const int src_dtype_tag = static_cast(src); + CAPTURE(src_dtype_tag); + std::vector kb(kf.size()), vb(vf.size()); + for (size_t i = 0; i < kf.size(); ++i) { + kb[i] = src == DType::kBF16 ? vt::F32ToBF16(kf[i]) : vt::F32ToF16(kf[i]); + vb[i] = src == DType::kBF16 ? vt::F32ToBF16(vf[i]) : vt::F32ToF16(vf[i]); + } + + std::vector kc_ref(cache_elems, 0); + std::vector vc_ref(cache_elems, 0); + Tensor ck = Host(kb.data(), src, {nt, H, D}); + Tensor cv = Host(vb.data(), src, {nt, H, D}); + Tensor ckc = Host(kc_ref.data(), DType::kI8, {nb, bs, H, D}); + Tensor cvc = Host(vc_ref.data(), DType::kI8, {nb, bs, H, D}); + Tensor cs = Host(slots.data(), DType::kI64, {nt}); + vt::ReshapeAndCacheFp8(cq, ck, cv, ckc, cvc, cs, Fp8KVCacheDataType::kFp8E4M3, k_scale, + v_scale); + + void* dk = gpu.Alloc(kb.size() * sizeof(uint16_t)); + void* dv = gpu.Alloc(vb.size() * sizeof(uint16_t)); + void* dkc = gpu.Alloc(cache_elems); + void* dvc = gpu.Alloc(cache_elems); + void* ds = gpu.Alloc(slots.size() * sizeof(int64_t)); + std::vector zero(cache_elems, 0); + gpu.Copy(gq, dk, kb.data(), kb.size() * sizeof(uint16_t)); + gpu.Copy(gq, dv, vb.data(), vb.size() * sizeof(uint16_t)); + gpu.Copy(gq, dkc, zero.data(), cache_elems); + gpu.Copy(gq, dvc, zero.data(), cache_elems); + gpu.Copy(gq, ds, slots.data(), slots.size() * sizeof(int64_t)); + Tensor gk = Dev(dk, src, {nt, H, D}); + Tensor gv = Dev(dv, src, {nt, H, D}); + Tensor gkc = Dev(dkc, DType::kI8, {nb, bs, H, D}); + Tensor gvc = Dev(dvc, DType::kI8, {nb, bs, H, D}); + Tensor gs = Dev(ds, DType::kI64, {nt}); + vt::ReshapeAndCacheFp8(gq, gk, gv, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E4M3, k_scale, + v_scale); + + std::vector kc_got(cache_elems, 0); + std::vector vc_got(cache_elems, 0); + gpu.Copy(gq, kc_got.data(), dkc, cache_elems); + gpu.Copy(gq, vc_got.data(), dvc, cache_elems); + gpu.Synchronize(gq); + CHECK(kc_got == kc_ref); + CHECK(vc_got == vc_ref); + // The CPU oracle must have WRITTEN something, or the equality above is + // between two all-zero buffers and holds for any kernel. + CHECK(std::any_of(kc_ref.begin(), kc_ref.end(), [](uint8_t b) { return b != 0; })); + + gpu.Free(dk); + gpu.Free(dv); + gpu.Free(dkc); + gpu.Free(dvc); + gpu.Free(ds); + } + gpu.DestroyQueue(gq); +} + +// ─── G4 ───────────────────────────────────────────────────────────────────── +// READ parity: paged attention over the SAME fp8 cache bytes, ROCm vs CPU, in +// BOTH shapes the fp8 arm routes to — pure decode (the generic block kernel) +// and prefill (the tiled flash kernel). The cache is built once on the host so +// this case measures the READ alone; G3 already measures the store. +// +// The dequant itself is bit-identical by construction: the ROCm kernel decodes +// e4m3 with the same arithmetic as vt::F8E4M3ToF32 and multiplies by the same +// per-tensor scale (quant_utils.cuh:419-429). The only divergence available is +// the softmax REDUCTION ORDER (block-cooperative on ROCm, sequential on the +// CPU), so the band is tight. A wrong scale, a missing dequant, a swapped +// k_scale/v_scale or a dropped sign blows it by orders of magnitude. +TEST_CASE("rocm fp8 KV paged-attention read matches the CPU read") { + if (!HasRocm()) { + MESSAGE("SKIPPED: no ROCm backend in this build/host — the ROCm fp8 KV " + "paged-attention read parity gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + Queue cq{Cpu(), nullptr}; + + // 2 requests, 2 q-heads over 1 kv-head (GQA), head_size 16, block_size 4. + const int64_t nb = 4, bs = 4, H = 1, D = 16, hq = 2, num_reqs = 2; + const size_t cache_elems = static_cast(nb * bs * H * D); + auto raw = RandF32(cache_elems, 77); + const float k_scale = 0.005f, v_scale = 0.009f; + std::vector kc(cache_elems), vc(cache_elems); + for (size_t i = 0; i < cache_elems; ++i) { + kc[i] = vt::StoreKvFp8E4M3(raw[i], k_scale); + vc[i] = vt::StoreKvFp8E4M3(raw[cache_elems - 1 - i], v_scale); + } + std::vector bt = {0, 1, 2, 3}; // [num_reqs, max_blocks] + std::vector seq = {5, 3}; + + void* dkc = gpu.Alloc(cache_elems); + void* dvc = gpu.Alloc(cache_elems); + void* dbt = gpu.Alloc(bt.size() * sizeof(int32_t)); + void* dseq = gpu.Alloc(seq.size() * sizeof(int32_t)); + gpu.Copy(gq, dkc, kc.data(), cache_elems); + gpu.Copy(gq, dvc, vc.data(), cache_elems); + gpu.Copy(gq, dbt, bt.data(), bt.size() * sizeof(int32_t)); + gpu.Copy(gq, dseq, seq.data(), seq.size() * sizeof(int32_t)); + + struct Shape { + const char* name; + int64_t nt; + std::vector qsl; + }; + // nt == num_reqs -> pure decode; nt > num_reqs -> prefill. + const std::vector shapes = {{"decode", 2, {0, 1, 2}}, {"prefill", 4, {0, 3, 4}}}; + + for (const Shape& sh : shapes) { + CAPTURE(std::string(sh.name)); + auto qh = RandF32(static_cast(sh.nt * hq * D), 88); + std::vector qsl = sh.qsl; + + PagedAttentionArgs args; + args.scale = 0.25f; + args.causal = true; + args.kv_cache_dtype = Fp8KVCacheDataType::kFp8E4M3; + args.k_scale = k_scale; + args.v_scale = v_scale; + + std::vector cpu_out(static_cast(sh.nt * hq * D), 0.0f); + Tensor cqt = Host(qh.data(), DType::kF32, {sh.nt, hq, D}); + Tensor cot = Host(cpu_out.data(), DType::kF32, {sh.nt, hq, D}); + Tensor ckc = Host(kc.data(), DType::kI8, {nb, bs, H, D}); + Tensor cvc = Host(vc.data(), DType::kI8, {nb, bs, H, D}); + Tensor cbt = Host(bt.data(), DType::kI32, {num_reqs, 2}); + Tensor cseq = Host(seq.data(), DType::kI32, {num_reqs}); + Tensor cqsl = Host(qsl.data(), DType::kI32, {num_reqs + 1}); + vt::PagedAttention(cq, cot, cqt, ckc, cvc, cbt, cseq, cqsl, args); + + void* dq = gpu.Alloc(qh.size() * sizeof(float)); + void* dout = gpu.Alloc(qh.size() * sizeof(float)); + void* dqsl = gpu.Alloc(qsl.size() * sizeof(int32_t)); + gpu.Copy(gq, dq, qh.data(), qh.size() * sizeof(float)); + gpu.Copy(gq, dqsl, qsl.data(), qsl.size() * sizeof(int32_t)); + Tensor gqt = Dev(dq, DType::kF32, {sh.nt, hq, D}); + Tensor got = Dev(dout, DType::kF32, {sh.nt, hq, D}); + Tensor gkc = Dev(dkc, DType::kI8, {nb, bs, H, D}); + Tensor gvc = Dev(dvc, DType::kI8, {nb, bs, H, D}); + Tensor gbt = Dev(dbt, DType::kI32, {num_reqs, 2}); + Tensor gseq = Dev(dseq, DType::kI32, {num_reqs}); + Tensor gqsl = Dev(dqsl, DType::kI32, {num_reqs + 1}); + vt::PagedAttention(gq, got, gqt, gkc, gvc, gbt, gseq, gqsl, args); + + std::vector gpu_out(qh.size(), 0.0f); + gpu.Copy(gq, gpu_out.data(), dout, gpu_out.size() * sizeof(float)); + gpu.Synchronize(gq); + + double num = 0.0, den = 0.0, worst = 0.0; + for (size_t i = 0; i < gpu_out.size(); ++i) { + const double d0 = static_cast(gpu_out[i]) - static_cast(cpu_out[i]); + num += d0 * d0; + den += static_cast(cpu_out[i]) * static_cast(cpu_out[i]); + worst = std::max(worst, std::fabs(d0)); + } + // The CPU arm must have produced a non-degenerate output, or the comparison + // above is between two fields of zeros and would pass on any kernel. + CHECK(den > 0.0); + const double nmse = den > 0.0 ? num / den : 1.0; + CAPTURE(nmse); + CAPTURE(worst); + CHECK(nmse < 1e-6); + CHECK(worst < 1e-3); + + gpu.Free(dq); + gpu.Free(dout); + gpu.Free(dqsl); + } + + gpu.Free(dkc); + gpu.Free(dvc); + gpu.Free(dbt); + gpu.Free(dseq); + gpu.DestroyQueue(gq); +} + +// ─── G4b ──────────────────────────────────────────────────────────────────── +// THE INSTANTIATION PRODUCTION WILL USE. G4 above runs an f32 query into an f32 +// output, which resolves `PagedAttnOnline` +// (src/vt/rocm/rocm_paged_attn.hip). That is not the arm a served model takes: +// vLLM resolves ONE model dtype and every layer inherits it (AGENTS.md "Inherit +// vLLM defaults"), the gate models are bf16, and the bf16 query/output arm is +// what production runs. Without this case +// `PagedAttnOnline<__hip_bfloat16, uint8_t, __hip_bfloat16>` compiles, ships, +// and is never once executed against the oracle. +// +// The band is looser than G4's and deliberately so: both arms round an f32 +// accumulator to bf16 on the store, and bf16 carries 8 mantissa bits, so two +// accumulators that differ only in softmax reduction order can land on +// opposite sides of one rounding boundary. The output is a convex combination +// of V rows and every V here is inside [-2, 2], so |x| < 2 and one bf16 ulp is +// at most 2^1 * 2^-7 = 1.56e-2; even if EVERY element were a full ulp out the +// NMSE would be (2^-8)^2 = 1.5e-5. The band below admits that and nothing else +// — a missing dequant, a swapped k_scale/v_scale or a dropped sign moves the +// output by orders of magnitude, not by an ulp. +TEST_CASE("rocm fp8 KV paged-attention read matches the CPU read (bf16 query, bf16 out)") { + if (!HasRocm()) { + MESSAGE("SKIPPED: no ROCm backend in this build/host — the bf16-query/bf16-out " + "fp8 KV paged-attention read parity gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + Queue cq{Cpu(), nullptr}; + + const int64_t nb = 4, bs = 4, H = 1, D = 16, hq = 2, num_reqs = 2; + const size_t cache_elems = static_cast(nb * bs * H * D); + auto raw = RandF32(cache_elems, 77); + const float k_scale = 0.005f, v_scale = 0.009f; + std::vector kc(cache_elems), vc(cache_elems); + for (size_t i = 0; i < cache_elems; ++i) { + kc[i] = vt::StoreKvFp8E4M3(raw[i], k_scale); + vc[i] = vt::StoreKvFp8E4M3(raw[cache_elems - 1 - i], v_scale); + } + std::vector bt = {0, 1, 2, 3}; + std::vector seq = {5, 3}; + + void* dkc = gpu.Alloc(cache_elems); + void* dvc = gpu.Alloc(cache_elems); + void* dbt = gpu.Alloc(bt.size() * sizeof(int32_t)); + void* dseq = gpu.Alloc(seq.size() * sizeof(int32_t)); + gpu.Copy(gq, dkc, kc.data(), cache_elems); + gpu.Copy(gq, dvc, vc.data(), cache_elems); + gpu.Copy(gq, dbt, bt.data(), bt.size() * sizeof(int32_t)); + gpu.Copy(gq, dseq, seq.data(), seq.size() * sizeof(int32_t)); + + struct Shape { + const char* name; + int64_t nt; + std::vector qsl; + }; + const std::vector shapes = {{"decode", 2, {0, 1, 2}}, {"prefill", 4, {0, 3, 4}}}; + + for (const Shape& sh : shapes) { + const std::string shape_name(sh.name); + CAPTURE(shape_name); + auto qf = RandF32(static_cast(sh.nt * hq * D), 88); + std::vector qb(qf.size()); + for (size_t i = 0; i < qf.size(); ++i) qb[i] = vt::F32ToBF16(qf[i]); + std::vector qsl = sh.qsl; + + PagedAttentionArgs args; + args.scale = 0.25f; + args.causal = true; + args.kv_cache_dtype = Fp8KVCacheDataType::kFp8E4M3; + args.k_scale = k_scale; + args.v_scale = v_scale; + + std::vector cpu_out(qf.size(), 0); + Tensor cqt = Host(qb.data(), DType::kBF16, {sh.nt, hq, D}); + Tensor cot = Host(cpu_out.data(), DType::kBF16, {sh.nt, hq, D}); + Tensor ckc = Host(kc.data(), DType::kI8, {nb, bs, H, D}); + Tensor cvc = Host(vc.data(), DType::kI8, {nb, bs, H, D}); + Tensor cbt = Host(bt.data(), DType::kI32, {num_reqs, 2}); + Tensor cseq = Host(seq.data(), DType::kI32, {num_reqs}); + Tensor cqsl = Host(qsl.data(), DType::kI32, {num_reqs + 1}); + vt::PagedAttention(cq, cot, cqt, ckc, cvc, cbt, cseq, cqsl, args); + + void* dq = gpu.Alloc(qb.size() * sizeof(uint16_t)); + void* dout = gpu.Alloc(qb.size() * sizeof(uint16_t)); + void* dqsl = gpu.Alloc(qsl.size() * sizeof(int32_t)); + gpu.Copy(gq, dq, qb.data(), qb.size() * sizeof(uint16_t)); + gpu.Copy(gq, dqsl, qsl.data(), qsl.size() * sizeof(int32_t)); + Tensor gqt = Dev(dq, DType::kBF16, {sh.nt, hq, D}); + Tensor got = Dev(dout, DType::kBF16, {sh.nt, hq, D}); + Tensor gkc = Dev(dkc, DType::kI8, {nb, bs, H, D}); + Tensor gvc = Dev(dvc, DType::kI8, {nb, bs, H, D}); + Tensor gbt = Dev(dbt, DType::kI32, {num_reqs, 2}); + Tensor gseq = Dev(dseq, DType::kI32, {num_reqs}); + Tensor gqsl = Dev(dqsl, DType::kI32, {num_reqs + 1}); + vt::PagedAttention(gq, got, gqt, gkc, gvc, gbt, gseq, gqsl, args); + + std::vector gpu_out(qb.size(), 0); + gpu.Copy(gq, gpu_out.data(), dout, gpu_out.size() * sizeof(uint16_t)); + gpu.Synchronize(gq); + + double num = 0.0, den = 0.0, worst = 0.0; + for (size_t i = 0; i < gpu_out.size(); ++i) { + const double g = static_cast(vt::BF16ToF32(gpu_out[i])); + const double c = static_cast(vt::BF16ToF32(cpu_out[i])); + num += (g - c) * (g - c); + den += c * c; + worst = std::max(worst, std::fabs(g - c)); + } + // The CPU arm must have produced a non-degenerate output, or the comparison + // is between two fields of zeros and would pass on any kernel. + CHECK(den > 0.0); + const double nmse = den > 0.0 ? num / den : 1.0; + CAPTURE(nmse); + CAPTURE(worst); + CHECK(nmse < 1e-4); + CHECK(worst < 2e-2); + + gpu.Free(dq); + gpu.Free(dout); + gpu.Free(dqsl); + } + + gpu.Free(dkc); + gpu.Free(dvc); + gpu.Free(dbt); + gpu.Free(dseq); + gpu.DestroyQueue(gq); +} + +// ─── G5 ───────────────────────────────────────────────────────────────────── +// fp8_e5m2 stays a NAMED later brick (spec W5) on ROCm exactly as on CPU and +// CUDA — it must be refused, never silently mis-stored through the e4m3 +// converter. There are THREE refusals on that path and only one is a ROCm-side +// guarantee: +// +// * the op wrapper, `src/vt/ops.cpp` `ReshapeAndCacheFp8` — device-independent, +// evaluated ABOVE the device checks and above GetOp, so it fires identically +// on a CPU queue and cannot be a ROCm guarantee. +// * the CPU kernel, `src/vt/cpu/cpu_cache.cpp` `ReshapeAndCacheFp8Kernel`. +// * the ROCm kernel's own guard, `src/vt/rocm/rocm_dense_basic.hip` +// `ReshapeAndCacheFp8KernelRocm`, which is defence in depth for any future +// caller that reaches the registered provider without going through the +// wrapper. +// +// A layered refusal needs an assertion that NAMES its layer. This version +// reaches the kernel guard the only way anything can — through the registered +// provider — and requires the message to carry both `rocm reshape_and_cache_fp8` +// and `fp8_e5m2`, which no other layer produces. +TEST_CASE("the ROCm fp8 KV store kernel refuses e5m2 (later brick)") { + if (!HasRocm()) { + MESSAGE("SKIPPED: no ROCm backend in this build/host — the ROCm-kernel e5m2 " + "refusal gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kROCM); + Queue gq = gpu.CreateQueue(); + const int64_t nb = 1, bs = 4, H = 1, D = 16, page = H * D; + std::vector k(static_cast(page), 1.0f); + std::vector slots = {0}; + void* dk = gpu.Alloc(k.size() * sizeof(float)); + void* dkc = gpu.Alloc(static_cast(nb * bs * page)); + void* dvc = gpu.Alloc(static_cast(nb * bs * page)); + void* ds = gpu.Alloc(sizeof(int64_t)); + gpu.Copy(gq, dk, k.data(), k.size() * sizeof(float)); + gpu.Copy(gq, ds, slots.data(), sizeof(int64_t)); + gpu.Synchronize(gq); + Tensor gk = Dev(dk, DType::kF32, {1, H, D}); + Tensor gkc = Dev(dkc, DType::kI8, {nb, bs, H, D}); + Tensor gvc = Dev(dvc, DType::kI8, {nb, bs, H, D}); + Tensor gs = Dev(ds, DType::kI64, {1}); + + // The registered ROCm provider, resolved exactly as vt::ReshapeAndCacheFp8 + // resolves it, then called directly so the wrapper's own e5m2 check is not in + // the way. Anything that reaches this kernel reaches it through this pointer. + auto* fn = reinterpret_cast( + vt::GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kROCM)); + REQUIRE(fn != nullptr); + std::string msg; + try { + fn(gq, gk, gk, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E5M2, 0.01f, 0.01f); + FAIL("the ROCm fp8 KV store kernel must refuse e5m2, not store it as e4m3"); + } catch (const std::runtime_error& e) { + msg = e.what(); + } + CAPTURE(msg); + // The refusal must come from the ROCm KERNEL and name the missing part, not + // from the device-independent wrapper this call deliberately bypassed. + CHECK(msg.find("rocm reshape_and_cache_fp8") != std::string::npos); + CHECK(msg.find("fp8_e5m2") != std::string::npos); + + // e4m3 through the SAME pointer still runs: the guard above refuses one kind, + // it does not disable the kernel. + fn(gq, gk, gk, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E4M3, 0.01f, 0.01f); + gpu.Synchronize(gq); + + gpu.Free(dk); + gpu.Free(dkc); + gpu.Free(dvc); + gpu.Free(ds); + gpu.DestroyQueue(gq); +}