From 4d71e776efc18cb5e61a26e642ddad8de5339134 Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 21 Aug 2026 14:49:40 +0000 Subject: [PATCH 1/2] feat(KV-FP8): the CUDA fp8 KV store and read, reached by removing the W1 guard that refused CUDA before the provider table (#1593) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `KV-FP8` W1 landed the CPU half in 2026-07 and left the CUDA arm as a named later brick. It is now the critical path of benchmark campaign #1574, whose subject `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` declares `kv_cache_quant_algo: "FP8"` in `hf_quant_config.json`, so no cell of that three-way can be served correctly without it. W1 is the ORACLE for this wave. Every gate here compares CUDA to the landed CPU kernels; nothing re-ports the numerics. THE STORE is a 1:1 port of the fp8 branch of vLLM's `reshape_and_cache_flash_kernel` (`csrc/libtorch_stable/cache_kernels.cu:314-401`) plus `CopyWithScaleOp` (`:241-252`) at pin `555967922`, restricted to upstream's `is_contiguous_heads && kv_scale_stride == 0` arm (`:352-366`) — the only arm the op's wrapper admits, because the vt cache is the NHD unbind slice and `ReshapeAndCacheFp8` takes two scalar scales. The converter is upstream's own `__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)` (`quant_utils.cuh:497-503`), a true DIVIDE rather than the activation path's hoisted reciprocal multiply, and its byte-for-byte equality to the CPU software codec `vt::F32ToF8E4M3` is already MEASURED at zero tolerance on sm_110 and sm_121a (`.agents/specs/vt-fp8-quant-arch-gate.md` G2). THE READ adds `LoadKv(ptr, i, scale)` beside `Load`. It is INERT on the f32 and bf16 arms — they forward to `Load` unchanged, so every existing caller reads the same bytes in the same order — and on `uint8_t` it is upstream's `scaled_vec_conversion` (`quant_utils.cuh:302-308`), written as the SAME ARITHMETIC as `vt::F8E4M3ToF32` so that CUDA==CPU on the read is a property of the source rather than of a measurement. Only the two correctness-grade kernels serve fp8, which is what the existing ladder already implies: the WMMA prefill kernels stage bf16 fragments, the vendored FA-2 launchers take bf16 pointers, and the vectorized decode-opt/GQA kernels read through `LoadRowN`/`LoadRow8`, 128-bit `uint4` loads specialized for bf16 and f32 only. Upstream draws the same line from the other side (`flash_attn.py:181-187,796-805`). WHAT ACTUALLY MADE THE ARM UNREACHABLE was neither kernel. Both W1 wrappers carried `VT_CHECK(q.device.type == DeviceType::kCPU, ... "a named later brick")` evaluated BEFORE provider lookup, so no CUDA kernel could ever have been reached however well it was registered. That is the RED this change was written against. The STORE now resolves through the provider table like every other op, because `kReshapeAndCacheFp8` is its own `OpId` that only CPU and CUDA register and an unimplemented backend refuses BY NAME inside `GetOp`. The READ cannot: it rides ADDITIVE fields on `PagedAttentionArgs` of an op `kMETAL` and `kROCM` already register for the FLOAT path, and nothing in the provider table separates the two arms, so an fp8 cache would reach a float kernel and return silent garbage. It keeps an explicit CPU-or-CUDA list whose message names the missing part. UNREACHED, DELIBERATELY, AND NOT NEW. Nothing calls the fp8 KV path from a production entry point on either backend: `vt::ReshapeAndCacheFp8` and `PagedAttentionArgs::kv_cache_dtype` have no caller outside their tests. W1 landed in that state and this does not change it. `KV-FP8` W3 owns the wiring — half-sized KV blocks in the runner, `--kv-cache-dtype` threaded from the CLI, and the checkpoint `k_scale`/`v_scale` path — and #1593 tracks it. Listed under `## Owed` in `.agents/specs/fp8-kv-cache.md`. THE DEVICE GATES ARE UNEXECUTED AND THE CUDA TRANSLATION UNITS ARE UNCOMPILED. The implementing session had no CUDA toolkit (`nvcc` is absent on `mudler-ubuntu-box`) and no device (the fleet was leased for #1574), so G2 (provider registration), G3 (store byte parity, f32 and bf16), G4 (paged-read parity, decode and prefill) and G5 (the e5m2 refusal) all skip with a MESSAGE naming what did not run. G1 and G1b — provider routing and the Metal/ROCm refusal — are the only cases that ran, and they run on the CPU leg. The first CUDA build or `rc` lease that touches this row must run `ctest -R test_cuda_fp8_kv_cache` before W2 counts as measured; the spec's `## Owed` says so. EVIDENCE. RED first: `test_cuda_fp8_kv_cache` 6 assertions / 6 failed, on `reshape_and_cache_fp8: only the CPU fp8-KV store is implemented in W1` at `ops.cpp:3478` and `paged_attention: only the CPU fp8-KV read is implemented in W1` at `ops.cpp:3811`. Green after: 6 cases / 10 assertions. Three negative mutations, each rebuilt, run and restored against a pre-taken sha256 — reinstate the store guard (3 failed), reinstate the read guard (7 failed), delete the Metal/ROCm refusal (4 failed). No sibling regressions: `test_ops_fp8_kv_cache` 8/511, `test_ops_reshape_cache` 12/192, `test_ops_paged_attn` 14/1646, `test_ops_paged_attn_dtype` 3/172. Issue: https://github.com/mudler/vllm.cpp/issues/1593 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/engine-matrix.md | 2 +- .agents/issue-index.md | 1 + .agents/quantization-matrix.md | 2 +- .agents/specs/fp8-kv-cache.md | 117 +++++- docs/FEATURES.md | 2 +- include/vt/ops.h | 12 +- src/vt/cuda/cuda_cache.cu | 125 ++++++ src/vt/cuda/cuda_paged_attn.cu | 173 +++++++-- src/vt/ops.cpp | 32 +- tests/CMakeLists.txt | 5 + tests/vt/test_cuda_fp8_kv_cache.cpp | 566 ++++++++++++++++++++++++++++ 11 files changed, 986 insertions(+), 51 deletions(-) create mode 100644 tests/vt/test_cuda_fp8_kv_cache.cpp diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 3463e7146..b330ae7ab 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`). **Residuals (honest, named):** the CUDA fp8 store + fp8 paged-attention read (the GPU memory-halving path, DGX-blocked), the runner/spec integration (half-sized KV blocks + checkpoint-scale threading + `--kv-cache-dtype`/`--calculate-kv-scales`), fp8_e5m2 CPU compute + per-head scales — all W2-W5 in the spec | 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) | [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 and the CUDA TUs UNCOMPILED** (no toolkit, no device in the implementing session) -- see the spec's `## Owed`. **Residuals (honest, named):** the runner/spec integration (half-sized KV blocks + checkpoint-scale threading + `--kv-cache-dtype`/`--calculate-kv-scales`), fp8_e5m2 CPU compute + per-head scales — all W2-W5 in the spec | 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) | [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/issue-index.md b/.agents/issue-index.md index 9a66f7d73..c6caaa3f5 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -527,3 +527,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1563](https://github.com/mudler/vllm.cpp/issues/1563) | `GATE-SQUASH-SEPARATOR` | **A markdown `---` horizontal rule anywhere in a pull request body silently voids the trailer block, and `check-commit-trailers.py` blames the trailers instead of the framing.** Found 2026-08-21 writing the body for PR [#1550](https://github.com/mudler/vllm.cpp/pull/1550) ([#1542](https://github.com/mudler/vllm.cpp/issues/1542)). `parsed_trailers()` shells out to git's trailer parser, and **git treats a line of exactly `---` as the start of the patch section**, so everything after the first one is not part of the message and a trailer block below it is invisible. Reproduced with no repository state: a body of `subject / prose / --- / more prose / FOLLOWING_AGENTS_PROTOCOL / the three trailers` reports `[trailers] Following-Agents-Protocol must appear exactly once` and `[attribution] AI-Assisted must appear exactly once`; `sed -i '/^---$/d'` on that same file reports `OK: commit trailer contract`, and the `---` is the only difference. **The MESSAGE is the defect, not only the behaviour**: `Following-Agents-Protocol` appears EXACTLY ONCE in the body while the checker says it must appear exactly once, so a reader counts occurrences, finds one, counts again and dumps bytes before thinking to test the parser's own framing. `_strict_errors` already computes `_paragraphs(body)[-1]` correctly as the three trailers verbatim, so the checker holds the information needed to say "the trailer paragraph is present but git could not parse it; a `---` line at line N ends the message". Worse, the neighbouring `FOLLOWING_AGENTS_PROTOCOL must appear exactly once as a separate paragraph before the trailer paragraph` check stays SILENT, so the two errors that fire both point away from the cause. **Beyond one confusing message**: the repository sets `squash_merge_commit_message = PR_BODY`, so the body IS the landed commit message, and a body carrying a `---` lands a commit whose trailers `git interpret-trailers` cannot see, on a branch that is never force-pushed. Same permanent-damage shape AGENTS.md records for the `---------` separator GitHub wrote under `COMMIT_MESSAGES`, arriving from the AUTHOR side rather than the forge side. `scripts/agent-pr-body.py --pr ` DOES catch it and caught it here before the merge; the exposure is a body never passed through that command, which AGENTS.md notes is not a gate and cannot be one because it reaches the network, while the CI guard reads the frozen `pull_request` payload and so does not re-read a body edited after the final push. NOT FIXED HERE: it changes a checker's semantics and its message, so under `## Changing the rules or a checker` it needs its own row, a red-before test and green-after evidence. Two candidate repairs, neither chosen: name the `---` line, or strip patch-section framing before parsing so a markdown rule is inert -- the second changes what the contract accepts and is the larger decision. Suggested minimum: `tests/scripts/test_check_commit_trailers.py` gains a case pinning the reproduction above | bug | | [#1454](https://github.com/mudler/vllm.cpp/issues/1454) | `SPEC-MTP-GGUF` | **`test_qwen3_5_gguf_mtp.cpp` reported `Status: SUCCESS!` with `assertions: 0` on every CI run, and its one arithmetic guarantee was a tautology.** Both cases opened `if (path == nullptr) return;` on `VLLM_MTP_GGUF_MODEL`, and a bare `return` from a doctest case is a PASS: re-derived on a clean Release build at `947e5f648`, unset, the file printed `test cases: 2 \| 2 passed \| 0 failed \| 0 skipped`, `assertions: 0`, `Status: SUCCESS!`, exit 0, and printed nothing else. The variable is set nowhere in `.github/workflows/`, so that was the state of every run. Second defect in the same file: the comment at `:52` stated `num_hidden_layers + depth == block_count` and the line under it asserted `CHECK(c.num_hidden_layers > 0)`, true of every valid model. MEASURED, not argued: mutating `src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp:889` to `c.num_hidden_layers = block_count;` compiled clean and left the file at 2/2 cases, 0 assertions, `SUCCESS!`, exit 0. FIXED IN FLOW. The invariant is now pinned **HERMETICALLY** on KV-only synthetic GGUFs carrying no weight bytes, so CI checks it every run rather than never - 65/1 (the shipped Qwen3.8-27B pair), 25/1 (the Qwen3.5-2B reference this suite was developed against) and 28/3, the third arm separating `- nextn` from `- 1` - plus a head-less arm asserting the key is NOT published, which is the half `NumMtpLayers` cannot express because it answers 1 for an absent key. The two env-gated cases stay, now skipping with a `MESSAGE` naming the variable as `test_gguf_mmproj_reach.cpp` does, and the live one re-derives the invariant from the file's own `block_count` kv. Unset 4 cases / 18 assertions / `SUCCESS!` / rc 0; live on `Qwen3.8-27B-Q4_K_M.gguf` 4 / 38 / `SUCCESS!` / rc 0. Both mutants now red (9/18 and 5/18, exit 1), compiled clean, restored against a pre-taken sha256. **The production line is CORRECT and was not touched**: `block_count - nextn` landed `1a4db5c3c`, the `mtp_num_hidden_layers` republication `493327b4e`. Related but distinct: [#821](https://github.com/mudler/vllm.cpp/issues/821) W2 (`0adeb8b0e`) pins the same arithmetic for the 27B artifact on a committed manifest in `tests/vllm/models/test_qwen38_27b_gguf_manifest.cpp`, and that gate DOES catch both mutants - so the invariant was not globally unpinned, it was unpinned in this row's own file | bug | | [#1434](https://github.com/mudler/vllm.cpp/issues/1434) | `GATE-DOC-CHECKPOINT-STATES` | **`scripts/check-doc-checkpoint.py` could not see `PARTIAL`, so 118 state cells could move with no gate observing them.** `STATES` (`:56-66`) is the whole definition of what a lifecycle state IS for the gate that enforces AGENTS.md's `docs/STATUS.md` / `docs/BENCHMARKS.md` / spec `## Now` triple, and `row_states` drops any row it cannot match. `lifecycle_moves` and `moved_rows` then iterate the AFTER map, so leaving the matched set is silent by construction. Re-derived at `947e5f648` (the report measured `63d87805c`): `PARTIAL` **118** cells and `ANCHOR-BACKFILL` **73**, against `DONE` 77 and `BLOCKED` 9 — `PARTIAL` is the second most used state in the matrices and the gate was blind to it. Over the seven tables `ROW_TABLES` actually reads, the resolved population goes from **153 rows to 226**, a 47.7 % widening. Two of the transitions the report names behave differently from its description, measured with scratch commits at `947e5f648` on an unmodified checker: `READY -> PARTIAL` rc **0** and `PARTIAL -> READY` rc **0** are the real blind spots, while the report's suggested `PARTIAL -> ACTIVE` already reds — by accident, reporting **`added as ACTIVE`** for a row that has existed for months, because it is absent from the BEFORE map. FIXED IN FLOW for `PARTIAL` only. **`ANCHOR-BACKFILL` is deliberately excluded**: `.agents/feature-matrix.md:14-17` defines it as a property of the RECORD (*a legacy implemented row without exact code, test and real-spec anchors*), `docs/STATUS.md` carries no such term and would have nothing true to write on a `DONE <-> ANCHOR-BACKFILL` move, and `REQUIRED["lifecycle"]` cannot demand the spec's `## Now` alone — so admitting it would demand a public-document edit with nothing to say, which is the exact shape `check-doc-checkpoint.py:4-17` records as the reason the file was rewritten (16 of 20 red CI runs, six hardcoded escape hatches). One row's resolved state moves and the move is a REPAIR: `KV-BLOCK-POOL` says `` `PARTIAL` (not `DONE`) `` in its prose and the last-match heuristic believed the parenthesis, resolving `DONE`. No pinned counter moves — `check-gate-commands.py` has its own `GATED_STATES` and `RUNNABLE_BASELINE` is keyed on matrix rows, `UNOWNED_HIGH_WATER` is unmoved because this row names an owner, and no matrix row or public document changes — which was measured, not assumed, because this is the [#1376](https://github.com/mudler/vllm.cpp/issues/1376) ratchet shape. Remainder listed under `## Owed` in [doc-checkpoint-lifecycle-states.md](specs/doc-checkpoint-lifecycle-states.md): `ANCHOR-BACKFILL` moves, `.agents/sglang-matrix.md` never entering `ROW_TABLES`, a row that leaves the matched set entirely, and a new row added directly as `PARTIAL` | bug | +| [#1593](https://github.com/mudler/vllm.cpp/issues/1593) | `KV-FP8` | **`KV-FP8` W2 and W3: the CUDA fp8 KV store, its paged-attention read, and the runner integration.** W1 landed the CPU half (`vt::ReshapeAndCacheFp8`, the read dequant in CPU paged attention, `vllm::v1::ParseCacheDType`) and left W2/W3/W4 `later`. The issue is now the critical path of benchmark campaign [#1574](https://github.com/mudler/vllm.cpp/issues/1574), whose subject `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` declares `kv_cache_quant_algo: "FP8"` and carries ZERO `k_scale`/`v_scale` tensors, so every published profile serves it with `--kv-cache-dtype fp8` and no cell can be served correctly without this. **W2 IS LANDED HERE**: the CUDA fp8-e4m3 store (`src/vt/cuda/cuda_cache.cu`), the fp8 dequant on the CUDA paged-attention read (`src/vt/cuda/cuda_paged_attn.cu` `LoadKv` + `LaunchPagedFp8`), the removal of the two W1 device-class refusals that made the CUDA arm unreachable however well it was registered, and a named CPU-or-CUDA refusal for the READ because it rides ADDITIVE `PagedAttentionArgs` fields on an op `kMETAL`/`kROCM` already register for the FLOAT path — without which an fp8 cache would be read as that backend's float dtype and return silent garbage. Gate `tests/vt/test_cuda_fp8_kv_cache.cpp`, RED-first on the provider-routing case. **The device half of that gate is UNEXECUTED and the CUDA TUs are UNCOMPILED**: the implementing session had no `nvcc` and no device, and says so under `## Owed` in [fp8-kv-cache.md](specs/fp8-kv-cache.md) together with the reachability debt — nothing calls the fp8 KV path from a production entry point on either backend, which is **W3's** wiring (half-sized KV blocks, `--kv-cache-dtype` threading, the checkpoint scale path including this checkpoint's scales-absent case). W3, W4, the Metal/ROCm arms and fp8_e5m2 remain owed | feature | diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index e9bebe433..74b39e7c6 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). e5m2 CPU compute + per-head scales + CUDA + runner integration 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) | [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 and whose CUDA TUs are UNCOMPILED (spec `## Owed`). e5m2 compute, per-head scales, the Metal/ROCm arms and the runner integration 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) | [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 694ef5535..a61187db7 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 (`KV-FP8`, `QUANT-KV-FP8`) +# fp8 KV cache (`cache_dtype=fp8*`) — spike + W1 + W2 (`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 @@ -19,9 +19,12 @@ re-port). CPU-buildable brick: an fp8-e4m3 K/V **store** (`vt::ReshapeAndCacheFp8`) + the fp8 **read** dequant in CPU paged attention + the `cache_dtype` config parse (`vllm::v1::ParseCacheDType`), all unit-gated RED-first. -- **Out (named later bricks):** the CUDA fp8-KV store kernel and the CUDA - fp8-KV paged-attention read (the GPU is the memory-halving e2e), fp8_e5m2 CPU - compute, per-attention-head scales, the full engine-runner integration +- **In (W2, `## W2 — the CUDA arm` below):** the CUDA fp8-e4m3 K/V store kernel + and the fp8 dequant on the CUDA paged-attention read, gated for parity against + the W1 CPU reference. +- **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 + — see `## W2` below), the full engine-runner integration (half-sized KV blocks in the real runner + checkpoint `k_scale`/`v_scale` threading + `--kv-cache-dtype`/`--calculate-kv-scales` CLI), and the vendor KV dtypes (`fp8_inc`, `fp8_ds_mla` — `QUANT-KV-FP8-VENDOR`) and turboquant / @@ -99,9 +102,15 @@ W1 (this change; CPU-only, `-Werror`): (mirror `CacheDType` + `is_quantized_kv_cache`). - `tests/vt/test_ops_fp8_kv_cache.cpp` (NEW) + its `tests/CMakeLists.txt` line. -Later bricks: the CUDA fp8 store + fp8 paged-attention read (GPU memory-halving -e2e); the runner/spec integration (half-sized blocks + checkpoint scale -threading + CLI); fp8_e5m2 CPU compute; per-head scales. +W2 (`## W2 — the CUDA arm` below): `src/vt/cuda/cuda_cache.cu` (the fp8 store +kernel + its kCUDA registration), `src/vt/cuda/cuda_paged_attn.cu` (`LoadKv`, +the two scale parameters on `PagedAttentionKernel`/`PagedFlashKernel`, +`LaunchPagedBlock`, `LaunchPagedFp8`), `src/vt/ops.cpp` (the device-class guards +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. ## Tests to port @@ -124,9 +133,17 @@ threading + CLI); fp8_e5m2 CPU compute; per-head scales. a wrong store direction (`hp * scale`) fails 3 cases / 480 assertions; a wrong read `v_scale` diverges > 0.05 from the baseline; an auto (no-dequant) read of an fp8 cache is refused. No sibling regressions (reshape 12/12, paged 14/14). -- **Later:** the CUDA fp8 store + read parity vs this CPU reference; the real - memory-halving e2e (KV blocks ~2× on a gate model at token parity) is the - binding gate and is DGX-blocked (docs/BENCHMARKS PENDING). +- **Correctness (W2, provider routing — the CPU leg):** `test_cuda_fp8_kv_cache` + 6 cases / 10 assertions GREEN on a CPU-only build. RED-first proven: with the + W1 device-class guards in place the suite reports 10 assertions / 6 failed for + the store guard plus the read guard, naming both refusal strings. +- **Correctness (W2, device — UNEXECUTED, see `## Owed`):** the store byte gate + (zero tolerance, f32 + bf16 source, a padded slot), the paged-read parity gate + (decode + prefill, NMSE < 1e-6 and worst < 1e-3 vs the CPU arm) and the + registration gate need a CUDA build and a device. Neither was available to the + implementing session, so **the CUDA TUs in this change are UNCOMPILED**. +- **Later:** the real memory-halving e2e (KV blocks ~2× on a gate model at token + parity) is the binding gate and is DGX-blocked (docs/BENCHMARKS PENDING). ## Dependencies @@ -141,11 +158,89 @@ vendor/turbo/nvfp4 KV dtypes are separate rows. |---|---|---| | W0 | this spike | DONE (this commit) | | W1 | CPU fp8-e4m3 store + read dequant + config parse + unit gate | DONE (this commit) | -| W2 | CUDA fp8-e4m3 store + fp8 paged-attention read (parity vs W1) | later | +| W2 | CUDA fp8-e4m3 store + fp8 paged-attention read (parity vs W1) | DONE (code + gate landed; the DEVICE cases are UNEXECUTED — see `## Owed`) | | W3 | runner/spec integration: half-sized KV blocks + checkpoint k/v_scale threading + `--kv-cache-dtype`/`--calculate-kv-scales` | later | | W4 | memory-halving e2e on a gate model (the binding gate, DGX) | later | | W5 | fp8_e5m2 CPU+CUDA compute; per-attention-head scales | later | +## W2 — the CUDA arm (#1593) + +Issue: [#1593](https://github.com/mudler/vllm.cpp/issues/1593). W1 is the +ORACLE: every W2 gate compares CUDA to the landed CPU kernels, never to a fresh +reference. + +**Store** (`src/vt/cuda/cuda_cache.cu`, `ReshapeAndCacheFp8KernelCuda` + +`ReshapeAndCacheFp8Kernel`, registered for `DeviceType::kCUDA`). A 1:1 port +of the fp8 branch of `reshape_and_cache_flash_kernel` +(`cache_kernels.cu:314-401`) + `CopyWithScaleOp` (`:241-252`), restricted to +upstream's `is_contiguous_heads && kv_scale_stride == 0` arm (`:352-366`) — +which is the only arm the op's wrapper admits, because the vt cache is the NHD +unbind slice and `ReshapeAndCacheFp8` takes two scalar scales. The converter is +upstream's own `__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)` +(`quant_utils.cuh:497-503`) — a true DIVIDE, not the activation path's hoisted +reciprocal multiply — and its byte-for-byte equality to the CPU software codec +`vt::F32ToF8E4M3` is already MEASURED at zero tolerance on sm_110 and sm_121a +([vt-fp8-quant-arch-gate.md](vt-fp8-quant-arch-gate.md) G2). Source dtypes +f32/f16/bf16, the same set the CPU `LoadSrcF32` serves. + +**Read** (`src/vt/cuda/cuda_paged_attn.cu`). `LoadKv(ptr, i, scale)` joins +`Load`: inert on the f32/bf16 arms (they forward to `Load` unchanged, so every +existing caller reads the same bytes in the same order), and on `uint8_t` it is +`Fp8E4M3ToF32Dev(byte) * scale` — upstream's `scaled_vec_conversion` (`quant_utils.cuh:302-308`), written as the SAME ARITHMETIC as +`vt::F8E4M3ToF32` so CUDA==CPU on the read is a property of the source rather +than of a measurement this session could not take. `PagedAttentionKernel` and +`PagedFlashKernel` gain `k_scale`/`v_scale`; `LaunchPagedByKv` keys on +`args.kv_cache_dtype` (never on the storage dtype, which is a bare `kI8` byte) +and routes to `LaunchPagedFp8`. + +**Scope of the read, argued.** Only the two correctness-grade kernels serve fp8: +the tiled flash prefill and the block decode. That is what the existing ladder +already implies — the WMMA prefill kernels stage `__nv_bfloat16` fragments, the +vendored FA-2 launchers take bf16 pointers, and the vectorized decode-opt/GQA +kernels read through `LoadRowN`/`LoadRow8`, 128-bit `uint4` loads specialized +for bf16 and f32 only. Upstream draws the same line from the other side: +FlashAttention serves a quantized KV cache only where +`flash_attn_supports_kv_cache_dtype` says so (`flash_attn.py:181-187,796-805`). +A tensor-core fp8 read is a PERFORMANCE brick; W2's gate is parity, and W4 owns +the memory/throughput measurement. + +**The device-class guards are gone, but not the refusal.** W1 hard-refused every +non-CPU queue inside both op wrappers, before provider lookup — that is what kept +the CUDA arm unreachable. The STORE now resolves through the provider table like +every other op, because `kReshapeAndCacheFp8` is its own `OpId` that only CPU and +CUDA register, so an unimplemented backend refuses BY NAME inside `GetOp`. The +READ cannot: it rides ADDITIVE fields on `PagedAttentionArgs` of an op `kMETAL` +and `kROCM` already register for the FLOAT path, and nothing in the provider +table can tell the two arms apart, so an fp8 cache would reach a float kernel and +return silent garbage. `src/vt/ops.cpp` therefore keeps an explicit CPU-or-CUDA +list there whose message names the missing part, and +`tests/vt/test_cuda_fp8_kv_cache.cpp` gates it on both `kMETAL` and `kROCM`. + +## Owed + +- **The W2 device gates are UNEXECUTED** (#1593). `tests/vt/test_cuda_fp8_kv_cache.cpp` + G2 (provider registration, CUDA build), G3 (store byte parity, f32 + bf16), G4 + (paged-read parity, decode + prefill) and G5 (e5m2 refusal) all need a CUDA + toolkit and a device; the implementing session had NEITHER — `nvcc` is absent + on `mudler-ubuntu-box` and the GPU fleet was leased for the #1574 campaign. + **The CUDA translation units in this change have therefore never been + compiled**, let alone run. G1/G1b (provider routing and the Metal/ROCm refusal) + are the only cases that executed, and they run on the CPU leg. The first CUDA + build or `rc` lease that touches this row must run + `ctest -R test_cuda_fp8_kv_cache` and record the result here before W2 counts + as measured. Until then the wave table's `DONE` means "landed and gated", not + "measured on hardware". +- **Nothing reaches the fp8 KV path from a production entry point yet**, on + either backend. `vt::ReshapeAndCacheFp8` and `PagedAttentionArgs::kv_cache_dtype` + have no caller outside their tests; W1 landed in that state and W2 does not + change it. **`KV-FP8` W3 owns the wiring** — half-sized KV blocks in the runner, + `--kv-cache-dtype` threaded from the CLI, and the checkpoint `k_scale`/`v_scale` + path — and it is tracked by #1593 alongside W2. +- **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. +- fp8_e5m2 and per-attention-head scales stay refused on both backends (W5). + ## Risks/decisions - **Storage as `DType::kI8`, interpretation as a separate enum.** vLLM's kernel diff --git a/docs/FEATURES.md b/docs/FEATURES.md index 80c827925..1c4655fc7 100644 --- a/docs/FEATURES.md +++ b/docs/FEATURES.md @@ -52,7 +52,7 @@ are our reading of their documented behavior, not measurements. | Block-paged KV with refcount and LRU evict | ✅ | ✅ | ✅ | ◐ | | Hybrid KV groups (full attention + GDN/Mamba) | ◐ GDN gate activation resolved from the checkpoint's `output_gate_type` (silu/swish/sigmoid; anything else refused at load, #489) | ✅ | ◐ | ◐ | | Sliding-window and chunked-local attention | ◐ | ✅ | ✅ | ✅ | -| fp8 KV cache | ◐ CPU only | ✅ | ✅ | ✅ | +| fp8 KV cache | ◐ e4m3 store + read dequant on CPU and CUDA (#1593); nothing serves it yet: no runner block sizing and no `--kv-cache-dtype`. Metal/ROCm refused by name. CUDA gate UNRUN ([spec](../.agents/specs/fp8-kv-cache.md)) | ✅ | ✅ | ✅ | | KV offload to host memory | ✅ | ✅ | ✅ | ☐ | | External KV provider ABI (LMCache) | ☐ | ✅ | ◐ | ☐ | | KV events (block create / evict publish) | ◐ no transport | ✅ | ☐ | ☐ | diff --git a/include/vt/ops.h b/include/vt/ops.h index 70ae1799d..aa737b63f 100644 --- a/include/vt/ops.h +++ b/include/vt/ops.h @@ -1121,7 +1121,7 @@ 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). kAuto (default) => the cache holds + // 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 @@ -1129,6 +1129,9 @@ struct PagedAttentionArgs { // (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. Fp8KVCacheDataType kv_cache_dtype = Fp8KVCacheDataType::kAuto; float k_scale = 1.0f; float v_scale = 1.0f; @@ -3446,9 +3449,10 @@ void ReshapeAndCache(Queue& q, const Tensor& k, const Tensor& v, Tensor& k_cache // the fp8::scaled_convert scale convention, quant_utils.cuh:296-308) @ pin // 555967922. k_scale/v_scale are the per-tensor scales BaseKVCacheMethod loads // from the checkpoint (kv_cache.py:108-191); both must be > 0. Same shape/stride -// contract as ReshapeAndCache; the ONLY difference is the fp8 store. CPU-only in -// W1 (the CUDA fp8-KV store kernel is a named later brick); kFp8E5M2 CPU compute -// is likewise a later brick. +// contract as ReshapeAndCache; the ONLY difference is the fp8 store. Implemented +// on CPU (W1, src/vt/cpu/cpu_cache.cpp) and CUDA (W2, src/vt/cuda/cuda_cache.cu, +// gated byte-for-byte against the CPU arm); a backend that registers no provider +// refuses by name in GetOp. kFp8E5M2 is a named later brick (spec W5). void ReshapeAndCacheFp8(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); diff --git a/src/vt/cuda/cuda_cache.cu b/src/vt/cuda/cuda_cache.cu index 06241b5c5..764110b2c 100644 --- a/src/vt/cuda/cuda_cache.cu +++ b/src/vt/cuda/cuda_cache.cu @@ -5,6 +5,9 @@ // layout — see the M1.6 Task-2 layout trap note). // Correctness-grade (M1.6): one block per token, threads stride over the page // (num_kv_heads*head_size). The perf kernel (vectorized / fp8) is M2.4. +#include +#include +#include #include #include @@ -94,6 +97,125 @@ void ReshapeAndCacheKernelCuda(Queue& q, const Tensor& k, const Tensor& v, Tenso Check(cudaGetLastError(), "reshape_and_cache launch"); } +// ─── fp8 KV-cache write (KV-FP8 W2, #1593) ───────────────────────────────── +// The CUDA arm of vt::ReshapeAndCacheFp8, and the CUDA sibling of the CPU kernel +// in src/vt/cpu/cpu_cache.cpp that is its ORACLE. +// +// 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. Upstream's `is_contiguous_heads && kv_scale_stride == 0` fast +// path (`:352-366`) is the ONLY one this op's wrapper admits: the vt paged cache +// is the NHD unbind slice (head_stride == head_size) and the scales are +// per-TENSOR (`kv_scale_stride == 0`). The HND / per-attention-head arm +// (`:367-400`) is a named later brick (spec W5), and per-head scales cannot +// reach here because ReshapeAndCacheFp8 takes two scalars. +// +// THE CONVERTER IS UPSTREAM'S OWN, and its equality to the CPU codec is already +// MEASURED. `fp8::scaled_convert` is +// `__nv_cvt_float_to_fp8(a / scale, __NV_SATFINITE, __NV_E4M3)` +// (quant_utils.cuh:497-503) — a true DIVIDE, unlike the activation-quant path's +// hoisted reciprocal multiply (cuda_quant_fp8.cu:56-63), which matters because +// the two differ by up to one f32 ulp before the round. That same intrinsic is +// gated byte-for-byte at zero tolerance against the CPU software codec +// vt::F32ToF8E4M3 on sm_110 and sm_121a +// (.agents/specs/vt-fp8-quant-arch-gate.md G2), and +// tests/vt/test_cuda_fp8_kv_cache.cpp re-takes that equality on this KV path. +// +// Destination arithmetic is the auto path's, with element size 1: the cache is +// DType::kI8 (the "byte never guesses its semantic type" rule, include/vt/dtype.h) +// and the fp8 INTERPRETATION travels as Fp8KVCacheDataType, exactly as upstream +// carries cache_t = uint8_t plus a KV_DTYPE template parameter. + +// Pointer overloads, not by-value ones: __half and __nv_bfloat16 both carry an +// implicit `operator float()`, so a by-value set would put a user conversion in +// the overload resolution for every call. Same shape as cuda_quant_fp8.cu's +// LoadIn and cuda_paged_attn.cu's Load. +__device__ __forceinline__ float Fp8SrcToF32(const float* p, int64_t i) { return p[i]; } +__device__ __forceinline__ float Fp8SrcToF32(const __nv_bfloat16* p, int64_t i) { + return __bfloat162float(p[i]); +} +__device__ __forceinline__ float Fp8SrcToF32(const __half* p, int64_t i) { + return __half2float(p[i]); +} + +// fp8 = Quantize(hp / scale) — quant_utils.cuh:296-300 "Convention of the scale". +__device__ __forceinline__ uint8_t StoreKvFp8E4M3Dev(float hp, float scale) { + return static_cast(__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)); +} + +template +__global__ void ReshapeAndCacheFp8Kernel( + const Tin* __restrict__ key, const Tin* __restrict__ value, + uint8_t* __restrict__ key_cache, uint8_t* __restrict__ value_cache, + const int64_t* __restrict__ 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; // element offset + 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(Fp8SrcToF32(key, ksrc + e), k_scale); + value_cache[vdst + e] = StoreKvFp8E4M3Dev(Fp8SrcToF32(value, vsrc + e), v_scale); + } +} + +void ReshapeAndCacheFp8KernelCuda(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, + "cuda 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; + const int64_t k_block_stride = k_cache.stride[0]; + const int64_t k_page_stride = k_cache.stride[1]; + const int64_t v_block_stride = v_cache.stride[0]; + const int64_t v_page_stride = v_cache.stride[1]; + const int64_t k_tok_stride = k.stride[0]; + const int64_t v_tok_stride = v.stride[0]; + const unsigned grid = static_cast(num_slots); + const unsigned block = static_cast(n_elems < 512 ? n_elems : 512); + const cudaStream_t s = AsStream(q); + const int64_t* slots = slot_mapping.Ptr(); + uint8_t* kc = k_cache.Ptr(); + uint8_t* vc = v_cache.Ptr(); + // The SOURCE dtype is the model float dtype and is typed here, unlike the auto + // path's raw-word copy: the fp8 store converts, so it must know what it reads. + // Same set the CPU LoadSrcF32 serves (cpu_cache.cpp). + switch (k.dtype) { + case DType::kF32: + ReshapeAndCacheFp8Kernel<<>>( + k.Ptr(), v.Ptr(), kc, vc, slots, block_size, n_elems, k_block_stride, + k_page_stride, v_block_stride, v_page_stride, k_tok_stride, v_tok_stride, k_scale, + v_scale); + break; + case DType::kBF16: + ReshapeAndCacheFp8Kernel<__nv_bfloat16><<>>( + k.Ptr<__nv_bfloat16>(), v.Ptr<__nv_bfloat16>(), kc, vc, slots, block_size, n_elems, + k_block_stride, k_page_stride, v_block_stride, v_page_stride, k_tok_stride, + v_tok_stride, k_scale, v_scale); + break; + case DType::kF16: + ReshapeAndCacheFp8Kernel<__half><<>>( + k.Ptr<__half>(), v.Ptr<__half>(), kc, vc, slots, block_size, n_elems, k_block_stride, + k_page_stride, v_block_stride, v_page_stride, k_tok_stride, v_tok_stride, k_scale, + v_scale); + break; + default: + VT_CHECK(false, "cuda reshape_and_cache_fp8: unsupported source dtype (f32/f16/bf16)"); + } + Check(cudaGetLastError(), "reshape_and_cache_fp8 launch"); +} + // ─── MLA cache write (W3) ────────────────────────────────────────────────── // Ported 1:1 from vllm/csrc/libtorch_stable/cache_kernels.cu:401-442 // `concat_and_cache_mla_kernel` @ e24d1b24 — ONE block per token, threads stride @@ -170,6 +292,9 @@ struct Registrar { RegisterOp( OpId::kReshapeAndCache, DeviceType::kCUDA, reinterpret_cast(static_cast(&ReshapeAndCacheKernelCuda))); + RegisterOp(OpId::kReshapeAndCacheFp8, DeviceType::kCUDA, + reinterpret_cast( + static_cast(&ReshapeAndCacheFp8KernelCuda))); RegisterOp( OpId::kConcatAndCacheMla, DeviceType::kCUDA, reinterpret_cast(static_cast(&ConcatAndCacheMlaKernelCuda))); diff --git a/src/vt/cuda/cuda_paged_attn.cu b/src/vt/cuda/cuda_paged_attn.cu index 87da049ee..c55b9df50 100644 --- a/src/vt/cuda/cuda_paged_attn.cu +++ b/src/vt/cuda/cuda_paged_attn.cu @@ -132,6 +132,47 @@ __device__ inline float Load(const __nv_bfloat16* p, int64_t i) { return __bfloa __device__ inline void Store(float* p, int64_t i, float v) { p[i] = v; } __device__ inline void Store(__nv_bfloat16* p, int64_t i, float v) { p[i] = __float2bfloat16(v); } +// ─── fp8 KV-cache READ (KV-FP8 W2, #1593) ────────────────────────────────── +// One K/V-CACHE element as f32, with the fp8 dequant folded in. `scale` is the +// per-tensor k_scale / v_scale BaseKVCacheMethod loads from the checkpoint +// (vllm/model_executor/layers/quantization/kv_cache.py:108-191) and is INERT on +// the float arms, which forward to Load() unchanged — so every existing bf16/f32 +// caller reads exactly the bytes, in exactly the order, it read before. +// +// The fp8 arm mirrors upstream's attention-side dequant +// `scaled_vec_conversion` (quant_utils.cuh:302-308): fp8 byte -> +// float, then multiply by the scale, i.e. `Dequant(FP8) * scale = HP` (the +// convention at :296-300). It is written as the SAME ARITHMETIC as the W1 CPU +// codec vt::F8E4M3ToF32 (include/vt/fp8_kv.h) rather than as the hardware +// `__nv_cvt_fp8_to_halfraw`, because W1 is this wave's oracle and sharing the +// decode makes CUDA==CPU on the read a property of the source rather than a +// measurement. The two agree in any case: every one of the 256 e4m3 codes is +// exactly representable in fp16, so upstream's fp8->half->float round trip is +// lossless. `std::ldexp(mantissa, exp - 7)` on a float IS `ldexpf`, so this is +// the CPU codec line for line. +__device__ __forceinline__ float Fp8E4M3ToF32Dev(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 CUDART_NAN_F; // e4m3fn NaN (0x7F/0xFF) + 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); +} + +__device__ inline float LoadKv(const float* p, int64_t i, float scale) { + (void)scale; + return Load(p, i); +} +__device__ inline float LoadKv(const __nv_bfloat16* p, int64_t i, float scale) { + (void)scale; + return Load(p, i); +} +__device__ inline float LoadKv(const uint8_t* p, int64_t i, float scale) { + return Fp8E4M3ToF32Dev(p[i]) * scale; +} + // FlashAttention local-mask bounds for one bottom-right-aligned absolute query // position p. Negative window values mean the corresponding full bound. Public // PagedAttentionArgs uses nullopt for full attention; launchers unwrap it to -1. @@ -189,7 +230,8 @@ __global__ void PagedAttentionKernel(Tout* out, const TQ* query, const TKV* k_ca 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) { + int window_left, int window_right, float k_scale, + float v_scale) { const int64_t t = blockIdx.x; // global query-token index const int64_t h = blockIdx.y; // q-head // Find request r with query_start_loc[r] <= t < query_start_loc[r+1]. @@ -232,7 +274,7 @@ __global__ void PagedAttentionKernel(Tout* out, const TQ* query, const TKV* k_ca const int64_t kbase = blk * kc_blk + off * kc_pg + g * kc_hd; float part = 0.0f; for (int64_t e = threadIdx.x; e < d; e += blockDim.x) - part += Load(query, qoff + e) * Load(k_cache, kbase + e); + part += Load(query, qoff + e) * LoadKv(k_cache, kbase + e, k_scale); red[threadIdx.x] = part; __syncthreads(); for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { @@ -248,7 +290,7 @@ __global__ void PagedAttentionKernel(Tout* out, const TQ* query, const TKV* k_ca const float pw = expf(s - m_new); const int64_t vbase = blk * vc_blk + off * vc_pg + g * vc_hd; for (int64_t e = threadIdx.x; e < d; e += blockDim.x) - acc[e] = acc[e] * corr + pw * Load(v_cache, vbase + e); + acc[e] = acc[e] * corr + pw * LoadKv(v_cache, vbase + e, v_scale); __syncthreads(); if (threadIdx.x == 0) { s_l = s_l * corr + pw; @@ -606,7 +648,7 @@ __global__ void PagedFlashKernel(Tout* out, const TQ* query, const TKV* k_cache, int 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, int bn) { + int window_right, int bn, float k_scale, float v_scale) { const int tile_idx = blockIdx.x; const int h = blockIdx.y; // q-head if (tile_idx >= num_tiles) return; @@ -675,12 +717,14 @@ __global__ void PagedFlashKernel(Tout* out, const TQ* query, const TKV* k_cache, const int j = j0 + kk; const int blk = block_table[static_cast(r) * bt_row + (j / block_size) * bt_col]; const int off = j % block_size; - ksm[idx] = Load(k_cache, static_cast(blk) * kc_blk + - static_cast(off) * kc_pg + - static_cast(g) * kc_hd + ee); - vsm[idx] = Load(v_cache, static_cast(blk) * vc_blk + - static_cast(off) * vc_pg + - static_cast(g) * vc_hd + ee); + ksm[idx] = LoadKv(k_cache, + static_cast(blk) * kc_blk + static_cast(off) * kc_pg + + static_cast(g) * kc_hd + ee, + k_scale); + vsm[idx] = LoadKv(v_cache, + static_cast(blk) * vc_blk + static_cast(off) * vc_pg + + static_cast(g) * vc_hd + ee, + v_scale); } __syncthreads(); @@ -2023,6 +2067,32 @@ bool DecodeD128Enabled() { // --- Launchers ------------------------------------------------------------- +// The correctness-grade block kernel, lifted out of LaunchDecode's tail so the +// fp8 KV arm (KV-FP8 W2) can reach it WITHOUT instantiating the vectorized +// decode-opt / GQA kernels above it, whose LoadRowN/LoadRow8 128-bit loaders are +// bf16/f32-only. Byte-for-byte the launch LaunchDecode always made: same kernel, +// same grid, same shared memory, same argument order. +template +void LaunchPagedBlock(cudaStream_t s, 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, const PagedAttentionArgs& args, + int64_t num_tokens, int64_t hq, int64_t d, int64_t num_reqs, + int64_t num_kv_heads, int64_t block_size) { + const dim3 grid(static_cast(num_tokens), static_cast(hq)); + const size_t shmem = (static_cast(d) + kPagedBlock) * sizeof(float); + auto* kernel = args.window_size.has_value() + ? PagedAttentionKernel + : PagedAttentionKernel; + kernel<<>>( + out.Ptr(), query.Ptr(), k_cache.Ptr(), v_cache.Ptr(), + block_table.Ptr(), seq_lens.Ptr(), query_start_loc.Ptr(), num_reqs, + hq, num_kv_heads, d, block_size, block_table.stride[0], block_table.stride[1], + k_cache.stride[0], k_cache.stride[1], k_cache.stride[2], v_cache.stride[0], v_cache.stride[1], + v_cache.stride[2], args.scale, args.logits_soft_cap, args.causal, WindowLeft(args), + WindowRight(args), args.k_scale, args.v_scale); + Check(cudaGetLastError(), "paged_attention decode launch"); +} + template void LaunchDecode(cudaStream_t s, Tensor& out, const Tensor& query, const Tensor& k_cache, const Tensor& v_cache, const Tensor& block_table, const Tensor& seq_lens, @@ -2098,18 +2168,9 @@ void LaunchDecode(cudaStream_t s, Tensor& out, const Tensor& query, const Tensor Check(cudaGetLastError(), "paged_attention decode-opt launch"); return; } - const size_t shmem = (static_cast(d) + kPagedBlock) * sizeof(float); - auto* kernel = args.window_size.has_value() - ? PagedAttentionKernel - : PagedAttentionKernel; - kernel<<>>( - out.Ptr(), query.Ptr(), k_cache.Ptr(), v_cache.Ptr(), - block_table.Ptr(), seq_lens.Ptr(), query_start_loc.Ptr(), num_reqs, - hq, num_kv_heads, d, block_size, block_table.stride[0], block_table.stride[1], - k_cache.stride[0], k_cache.stride[1], k_cache.stride[2], v_cache.stride[0], v_cache.stride[1], - v_cache.stride[2], args.scale, args.logits_soft_cap, args.causal, WindowLeft(args), - WindowRight(args)); - Check(cudaGetLastError(), "paged_attention decode launch"); + LaunchPagedBlock(s, out, query, k_cache, v_cache, block_table, seq_lens, + query_start_loc, args, num_tokens, hq, d, num_reqs, + num_kv_heads, block_size); } // Per-request query-tile layout, built DIRECTLY on the device from the device @@ -2230,7 +2291,8 @@ void LaunchPrefillFlash(cudaStream_t s, Tensor& out, const Tensor& query, const num_tiles, static_cast(hq), static_cast(num_kv_heads), static_cast(d), static_cast(block_size), block_table.stride[0], block_table.stride[1], k_cache.stride[0], k_cache.stride[1], k_cache.stride[2], v_cache.stride[0], v_cache.stride[1], v_cache.stride[2], - args.scale, args.logits_soft_cap, args.causal, WindowLeft(args), WindowRight(args), bn); + args.scale, args.logits_soft_cap, args.causal, WindowLeft(args), WindowRight(args), bn, + args.k_scale, args.v_scale); Check(cudaGetLastError(), "paged_attention prefill flash launch"); Check(cudaFreeAsync(d_tiles, s), "paged flash tiles free"); } @@ -2848,6 +2910,61 @@ void LaunchPaged(cudaStream_t s, Tensor& out, const Tensor& query, const Tensor& } } +// ─── fp8 KV-cache READ dispatch (KV-FP8 W2, #1593) ───────────────────────── +// TKV is `uint8_t`: the cache pages are 1-byte fp8-e4m3 (DType::kI8) and each +// read is dequantized as Dequant(fp8) * k_scale|v_scale inside LoadKv, mirroring +// upstream's `scaled_vec_conversion` (quant_utils.cuh:302-308). +// +// SCOPE, argued rather than assumed. Only the two CORRECTNESS-GRADE kernels are +// reachable from here — the tiled flash prefill and the block decode — and that +// is not a shortcut, it is what the ladder above already implies. Every faster +// arm is bf16-NATIVE by construction: the WMMA prefill ladder stages +// `__nv_bfloat16` fragments, the vendored FA-2 launchers take bf16 pointers, and +// the vectorized decode-opt/GQA kernels read the cache through LoadRowN/LoadRow8, +// which are 128-bit `uint4` loads specialized for bf16 and f32 only. Upstream +// draws the same line from the other side: FlashAttention only serves a +// quantized KV cache when `flash_attn_supports_kv_cache_dtype` says so +// (flash_attn.py:181-187,796-805) and otherwise the backend refuses. A tensor- +// core fp8 read is a PERFORMANCE brick, not this one; W2's gate is parity with +// the W1 CPU reference, and W4 owns the memory/throughput measurement. +template +void LaunchPagedFp8Out(cudaStream_t s, 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, const PagedAttentionArgs& args) { + const int64_t num_tokens = query.shape[0], hq = query.shape[1], d = query.shape[2]; + const int64_t num_reqs = seq_lens.shape[0]; + const int64_t num_kv_heads = k_cache.shape[2], block_size = k_cache.shape[1]; + if (num_tokens == 0 || hq == 0 || d == 0) return; + // Same predicate LaunchPaged uses to pick the tiled prefill kernel. + const bool is_prefill = num_tokens > num_reqs; + if (is_prefill && d <= kMaxEpl * 32 && PrefillFlashEnabled()) { + LaunchPrefillFlash(s, out, query, k_cache, v_cache, block_table, seq_lens, + query_start_loc, args, hq, d, num_reqs, num_kv_heads, + block_size); + return; + } + LaunchPagedBlock(s, out, query, k_cache, v_cache, block_table, seq_lens, + query_start_loc, args, num_tokens, hq, d, num_reqs, + num_kv_heads, block_size); +} + +template +void LaunchPagedFp8(cudaStream_t s, 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, const PagedAttentionArgs& args) { + switch (out.dtype) { + case DType::kF32: + LaunchPagedFp8Out(s, out, query, k_cache, v_cache, block_table, seq_lens, + query_start_loc, args); + break; + case DType::kBF16: + LaunchPagedFp8Out(s, out, query, k_cache, v_cache, block_table, seq_lens, + query_start_loc, args); + break; + default: VT_CHECK(false, "cuda paged_attention: unsupported out dtype (fp8 KV read)"); + } +} + // Dispatch on (query dtype, KV-cache dtype). Both f32 and bf16 caches are valid // (Phase-1 bf16 KV cache mirrors vLLM's bf16 flash_attn KV store); the query may // independently be f32 (Phase 1) or bf16. @@ -2855,6 +2972,16 @@ template void LaunchPagedByKv(cudaStream_t s, 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, const PagedAttentionArgs& args) { + // fp8 KV cache: the STORAGE dtype is a raw byte (kI8) and the INTERPRETATION + // travels in args.kv_cache_dtype, exactly as upstream carries cache_t=uint8_t + // plus a KV_DTYPE template parameter (dtype_fp8.cuh:9-13). Key on the + // interpretation, never on the storage dtype: a kI8 tensor with kAuto is not + // an fp8 cache, and the op wrapper already refuses that pair. + if (args.kv_cache_dtype == Fp8KVCacheDataType::kFp8E4M3) { + LaunchPagedFp8(s, out, query, k_cache, v_cache, block_table, seq_lens, query_start_loc, + args); + return; + } switch (k_cache.dtype) { case DType::kF32: LaunchPaged(s, out, query, k_cache, v_cache, block_table, seq_lens, diff --git a/src/vt/ops.cpp b/src/vt/ops.cpp index 14e8d72a0..092150546 100644 --- a/src/vt/ops.cpp +++ b/src/vt/ops.cpp @@ -3433,8 +3433,8 @@ void ReshapeAndCacheFp8(Queue& q, const Tensor& k, const Tensor& v, Tensor& k_ca VT_CHECK(kind != Fp8KVCacheDataType::kAuto, "reshape_and_cache_fp8: kind must be an fp8 dtype (use ReshapeAndCache for auto)"); VT_CHECK(kind == Fp8KVCacheDataType::kFp8E4M3, - "reshape_and_cache_fp8: only fp8_e4m3 is implemented on CPU in W1 " - "(fp8_e5m2 CPU compute is a named later brick)"); + "reshape_and_cache_fp8: only fp8_e4m3 is implemented " + "(fp8_e5m2 compute is a named later brick, spec W5)"); VT_CHECK(k.rank == 3 && v.rank == 3, "reshape_and_cache_fp8: k/v must be rank-3 [num_tokens,num_kv_heads,head_size]"); VT_CHECK(k_cache.rank == 4 && v_cache.rank == 4, @@ -3475,9 +3475,11 @@ void ReshapeAndCacheFp8(Queue& q, const Tensor& k, const Tensor& v, Tensor& k_ca VT_CHECK(k_cache.stride[2] == head_size && v_cache.stride[2] == head_size, "reshape_and_cache_fp8: k_cache/v_cache page must be head-contiguous " "(stride[2] == head_size) — the NHD unbind-slice layout"); - VT_CHECK(q.device.type == DeviceType::kCPU, - "reshape_and_cache_fp8: only the CPU fp8-KV store is implemented in W1 " - "(the CUDA fp8-KV store kernel is a named later brick)"); + // NO device-class guard. W1 hard-refused every non-CPU queue here, which is + // what kept the CUDA arm unreachable; W2 lands that arm (cuda_cache.cu), so + // the op resolves through the provider table like every other op and a device + // with no registered fp8-KV store refuses BY NAME in GetOp + // (src/vt/op_provider.cpp:563-567) instead of by device class. VT_CHECK(k.device == q.device && v.device == q.device && k_cache.device == q.device && v_cache.device == q.device && slot_mapping.device == q.device, "reshape_and_cache_fp8: device mismatch (k/v/k_cache/v_cache/slot_mapping/queue)"); @@ -3802,15 +3804,25 @@ void PagedAttention(Queue& q, Tensor& out, const Tensor& query, const Tensor& k_ "paged_attention: k_cache/v_cache must share one float dtype"); } else { VT_CHECK(args.kv_cache_dtype == Fp8KVCacheDataType::kFp8E4M3, - "paged_attention: only fp8_e4m3 KV read is implemented on CPU in W1 " - "(fp8_e5m2 CPU read is a named later brick)"); + "paged_attention: only the fp8_e4m3 KV read is implemented " + "(the fp8_e5m2 read is a named later brick, spec W5)"); VT_CHECK(k_cache.dtype == DType::kI8 && v_cache.dtype == DType::kI8, "paged_attention: fp8 KV read requires 1-byte fp8 cache (DType::kI8)"); VT_CHECK(args.k_scale > 0.0f && args.v_scale > 0.0f, "paged_attention: fp8 KV read requires k_scale/v_scale > 0"); - VT_CHECK(q.device.type == DeviceType::kCPU, - "paged_attention: only the CPU fp8-KV read is implemented in W1 " - "(the CUDA fp8-KV paged-attention kernel is a named later brick)"); + // 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"); } // metadata: block_table [num_reqs, max_blocks] i32, seq_lens [num_reqs] i32, // query_start_loc [num_reqs+1] i32. diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 62da9ec97..61418d028 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -2176,6 +2176,11 @@ vllm_cpp_add_test(test_ops_reshape_cache vt/test_ops_reshape_cache.cpp) # KV-FP8 W1: fp8 KV-cache store (ReshapeAndCacheFp8) + paged-attention read # dequant + the ParseCacheDType config wiring. vllm_cpp_add_test(test_ops_fp8_kv_cache vt/test_ops_fp8_kv_cache.cpp) +# KV-FP8 W2 (#1593): the CUDA fp8-e4m3 KV store + the fp8 dequant on the CUDA +# paged-attention read, gated for parity against the W1 CPU reference. The +# device cases skip with a MESSAGE when no CUDA backend is present; the +# provider-routing case runs on every leg. +vllm_cpp_add_test(test_cuda_fp8_kv_cache vt/test_cuda_fp8_kv_cache.cpp) # MLA campaign W3: the compressed-latent cache write + the grouped-topk router. vllm_cpp_add_test(test_ops_mla_cache vt/test_ops_mla_cache.cpp) # MLA campaign W4: the two-stage split-KV MQA decode over the compressed latent. diff --git a/tests/vt/test_cuda_fp8_kv_cache.cpp b/tests/vt/test_cuda_fp8_kv_cache.cpp new file mode 100644 index 000000000..1d9821b69 --- /dev/null +++ b/tests/vt/test_cuda_fp8_kv_cache.cpp @@ -0,0 +1,566 @@ +// CUDA fp8 KV-cache store + paged-attention read gate (KV-FP8 W2, #1593). +// +// 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 W2 — the CUDA arm is measured against it, never against a fresh +// reference — so this file only ever compares CUDA to the landed CPU kernels. +// +// 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:302-308 +// (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) +// +// FIVE gates, and they do not all run in the same build: +// +// G1 (runs in every build WITHOUT the CUDA backend, i.e. the x86 CI leg): the +// W1 device-class refusal is GONE. W1 hard-refused any non-CPU queue inside +// the op wrapper, BEFORE provider lookup ("the CUDA fp8-KV store kernel is a +// named later brick"). That guard is what W2 removes; while it stands no +// CUDA kernel can be reached however well it is registered, so this case is +// the RED-first assertion for the whole wave and the one gate a host with no +// CUDA toolkit can actually execute. +// G2 (CUDA build): the CUDA providers are REGISTERED for the fp8 store and the +// paged read — the shared-seam reach check. vt::ops.cpp dispatches through +// GetOp(OpId, DeviceType) and nothing else can select a kernel, so a +// registered provider IS the production path. +// G3 (CUDA device): STORE parity — the CUDA store writes the SAME BYTES as the +// CPU store, zero tolerance, over f32 and bf16 sources, with a padded (-1) +// slot and a strided unbind-slice cache. +// G4 (CUDA device): READ parity — paged attention over identical fp8 cache +// bytes, CUDA vs CPU, in both the decode and the prefill shape (the two +// kernels the fp8 arm routes to). +// G5 (CUDA device): fp8_e5m2 stays refused on CUDA as it is on CPU. +// +// G3/G4/G5 SKIP CLEANLY when no CUDA 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 HasCuda() { + try { + vt::GetBackend(DeviceType::kCUDA); + return true; + } catch (const std::runtime_error&) { + return false; + } +} + +Device Cpu() { return Device{DeviceType::kCPU, 0}; } +Device Gpu() { return Device{DeviceType::kCUDA, 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 ───────────────────────────────────────────────────────────────────── +// RED-first for the whole wave, and the only case here a CUDA-less host runs. +// +// Under W1 both wrappers carried `VT_CHECK(q.device.type == DeviceType::kCPU, +// ... "is a named later brick")`, evaluated BEFORE the provider table is +// consulted. W2 deletes it, so a non-CPU queue now resolves through GetOp like +// every other op and refuses BY NAME when nothing is registered +// (src/vt/op_provider.cpp:563-567, "no kernel for op ..."). +// +// Compiled only where the CUDA backend is absent: in a CUDA build the op IS +// registered, so these calls would dispatch a real kernel over host pointers. +// The CUDA build asserts the same property from the other side, in G2. +#ifndef VLLM_CPP_CUDA +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 CUDA 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); + CHECK(store_msg.find("only the CPU fp8-KV store") == 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 CUDA 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); + CHECK(read_msg.find("only the CPU fp8-KV read") == std::string::npos); +} +#endif // !VLLM_CPP_CUDA + +// ─── G1b ──────────────────────────────────────────────────────────────────── +// The other half of removing the device-class guard, and the reason it could not +// simply be deleted: the fp8 READ rides ADDITIVE fields on PagedAttentionArgs of +// an op kMETAL and kROCM already register for the FLOAT path (metal_ops.mm, +// rocm_ops.hip). The provider table cannot tell the two arms apart, so an fp8 +// cache reaching one of those kernels 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. +// +// Runs in every build: the check fires in the op wrapper, before any device or +// provider is touched, so no Metal/ROCm 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; + + for (DeviceType dt : {DeviceType::kMETAL, DeviceType::kROCM}) { + const Device dev{dt, 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 kCUDA 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 CUDA arm compiled and registered", which is exactly the question +// a `#ifdef`-elided kernel silently answers "no" to. +#ifdef VLLM_CPP_CUDA +TEST_CASE("the CUDA fp8 KV store and paged read are registered providers") { + CHECK(vt::GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kCUDA) != nullptr); + CHECK(vt::GetOp(OpId::kPagedAttention, DeviceType::kCUDA) != nullptr); +} +#endif // VLLM_CPP_CUDA + +// ─── G3 ───────────────────────────────────────────────────────────────────── +// STORE parity, byte for byte, zero tolerance. The CPU kernel is the oracle. +// +// The two arms are not 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 CUDA kernel is upstream's own +// `__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)`. That equality +// is already MEASURED in this tree at zero tolerance on sm_110 and sm_121a for +// the identical converter pair (.agents/specs/vt-fp8-quant-arch-gate.md G2, CPU +// vs CUDA QuantFp8Static); 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("cuda fp8 KV store is byte-identical to the CPU store") { + if (!HasCuda()) { + MESSAGE("SKIPPED: no CUDA backend in this build/host — the CUDA fp8 KV store " + "parity gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + Queue gq = gpu.CreateQueue(); + Queue cq{Cpu(), nullptr}; + + // Two blocks, block_size 4, 2 kv-heads, head_size 16 (upstream requires + // head_size % 16 == 0 on the fp8 path). 6 tokens, one PADDED (-1) so the skip + // branch is exercised on both arms. + const int64_t nb = 2, bs = 4, H = 2, D = 16, page = H * D, nt = 6; + const size_t cache_elems = static_cast(nb * bs * page); + auto k = RandF32(static_cast(nt * page), 11); + auto v = RandF32(static_cast(nt * page), 22); + std::vector slots = {0, 5, -1, 7, 2, 1}; + const float k_scale = 0.004f, v_scale = 0.011f; + + // CPU reference bytes, seeded with a recognisable fill so an untouched byte + // (the padded slot's, and every unwritten page) compares too. + std::vector kc_ref(cache_elems, 0xAB); + std::vector vc_ref(cache_elems, 0xCD); + Tensor ck = Host(k.data(), DType::kF32, {nt, H, D}); + Tensor cv = Host(v.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(k.size() * sizeof(float)); + void* dv = gpu.Alloc(v.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, k.data(), k.size() * sizeof(float)); + gpu.Copy(gq, dv, v.data(), v.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 CUDA ones: a quantized byte may + // legitimately equal the 0xAB fill, and counting CUDA'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); +} + +// bf16 source arm of the same store — the dtype vLLM actually resolves for a +// model (AGENTS.md "Inherit vLLM defaults"). Upstream widens bf16 to f32 BEFORE +// the divide (quant_utils.cuh:482-489, `__bfloat162float(a) / scale`) and the +// CPU LoadSrcF32 does the same, so the two must still agree byte for byte. +TEST_CASE("cuda fp8 KV store is byte-identical to the CPU store (bf16 source)") { + if (!HasCuda()) { + MESSAGE("SKIPPED: no CUDA backend in this build/host — the bf16-source fp8 KV " + "store parity gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + 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 kb(kf.size()), vb(vf.size()); + for (size_t i = 0; i < kf.size(); ++i) { + kb[i] = vt::F32ToBF16(kf[i]); + vb[i] = vt::F32ToBF16(vf[i]); + } + std::vector slots = {3, 0, 2, 1}; + const float k_scale = 0.007f, v_scale = 0.003f; + + std::vector kc_ref(cache_elems, 0); + std::vector vc_ref(cache_elems, 0); + Tensor ck = Host(kb.data(), DType::kBF16, {nt, H, D}); + Tensor cv = Host(vb.data(), DType::kBF16, {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, DType::kBF16, {nt, H, D}); + Tensor gv = Dev(dv, DType::kBF16, {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); + + 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, CUDA 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 CUDA kernel decodes +// e4m3 with the same arithmetic as vt::F8E4M3ToF32 and multiplies by the same +// per-tensor scale (quant_utils.cuh:302-308). The only divergence available is +// the softmax REDUCTION ORDER (block-cooperative on CUDA, 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("cuda fp8 KV paged-attention read matches the CPU read") { + if (!HasCuda()) { + MESSAGE("SKIPPED: no CUDA backend in this build/host — the CUDA fp8 KV " + "paged-attention read parity gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + 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); +} + +// ─── G5 ───────────────────────────────────────────────────────────────────── +// fp8_e5m2 stays a NAMED later brick (spec W5) on CUDA exactly as on CPU — it +// must be refused, never silently mis-stored through the e4m3 converter. +TEST_CASE("cuda fp8 KV store refuses e5m2 (later brick)") { + if (!HasCuda()) { + MESSAGE("SKIPPED: no CUDA backend in this build/host — the CUDA e5m2 refusal " + "gate did NOT run"); + return; + } + Backend& gpu = vt::GetBackend(DeviceType::kCUDA); + 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}); + CHECK_THROWS_AS(vt::ReshapeAndCacheFp8(gq, gk, gk, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E5M2, + 0.01f, 0.01f), + std::runtime_error); + gpu.Free(dk); + gpu.Free(dkc); + gpu.Free(dvc); + gpu.Free(ds); + gpu.DestroyQueue(gq); +} From 3e9c7712bd2a0ce3c399b4b74fc8d6235d4e2cba Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Fri, 21 Aug 2026 20:22:23 +0000 Subject: [PATCH 2/2] fix(KV-FP8): the W2 anchors point at an identity template, and G5 could not fail for any CUDA defect (#1593, #1636) Repairs the six findings of the fresh review of #1593 W2. Four are records, two are gates, and the two merge-blocking ones land in the commit message, which `squash_merge_commit_message = PR_BODY` makes permanent. THE ANCHOR. Four new sites and the pull request body cited `scaled_vec_conversion` at `quant_utils.cuh:302-308`. Verified against the pinned tree at `5559679229bc961848b121ccdeaa8fa5d79bec98`: lines 301-305 are the GENERIC primary template, whose body is `return x;`, and 307-314 are the `// fp8 -> half` `` specialization. The `` one is at `:419-429`. Same class, same change: `Fp8KVCacheDataType` was cited at `dtype_fp8.cuh:9-13`, the `#include ` guard; the enum is at `:15-19`. An anchor is how the next reader checks a port against the oracle, and one landing on an identity template invites the conclusion that the port is unfaithful. Three W1 copies of the same wrong anchor (`include/vt/fp8_kv.h:92`, `include/vt/ops.h:1129`, `src/vt/cpu/cpu_paged_attn.cpp:164`) are outside this change's authority and are filed as #1636, owned by `KV-FP8` and listed under the spec's `## Owed`. THE TRANSLATION UNITS COMPILE. The body said they do not. CI job `cuda-fat-build` built both changed files for `80;86;87;89;90a;100a;103a;110; 120a;121a` under `-Werror=all-warnings` and PASSED on `4d71e776e`, run 32495320287, job 96812232428. What stays true is that nothing has been EXECUTED on a device, because that job configures `-DVLLM_CPP_BUILD_TESTS=OFF`. The spec now separates the two states. G5 COULD NOT FAIL FOR ANY CUDA DEFECT, and the fix needed three mutations to state correctly. The e5m2 refusal exists in the op wrapper, the CPU kernel and the CUDA kernel. Deleting the CUDA one on a CPU build gives `ninja: no work to do` and leaves the file 7/10 SUCCESS. Deleting the wrapper's leaves W1's `test_ops_fp8_kv_cache` GREEN at 8/511, because execution falls through to the CPU kernel's check and W1's case asserts `CHECK_THROWS_AS` on `std::runtime_error` rather than a message; only deleting both turns it red (7/8, 510/511). A layered refusal needs an assertion that names its layer, so G5 now resolves the registered provider with `GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kCUDA)`, calls it directly, and requires both `cuda reshape_and_cache_fp8` and `fp8_e5m2` in the message. TWO UNGATED INSTANTIATIONS. `LaunchPagedFp8Out<__nv_bfloat16, __nv_bfloat16>` is what a served bf16 model takes and what #1574's subject runs, and G4 exercised only ``; the store's `DType::kF16 -> __half` arm was untouched while the wrapper's `IsFloat()` admits it. G4b and the f16 leg of G3 close both. Neither has a red-first mutation: they are device cases and this session has no `nvcc` and no device, which the first mutation above measures rather than assumes. They are listed under `## Owed`. TWO PROSE OVERSTATEMENTS. The store is elementwise-identical, not 1:1: upstream vectorizes the contiguous-heads arm through `vectorize_with_alignment` (`cache_kernels.cu:360-363`) and this is a scalar strided loop over the same elements in the same order, which is a bandwidth difference W4 owns. And the read is the CPU codec for every one of the 254 finite e4m3 codes, not "line for line": on `0x7F` and `0xFF` the CPU returns `quiet_NaN()` (`0x7FC00000`) and the device returns `CUDART_NAN_F` (`0x7FFFFFFF`), which no gate here can see because a NaN compares unequal to itself. Focused gate after: `test_cuda_fp8_kv_cache` 7 cases / 10 assertions SUCCESS, Release CPU build with `-Wall -Wextra -Werror`. Siblings unmoved: `test_ops_fp8_kv_cache` 8/511, `test_ops_reshape_cache` 12/192, `test_ops_paged_attn` 14/1646, `test_ops_paged_attn_dtype` 3/172. Every mutation restored against a pre-taken sha256. Issue: https://github.com/mudler/vllm.cpp/issues/1593 Anchor debt: https://github.com/mudler/vllm.cpp/issues/1636 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/issue-index.md | 1 + .agents/specs/fp8-kv-cache.md | 127 ++++++++--- src/vt/cuda/cuda_cache.cu | 9 + src/vt/cuda/cuda_paged_attn.cu | 25 ++- tests/vt/test_cuda_fp8_kv_cache.cpp | 337 ++++++++++++++++++++++------ 5 files changed, 400 insertions(+), 99 deletions(-) diff --git a/.agents/issue-index.md b/.agents/issue-index.md index c6caaa3f5..59ccd6586 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -528,3 +528,4 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1454](https://github.com/mudler/vllm.cpp/issues/1454) | `SPEC-MTP-GGUF` | **`test_qwen3_5_gguf_mtp.cpp` reported `Status: SUCCESS!` with `assertions: 0` on every CI run, and its one arithmetic guarantee was a tautology.** Both cases opened `if (path == nullptr) return;` on `VLLM_MTP_GGUF_MODEL`, and a bare `return` from a doctest case is a PASS: re-derived on a clean Release build at `947e5f648`, unset, the file printed `test cases: 2 \| 2 passed \| 0 failed \| 0 skipped`, `assertions: 0`, `Status: SUCCESS!`, exit 0, and printed nothing else. The variable is set nowhere in `.github/workflows/`, so that was the state of every run. Second defect in the same file: the comment at `:52` stated `num_hidden_layers + depth == block_count` and the line under it asserted `CHECK(c.num_hidden_layers > 0)`, true of every valid model. MEASURED, not argued: mutating `src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp:889` to `c.num_hidden_layers = block_count;` compiled clean and left the file at 2/2 cases, 0 assertions, `SUCCESS!`, exit 0. FIXED IN FLOW. The invariant is now pinned **HERMETICALLY** on KV-only synthetic GGUFs carrying no weight bytes, so CI checks it every run rather than never - 65/1 (the shipped Qwen3.8-27B pair), 25/1 (the Qwen3.5-2B reference this suite was developed against) and 28/3, the third arm separating `- nextn` from `- 1` - plus a head-less arm asserting the key is NOT published, which is the half `NumMtpLayers` cannot express because it answers 1 for an absent key. The two env-gated cases stay, now skipping with a `MESSAGE` naming the variable as `test_gguf_mmproj_reach.cpp` does, and the live one re-derives the invariant from the file's own `block_count` kv. Unset 4 cases / 18 assertions / `SUCCESS!` / rc 0; live on `Qwen3.8-27B-Q4_K_M.gguf` 4 / 38 / `SUCCESS!` / rc 0. Both mutants now red (9/18 and 5/18, exit 1), compiled clean, restored against a pre-taken sha256. **The production line is CORRECT and was not touched**: `block_count - nextn` landed `1a4db5c3c`, the `mtp_num_hidden_layers` republication `493327b4e`. Related but distinct: [#821](https://github.com/mudler/vllm.cpp/issues/821) W2 (`0adeb8b0e`) pins the same arithmetic for the 27B artifact on a committed manifest in `tests/vllm/models/test_qwen38_27b_gguf_manifest.cpp`, and that gate DOES catch both mutants - so the invariant was not globally unpinned, it was unpinned in this row's own file | bug | | [#1434](https://github.com/mudler/vllm.cpp/issues/1434) | `GATE-DOC-CHECKPOINT-STATES` | **`scripts/check-doc-checkpoint.py` could not see `PARTIAL`, so 118 state cells could move with no gate observing them.** `STATES` (`:56-66`) is the whole definition of what a lifecycle state IS for the gate that enforces AGENTS.md's `docs/STATUS.md` / `docs/BENCHMARKS.md` / spec `## Now` triple, and `row_states` drops any row it cannot match. `lifecycle_moves` and `moved_rows` then iterate the AFTER map, so leaving the matched set is silent by construction. Re-derived at `947e5f648` (the report measured `63d87805c`): `PARTIAL` **118** cells and `ANCHOR-BACKFILL` **73**, against `DONE` 77 and `BLOCKED` 9 — `PARTIAL` is the second most used state in the matrices and the gate was blind to it. Over the seven tables `ROW_TABLES` actually reads, the resolved population goes from **153 rows to 226**, a 47.7 % widening. Two of the transitions the report names behave differently from its description, measured with scratch commits at `947e5f648` on an unmodified checker: `READY -> PARTIAL` rc **0** and `PARTIAL -> READY` rc **0** are the real blind spots, while the report's suggested `PARTIAL -> ACTIVE` already reds — by accident, reporting **`added as ACTIVE`** for a row that has existed for months, because it is absent from the BEFORE map. FIXED IN FLOW for `PARTIAL` only. **`ANCHOR-BACKFILL` is deliberately excluded**: `.agents/feature-matrix.md:14-17` defines it as a property of the RECORD (*a legacy implemented row without exact code, test and real-spec anchors*), `docs/STATUS.md` carries no such term and would have nothing true to write on a `DONE <-> ANCHOR-BACKFILL` move, and `REQUIRED["lifecycle"]` cannot demand the spec's `## Now` alone — so admitting it would demand a public-document edit with nothing to say, which is the exact shape `check-doc-checkpoint.py:4-17` records as the reason the file was rewritten (16 of 20 red CI runs, six hardcoded escape hatches). One row's resolved state moves and the move is a REPAIR: `KV-BLOCK-POOL` says `` `PARTIAL` (not `DONE`) `` in its prose and the last-match heuristic believed the parenthesis, resolving `DONE`. No pinned counter moves — `check-gate-commands.py` has its own `GATED_STATES` and `RUNNABLE_BASELINE` is keyed on matrix rows, `UNOWNED_HIGH_WATER` is unmoved because this row names an owner, and no matrix row or public document changes — which was measured, not assumed, because this is the [#1376](https://github.com/mudler/vllm.cpp/issues/1376) ratchet shape. Remainder listed under `## Owed` in [doc-checkpoint-lifecycle-states.md](specs/doc-checkpoint-lifecycle-states.md): `ANCHOR-BACKFILL` moves, `.agents/sglang-matrix.md` never entering `ROW_TABLES`, a row that leaves the matched set entirely, and a new row added directly as `PARTIAL` | bug | | [#1593](https://github.com/mudler/vllm.cpp/issues/1593) | `KV-FP8` | **`KV-FP8` W2 and W3: the CUDA fp8 KV store, its paged-attention read, and the runner integration.** W1 landed the CPU half (`vt::ReshapeAndCacheFp8`, the read dequant in CPU paged attention, `vllm::v1::ParseCacheDType`) and left W2/W3/W4 `later`. The issue is now the critical path of benchmark campaign [#1574](https://github.com/mudler/vllm.cpp/issues/1574), whose subject `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` declares `kv_cache_quant_algo: "FP8"` and carries ZERO `k_scale`/`v_scale` tensors, so every published profile serves it with `--kv-cache-dtype fp8` and no cell can be served correctly without this. **W2 IS LANDED HERE**: the CUDA fp8-e4m3 store (`src/vt/cuda/cuda_cache.cu`), the fp8 dequant on the CUDA paged-attention read (`src/vt/cuda/cuda_paged_attn.cu` `LoadKv` + `LaunchPagedFp8`), the removal of the two W1 device-class refusals that made the CUDA arm unreachable however well it was registered, and a named CPU-or-CUDA refusal for the READ because it rides ADDITIVE `PagedAttentionArgs` fields on an op `kMETAL`/`kROCM` already register for the FLOAT path — without which an fp8 cache would be read as that backend's float dtype and return silent garbage. Gate `tests/vt/test_cuda_fp8_kv_cache.cpp`, RED-first on the provider-routing case. **The device half of that gate is UNEXECUTED and the CUDA TUs are UNCOMPILED**: the implementing session had no `nvcc` and no device, and says so under `## Owed` in [fp8-kv-cache.md](specs/fp8-kv-cache.md) together with the reachability debt — nothing calls the fp8 KV path from a production entry point on either backend, which is **W3's** wiring (half-sized KV blocks, `--kv-cache-dtype` threading, the checkpoint scale path including this checkpoint's scales-absent case). W3, W4, the Metal/ROCm arms and fp8_e5m2 remain owed | feature | +| [#1636](https://github.com/mudler/vllm.cpp/issues/1636) | `KV-FP8` | **`KV-FP8` W1's three read-side comments anchor `scaled_vec_conversion` at `quant_utils.cuh:302-308`, which at pin `555967922` is the IDENTITY primary template plus the header of the fp8->HALF specialization.** Lines 301-305 are `template ... { return x; }` and 307-314 are the `` conversion; the `` one the comments describe is at `:419-429` under the `// fp8 -> float` label at `:418`. Sites, all landed by W1 and all outside the W2 change's authority: `include/vt/fp8_kv.h:92`, `include/vt/ops.h:1129`, `src/vt/cpu/cpu_paged_attn.cpp:164`. W2 ([#1593](https://github.com/mudler/vllm.cpp/issues/1593), PR [#1606](https://github.com/mudler/vllm.cpp/pull/1606)) copied the same wrong anchor into four new places and CORRECTED all four there; these three are filed rather than fixed in flow. Same shape, second anchor: `Fp8KVCacheDataType` is cited at `dtype_fp8.cuh:9-13`, which is the `#include ` guard -- the enum is at `:15-19` (`include/vt/fp8_kv.h:5`, `:30`). Third, a different kind: `.agents/engine-matrix.md` and `.agents/quantization-matrix.md` both say the W2 CUDA translation units are UNCOMPILED, and CI job `cuda-fat-build` built them for ten architectures under `-Werror=all-warnings` and PASSED on `4d71e776efc18cb5e61a26e642ddad8de5339134` (run 32495320287, job 96812232428). What stays true is that nothing has been EXECUTED on a device, because that job configures `-DVLLM_CPP_BUILD_TESTS=OFF`; both clauses need the narrower statement. An upstream anchor is how the next reader checks a port against the oracle, and one that lands on a `return x;` primary template invites the conclusion that the port is unfaithful. Listed under `## Owed` in [fp8-kv-cache.md](specs/fp8-kv-cache.md) | bug | diff --git a/.agents/specs/fp8-kv-cache.md b/.agents/specs/fp8-kv-cache.md index a61187db7..c75306498 100644 --- a/.agents/specs/fp8-kv-cache.md +++ b/.agents/specs/fp8-kv-cache.md @@ -57,9 +57,9 @@ store/read kernels are vLLM's own csrc: **`FP8 = Quantize(HP / scale)`; `Dequant(FP8) * scale = HP`.** `k_scale`/ `v_scale` are `[1]` (per-tensor) or `[num_heads]` (per-head, `kv_scale_stride`, `:365-401`). Store dtype `cache_t = uint8_t`; the fp8 *interpretation* is the - `Fp8KVCacheDataType` template param (`csrc/attention/dtype_fp8.cuh:9-13`). + `Fp8KVCacheDataType` template param (`csrc/attention/dtype_fp8.cuh:15-19`). - **Read (dequant).** The fp8 attention read multiplies back by the scale: - `scaled_vec_conversion` (`quant_utils.cuh:302-308`) = + `scaled_vec_conversion` (`quant_utils.cuh:419-429`) = `half_to_float(fp8_to_half(byte)) * scale`. Consumed by the FA/flashinfer fp8 paths and the reference `test_cache.py`'s `convert_fp8`. - **Memory accounting.** An fp8 KV element is 1 byte vs bf16's 2 → the KV block @@ -85,9 +85,9 @@ existing copies), the fp8 store op, the read dequant, and the config parse. W1 (this change; CPU-only, `-Werror`): - `include/vt/fp8_kv.h` (NEW) — `Fp8KVCacheDataType` enum (mirror - `dtype_fp8.cuh:9-13`) + `F8E4M3ToF32`/`F32ToF8E4M3`/`StoreKvFp8E4M3`/ + `dtype_fp8.cuh:15-19`) + `F8E4M3ToF32`/`F32ToF8E4M3`/`StoreKvFp8E4M3`/ `LoadKvFp8E4M3` (bit-match the landed codecs; the store/load scale convention - from `quant_utils.cuh:296-308`). + from `quant_utils.cuh:296-300`). - `include/vt/ops.h` — `OpId::kReshapeAndCacheFp8`, `ReshapeAndCacheFp8Fn`, `vt::ReshapeAndCacheFp8` decl, and the additive `PagedAttentionArgs` `kv_cache_dtype`/`k_scale`/`v_scale` fields (default kAuto/1.0 → every existing @@ -134,14 +134,19 @@ threading + CLI); fp8_e5m2 compute; per-head scales; the Metal and ROCm arms. read `v_scale` diverges > 0.05 from the baseline; an auto (no-dequant) read of an fp8 cache is refused. No sibling regressions (reshape 12/12, paged 14/14). - **Correctness (W2, provider routing — the CPU leg):** `test_cuda_fp8_kv_cache` - 6 cases / 10 assertions GREEN on a CPU-only build. RED-first proven: with the - W1 device-class guards in place the suite reports 10 assertions / 6 failed for - the store guard plus the read guard, naming both refusal strings. + 7 cases / 10 assertions GREEN on a CPU-only build. Only G1 and G1b assert + there; the four device cases skip with a MESSAGE naming what did not run. + RED-first proven: with the W1 device-class guards in place the suite reports + 10 assertions / 6 failed for the store guard plus the read guard, naming both + refusal strings. - **Correctness (W2, device — UNEXECUTED, see `## Owed`):** the store byte gate - (zero tolerance, f32 + bf16 source, a padded slot), the paged-read parity gate - (decode + prefill, NMSE < 1e-6 and worst < 1e-3 vs the CPU arm) and the - registration gate need a CUDA build and a device. Neither was available to the - implementing session, so **the CUDA TUs in this change are UNCOMPILED**. + (zero tolerance; f32, bf16 and f16 sources; a padded slot), the paged-read + parity gate in both the f32 and the bf16 query/output instantiation (decode + + prefill), the CUDA-kernel e5m2 refusal and the registration gate all need a + device. **The CUDA translation units DO compile** — CI job `cuda-fat-build` + builds them for ten architectures under `-Werror=all-warnings` — but that job + configures `-DVLLM_CPP_BUILD_TESTS=OFF`, so nothing in this file has ever + been executed on a device. - **Later:** the real memory-halving e2e (KV blocks ~2× on a gate model at token parity) is the binding gate and is DGX-blocked (docs/BENCHMARKS PENDING). @@ -170,12 +175,17 @@ ORACLE: every W2 gate compares CUDA to the landed CPU kernels, never to a fresh reference. **Store** (`src/vt/cuda/cuda_cache.cu`, `ReshapeAndCacheFp8KernelCuda` + -`ReshapeAndCacheFp8Kernel`, registered for `DeviceType::kCUDA`). A 1:1 port -of the fp8 branch of `reshape_and_cache_flash_kernel` +`ReshapeAndCacheFp8Kernel`, registered for `DeviceType::kCUDA`). An +ELEMENTWISE-IDENTICAL port of the fp8 branch of `reshape_and_cache_flash_kernel` (`cache_kernels.cu:314-401`) + `CopyWithScaleOp` (`:241-252`), restricted to upstream's `is_contiguous_heads && kv_scale_stride == 0` arm (`:352-366`) — which is the only arm the op's wrapper admits, because the vt cache is the NHD -unbind slice and `ReshapeAndCacheFp8` takes two scalar scales. The converter is +unbind slice and `ReshapeAndCacheFp8` takes two scalar scales. It is NOT +instruction-identical, and calling it a 1:1 port would overstate it: upstream's +contiguous-heads arm moves the row through `vectorize_with_alignment` +(`:360-363`, `VEC_SIZE` 8 for a 2-byte source and 4 for f32), and ours is a +scalar strided loop over the same elements in the same order. Same bytes out, +fewer bytes per instruction; W4 owns closing the bandwidth gap. The converter is upstream's own `__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)` (`quant_utils.cuh:497-503`) — a true DIVIDE, not the activation path's hoisted reciprocal multiply — and its byte-for-byte equality to the CPU software codec @@ -187,13 +197,29 @@ f32/f16/bf16, the same set the CPU `LoadSrcF32` serves. `Load`: inert on the f32/bf16 arms (they forward to `Load` unchanged, so every existing caller reads the same bytes in the same order), and on `uint8_t` it is `Fp8E4M3ToF32Dev(byte) * scale` — upstream's `scaled_vec_conversion` (`quant_utils.cuh:302-308`), written as the SAME ARITHMETIC as +uint8_t>` (`quant_utils.cuh:419-429`), written as the SAME ARITHMETIC as `vt::F8E4M3ToF32` so CUDA==CPU on the read is a property of the source rather than of a measurement this session could not take. `PagedAttentionKernel` and `PagedFlashKernel` gain `k_scale`/`v_scale`; `LaunchPagedByKv` keys on `args.kv_cache_dtype` (never on the storage dtype, which is a bare `kI8` byte) and routes to `LaunchPagedFp8`. +Same-arithmetic holds for all 254 FINITE e4m3 codes and NOT for the two NaN +ones. On `0x7F`/`0xFF` the CPU returns `std::numeric_limits::quiet_NaN()` +(`0x7FC00000`) and the device returns `CUDART_NAN_F` (`0x7FFFFFFF`): both quiet, +both propagating, different payload. **No gate in this file can see that**, +because a NaN compares unequal to itself, so a byte or NMSE comparison fails +on any payload rather than on the wrong one. It is recorded rather than +measured, and reaching it needs a non-finite input in the first place — `__NV_SATFINITE` +clamps an out-of-range magnitude to `0x7E`/`0xFE`, so a store of a finite +`hp / scale` never writes a NaN code. + +`LaunchPagedFp8`/`LaunchPagedFp8Out` are templated on `TQ, Tout` over +`{float, __nv_bfloat16}`, so the fp8 read has FOUR instantiations and the gate +exercises two of them: `` (G4) and `<__nv_bfloat16,__nv_bfloat16>` +(G4b), the one a served model takes. The two mixed-dtype instantiations are +compiled and ungated. + **Scope of the read, argued.** Only the two correctness-grade kernels serve fp8: the tiled flash prefill and the block decode. That is what the existing ladder already implies — the WMMA prefill kernels stage `__nv_bfloat16` fragments, the @@ -217,20 +243,69 @@ return silent garbage. `src/vt/ops.cpp` therefore keeps an explicit CPU-or-CUDA list there whose message names the missing part, and `tests/vt/test_cuda_fp8_kv_cache.cpp` gates it on both `kMETAL` and `kROCM`. +**Where a refusal is asserted decides what it proves.** The e5m2 refusal exists +in THREE places — the `ReshapeAndCacheFp8` op wrapper (`src/vt/ops.cpp`), the +CPU kernel (`src/vt/cpu/cpu_cache.cpp`) and `ReshapeAndCacheFp8KernelCuda` — and +only the third is a CUDA guarantee. The wrapper check is device-independent and +sits ABOVE both the device checks and `GetOp`, so a case that calls +`vt::ReshapeAndCacheFp8` with device tensors and asserts a bare throw cannot +distinguish any of the three. G5 was written that way. MEASURED, not argued: +deleting the CUDA kernel's `VT_CHECK` on a CPU build produces `ninja: no work to +do` and leaves the file 7 cases / 10 assertions `SUCCESS!`; deleting the op +wrapper's leaves W1's `test_ops_fp8_kv_cache` GREEN at 8/511, because execution +then falls through to the CPU kernel's own check and W1's case asserts +`CHECK_THROWS_AS(..., std::runtime_error)` rather than a message. Only deleting +the wrapper AND the CPU kernel check together turns W1 red (7/8 cases, 510/511 +assertions). So what W1 pins is "e5m2 is refused somewhere on the CPU path", and +a layered refusal needs an assertion that NAMES its layer. + +G5 is now written the only way that reaches the kernel guard: resolve the +registered provider with `GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kCUDA)`, +call it directly, and require the message to contain both +`cuda reshape_and_cache_fp8` and `fp8_e5m2`, which no other layer produces. It +then calls the same pointer with e4m3 to prove the guard refuses one kind rather +than disabling the kernel. Bypassing the wrapper is deliberate and is the point +of the case; the production path still goes through it, and G1/G1b/G2 gate that. + ## Owed - **The W2 device gates are UNEXECUTED** (#1593). `tests/vt/test_cuda_fp8_kv_cache.cpp` - G2 (provider registration, CUDA build), G3 (store byte parity, f32 + bf16), G4 - (paged-read parity, decode + prefill) and G5 (e5m2 refusal) all need a CUDA - toolkit and a device; the implementing session had NEITHER — `nvcc` is absent - on `mudler-ubuntu-box` and the GPU fleet was leased for the #1574 campaign. - **The CUDA translation units in this change have therefore never been - compiled**, let alone run. G1/G1b (provider routing and the Metal/ROCm refusal) - are the only cases that executed, and they run on the CPU leg. The first CUDA - build or `rc` lease that touches this row must run - `ctest -R test_cuda_fp8_kv_cache` and record the result here before W2 counts - as measured. Until then the wave table's `DONE` means "landed and gated", not - "measured on hardware". + G2 (provider registration, CUDA build), G3 (store byte parity over the f32, + bf16 and f16 sources), G4 and G4b (paged-read parity, decode + prefill, for + the f32 and the bf16 query/output instantiation) and G5 (the CUDA kernel's own + e5m2 refusal, reached through the registered provider) all need a device. + Neither the implementing session nor the repair session had one — `nvcc` is + absent on `mudler-ubuntu-box` and the GPU fleet was leased for the #1574 + campaign — so none of them has run and none has a red-first mutation. + **They do COMPILE.** CI job `cuda-fat-build` (`.github/workflows/ci.yml:773`) + built both changed CUDA translation units for `80;86;87;89;90a;100a;103a;110; + 120a;121a` under `-Werror=all-warnings` and passed on `4d71e776e` + ([run 32495320287](https://github.com/mudler/vllm.cpp/actions/runs/32495320287/job/96812232428)), + but it configures `-DVLLM_CPP_BUILD_TESTS=OFF`, so it never links or runs this + file. G1/G1b (provider routing and the Metal/ROCm refusal) are the only cases + that executed, and they run on the CPU leg. The first `rc` lease that touches + this row must run `ctest -R test_cuda_fp8_kv_cache` and record the result here + before W2 counts as measured. Until then the wave table's `DONE` means "landed + and gated", not "measured on hardware". +- **A CPU-only gate cannot see a CUDA defect, and this one was PROVEN blind.** + The W2 review deleted the entire production call site for the CUDA fp8 read + and the focused gate stayed 100% green; inverting the CUDA store's scale + direction produced `ninja: no work to do`, because a CPU build's + `compile_commands.json` contains ZERO `.cu` translation units — 1046 entries + in the review's configure and 1030 in the repair session's, none of them CUDA. + That is the honest consequence of the state above, not a defect in the gate, + and it is why the device run is the binding evidence. +- **Three W1 read-side citations point at the WRONG upstream lines** and are + outside the W2 change's authority to edit: `include/vt/fp8_kv.h:92`, + `include/vt/ops.h:1129` and `src/vt/cpu/cpu_paged_attn.cpp:164` each cite + `scaled_vec_conversion` at `quant_utils.cuh:302-308`, which at + pin `555967922` is the generic primary template plus the header of the + `` (fp8 -> half) specialization. The `` one + is at `:419-429`. W2 corrected its own four copies; these three are owed and + tracked by [#1636](https://github.com/mudler/vllm.cpp/issues/1636), which also + owes the "CUDA TUs are UNCOMPILED" clause in + [`.agents/engine-matrix.md`](../engine-matrix.md) and + [`.agents/quantization-matrix.md`](../quantization-matrix.md). - **Nothing reaches the fp8 KV path from a production entry point yet**, on either backend. `vt::ReshapeAndCacheFp8` and `PagedAttentionArgs::kv_cache_dtype` have no caller outside their tests; W1 landed in that state and W2 does not diff --git a/src/vt/cuda/cuda_cache.cu b/src/vt/cuda/cuda_cache.cu index 764110b2c..e6e069b14 100644 --- a/src/vt/cuda/cuda_cache.cu +++ b/src/vt/cuda/cuda_cache.cu @@ -110,6 +110,15 @@ void ReshapeAndCacheKernelCuda(Queue& q, const Tensor& k, const Tensor& v, Tenso // (`:367-400`) is a named later brick (spec W5), and per-head scales cannot // reach here because ReshapeAndCacheFp8 takes two scalars. // +// ELEMENTWISE-IDENTICAL, NOT INSTRUCTION-IDENTICAL. Upstream's contiguous-heads +// arm moves the row through `vectorize_with_alignment` (`:360-363`, +// VEC_SIZE 8 for a 2-byte source and 4 for f32), which converts the same +// elements in the same order under a vectorized load/store. The loop below is a +// SCALAR strided one, so it writes the same bytes and reads the same inputs +// while moving them one at a time. That is a bandwidth difference, not a +// numerical one, and W4 — which owns the memory/throughput measurement — owns +// closing it. Do not read "ported" here as "the same instructions". +// // THE CONVERTER IS UPSTREAM'S OWN, and its equality to the CPU codec is already // MEASURED. `fp8::scaled_convert` is // `__nv_cvt_float_to_fp8(a / scale, __NV_SATFINITE, __NV_E4M3)` diff --git a/src/vt/cuda/cuda_paged_attn.cu b/src/vt/cuda/cuda_paged_attn.cu index c55b9df50..dad948917 100644 --- a/src/vt/cuda/cuda_paged_attn.cu +++ b/src/vt/cuda/cuda_paged_attn.cu @@ -140,16 +140,27 @@ __device__ inline void Store(__nv_bfloat16* p, int64_t i, float v) { p[i] = __fl // caller reads exactly the bytes, in exactly the order, it read before. // // The fp8 arm mirrors upstream's attention-side dequant -// `scaled_vec_conversion` (quant_utils.cuh:302-308): fp8 byte -> +// `scaled_vec_conversion` (quant_utils.cuh:419-429): fp8 byte -> // float, then multiply by the scale, i.e. `Dequant(FP8) * scale = HP` (the // convention at :296-300). It is written as the SAME ARITHMETIC as the W1 CPU // codec vt::F8E4M3ToF32 (include/vt/fp8_kv.h) rather than as the hardware // `__nv_cvt_fp8_to_halfraw`, because W1 is this wave's oracle and sharing the // decode makes CUDA==CPU on the read a property of the source rather than a -// measurement. The two agree in any case: every one of the 256 e4m3 codes is -// exactly representable in fp16, so upstream's fp8->half->float round trip is -// lossless. `std::ldexp(mantissa, exp - 7)` on a float IS `ldexpf`, so this is -// the CPU codec line for line. +// measurement. The two agree in any case: every one of the 254 FINITE e4m3 codes +// is exactly representable in fp16, so upstream's fp8->half->float round trip is +// lossless. `std::ldexp(mantissa, exp - 7)` on a float IS `ldexpf`, so each of +// those 254 decodes to the same f32 bits as the CPU codec. +// +// THE TWO NaN CODES are not identical, and the suite cannot see it. On 0x7F and +// 0xFF the CPU returns `std::numeric_limits::quiet_NaN()` (0x7FC00000) +// and this returns `CUDART_NAN_F` (0x7FFFFFFF): both are quiet NaNs and both +// propagate the same way, but the PAYLOAD differs. No gate here +// distinguishes them, because a NaN compares unequal to everything including +// itself, so a byte or NMSE comparison fails on ANY payload rather than on the +// wrong one. It is recorded here rather than measured. Reaching it also needs a +// non-finite input: `__NV_SATFINITE` clamps an out-of-range magnitude to the max +// finite code (0x7E/0xFE), so a store of a finite `hp / scale` never writes +// 0x7F/0xFF. __device__ __forceinline__ float Fp8E4M3ToF32Dev(uint8_t byte) { const uint32_t sign = static_cast(byte >> 7) & 0x1U; const uint32_t exp = static_cast(byte >> 3) & 0xFU; @@ -2913,7 +2924,7 @@ void LaunchPaged(cudaStream_t s, Tensor& out, const Tensor& query, const Tensor& // ─── fp8 KV-cache READ dispatch (KV-FP8 W2, #1593) ───────────────────────── // TKV is `uint8_t`: the cache pages are 1-byte fp8-e4m3 (DType::kI8) and each // read is dequantized as Dequant(fp8) * k_scale|v_scale inside LoadKv, mirroring -// upstream's `scaled_vec_conversion` (quant_utils.cuh:302-308). +// upstream's `scaled_vec_conversion` (quant_utils.cuh:419-429). // // SCOPE, argued rather than assumed. Only the two CORRECTNESS-GRADE kernels are // reachable from here — the tiled flash prefill and the block decode — and that @@ -2974,7 +2985,7 @@ void LaunchPagedByKv(cudaStream_t s, Tensor& out, const Tensor& query, const Ten const Tensor& query_start_loc, const PagedAttentionArgs& args) { // fp8 KV cache: the STORAGE dtype is a raw byte (kI8) and the INTERPRETATION // travels in args.kv_cache_dtype, exactly as upstream carries cache_t=uint8_t - // plus a KV_DTYPE template parameter (dtype_fp8.cuh:9-13). Key on the + // plus a KV_DTYPE template parameter (dtype_fp8.cuh:15-19). Key on the // interpretation, never on the storage dtype: a kI8 tensor with kAuto is not // an fp8 cache, and the op wrapper already refuses that pair. if (args.kv_cache_dtype == Fp8KVCacheDataType::kFp8E4M3) { diff --git a/tests/vt/test_cuda_fp8_kv_cache.cpp b/tests/vt/test_cuda_fp8_kv_cache.cpp index 1d9821b69..ff69974a7 100644 --- a/tests/vt/test_cuda_fp8_kv_cache.cpp +++ b/tests/vt/test_cuda_fp8_kv_cache.cpp @@ -8,14 +8,14 @@ // 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:302-308 +// 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) // -// FIVE gates, and they do not all run in the same build: +// The gates, and they do not all run in the same build: // // G1 (runs in every build WITHOUT the CUDA backend, i.e. the x86 CI leg): the // W1 device-class refusal is GONE. W1 hard-refused any non-CPU queue inside @@ -24,19 +24,25 @@ // CUDA kernel can be reached however well it is registered, so this case is // the RED-first assertion for the whole wave and the one gate a host with no // CUDA toolkit can actually execute. +// G1b (every build): the fp8 READ is refused by name on kMETAL and kROCM. The +// check fires in the op wrapper, so no Metal or ROCm backend need be linked. // G2 (CUDA build): the CUDA providers are REGISTERED for the fp8 store and the // paged read — the shared-seam reach check. vt::ops.cpp dispatches through // GetOp(OpId, DeviceType) and nothing else can select a kernel, so a // registered provider IS the production path. // G3 (CUDA device): STORE parity — the CUDA store writes the SAME BYTES as the -// CPU store, zero tolerance, over f32 and bf16 sources, with a padded (-1) -// slot and a strided unbind-slice cache. +// 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 (CUDA device): READ parity — paged attention over identical fp8 cache // bytes, CUDA vs CPU, in both the decode and the prefill shape (the two -// kernels the fp8 arm routes to). -// G5 (CUDA device): fp8_e5m2 stays refused on CUDA as it is on CPU. +// kernels the fp8 arm routes to), for an f32 query/output... +// G4b (CUDA device): ...and for the bf16 query/output a served model actually +// runs, which is a DIFFERENT template instantiation of the same launcher. +// G5 (CUDA device): fp8_e5m2 stays refused BY THE CUDA KERNEL, reached through +// the registered provider. The op wrapper's own e5m2 refusal is device- +// independent and is gated by W1 at tests/vt/test_ops_fp8_kv_cache.cpp:342. // -// G3/G4/G5 SKIP CLEANLY when no CUDA backend is present, which is the house +// G3/G4/G4b/G5 SKIP CLEANLY when no CUDA 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 @@ -341,14 +347,24 @@ TEST_CASE("cuda fp8 KV store is byte-identical to the CPU store") { gpu.DestroyQueue(gq); } -// bf16 source arm of the same store — the dtype vLLM actually resolves for a -// model (AGENTS.md "Inherit vLLM defaults"). Upstream widens bf16 to f32 BEFORE -// the divide (quant_utils.cuh:482-489, `__bfloat162float(a) / scale`) and the -// CPU LoadSrcF32 does the same, so the two must still agree byte for byte. -TEST_CASE("cuda fp8 KV store is byte-identical to the CPU store (bf16 source)") { +// 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 `ReshapeAndCacheFp8KernelCuda` 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 CUDA +// kernel through `Fp8SrcToF32` 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("cuda fp8 KV store is byte-identical to the CPU store (bf16 and f16 sources)") { if (!HasCuda()) { - MESSAGE("SKIPPED: no CUDA backend in this build/host — the bf16-source fp8 KV " - "store parity gate did NOT run"); + MESSAGE("SKIPPED: no CUDA backend in this build/host — the bf16/f16-source fp8 " + "KV store parity gate did NOT run"); return; } Backend& gpu = vt::GetBackend(DeviceType::kCUDA); @@ -359,54 +375,65 @@ TEST_CASE("cuda fp8 KV store is byte-identical to the CPU store (bf16 source)") 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 kb(kf.size()), vb(vf.size()); - for (size_t i = 0; i < kf.size(); ++i) { - kb[i] = vt::F32ToBF16(kf[i]); - vb[i] = vt::F32ToBF16(vf[i]); - } std::vector slots = {3, 0, 2, 1}; const float k_scale = 0.007f, v_scale = 0.003f; - std::vector kc_ref(cache_elems, 0); - std::vector vc_ref(cache_elems, 0); - Tensor ck = Host(kb.data(), DType::kBF16, {nt, H, D}); - Tensor cv = Host(vb.data(), DType::kBF16, {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, DType::kBF16, {nt, H, D}); - Tensor gv = Dev(dv, DType::kBF16, {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); + // 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]); + } - gpu.Free(dk); - gpu.Free(dv); - gpu.Free(dkc); - gpu.Free(dvc); - gpu.Free(ds); + 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); } @@ -418,7 +445,7 @@ TEST_CASE("cuda fp8 KV store is byte-identical to the CPU store (bf16 source)") // // The dequant itself is bit-identical by construction: the CUDA kernel decodes // e4m3 with the same arithmetic as vt::F8E4M3ToF32 and multiplies by the same -// per-tensor scale (quant_utils.cuh:302-308). The only divergence available is +// per-tensor scale (quant_utils.cuh:419-429). The only divergence available is // the softmax REDUCTION ORDER (block-cooperative on CUDA, 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. @@ -530,13 +557,169 @@ TEST_CASE("cuda fp8 KV paged-attention read matches the CPU read") { gpu.DestroyQueue(gq); } +// ─── G4b ──────────────────────────────────────────────────────────────────── +// THE INSTANTIATION PRODUCTION WILL USE. G4 above runs an f32 query into an f32 +// output, which resolves `LaunchPagedFp8Out` +// (src/vt/cuda/cuda_paged_attn.cu). 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 #1574's subject +// `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` — the campaign that makes this row the +// critical path — runs a bf16 query and a bf16 output. Without this case +// `LaunchPagedFp8Out<__nv_bfloat16, __nv_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("cuda fp8 KV paged-attention read matches the CPU read (bf16 query, bf16 out)") { + if (!HasCuda()) { + MESSAGE("SKIPPED: no CUDA 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::kCUDA); + 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 CUDA exactly as on CPU — it -// must be refused, never silently mis-stored through the e4m3 converter. -TEST_CASE("cuda fp8 KV store refuses e5m2 (later brick)") { +// must be refused, never silently mis-stored through the e4m3 converter. There +// are THREE refusals on that path and only one is a CUDA-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 CUDA guarantee. +// * the CPU kernel, `src/vt/cpu/cpu_cache.cpp` `ReshapeAndCacheFp8Kernel`. +// * the CUDA kernel's own guard, `src/vt/cuda/cuda_cache.cu` +// `ReshapeAndCacheFp8KernelCuda`, which is defence in depth for any future +// caller that reaches the registered provider without going through the +// wrapper. +// +// The FIRST version of this case called `vt::ReshapeAndCacheFp8` with device +// tensors and asserted a bare throw. That reads like a device gate and is not +// one, and both halves were MEASURED rather than argued. Deleting the CUDA +// kernel's VT_CHECK on a CPU build gives `ninja: no work to do` and leaves this +// file 7/10 SUCCESS. Deleting the op wrapper's leaves W1's +// `test_ops_fp8_kv_cache` GREEN at 8/511, because execution falls through to +// the CPU kernel's check and W1's `refuses e5m2` case +// (`tests/vt/test_ops_fp8_kv_cache.cpp:342`) asserts CHECK_THROWS_AS on +// std::runtime_error, not a message; only deleting BOTH turns it red (7/8, +// 510/511). What W1 pins is therefore "refused somewhere on the CPU path". +// +// 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 `cuda reshape_and_cache_fp8` +// and `fp8_e5m2`, which no other layer produces. +TEST_CASE("the CUDA fp8 KV store kernel refuses e5m2 (later brick)") { if (!HasCuda()) { - MESSAGE("SKIPPED: no CUDA backend in this build/host — the CUDA e5m2 refusal " - "gate did NOT run"); + MESSAGE("SKIPPED: no CUDA backend in this build/host — the CUDA-kernel e5m2 " + "refusal gate did NOT run"); return; } Backend& gpu = vt::GetBackend(DeviceType::kCUDA); @@ -555,9 +738,31 @@ TEST_CASE("cuda fp8 KV store refuses e5m2 (later brick)") { 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}); - CHECK_THROWS_AS(vt::ReshapeAndCacheFp8(gq, gk, gk, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E5M2, - 0.01f, 0.01f), - std::runtime_error); + + // The registered CUDA 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::kCUDA)); + REQUIRE(fn != nullptr); + std::string msg; + try { + fn(gq, gk, gk, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E5M2, 0.01f, 0.01f); + FAIL("the CUDA 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 CUDA KERNEL and name the missing part, not + // from the device-independent wrapper this call deliberately bypassed. + CHECK(msg.find("cuda 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);