diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md index 773f06b34..68228904f 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** (no device in the implementing session), though the CUDA TUs DO COMPILE: CI `cuda-fat-build` built them for ten architectures under `-Werror=all-warnings` on `4d71e776e` (run 32495320287). That job sets `-DVLLM_CPP_BUILD_TESTS=OFF`, so nothing has EXECUTED them -- see the spec's `## Owed`. **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 7c29bfb76..074ac1aeb 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -543,6 +543,8 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1575](https://github.com/mudler/vllm.cpp/issues/1575) | `SPEC-DFLASH2` | **`build-newest-gcc` has been RED on `main` since `5702d8f83`, and it is a recurrence of the class `tests/support/process_id.h` was created to close, not a missing include.** `tests/vllm/models/test_qwen3_dflash2_gguf.cpp:547` calls `::getpid()` while including nothing that declares it; the lane builds in a `gcc:16` container where libstdc++ no longer supplies the transitive ``, so it fails to COMPILE (`error: '::getpid' has not been declared; did you mean 'getpt'?`) while every local toolchain at gcc 13 stays green. The seam's own header predicts this: it says the spelling "was fixed once in three files and came back in five more, because each new loader test copies the temp-directory helper from the last one", which is exactly what W5's GGUF test did. The convention is otherwise unanimous — every other file in the tree naming `getpid` either includes `` or uses the helper, and the sibling `tests/vllm/entrypoints/test_dspark_draft_routing.cpp` uses the helper AND carries a comment warning not to reintroduce this class. ATTRIBUTED, not inferred: `build-newest-gcc` is absent from the baseline's failed list at `92406c620` and present at `5702d8f83`, the commit that added the file. MEASURED red-first in the lane's own container at base `947e5f648`, file sha256 `83bba319…`: unmodified `RED_RC=1` with the error byte-identical to CI; with the seam applied `GREEN_RC=0`, `compile_err=0`, `git diff --stat` confirming the edit applied; local gcc 13 `GCC13_RC=0`, so no regression on the shipped toolchain. FIXED IN FLOW with the portable spelling (`#include "support/process_id.h"`, `vllm_test::ProcessId()`) rather than ``, which would work on POSIX but re-copies the idiom the helper centralises and does not compile on MSVC. Found while measuring [#1464](https://github.com/mudler/vllm.cpp/issues/1464) at `origin/main` | 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 | +| [#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 | | [#1594](https://github.com/mudler/vllm.cpp/issues/1594) | `BACKEND-BENCH-CUDA-SGLANG-PREFLIGHT` | **`run_serve_low.py` could not name a third benchmark subject, and the flag that would name one is the key of the evidence tree.** `--model-key` was `choices=("27", "35")` at `:663`, and the value is a directory name at `:75` (`evidence/corpus//`), `:81` (`evidence/raw///`) and again at `:350`/`:352` for the in-container paths the pinned `sglang.bench_serving` reads and writes. So [#1574](https://github.com/mudler/vllm.cpp/issues/1574)'s subject, `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` @ `36f717a2`, could only be measured by mislabelling its evidence under an existing key. **The failure to prevent is not a crash**: it is a run that completes, writes a plausible raw result, and files it under another subject's key, which a `choices` edit alone would not catch. FIXED IN FLOW. `MODEL_KEYS` is now one definition the parser's `choices` and a new fail-closed `build_bench_command` refusal both read, and the third key is `q38mtp` -- deliberately not digit-prefixed, because `27` already names a DIFFERENT 27B checkpoint (`unsloth/Qwen3.6-27B-NVFP4`) and `online_gate.py`'s `27`/`27n` pair records that two 27B checkpoints share no goldens and no comparable ratios. The test enters through `main()` on a real argv with a stubbed `subprocess.run` that resolves the container `--output-file` back through the command's OWN bind mount, so the host derivation and the container derivation must agree or the run cannot read its own result back; it asserts, for every admitted key, that all four derived paths carry that key and that no other admitted key's tree was named or created. Three negative mutations recorded red and restored: the corpus path hardcoded to `27` (`missing corpus partition`), the container raw path alone hardcoded to `27` (host read `FileNotFoundError` -- the insidious one the corpus precondition cannot see), and the key guard deleted. No workload shape, concurrency ladder or timed semantic changed, and no server was started. `docs/USAGE.md` is deliberately unedited: it documents no `tools/bench` harness, this flag reaches no shipped command, and adding the first `tools/bench` entry for one choice value would be a new obligation rather than a projection of a changed one | gap | | [#1605](https://github.com/mudler/vllm.cpp/issues/1605) | `BACKEND-BENCH-CUDA-SGLANG-PREFLIGHT` | **`test_serve_low_client.py`'s concurrency-cap case reds at random, and the number it reports belongs to the fixture rather than to the client.** `tests/tools/test_serve_low_client.py:135` asserts `_CompletionHandler.peak <= 2` after `run_usage_batch(..., max_concurrency=2)`. Measured at `f7d517c09`: **1 failure in 30** standalone runs of that one case, `AssertionError: 3 not less than or equal to 2`; the fresh review of [#1596](https://github.com/mudler/vllm.cpp/pull/1596) measured 2 in 8 on the same head. `tests/CMakeLists.txt:12` registers all of `tests/tools` as the CTest target `test_serve_low_tools`, so it reds CI at random on whatever diff is in flight, and it has done so since `e58858a91` with no issue naming it. The cap itself is structural and is NOT what fails: `run_usage_batch` submits into a `ThreadPoolExecutor(max_workers=max_concurrency)`. `peak` counts handler threads inside the fixture's `do_POST`, between `active += 1` and the `finally: active -= 1`, and that window is strictly WIDER than the client's -- `post_json` returns once `response.read()` has the body and closes the connection in its own `finally`, while the handler thread has yet to leave its `try` and retake the lock. A worker that returns and submits its next request in that gap makes a third handler live with two client threads. So the reported 3 is the instrument's own overlap, and reporting it as a breached concurrency cap points the reader at `run_usage_batch`, where the 3 did not come from. **NOT FIXED IN FLOW, deliberately**: the repair changes what the case measures (count the client's in-flight window, or hold the handler open until the client has read), which changes the guarantee it asserts, so under `## Every change starts from an issue` it takes the row, spec and fresh-review path rather than riding inside a model-key change. The lower bound `assertGreaterEqual(peak, 2)` is the half that proves parallelism happened and must survive any repair: widening the bound to `<= max_concurrency + 1` would make the case green and measure nothing. Found by the fresh review of [#1596](https://github.com/mudler/vllm.cpp/pull/1596) (F4), reproduced independently before filing | bug | | [#1595](https://github.com/mudler/vllm.cpp/issues/1595) | `BACKEND-TENSTORRENT-HOST-FREE-FORWARD` | `tests/vllm/models/test_qwen3_dflash2_gguf.cpp:547` calls `::getpid()` with no `` in its include block, so the TU does not compile under clang-20 (`no member named 'getpid' in the global namespace`); landed on main in the SPEC-DFLASH2 GGUF arm and blocks every post-merge rebuild on a clang host, including the TT card gates; fixed in-flow with the one-line missing include, matching the sibling-test convention | bug | diff --git a/.agents/quantization-matrix.md b/.agents/quantization-matrix.md index e9bebe433..1a627c0f7 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; the CUDA TUs COMPILE (CI `cuda-fat-build`, ten architectures, run 32495320287 on `4d71e776e`) but that job builds with tests OFF, so none has been executed (spec `## Owed`). 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..c75306498 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 / @@ -54,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 @@ -82,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 @@ -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,22 @@ 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` + 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 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). ## Dependencies @@ -141,11 +163,159 @@ 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`). 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. 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 +`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: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 +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`. + +**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 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 + 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 e51834fad..b2358b7d1 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..e6e069b14 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,134 @@ 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. +// +// 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)` +// (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 +301,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..dad948917 100644 --- a/src/vt/cuda/cuda_paged_attn.cu +++ b/src/vt/cuda/cuda_paged_attn.cu @@ -132,6 +132,58 @@ __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: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 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; + 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 +241,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 +285,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 +301,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 +659,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 +728,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 +2078,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 +2179,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 +2302,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 +2921,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: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 +// 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 +2983,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: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) { + 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 5762448a7..856722a59 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -2179,6 +2179,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..ff69974a7 --- /dev/null +++ b/tests/vt/test_cuda_fp8_kv_cache.cpp @@ -0,0 +1,771 @@ +// 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:419-429 +// (scaled_vec_conversion) +// scale quant_utils.cuh:296-300 — FP8 = Quantize(HP / scale); +// Dequant(FP8) * scale = HP +// scales vllm/model_executor/layers/quantization/kv_cache.py:108-191 +// (BaseKVCacheMethod: per-TENSOR k_scale/v_scale, 1.0 uncalibrated) +// +// The gates, and they do not all run in the same build: +// +// G1 (runs in every build WITHOUT the 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. +// 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 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), 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/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 + +#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); +} + +// 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/f16-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 slots = {3, 0, 2, 1}; + const float k_scale = 0.007f, v_scale = 0.003f; + + // Both narrow dtypes are 2-byte, so one uint16_t staging buffer serves each. + for (DType src : {DType::kBF16, DType::kF16}) { + const int src_dtype_tag = static_cast(src); + CAPTURE(src_dtype_tag); + std::vector kb(kf.size()), vb(vf.size()); + for (size_t i = 0; i < kf.size(); ++i) { + kb[i] = src == DType::kBF16 ? vt::F32ToBF16(kf[i]) : vt::F32ToF16(kf[i]); + vb[i] = src == DType::kBF16 ? vt::F32ToBF16(vf[i]) : vt::F32ToF16(vf[i]); + } + + std::vector kc_ref(cache_elems, 0); + std::vector vc_ref(cache_elems, 0); + Tensor ck = Host(kb.data(), src, {nt, H, D}); + Tensor cv = Host(vb.data(), src, {nt, H, D}); + Tensor ckc = Host(kc_ref.data(), DType::kI8, {nb, bs, H, D}); + Tensor cvc = Host(vc_ref.data(), DType::kI8, {nb, bs, H, D}); + Tensor cs = Host(slots.data(), DType::kI64, {nt}); + vt::ReshapeAndCacheFp8(cq, ck, cv, ckc, cvc, cs, Fp8KVCacheDataType::kFp8E4M3, k_scale, + v_scale); + + void* dk = gpu.Alloc(kb.size() * sizeof(uint16_t)); + void* dv = gpu.Alloc(vb.size() * sizeof(uint16_t)); + void* dkc = gpu.Alloc(cache_elems); + void* dvc = gpu.Alloc(cache_elems); + void* ds = gpu.Alloc(slots.size() * sizeof(int64_t)); + std::vector zero(cache_elems, 0); + gpu.Copy(gq, dk, kb.data(), kb.size() * sizeof(uint16_t)); + gpu.Copy(gq, dv, vb.data(), vb.size() * sizeof(uint16_t)); + gpu.Copy(gq, dkc, zero.data(), cache_elems); + gpu.Copy(gq, dvc, zero.data(), cache_elems); + gpu.Copy(gq, ds, slots.data(), slots.size() * sizeof(int64_t)); + Tensor gk = Dev(dk, src, {nt, H, D}); + Tensor gv = Dev(dv, src, {nt, H, D}); + Tensor gkc = Dev(dkc, DType::kI8, {nb, bs, H, D}); + Tensor gvc = Dev(dvc, DType::kI8, {nb, bs, H, D}); + Tensor gs = Dev(ds, DType::kI64, {nt}); + vt::ReshapeAndCacheFp8(gq, gk, gv, gkc, gvc, gs, Fp8KVCacheDataType::kFp8E4M3, k_scale, + v_scale); + + std::vector kc_got(cache_elems, 0); + std::vector vc_got(cache_elems, 0); + gpu.Copy(gq, kc_got.data(), dkc, cache_elems); + gpu.Copy(gq, vc_got.data(), dvc, cache_elems); + gpu.Synchronize(gq); + CHECK(kc_got == kc_ref); + CHECK(vc_got == vc_ref); + // The CPU oracle must have WRITTEN something, or the equality above is + // between two all-zero buffers and holds for any kernel. + CHECK(std::any_of(kc_ref.begin(), kc_ref.end(), [](uint8_t b) { return b != 0; })); + + gpu.Free(dk); + gpu.Free(dv); + gpu.Free(dkc); + gpu.Free(dvc); + gpu.Free(ds); + } + gpu.DestroyQueue(gq); +} + +// ─── G4 ───────────────────────────────────────────────────────────────────── +// READ parity: paged attention over the SAME fp8 cache bytes, 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: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. +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); +} + +// ─── 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. 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-kernel 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}); + + // 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); + gpu.Free(ds); + gpu.DestroyQueue(gq); +}