diff --git a/benchmarks/README.md b/benchmarks/README.md
index 6218903f2..5cbb0b34e 100644
--- a/benchmarks/README.md
+++ b/benchmarks/README.md
@@ -29,3 +29,12 @@ python benchmarks/bench_offload_cache_copy.py
 
 For host RAM vs PCIe bandwidth and the offload/hybrid backend pick, use `ft bench bw`
 instead — it writes the JSON profile the engine reads.
+
+**`bench_kv_quant.py`** compares BF16, FP8 and NVFP4 KV storage bytes, one-step
+scatter latency and paged decode latency on synthetic inputs. No checkpoint is
+required. Keep the GPU idle and use identical arguments for A/B comparisons;
+this does not measure model quality or end-to-end serving throughput.
+
+```bash
+PYTHONPATH=python:. uv run python benchmarks/bench_kv_quant.py --lengths 1024,8192,32768
+```
diff --git a/benchmarks/bench_kv_quant.py b/benchmarks/bench_kv_quant.py
new file mode 100644
index 000000000..18707eb07
--- /dev/null
+++ b/benchmarks/bench_kv_quant.py
@@ -0,0 +1,77 @@
+"""Paged KV storage/scatter/decode microbenchmark, independent of model weights.
+
+Run with PYTHONPATH=python:. uv run python benchmarks/bench_kv_quant.py.
+Compare identical arguments on the baseline and candidate; this does not measure
+end-to-end model quality, TTFT, or serving throughput.
+"""
+
+import argparse
+import json
+
+import torch
+import triton.testing
+
+from freetoken.distributed import set_tp_info
+from freetoken.kernel.triton.attention import decode_paged_attention
+from freetoken.kvcache.mha_pool import MHAKVCache
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--formats", default="none,fp8,nvfp4")
+    parser.add_argument("--lengths", default="1024,8192,32768")
+    parser.add_argument("--batch", type=int, default=1)
+    parser.add_argument("--heads", type=int, default=4)
+    parser.add_argument("--group", type=int, default=4)
+    parser.add_argument("--dim", type=int, default=128)
+    args = parser.parse_args()
+    set_tp_info(rank=0, size=1)
+    torch.manual_seed(42)
+    batch, heads, dim = args.batch, args.heads, args.dim
+    qheads = heads * args.group
+    device = torch.device("cuda")
+    results = []
+    for length in map(int, args.lengths.split(",")):
+        slots = batch * length
+        k, v = [torch.randn(slots, heads * dim, device=device, dtype=torch.bfloat16)
+                for _ in range(2)]
+        loc = torch.arange(slots, device=device, dtype=torch.int32)
+        q = torch.randn(batch, qheads, dim, device=device, dtype=torch.bfloat16)
+        indptr = torch.arange(batch + 1, device=device, dtype=torch.int32) * length
+        pos = torch.full((batch,), length - 1, device=device, dtype=torch.int32)
+        scratch = torch.empty(batch, qheads, 8, dim, device=device)
+        lse = torch.empty(batch, qheads, 8, device=device)
+        splits = torch.full((batch,), 8, device=device, dtype=torch.int32)
+        out = torch.empty_like(q)
+        for quant in args.formats.split(","):
+            pool = MHAKVCache(heads, 1, dim, slots, 1, q.dtype, device, kv_quant=quant)
+            pool.store_kv(k, v, loc, 0)
+            extra = {}
+            if quant == "nvfp4":
+                extra = dict(kv_quant=quant, k_block_scale=pool.k_block_scale(0),
+                             v_block_scale=pool.v_block_scale(0))
+            kc, vc = [getattr(pool, name)(0).flatten(0, 1) for name in ("k_cache", "v_cache")]
+
+            def decode():
+                return decode_paged_attention(q, kc, vc, indptr, loc, pos,
+                    scratch, lse, splits, 8, dim ** -.5, out=out,
+                    k_scale=pool.k_scale(0), v_scale=pool.v_scale(0), **extra)
+
+            decode()
+            decode_ms = triton.testing.do_bench(decode, warmup=100, rep=300)
+            store_ms = triton.testing.do_bench(
+                lambda: pool.store_kv(k[-batch:], v[-batch:], loc[-batch:], 0),
+                warmup=100, rep=300)
+            record = dict(format=quant, length=length, batch=batch,
+                          kv_bytes=pool.unit_bytes()[0] * slots,
+                          decode_ms=decode_ms, store_ms=store_ms)
+            results.append(record)
+            print(json.dumps(record), flush=True)
+            del pool, kc, vc
+        del k, v
+    print(json.dumps(dict(gpu=torch.cuda.get_device_name(), torch=torch.__version__,
+                         heads=heads, group=args.group, dim=dim, results=results)))
+
+
+if __name__ == "__main__":
+    main()
diff --git a/docs/cli.md b/docs/cli.md
index 0c3394398..2733c7ea8 100644
--- a/docs/cli.md
+++ b/docs/cli.md
@@ -71,8 +71,32 @@ ft serve --model ... --gpu GPU-9e8d7c6b  # the same card by UUID (a unique prefi
 | `--num-pages` / `--num-tokens` | auto | KV capacity override in pages / tokens (mutually exclusive; auto sizes from VRAM left after weights and MoE cache) |
 | `--page-size` | 1 | KV page size; DSV4 forces 128, the TRTLLM backend needs 16/32/64, SWA models require 1 |
 | `--cache-type` | radix | `radix` (prefix reuse; SWA/GDN-aware variants picked automatically) or `naive` |
+| `--kv-cache-dtype` | bf16 | `bf16`, `fp8`, `nvfp4`, or `fp8-fp4`. Generic [FP8](#fp8-kv-cache) and [NVFP4](#nvfp4-kv-cache) use their own scale layouts; [DeepSeek-V4.1](deepseek-v41.md#native-fp8-fp4-kv-storage) uses `fp8-fp4` to preserve its native FP8 window, FP4 compressed KV and MXFP4 index keys. |
 | `--attention-backend`, `--attn` | auto | `trtllm`/`fi`/`fa`/`triton`/`dsv4_sparse`/`dsa`; `prefill,decode` pair allowed; auto picks per model + GPU |
 
+### FP8 KV cache
+
+`ft serve --kv-cache-dtype fp8` halves the bytes per cached token (8-bit codes instead
+of 16), so a card that held N tokens holds close to 2N. Each `(token, kv head)` row
+keeps its own fp32 scale, which costs ~3% back at `head_dim=128`. Requirements and
+trade-offs:
+
+- Needs the **triton** attention backend; `--attn auto` selects it (and refuses an
+  explicit `fi`/`fa`/`trtllm`, which cannot be shown to apply these scales).
+- Works on the plain paged, hybrid-SWA and QSA sparse (Qwen3.8-Flash-Next) KV pools.
+  On QSA the block-selection index keys stay 16-bit; only the selected K/V rows are
+  read back as codes. MLA/DSA latent KV, DeepSeek-V4's tiered pool and the block-sparse
+  MiniMax-M3 pool stay 16-bit; asking for fp8 there fails at startup rather than
+  silently ignoring the flag.
+- The same bytes on every GPU FreeToken targets: the codes sit in a plain byte buffer
+  and are decoded in software, so the cache holds identical data and produces identical
+  numbers on any card (the fp8 type is deliberately kept out of the kernels, which is
+  also what makes the feature work on the RTX 30 series).
+- Accuracy is checkpoint-dependent. Expect it to matter most on long contexts and on
+  models with outlier key channels; keep `bf16` when a run must be bit-reproducible.
+- `ft ctl stats` / `/v1/cache/status` report the smaller `kv_bytes_per_token`, and
+  `ft ctl cache --kv N` moves the same (now cheaper) pool.
+
 ### MoE offload
 
 See [models.md](models.md#moe-strategies) for what each strategy does.
@@ -101,12 +125,17 @@ See [models.md](models.md#moe-strategies) for what each strategy does.
 
 ### Image input
 
-Experimental. Needs a checkpoint whose family registers a vision encoder ([models.md](models.md#image-input) lists them and how each one
-maps the flags below); a request carrying images is rejected otherwise. Images are accepted on all three protocols (OpenAI `image_url`,
-Anthropic `image` blocks, Responses `input_image`) as an http(s) URL or base64. Images inside a tool
-result (an Anthropic `tool_result` block from Claude Code's Read, a Responses `function_call_output`
-from Codex's view_image) are moved to the user turn that follows the tool message, as vLLM does,
-because chat templates render tool messages as plain text.
+Experimental. Needs a checkpoint whose family registers a vision encoder;
+[models.md](models.md#image-input) lists the families and their image options.
+A request carrying images is rejected otherwise. Images are accepted on all three protocols (OpenAI
+`image_url`, Anthropic `image` blocks, Responses `input_image`) as an http(s) URL
+or base64. Tool-result image support depends on the model's chat encoding:
+DeepSeek-V4.1 accepts ordered text and images inside Anthropic `tool_result`
+blocks and Responses `function_call_output` without moving them. For Qwen VL and
+other chat templates that render tool messages as plain text, images are moved
+to a user turn after the tool message, as vLLM does. This supports tool images
+from Claude Code's Read and Codex's view_image while keeping the tool's text
+under its original call ID.
 `GET /v1/stats` reports what the server accepts as `model.input_modalities` (`["text"]` or `["text", "image"]`),
 so a client can gate its attachment controls without reading the checkpoint config.
 
@@ -115,12 +144,19 @@ so a client can gate its attachment controls without reading the checkpoint conf
 | `--text-model-only` | off | Serve a multimodal checkpoint text-only: no encoder tower is built (its VRAM goes to the KV/expert pools) and every multimodal input is rejected. Same as `--mm-disable` with every encoder kind |
 | `--mm-disable` | none | Encoder towers to leave unbuilt (`vision`, `audio`); every input they would serve is rejected |
 | `--mm-encoder-weights` | host | Where the encoder tower's block weights live. `host` streams them from pinned host banks two blocks at a time behind the compute, so the GPU holds two blocks instead of the whole tower; small images pay the copy time, large ones hide it behind the compute. `gpu` keeps them resident. An encoder without a block stack stays resident either way |
-| `--image-min-tokens`, `--image-max-tokens` | processor defaults | Per-image token budget: the image processor resizes every image to take between these many tokens, converted to the family's own units by its processor. A family with fixed budgets honors the maximum only and refuses one below its smallest budget at start-up |
+| `--image-min-tokens`, `--image-max-tokens` | processor defaults | Per-image token budget, converted to the family's own units by its processor. A family with fixed budgets honors the maximum only and refuses one below its smallest budget at start-up. DeepSeek-V4.1 supports the maximum and rejects the minimum flag, as described below |
 | `--mm-processor-kwargs` | none | JSON object of extra keyword arguments for the checkpoint's image processor call, for knobs the token budget does not cover; applied after the budget, so an explicit key wins |
 | `--mm-embed-cache-device` | cpu | Where encoded image embeddings live between prefill chunks. `cpu` keeps them out of the VRAM budget; `cuda` skips the copy back |
 | `--allowed-media-domains` | any | Comma-separated hostname allowlist for image URLs; requests for other domains are rejected with a 400. Empty allows any domain |
 | `--allowed-local-media-path` | off | Directory `file://` image refs may be read from; unset rejects local files |
 
+DeepSeek-V4.1's `--image-max-tokens` limits the entire image span, including start,
+row-newline and end tokens. It rejects `--image-min-tokens`; use
+`--mm-processor-kwargs '{"vision_min_pixels": 295936}'` to set the minimum pixel
+area before the maximum-token resize. Its supported processor options are
+`vision_min_pixels`, `vision_max_n_token` and `vision_max_wh_ratio`;
+`vision_max_n_token` takes precedence over `--image-max-tokens` when both are set.
+
 ## ft shell
 
 ```bash
@@ -202,3 +238,33 @@ profile that `ft serve --moe-strategy auto` and `--moe-hybrid-max-fetch -1` then
 - `--threshold` (default 2.0) sets the call: recommend hybrid when CPU bandwidth beats PCIe
   by that factor.
 
+
+### NVFP4 KV cache
+
+`ft serve --model <checkpoint> --kv-cache-dtype nvfp4 --attention-backend triton`
+opts into packed E2M1 KV storage. The initial implementation supports plain paged
+FULL attention (MHA/GQA), hybrid-SWA, and the full-attention portion of hybrid-linear
+models, QSA, and MLA/DSA (including GLM-5.3-Flash). Head dimensions must be
+divisible by 16. DSV4 and BSA pools are rejected at startup. `auto` selects
+Triton, QSA sparse, or DSA attention for supported models. For MLA/DSA use
+`--attention-backend auto` or `--attention-backend dsa`. Only the latent slab is
+quantized; indexer keys, kpool tails/gates, and recurrent states retain their
+existing precision. A 512-element latent row occupies 292 bytes instead of 1024
+bytes in BF16, excluding those other tiers.
+
+Each K or V row stores `head_dim / 2` packed bytes, `head_dim / 16` E4M3 block-scale
+bytes, and one FP32 row scale. At head_dim 128 this is 76 bytes, versus 256 for
+BF16 and 132 for the existing FP8 format. Pool management, recurrent states,
+attention workspace and model weights consume additional memory.
+
+The second-level scale is dynamic per token/head, so appending a token never
+rescales an existing prefix. This is a FreeToken KV layout, not an external
+NVFP4 checkpoint or attention-library ABI. K/V are restored inside attention;
+Q and attention arithmetic retain their compute precision. The MoE weight option
+`--nvfp4-backend` is independent. Paged MHA prefill uses fresh compute-dtype K/V
+while cached prefixes are restored, as in the FP8 path. MLA/DSA stores fresh
+latent rows first and reads the quantized cache in both prefill and decode.
+
+NVFP4 is opt-in: assess quality on your checkpoint and workload before using it
+for long-context inference. Capacity savings do not guarantee faster decode;
+packing, reconstruction, and the selected attention backend affect throughput.
diff --git a/docs/deepseek-v41-kv-quantization-plan.md b/docs/deepseek-v41-kv-quantization-plan.md
new file mode 100644
index 000000000..2f188acf6
--- /dev/null
+++ b/docs/deepseek-v41-kv-quantization-plan.md
@@ -0,0 +1,253 @@
+# DeepSeek-V4.1 KV cache quantization proposal
+
+Status: design record, 2026-09-12. The accepted implementation names the native
+mixed mode `fp8-fp4`; see [current behavior and validation](deepseek-v41.md#native-fp8-fp4-kv-storage).
+The estimates below predate GPU validation. The experimental NVFP4 window mode
+is outside the accepted implementation scope.
+
+Target: one RTX 5090 (32 GiB), the current V4.1 branch, and
+`LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4` at revision
+`dfce15b92ed1fa76e80e2a46ba847e5b5451f12c`, including image inputs.
+
+## Recommendation
+
+Implement packed storage for the model's existing quantization first: FP8 for
+the sliding window, FP4 for compressed attention KV, and MXFP4 for index keys.
+The current implementation already rounds these values to those formats, then
+stores the reconstructed values in BF16. Keeping the original codes and scales
+can avoid additional quantization error relative to this branch's BF16 baseline.
+Readers must reconstruct the same BF16 values before the existing arithmetic;
+end-to-end equivalence still needs tests.
+
+Add NVFP4 window storage as a separate experimental option after that baseline
+works. It introduces an additional numerical change. Its incremental benefit is
+small with a small window pool, so it should not delay the first implementation.
+
+This reduces GPU KV memory. It does not shrink the host expert banks or Engram
+tables, and requires no model download or checkpoint conversion. Lower memory
+traffic may help attention, but decoding packed values also costs work. Overall
+generation speed, including CPU expert offload, must be measured.
+
+## Formats and proposed configuration
+
+The accepted CLI value is `fp8-fp4`, distinct from the generic `fp8` and `nvfp4`
+layouts. Keep BF16 as the default until validation completes.
+
+| Proposed option | Window | Compressed attention KV | Index keys |
+| --- | --- | --- | --- |
+| `--kv-cache-dtype bf16` | Existing BF16 storage | Existing BF16 storage | Existing BF16 storage |
+| `--kv-cache-dtype fp8-fp4` | Native FP8 packed storage | Native FP4 packed storage | Native MXFP4 packed storage |
+| `--kv-cache-dtype nvfp4` (future only) | Additional NVFP4 quantization | Same native FP4 storage | Same native MXFP4 storage |
+
+Report all three resolved formats and their bytes in startup diagnostics.
+
+| Tier | Width | Codes and scale representation | Bytes/row, including scales |
+| --- | ---: | --- | ---: |
+| BF16 window / compressed KV | 512 | BF16 values | 1,024 |
+| BF16 index keys | 128 | BF16 values | 256 |
+| Native window | 512 | E4M3 codes; one UE8M0 scale per 32 values | 528 |
+| Native compressed KV | 512 | Packed E2M1 codes; one E4M3 scale per 16 values; implicit global scale 1 | 288 |
+| Native index keys | 128 | Packed E2M1 codes; one UE8M0 scale per 32 values | 68 |
+| Experimental NVFP4 window | 512 | Packed E2M1 codes; E4M3 scale per 16 values; FP32 scale per row | 292 |
+
+V4.1 uses a shared K/V latent; do not multiply these figures by two for K and V.
+There are 40 window pools but only four compressed-KV/index-key owners, at layers
+2, 8, 14 and 20, with compression ratios 2, 2, 2 and 1. Eight layers compute index
+queries, which does not require eight index-key pools.
+
+The fixed checkpoint's [reference model](https://huggingface.co/LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4/blob/dfce15b92ed1fa76e80e2a46ba847e5b5451f12c/inference/model.py)
+uses these native formats: window at line 707, compressed KV at line 760, and
+index queries/keys at lines 546/552. Its
+[reference quantization kernel](https://huggingface.co/LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4/blob/dfce15b92ed1fa76e80e2a46ba847e5b5451f12c/inference/kernel.py)
+also exposes packed FP4 output. These files were inspected, not executed.
+
+Uniform FP8 storage is less attractive: it uses more bytes for already-FP4 tiers
+and can round compressed values again. For example, an E2M1 value of 1.5 times an
+E4M3 scale of 1.125 reconstructs to BF16 1.6875, which E4M3 cannot represent
+exactly. Preserve each tier's native representation instead.
+
+## Memory estimates
+
+These totals cover pool-owned tensors, including scales, compression state,
+scratch rows and the full-to-window map. They assume 128-token pages, two running
+requests, three scratch rows per source, and one extra dummy full page. K and M
+mean 1,024 and 1,048,576 tokens. Capacity is shared across requests, not available
+independently to each of the two requests.
+
+| Shared KV capacity | Window pool ratio | Current BF16, GiB | Native mixed storage, GiB | NVFP4 window plus native compressed/index, GiB |
+| --- | ---: | ---: | ---: | ---: |
+| 32K | 0.20 | 0.353405 | 0.159638 | 0.101120 |
+| 128K | 0.20 | 1.397678 | 0.630562 | 0.399868 |
+| 1M | 0.02 | 3.937756 | 1.293732 | 1.109177 |
+| 1M | 0.20 | 11.173664 | 5.041100 | 3.196675 |
+
+For 1M with a 0.02 window pool ratio, native storage saves about 2.64 GiB;
+quantizing the window further saves another 0.18 GiB. At the current nominal
+32K/0.20 setting, native storage saves about 198 MiB. The live automatic planner
+can round to a different page count, so these are not exact live-process totals.
+
+The window pool ratio controls retained window pages, not the model's 128-token
+attention window. Reducing this ratio is a separate cache-retention decision.
+Neither these estimates nor the model's context limit establishes usable 1M
+throughput on this machine.
+
+For reproduction, let `T` be usable shared token capacity, `F = T + 128`,
+`w = ceil(r * (T / 128 + 1))`, and `C = 2.5 * F`. For the configurations above,
+the window working-set floor is below `w`. With row byte sizes `bw`, `bc`, `bi`:
+
+```text
+window bytes     = 40 * 128 * w * bw
+compressed bytes = (C + 12) * bc
+index bytes      = (C + 12) * bi
+FP32 state bytes = 3 * (2 * w + 1) * 1024 * 4
+map bytes        = (F + 1) * 8
+```
+
+The BF16 totals were checked against the current pure-Python cost functions and
+an independent shape calculation; all four cases matched exactly. Packed totals
+use the same shapes with the proposed row sizes. Alignment or padding added by
+the implementation must be reflected in the estimates.
+
+Page tables, free lists and fixed position arrays add approximately 0.38, 1.51,
+12.04 and 12.05 MiB respectively. At 1M, a two-request eager decode snapshot adds
+16 MiB, with another 8 MiB temporary index-selection tensor during its creation.
+Query/selection/attention workspaces, prefix-cache metadata and CUDA allocator
+headroom are additional and workload dependent. Freed capacity can be used for
+longer contexts or more GPU expert caching; it is not automatically all spare
+VRAM if the planner reallocates it.
+
+## Implementation sequence
+
+### 1. Preserve quantization results and expose a storage layout
+
+Extend `kernel/triton/dsv41/quant.py` with packed encode/decode helpers and a CPU
+reference. Capture codes and scales where `_project`, `Indexer.keys` and
+`_publish` currently perform roundtrips in `models/deepseek_v41/attention.py`.
+Do not derive new scales from already reconstructed BF16 values in native mode.
+Keep query quantization and RoPE ordering unchanged, including the existing
+quantization of the RoPE channels.
+
+Keep a logical BF16 computation dtype and an explicit storage layout. Each packed
+uint8 row contains its code bytes followed by its scale bytes. The inherited BF16 dtype assertion cannot simply be
+changed to FP8: the three tiers have different storage types and strides. Centralize
+row-size and scale-shape definitions so allocation and byte accounting agree.
+
+### 2. Allocate and write every tier consistently
+
+Extend `kvcache/dsv41_paged_pool.py` to allocate the packed code/scale rows for
+each owner, including dummy and request-specific scratch rows. Keep the unfinished
+compression state rings in FP32; they hold reduction inputs and scores across
+partial pairs, not finished quantized KV values.
+
+Override the V4.1 backend's compressed write route. The inherited
+`attention/dsv4_compress.py::scatter_compressed` directly calls
+`index_copy_(..., kv.to(pool.dtype))`; changing only a pool writer would miss this
+path and cast values to integer bytes instead of encoding them. Route window,
+compressed attention and index-key writes through the appropriate codec in both
+prefill and decode. Preserve the current page addressing and request-isolated
+scratch destinations.
+
+### 3. Decode only selected tiles inside attention and the indexer
+
+Add V4.1-specific readers through `attention/dsv41_sparse.py`. The current shared
+`kernel/triton/dsv4/sparse_attn.py` assumes identical BF16 window/compressed
+strides. Both normal and split-K attention need mixed-format readers. Keep the
+V4 BF16 path unchanged; factor shared pieces only when their behavior is identical.
+
+Reconstruct a selected tile to the same BF16 values as the old roundtrip, then
+use the existing FP32 arithmetic. In particular, multiplying FP4 codes and scales
+directly into FP32 attention without the intermediate BF16 rounding changes the
+baseline. Do not simultaneously switch to FP8/FP4 tensor-core attention.
+
+Preserve the single-KV-tile staging design: the existing kernel documents an RTX
+5090 shared-memory constraint. Loading two complete window/compressed tiles and
+selecting afterwards can exceed it. Inspect compiled shared memory, registers and
+temporary allocations on sm120. Generic NVFP4 readers assume scalar base pointers;
+mixed per-column pool selection requires an adapted reader.
+
+In `kernel/triton/dsv41/indexer.py`, decode MXFP4 keys within the existing bounded
+tiles. Preserve score rounding, candidate reuse and deterministic tie ordering.
+Verify selected IDs, not just score closeness. Do not create a full-context BF16
+key copy, which would erase the memory benefit at long contexts.
+
+### 4. Complete accounting, cache lifecycle and configuration gates
+
+Update exact bytes, unit bytes, automatic budget estimates and the page solver in
+`kvcache/dsv41_cost_model.py` using the same layout definitions as allocation.
+Include every scale byte in reported bytes and rebuild validation. Keeping scales
+inside the original slabs lets inherited cleanup free codes and scales together;
+preserve the existing idle-only resize/rebind sequence.
+
+Codes and scales must share physical row identity through prefix reuse, window
+eviction, page reuse, dummy accesses and request cancellation. Reinitialize safe
+dummy values and do not permit stale scales to accompany new codes. Preserve
+graph-safe writes and recapture behavior wherever the baseline supports graphs.
+
+Only after writers, readers and accounting pass tests, enable V4.1's matching
+capability flags and pool checks in `attention/__init__.py`, `engine/engine.py`
+and `kvcache/__init__.py`. Keep rejection for other unsupported architectures.
+Update CLI help and `docs/deepseek-v41.md` with the actual mixed-tier semantics.
+
+### 5. Add the experimental NVFP4 window option
+
+Initially encode the existing FP8-roundtrip BF16 window values into NVFP4, so this
+option is explicitly one additional compression step relative to native mode.
+Use a separate FP32 scale per stored row; a changing cache-wide scale would
+invalidate previously stored prefixes. Existing `kernel/triton/kv_nvfp4.py`
+provides useful packing and row-scale patterns. Its quantizer must not replace
+the native compressed-KV codec, whose global scale is implicitly one.
+
+NVIDIA describes NVFP4 as E2M1 data with 16-value E4M3 block scales and a second
+FP32 scale in its [format explanation](https://developer.nvidia.com/blog/introducing-nvfp4-for-efficient-and-accurate-low-precision-inference/).
+As checked on 2026-09-12, TensorRT-LLM's
+[hardware matrix](https://nvidia.github.io/TensorRT-LLM/features/quantization.html#hardware-support-matrix)
+does not list NVFP4 KV support for sm120, though it lists it for sm100/103.
+This proposal therefore uses FreeToken's own packed storage and tile decoding;
+it does not assume a TensorRT-LLM attention kernel is available on the RTX 5090.
+
+## Validation and acceptance
+
+Extend existing tests rather than adding a parallel test framework:
+
+- `tests/kernels/test_dsv41_quant.py`: independent code/scale/reference checks,
+  zero and tiny values, outliers, rounding boundaries and FP4 tie cases; native
+  decode must reproduce the current BF16 roundtrip exactly.
+- `tests/kernels/test_dsv41_indexer.py`: score precision and exact selected IDs,
+  ties, candidate restriction, source sharing and long-key tile boundaries.
+- `tests/models/test_dsv41_attention.py`: identical-mode BF16 comparisons for
+  normal and split-K attention, sinks, all-masked selections, ratio 0/1/2,
+  odd chunk boundaries, 128-token page boundaries and two-request decode.
+- `tests/kvcache/test_dsv41_pool.py`: allocation bytes equal cost estimates,
+  solver maximal fit, fixed-window overrides, scratch/dummy rows, page reuse,
+  prefix reuse, eviction and rebuild replacing all code/scale pointers.
+- `tests/engine/test_kv_quant_config.py` and
+  `tests/engine/test_attention_backend_matrix.py`: accept completed V4.1 profiles
+  and retain rejection for unsupported pool/backend combinations. Cover graph
+  replay and resize recapture where supported by the baseline.
+- RTX 5090 integration: text, Japanese/code prompts, image prompts and mixed
+  multi-turn requests. Check independent requests cannot reuse each other's
+  scale or unfinished compression state. Confirm no full-context BF16 expansion.
+
+Benchmark the same pinned checkpoint, prompts, context lengths and batch sizes
+against this branch's pre-change BF16 implementation. Record TTFT, prefill rate,
+decode tokens/s, actual KV allocation, peak VRAM and temporary memory. Separate a
+fixed expert-cache comparison from a planner-retuned comparison that spends the
+saved KV bytes on experts. Warm kernels and Engram pages consistently.
+
+Repository policy requests performance comparisons against `main`; if `main`
+cannot run this V4.1 checkpoint, report that limitation explicitly rather than
+substitute a different model as an equivalent baseline. Test shared-kernel
+regressions on supported models if shared code changes.
+
+Native mode must first pass exact codec reconstruction and index-selection tests;
+then check model-level differences with matching launch modes and documented
+floating-point tolerances. NVFP4 window mode additionally needs teacher-forced
+logit/perplexity comparisons and long-context retrieval, Japanese/code and image
+quality evaluation. Establish acceptance thresholds before promoting it. Short
+arithmetic and image-caption smoke tests alone are insufficient for that decision.
+
+At the time of the initial proposal, no GPU quantized-KV tests, quality evaluations
+or speed benchmarks had been run. Current results belong in `deepseek-v41.md`.
+The first deliverable should be validated native mixed storage;
+the further NVFP4 window compression remains opt-in until its benefit and quality
+are measured.
diff --git a/docs/deepseek-v41.md b/docs/deepseek-v41.md
new file mode 100644
index 000000000..b1d235597
--- /dev/null
+++ b/docs/deepseek-v41.md
@@ -0,0 +1,465 @@
+# DeepSeek-V4.1 Flash NVFP4
+
+This branch adds experimental text and image inference for
+[`LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4`](https://huggingface.co/LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4),
+validated at revision `dfce15b92ed1fa76e80e2a46ba847e5b5451f12c`.
+The earlier `s-zaizen/DeepSeek-V4.1-Flash-NVFP4` checkpoint with FP8 Engram
+tables remains supported.
+The implementation uses the checkpoint's native names and shapes. It does not
+instantiate model code downloaded from the Hub.
+
+## Requirements on one RTX 5090
+
+The launch example targets one 32 GiB RTX 5090 and approximately 480 GiB of WSL
+memory on a 512 GiB host. The LibertAIDAI checkpoint occupies about 429 GB
+(400 GiB). Routed expert banks require 285.04 GiB of host memory; the 97.28 GiB
+Engram tables remain packed on disk and in the OS page cache. Only requested
+rows are decoded. Tables are required for text as well as images, and are never
+fully pinned or expanded to BF16.
+
+With every Engram page cached, expert banks plus tables total 382.32 GiB,
+leaving about 98 GiB of the WSL memory allocation for runtime overhead and other
+uses. The earlier FP8 Engram tables required 188.83 GiB; their FP4 replacements
+save 91.56 GiB. GPU-resident weight geometry is unchanged. These are byte-count
+estimates, not a guarantee for every context length or concurrent workload.
+
+The publisher describes the expert transcode as lossless but the Engram FP8 to
+FP4 conversion as lossy, and has not published end-to-end quality evaluations.
+Successful smoke tests cannot establish equivalent answer quality.
+
+Allow SSD space for the approximately 429 GB checkpoint and runtime caches, plus
+another approximately 429 GB if creating an FTW copy. Download the revision above
+with a Hugging Face client, then pass the completed local snapshot directory to
+`--model` in the launch command below.
+
+The launch example uses a 32,768-token context, 1,024-token prefill chunks and two
+concurrent requests. Adjust `--max-seq-len-override`, `--kv-reserve-tokens`,
+`--max-prefill-length` and `--max-running-requests` for your workload. Increasing
+the context limit alone does not reserve the corresponding KV capacity.
+Full-checkpoint long-context throughput and memory limits still need measurement.
+
+## Direct launch
+
+After [installing FreeToken](install.md), run this command with a downloaded
+snapshot or standalone FTW directory. The current default streams vision blocks
+from host memory (`--mm-encoder-weights host`). Historical validation below used
+resident vision weights, so its GPU allocation measurements do not describe this
+default placement:
+
+```bash
+ft serve --model /path/to/snapshot \
+  --served-model-name LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4 \
+  --moe-strategy offload --quant-backend moe.nvfp4=triton --expert-load parallel \
+  --moe-cache-auto --disable-moe-prefill-overlap \
+  --cuda-graph-max-bs 0 --kv-cache-dtype fp8-fp4 \
+  --max-seq-len-override 32768 --kv-reserve-tokens 32768 \
+  --max-prefill-length 1024 --max-running-requests 2 \
+  --swa-full-tokens-ratio 0.2 \
+  --host 0.0.0.0 --port 1919
+```
+
+On the validation host, the launch configuration explicitly loaded expert banks in
+parallel. The expert reader excludes the two large Engram shards; its banks and
+temporary shard buffers fit within this host's RAM. The generic automatic check
+counts the whole checkpoint and would choose serial loading. Use
+`--expert-load serial` for a lower-memory host. Prefill overlap is disabled so the
+GPU cache does not require two complete expert layers, and the planner sizes the
+cache from measured free GPU memory.
+
+The engine automatically selects `dsv41_sparse` attention. WSL's existing CPU
+expert fallback remains available when its CUDA pinned-memory quota cannot hold
+the expert banks. CPU and hybrid execution use the same NVFP4 expert values as
+the offload strategy. The Triton NVFP4 expert kernel supports this model's clamped
+SwiGLU; `moe.nvfp4=marlin` and `moe.nvfp4=b12x` are rejected.
+
+## Sizing a million-token context
+
+The following configuration is sized from tensor shapes and covered by budget
+tests. It has not been validated with a full million-token checkpoint run.
+
+```bash
+ft serve --model /path/to/snapshot \
+  --served-model-name LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4 \
+  --moe-strategy offload --quant-backend moe.nvfp4=triton --moe-cache-auto \
+  --disable-moe-prefill-overlap \
+  --cuda-graph-max-bs 0 --kv-cache-dtype bf16 \
+  --max-seq-len-override 1048576 --kv-reserve-tokens 1048576 \
+  --swa-full-tokens-ratio 0.02 \
+  --max-prefill-length 1024 --max-running-requests 2
+```
+
+Historical validation with GPU-resident vision weights measured 11.35 GiB of
+resident model parameters. The current default streams vision blocks from host
+memory through two GPU buffers, so its resident allocation differs. The BF16 KV
+pool for 1,048,576 tokens requires 3.94 GiB at a window ratio of 0.02; the default 0.2
+ratio requires 11.17 GiB. A ratio of 0.01 reduces KV storage to 3.54 GiB. These
+ratios control retained window pages and prefix reuse, while the model's attention
+window remains 128 tokens.
+
+One layer of 384 cached experts requires 7.13 GiB. Disabling prefill overlap permits
+this minimum; the automatic planner can assign additional expert slots from the
+measured free memory. With the historical resident-vision layout and ratio 0.02,
+the minimum persistent model, expert and KV allocation is 22.42 GiB, plus about
+24 MiB of GPU Engram staging for a 1,024-token chunk. Activations, CUDA workspaces
+and allocator overhead need additional memory.
+The two-buffer overlap path requires at least 768 expert slots and may exceed
+the default memory budget at this context length.
+
+Chunked prefill keeps Engram staging proportional to the chunk size. Increasing
+the context limit does not allocate a million rows of Engram staging. The host
+expert banks remain about 285 GiB; Engram tables occupy about 97 GiB on disk and
+are read through the reclaimable OS page cache. All figures describe allocation
+geometry, not measured full-model throughput or answer quality.
+
+## Image requests
+
+OpenAI chat accepts ordered `text` and `image_url` parts, Anthropic messages accept
+base64 or URL `image` blocks, and Responses accepts `input_image.image_url`.
+Use a base64 image data URL for local files; public HTTP(S) image URLs are also
+supported. Uploaded Responses file IDs are not supported.
+
+```python
+import base64
+from pathlib import Path
+from openai import OpenAI
+
+image = base64.b64encode(Path("example.png").read_bytes()).decode()
+client = OpenAI(base_url="http://localhost:1919/v1", api_key="local")
+response = client.chat.completions.create(
+    model="LibertAIDAI/DeepSeek-V4.1-Flash-NVFP4",
+    messages=[{"role": "user", "content": [
+        {"type": "text", "text": "Describe this image."},
+        {"type": "image_url", "image_url": {"url": f"data:image/png;base64,{image}"}},
+    ]}],
+    max_tokens=256,
+)
+print(response.choices[0].message.content)
+```
+
+Image spans include start, newline and end tokens and may cross prefill chunk
+boundaries. The shared multimodal path hashes the processed patches, grid and
+span layout into image-specific placeholder IDs, allowing matching image
+prefixes to reuse the prefix cache. Repeated images share encoded embeddings
+while those embeddings are needed by active requests; a complete prefix hit can
+skip image encoding. Legacy `media` and precomputed `mm_embeds` requests remain
+isolated from prefix reuse. Current input limits are 16 images per request,
+32 MiB per image and 64 megapixels per image.
+
+The native chat encoder also preserves ordered text and images in Anthropic
+`tool_result` blocks and Responses `function_call_output`. Image support is
+enabled by default; `--text-model-only` or `--mm-disable vision` disables the
+tower and rejects image inputs. Vision block weights use host streaming by
+default; select `--mm-encoder-weights gpu` to keep them resident.
+
+`--image-max-tokens` limits the entire image span, including start, row-newline
+and end tokens. `--image-min-tokens` is rejected; use
+`--mm-processor-kwargs '{"vision_min_pixels": 295936}'` to set the minimum pixel
+area before the maximum-token resize. The processor also accepts
+`vision_max_n_token` and `vision_max_wh_ratio`; `vision_max_n_token` overrides
+`--image-max-tokens` when both are supplied.
+
+## Native FP8-FP4 KV storage
+
+Use `--kv-cache-dtype fp8-fp4` to keep the model's existing quantized KV values
+packed in GPU memory. This mode uses `dsv41_sparse` attention with BF16 compute;
+`--attention-backend auto` selects it. Generic `fp8` and `nvfp4` KV modes remain
+unsupported for V4.1 because their layouts differ.
+
+The launch example selects `fp8-fp4` explicitly. Use `--kv-cache-dtype bf16` to
+select the previous representation, then restart the server. The CLI default
+remains `bf16`.
+
+| Stored tier | Native format | Bytes per row, including scales | BF16 bytes per row |
+| --- | --- | ---: | ---: |
+| Window, 512 values | E4M3 + UE8M0 scale per 32 values | 528 | 1,024 |
+| Compressed KV, 512 values | E2M1 + E4M3 scale per 16 values | 288 | 1,024 |
+| Index keys, 128 values | E2M1 + UE8M0 scale per 32 values | 68 | 256 |
+
+Each physical row contains codes followed by scale bytes, so page reuse and
+cache rebuilds move them together. The writer keeps codes from the existing
+quantization step, without recalculating scales from reconstructed BF16 values.
+Attention and indexer kernels restore selected tiles to BF16 at the original
+rounding point. Unfinished compression state remains FP32. The full cache is
+never expanded back to BF16 for attention.
+
+This mode avoids additional storage quantization relative to this branch's BF16
+baseline. It does not remove the model's existing quantization or establish
+equivalence to a differently quantized checkpoint. Model weights, Engram tables
+and host RAM requirements are unchanged; no new download is required.
+
+For a nominal shared 32K pool at window ratio 0.2 and two concurrent requests,
+the pool-owned tensors decrease from 379,465,736 to 171,409,848 bytes. A shared
+1M pool at ratio 0.02 decreases from 4,228,132,872 to 1,389,134,264 bytes. These
+include scales, pending state, scratch rows and the full-to-window map; query
+workspaces and allocator overhead are additional. The automatic planner may
+choose different page counts or spend the savings on more cached experts.
+Long-context throughput still needs separate measurement.
+
+The mode is selected at startup. Changing between `bf16` and `fp8-fp4` requires
+a server restart; resizing an idle cache preserves its storage format. Further
+NVFP4 quantization of the window is not part of this implementation.
+
+### Packed KV validation on RTX 5090
+
+On 2026-09-12 JST, 319 focused tests passed in the CUDA 13 development container
+with PyTorch 2.11.0+cu130, Triton 3.6.0 and the workspace sources installed.
+They covered codecs, sparse attention,
+index scores and selection, pool allocation/rebuild, model prefill/decode/images,
+engine configuration and the shared page manager. A subsequent focused CUDA
+test also passed for two requests crossing odd compression and page boundaries,
+with identical outputs and shared index IDs between BF16 and packed storage.
+
+The equivalent focused test command from the repository root is shown below.
+The recorded run invoked `python -m pytest` inside a private development
+container; the container configuration is not part of this repository.
+
+```bash
+uv run python -m pytest \
+  tests/kernels/test_dsv41_quant.py \
+  tests/kernels/test_dsv41_sparse.py \
+  tests/kernels/test_dsv41_indexer.py \
+  tests/models/test_dsv41_attention.py \
+  tests/models/test_deepseek_v41_model.py \
+  tests/engine/test_deepseek_v41_engine.py \
+  tests/kvcache/test_dsv41_pool.py \
+  tests/engine/test_kv_quant_config.py \
+  tests/engine/test_attention_backend_matrix.py \
+  tests/scheduler/test_dsv4_generic_manager.py -q -p no:cacheprovider
+```
+
+The native codecs reproduce the BF16 roundtrip values exactly. Index scores and
+selected IDs also matched at the model's 32 index heads. Attention with 64 heads
+showed small floating-point differences from the BF16 kernel: three of 32,768
+output elements in one decode case and 36 of 1,048,576 in one prefill case, both
+with maximum absolute difference `6.103515625e-5`. Packed storage therefore does
+not imply bit-identical model logits for every launch shape.
+
+An isolated attention comparison used batch 1, 64 heads, dimension 512, window
+128 and 640 selected rows, five warmups and 31 CUDA-event samples:
+
+| Attention call | BF16 median | Packed median | Packed / BF16 |
+| --- | ---: | ---: | ---: |
+| Decode, one query, five splits | 0.153664 ms | 0.143968 ms | 0.937 |
+| Prefill, 32 queries, no splits | 0.123488 ms | 0.275744 ms | 2.233 |
+
+These are CUDA-event elapsed times around Python attention-wrapper calls,
+including possible host enqueue gaps and output/scratch allocation effects;
+they are not isolated kernel execution times or end-to-end generation throughput.
+Prefill calls were slower despite using less KV storage. The packed kernel requires 100,352 bytes
+of shared memory against this GPU's 101,376-byte limit, using one pipeline stage.
+It retains BF16 tiles and promotes them separately for each FP32 dot operation;
+promoting too early or increasing pipeline stages exceeds the device limit.
+Rerun the 512-dimensional regressions after compiler/kernel changes.
+
+The baseline is this branch's existing V4.1 BF16 implementation. The repository's
+unmodified base at `04d4621` lacks this V4.1 model path and cannot serve as a
+same-checkpoint baseline.
+
+The private validation image was
+`sha256:d6f7429ed8939fd981d2e26359f377deddaf1e926839c8ccfe3e09403b2d4ad3`
+and used the model settings in [Direct launch](#direct-launch). Image digests
+identify local validation artifacts; they are not published images or build
+instructions. The private Dockerfile, Compose configuration and download helper
+are outside the scope of this repository.
+
+All 390 Python source files in the image matched the tested workspace. The
+server became ready at 21:27:26 UTC on September 11, using the pinned LibertAIDAI
+checkpoint and one RTX 5090. The automatic planner resolved the following
+allocations; pool byte counts include its owned tensors:
+
+| Allocation | Previous BF16 image | FP8-FP4 image |
+| --- | ---: | ---: |
+| KV pool | 363.06 MiB | 176.40 MiB |
+| Usable full pages, 128 tokens each | 259 | 279 |
+| Usable window pages | 51 | 55 |
+| Cached experts | 818 | 830 |
+| Free GPU memory after initialization | 2.98 GiB | 2.97 GiB |
+
+The allocated KV pool decreased by 51.4% while page and expert counts increased.
+This is the planner's actual allocation, not a comparison at fixed capacity;
+freed bytes were reassigned, so total free GPU memory did not increase.
+
+Before API checks in each mode, both compressed Engram tables were read once
+with a 64 MiB buffer and `mincore` reported 100% residency. Identical streaming
+chat requests used `temperature: 0`, `thinking: {"type": "disabled"}` and
+`stream_options: {"include_usage": true}`. The three requests followed a short
+`2 + 2` warmup. Each cell below is request wall time / first content-token time:
+
+| Request | BF16 | FP8-FP4 first pass | FP8-FP4 repeat |
+| --- | ---: | ---: | ---: |
+| Japanese, 22 input / 42 output tokens | 37.22 / 17.46 s | 49.69 / 18.56 s | 24.58 / 15.08 s |
+| Corn image, 198 input tokens | 31.15 / 18.60 s | 52.43 / 21.41 s | 24.71 / 15.48 s |
+| Multi-turn arithmetic, 27 input / 2 output tokens | 17.93 / 17.69 s | 18.66 / 18.42 s | 15.37 / 15.13 s |
+
+The Japanese prompt was `日本語で、空が青く見える理由を短い一文で説明してください。`
+with `max_tokens: 48`. Its answer was identical in all three runs. The image
+request used the same local `corn.jpeg` and `Describe this image briefly.`, also
+with `max_tokens: 48`. Both modes correctly described three ears of corn, one
+partially husked; wording differed, producing 42 BF16 tokens versus 41 packed
+tokens. Multi-turn history was user `What is 4 + 5?`, assistant `9`, then user
+`Double that number. Reply with only the number.` with `max_tokens: 16`. All runs
+answered `18` and ended normally. Both packed passes produced identical text.
+Health remained `ok` after the checks.
+
+These single samples do not establish a throughput improvement or equivalent
+model quality. Initial packed requests were slower, and repeated requests were
+faster. The indexer specializes its key width during eager decode, so new
+positions can compile new kernels; packed mode uses separate specializations
+from BF16. This is a likely contributor to first-pass latency, not an isolated
+measurement of compilation cost. Different expert/page allocations and cache
+warming also affect these results. Prefill attention remains slower in the
+isolated comparison above, and unseen context lengths can incur more compilation.
+
+## Implementation and current limits
+
+- V4.1 has a separate model registration and nested-config parser. It implements
+  CSA2 ratio-1/ratio-2 compressed attention, shared source layers, two-stage index
+  selection, shifted mHC mixing, image-aware routing and Engram lookup.
+- Resident projections use FP8 weights with 32x32 UE8M0 scales. Routed experts use
+  the checkpoint's NVFP4 values, per-16 E4M3 scales and global scales in W4A16
+  execution. This activation precision differs from the donor reference and is
+  not a claim of bit-identical logits or unchanged model quality.
+- Window, compressed and index keys retain their required FP8/FP4 quantization.
+  `--kv-cache-dtype bf16` stores reconstructed values; `fp8-fp4` stores their
+  original codes and scales. Additional NVFP4 window quantization is not implemented.
+- Eager execution is required. Explicit CUDA graph capture is rejected because
+  the variable-length index search is not ready for practical long-context graphs.
+- Equal index scores select earlier positions deterministically. This preserves
+  chunk consistency but can differ from the reference's unspecified top-k ties.
+- MTP weights are excluded; speculative decoding is not implemented.
+- Engram supports FP8 E4M3 and packed FP4 E2M1, both with block-32 UE8M0 scales.
+  FP4 values use the low nibble first. This is distinct from the experts'
+  NVFP4 block-16 E4M3 scales and global scales.
+- FTW conversion streams Engram tables to a standalone `engram/` directory and
+  retains their format in a version-2 manifest and nested configuration. Legacy
+  version-1 FP8 manifests remain readable. The source snapshot is not
+  needed when serving the completed FTW directory.
+
+## Validation
+
+### LibertAIDAI FP4 Engram checkpoint
+
+The pinned checkpoint has 143,317 tensors across 48 shards. All 138,240 backbone
+expert components passed the native header validator, and all 1,480 resident
+parameters matched the meta model without missing names, extra names or shape
+differences. Its 46,080 backbone expert global scales range from 2^-7 to 2^-3;
+all are positive, finite and exactly representable in the loader's FP16 storage.
+An independent arithmetic E2M1 decoder matched the native BF16 lookup exactly
+for 1,002 sampled rows in each real Engram table, including the first/last rows,
+all 16 E2M1 codes, block scales, duplicate IDs and reordered queries. These tests
+ran against the installed code in image
+`sha256:a9525a2fdd6f7c8217266c15a6bb93c1bc85a64e86113f01c06b2607de86e7f2`;
+all 388 Python source files in that image matched the tested workspace.
+
+All 48 shards (429,406,627,896 bytes) matched their SHA256 download etags during
+Docker import. The final cache passed size checks for all 90 repository files
+and safetensors/index checks for all 143,317 tensors. This was checkpoint
+verification in the private validation environment, not a repository download
+or import workflow.
+
+On 2026-09-11 UTC, the private container setup started the image on the same
+RTX 5090/driver 591.86 host, using the [Direct launch](#direct-launch) settings
+with `--kv-cache-dtype bf16`. The API became ready at 20:23:03 UTC, 169 seconds
+after process startup; parallel expert reading took 144 seconds. The resolved
+cache retained 818 experts and 259 KV pages, with 14 CPU-decode layers and 26
+GPU-offload layers. The engine reported 2.98 GiB of free GPU memory after
+initialization. The model process used 289.95 GiB RSS with no process swap after
+the first text request; this excludes unmapped Engram file-cache pages.
+
+With `thinking: {"type": "disabled"}` and `temperature: 0`, the first text
+request (`What is 2 + 2? Reply with only the number.`, `max_tokens: 16`) returned
+`4` in 64.27 seconds, with 18 input and two output tokens. The corn image request
+shown below (`max_tokens: 32`) returned the same correct description as the
+earlier checkpoint in 48.30 seconds, with 201 input and 29 output tokens.
+These two prompts are functional checks, not a general quality comparison or
+a throughput benchmark.
+
+After these requests, both Engram shards were read once with a 64 MiB buffer
+while the model remained running. This took 72.18 seconds. `mincore` then
+reported 100% residency for both tables: all 104,451,107,600 compressed bytes
+(97.28 GiB) were in Linux page cache. Container `memory.current` was 390.73 GiB,
+and Windows still reported 73.76 GiB of free physical memory. WSL swap usage
+remained zero before and after warming.
+
+WSL `MemAvailable` was 180.99 GiB, which includes reclaimable Engram pages;
+subtracting the entire Engram cache gives approximately 83.71 GiB as a
+conservative remaining-memory estimate with those pages retained. Docker's usual
+stats display was about 292 GiB because it subtracts inactive file cache. Do not
+interpret that display as the total with all Engram pages cached. Page residency
+is an observation at measurement time; the OS can reclaim cached pages later.
+The same text prompt still returned `4` after warming (18 input/two output
+tokens, 25.48 seconds). This measurement does not isolate cache-warming benefits
+or establish improved generation throughput.
+After that generation, both tables still had 100% page residency, container
+memory was 390.67 GiB, model RSS was 290.28 GiB, and both process and WSL swap
+usage were zero. The API remained healthy.
+
+The compatibility change passed 37 configuration/weight tests and a separate
+37-test Engram/model/Engine regression run. After the final CPU-device guard,
+the Engram suite passed 28 tests. These runs overlap. They cover independent
+E2M1 decoding, requested-row-only reads, packed FTW roundtrips, legacy FP8
+manifests, CUDA staging and image prefill/decode in the small engine.
+
+### Earlier s-zaizen FP8 Engram checkpoint
+
+The earlier s-zaizen checkpoint at revision
+`179b7cda25486efbaaf8637d696759d9a791d8bd` was loaded and served successfully on
+2026-09-11 UTC with one RTX 5090, NVIDIA driver 591.86, CUDA 13 and
+PyTorch `2.11.0+cu130`. The tested Docker image was
+`sha256:87805e7660460773ef16d6018bc547fafe783b98d019d7c8b5158ac08da7a18f`.
+All 48 checkpoint shards, totaling 527,293,384,648 bytes, matched their SHA256
+etags; all 188,245 tensor entries matched the checkpoint index.
+
+The private container setup used the [Direct launch](#direct-launch) settings
+with `--kv-cache-dtype bf16`, the local s-zaizen snapshot as `--model`, and
+`s-zaizen/DeepSeek-V4.1-Flash-NVFP4` as `--served-model-name`.
+
+This run used a 32,768-token context, 1,024-token prefill chunks, two concurrent
+request slots, parallel expert loading, automatic MoE cache sizing and disabled
+prefill overlap. Startup ran from 19:11:16 to API readiness at 19:14:14 UTC
+(178 seconds); the parallel expert read took 151 seconds. The resolved cache had
+818 expert slots and 259 KV pages. RAM use after initialization was 290.7 GiB;
+the engine reported 2.94 GiB of free GPU memory. The 285.04 GiB expert banks used
+14 CPU-decode layers and 26 GPU-offload layers, matching the WSL pin-budget plan.
+
+HTTP chat requests used `thinking: {"type": "disabled"}` and `temperature: 0`:
+
+| Request | Input / output tokens | Wall time | Result |
+|---|---:|---:|---|
+| First text request, `max_tokens: 16` | 18 / 2 | 73.15 s | `4` |
+| Image request, `max_tokens: 32` | 201 / 29 | 49.14 s | Description below |
+| Same text request after the image, `max_tokens: 16` | 18 / 2 | 17.78 s | `4` |
+
+The text prompt was `What is 2 + 2? Reply with only the number.` The image request
+used the model repository's `corn.jpeg` with
+`Describe this image in one short sentence.` Its response was:
+
+> Three ears of fresh corn with green husks are arranged on a white background, with the central ear partially peeled to reveal bright yellow kernels.
+
+Health checks succeeded, and `/v1/models` reported the configured served model ID
+and context length 32768. These are limited functional smoke tests. Request wall
+times include all request processing; the cause of the first-request delay was
+not isolated. General text/image quality, full million-token inference, sustained
+performance and FTW startup with this complete checkpoint remain unvalidated.
+
+Automated tests compare operations with independent CPU/dequantization references
+and exercise a complete small V4.1 model, including NVFP4 experts, Engram, vision,
+chunked prefill, decode and request isolation. A separate 1,048,576-key index
+search test checks bounded temporary GPU allocation.
+
+```bash
+uv run pytest tests/models/test_deepseek_v41* tests/models/test_dsv41_attention.py \
+  tests/kernels/test_dsv41* tests/kvcache/test_dsv41_pool.py \
+  tests/scheduler/test_multimodal_chunks.py -q
+```
+
+The broader suite run reported 1,858 passes and six failures. One missing AOT
+registration introduced by this change was fixed, followed by 30 passing targeted
+tests. Four GLM failures also reproduced at the unchanged HEAD (`04d462149e21`). An intermittent
+device-to-device copy test failed in the broader run and passed five isolated
+reruns; this does not establish a clean full-suite result. Separate focused runs
+passed nine tests in the private validation image.
+These counts describe separate runs and are not added together.
+
+The adapted reference components retain their upstream MIT attribution in
+`python/freetoken/models/deepseek_v41/NOTICE`.
diff --git a/docs/models.md b/docs/models.md
index 548a363c6..0b09f4c5d 100644
--- a/docs/models.md
+++ b/docs/models.md
@@ -20,6 +20,9 @@ for them; other checkpoints of the same architectures work too.
 | MiniMax-M3 | [nvidia/MiniMax-M3-NVFP4](https://huggingface.co/nvidia/MiniMax-M3-NVFP4) |
 | Muse-Glimmer | [meta-models/Muse-Glimmer-30B](https://huggingface.co/meta-models/Muse-Glimmer-30B), [RedHatAI/Muse-Glimmer-30B-NVFP4](https://huggingface.co/RedHatAI/Muse-Glimmer-30B-NVFP4) |
 
+This branch also includes experimental DeepSeek-V4.1 Flash NVFP4 text and image
+support. See [the V4.1 setup and validation notes](deepseek-v41.md).
+
 ### Image input
 
 These families accept image input by default; pass `--text-model-only` to skip the vision encoder. The flags are described in the
@@ -27,6 +30,7 @@ These families accept image input by default; pass `--text-model-only` to skip t
 
 | Family | Image tokens | `--image-min-tokens` / `--image-max-tokens` | `--mm-processor-kwargs` example |
 | --- | --- | --- | --- |
+| DeepSeek-V4.1 Flash (experimental, native ViT tower) | resized patch grid with start, row-newline and end tokens | maximum covers the entire image span; minimum is rejected; use `vision_min_pixels` instead | `{"vision_min_pixels": 295936, "vision_max_n_token": 2048}` |
 | Qwen3.6 (both variants, every listed weight format), Qwen3.8-Flash-Next, Qwen3-VL | one token per 32x32 pixels of the resized image, dynamic resolution | pixel areas in `size.shortest_edge` / `longest_edge`; checkpoint defaults 64 to 16384 tokens | `{"size": {"longest_edge": 1048576}}` |
 | Gemma-4 26B-A4B, 31B (`gemma4`: ViT tower, streamed under `--mm-encoder-weights host`) | one of the soft-token budgets 70 / 140 / 280 / 560 / 1120, every image scaled to its budget as far as the aspect ratio allows | the maximum picks the largest budget within it, below 70 is refused at start-up; the minimum has no effect | `{"max_soft_tokens": 1120}` |
 | Gemma-4 12B (`gemma4_unified`: linear patch embedder, resident under either placement) | same budgets, one 48x48 super-patch per soft token | same as the tower releases | same |
@@ -38,7 +42,7 @@ These families accept image input by default; pass `--text-model-only` to skip t
 
 `ft serve --moe-strategy {auto,fused,offload,cpu,hybrid}` (`--moe-backend` is the deprecated old spelling):
 
-- **fused** — experts resident on GPU (needs the VRAM); never auto-selected.
+- **fused** — experts resident on GPU (needs the memory); auto-selected on unified-memory GPUs for supported BF16 and block-FP8 formats.
 - **offload** — experts live in host RAM, an LRU cache of expert slots on GPU;
   misses stream over PCIe.
 - **cpu** — misses are computed on the CPU instead of fetched.
@@ -46,7 +50,8 @@ These families accept image input by default; pass `--text-model-only` to skip t
   CPU, overlapped. Run `ft bench bw` once per machine to calibrate the split.
 - **auto** — dense models always resolve to `fused`; MoE models resolve to
   `offload`, upgraded to `hybrid` when a cached `ft bench bw` profile
-  recommends it.
+  recommends it. On unified-memory GPUs, supported resident formats resolve to
+  `fused`; other formats keep `offload` and skip the hybrid upgrade.
 
 ## Notes
 
@@ -60,3 +65,9 @@ These families accept image input by default; pass `--text-model-only` to skip t
 - DeepSeek-V4 checkpoints must keep the `inference/config.json` subdir — the
   authoritative model args are read from there.
 - Qwen3.8-Flash-Next keeps a 47.7 GiB PLE n-gram table pinned in host RAM.
+- `--kv-cache-dtype fp8` (see [cli.md](cli.md#fp8-kv-cache)) covers the plain paged,
+  hybrid-SWA and QSA sparse KV pools — gpt-oss, Qwen3/3.5/3.6, GLM-4.x, Gemma-4,
+  MiniMax-M2.5, Muse-Glimmer, Llama/Qwen2/Mistral, Qwen3.8-Flash-Next (on QSA only the
+  selected K/V rows are read back as codes; block selection keeps 16-bit index keys).
+  MLA/DSA (GLM-5.2), DeepSeek-V4's tiered pool and MiniMax-M3's block-sparse pool stay
+  16-bit and reject it.
diff --git a/pyproject.toml b/pyproject.toml
index 5511c1f2b..f7120179a 100644
--- a/pyproject.toml
+++ b/pyproject.toml
@@ -47,7 +47,7 @@ dependencies = [
     "numpy>=2.0,<2.5",
     "openai>=2.0,<3",
     "partial-json-parser>=0.2,<1",
-    "pillow>=10,<13",
+    "pillow>=11,<13",
     "prompt_toolkit>=3.0,<4",
     "pydantic>=2.9,<3",
     "pyzmq>=27,<28",
@@ -117,6 +117,7 @@ where = ["python"]
 
 [tool.setuptools.package-data]
 "*" = ["csrc/**/*", "moe/configs/**/*.json"]
+"freetoken.models.deepseek_v41" = ["NOTICE"]
 
 [tool.pytest.ini_options]
 testpaths = ["tests"]
diff --git a/python/freetoken/attention/__init__.py b/python/freetoken/attention/__init__.py
index 8a410ffe4..08b7ab5fd 100644
--- a/python/freetoken/attention/__init__.py
+++ b/python/freetoken/attention/__init__.py
@@ -33,6 +33,13 @@ class BackendInfo:
     # Whether forward() honors a per-call AttentionSpec (window/sm_scale/sinks).
     # Non-consumers raise on a non-None spec instead of silently dropping it.
     consumes_attn_spec: bool = False
+    # Whether forward() reads an fp8 KV pool (codes + per-token/per-head scales).
+    # Backends that hand the cache to an external kernel must opt out until that
+    # kernel is proven to apply our scale layout; the engine then refuses (or auto-
+    # avoids) them for --kv-cache-dtype fp8.
+    supports_fp8_kv: bool = False
+    supports_nvfp4_kv: bool = False
+    supports_fp8_fp4_kv: bool = False
 
 
 SUPPORTED_ATTENTION_BACKENDS = Registry[BackendCreator]("Attention Backend")
@@ -84,6 +91,8 @@ def create_fa_backend(config: ModelConfig):
     BackendInfo(
         supported_types=frozenset({AttnType.FULL, AttnType.SWA}),
         consumes_attn_spec=True,
+        supports_fp8_kv=True,
+        supports_nvfp4_kv=True,
     ),
 )
 def create_triton_backend(config: ModelConfig):
@@ -102,9 +111,23 @@ def create_dsv4_sparse_backend(config: ModelConfig):
     return DSV4SparseAttnBackend(config)
 
 
+@SUPPORTED_ATTENTION_BACKENDS.register(
+    "dsv41_sparse",
+    BackendInfo(supported_types=frozenset({AttnType.DSV41}), supports_fp8_fp4_kv=True),
+)
+def create_dsv41_sparse_backend(config: ModelConfig):
+    from .dsv41_sparse import DSV41SparseAttnBackend
+
+    return DSV41SparseAttnBackend(config)
+
+
 @SUPPORTED_ATTENTION_BACKENDS.register(
     "dsa",
-    BackendInfo(supported_types=frozenset({AttnType.MLA, AttnType.DSA})),
+    BackendInfo(
+        supported_types=frozenset({AttnType.MLA, AttnType.DSA}),
+        supports_fp8_kv=True,
+        supports_nvfp4_kv=True,
+    ),
 )
 def create_dsa_backend(config: ModelConfig):
     # MLA with a grouped index (index_ratio > 1) is the kpool indexer layout.
@@ -137,6 +160,10 @@ def create_m3_sparse_backend(config: ModelConfig):
     "qsa_sparse",
     BackendInfo(
         supported_types=frozenset({AttnType.QSA}),
+        # The attend kernel dequantizes on load (kernel/triton/qsa/attend.py); the
+        # compressed index keys it scores against are a separate, always-16-bit tier.
+        supports_fp8_kv=True,
+        supports_nvfp4_kv=True,
         # 64-token pages: a 4-token compress group never straddles a page, so the
         # compressed row of a group is page_base // 4 + block-in-page.
         page_sizes=(64,),
diff --git a/python/freetoken/attention/base.py b/python/freetoken/attention/base.py
index 9115b8092..344e2b517 100644
--- a/python/freetoken/attention/base.py
+++ b/python/freetoken/attention/base.py
@@ -20,6 +20,7 @@ class AttnType(str, Enum):
     MLA = "mla"  # plain latent-KV MLA -> MLAKVCache
     DSA = "dsa"  # latent-KV MLA + DSA sparse indexer -> DSAKVCache
     DSV4 = "dsv4"  # DSV4 window+compressed sparse -> DSV4PagedKVCache
+    DSV41 = "dsv41"
     LINEAR = "linear"  # GDN/mamba state layers -> LinearStatePool
     # GQA block-sparse (MiniMax-M3): paged GQA K/V + a per-sparse-layer index-key
     # slab; the indexer picks top-k 128-token blocks per query -> BSAKVCache
diff --git a/python/freetoken/attention/dsa.py b/python/freetoken/attention/dsa.py
index 398d17002..2d6781da5 100644
--- a/python/freetoken/attention/dsa.py
+++ b/python/freetoken/attention/dsa.py
@@ -201,6 +201,9 @@ def _attend(
         return glm_dsa_sparse_attn(
             q_cat, self.kvcache.latent_rows(layer_id), sel, self.sm_scale,
             counts=cnt, d_v=self.kv_lora_rank,
+            pool_scale=self.kvcache.latent_scale(layer_id),
+            kv_quant=self.kvcache.kv_quant,
+            pool_block_scale=self.kvcache.latent_block_scale(layer_id),
         )
 
     def mla_forward(
diff --git a/python/freetoken/attention/dsv41_sparse.py b/python/freetoken/attention/dsv41_sparse.py
new file mode 100644
index 000000000..208e17fba
--- /dev/null
+++ b/python/freetoken/attention/dsv41_sparse.py
@@ -0,0 +1,65 @@
+"""V4.1 sparse attention with explicit compressed-KV source ownership."""
+
+from __future__ import annotations
+
+import torch
+
+from freetoken.core import get_global_ctx
+from .dsv4_sparse import DSV4SparseAttnBackend
+
+
+class DSV41SparseAttnBackend(DSV4SparseAttnBackend):
+    def __init__(self, config):
+        self.config = config
+        self.device = get_global_ctx().kv_cache.device
+        self.window_size = config.dsv41_args.window_size
+        self.capture = None
+        self.capture_bs = []
+        self.max_graph_bs = 0
+        self._window_ar = torch.arange(self.window_size, device=self.device)
+        self.begin_forward()
+
+    def begin_forward(self):
+        self.shared_indices = {}
+        self.shared_candidates = {}
+
+    def scatter_compressed(self, layer_id, tier, rows, kv):
+        if tier == "attn":
+            self.pool.store_compressed(kv, layer_id, rows)
+        elif tier == "idx":
+            self.pool.store_indexer(kv, layer_id, rows)
+        else:
+            raise ValueError(f"Unknown V4.1 compressed tier: {tier}")
+
+    def attend(self, q, layer_id, topk_idxs, n_window, attn_sink, softmax_scale,
+               cmp_counts=None, has_compression=True):
+        pool = self.pool
+        source = pool.kv_sources[layer_id]
+        compressed = pool.cmp_pool[source] if source is not None else pool.window_pool[layer_id]
+        packed = getattr(pool, "kv_quant", "none") == "fp8-fp4"
+        if q.is_cuda:
+            if packed:
+                from freetoken.kernel.triton.dsv41.sparse_attn import sparse_attn_paged
+            else:
+                from freetoken.kernel.triton.dsv4.sparse_attn import sparse_attn_paged
+            return sparse_attn_paged(q, pool.window_pool[layer_id], compressed, attn_sink,
+                                     topk_idxs.int(), n_window, softmax_scale, cmp_counts)
+        flat = q.reshape(-1, *q.shape[-2:])
+        ids = topk_idxs.reshape(flat.shape[0], -1).long()
+        window = pool.window_pool[layer_id][ids[:, :n_window].clamp_min(0)]
+        cmps = compressed[ids[:, n_window:].clamp_min(0)]
+        if packed:
+            from freetoken.kernel.triton.dsv41.quant import unpack_fp4, unpack_fp8
+            window = unpack_fp8(window, block_size=32, dtype=torch.bfloat16)
+            cmps = (unpack_fp4(cmps, block_size=16, scale_format="e4m3", dtype=torch.bfloat16)
+                    if source is not None else window[:, :0])
+        kv = torch.cat((window, cmps), 1)
+        logits = torch.einsum("qhd,qkd->qhk", flat.float(), kv.float()) * softmax_scale
+        live = ids >= 0
+        if cmp_counts is not None:
+            columns = torch.arange(ids.shape[1], device=ids.device)
+            live &= columns[None] < n_window + cmp_counts.reshape(-1, 1)
+        logits.masked_fill_(~live[:, None, :], -torch.inf)
+        sink = attn_sink.float().view(1, -1, 1).expand(flat.shape[0], -1, 1)
+        prob = torch.cat((logits, sink), -1).softmax(-1)[..., :-1]
+        return torch.einsum("qhk,qkd->qhd", prob, kv.float()).to(q.dtype).view_as(q)
diff --git a/python/freetoken/attention/qsa_sparse.py b/python/freetoken/attention/qsa_sparse.py
index b07c95232..842ff4651 100644
--- a/python/freetoken/attention/qsa_sparse.py
+++ b/python/freetoken/attention/qsa_sparse.py
@@ -112,7 +112,17 @@ def __init__(self, config: ModelConfig) -> None:
             f"qsa_sparse backend needs a QSA pool, got {type(self.kvcache).__name__}"
         )
         self.device = self.kvcache.device
+        # The pool's COMPUTE dtype, never its store dtype (the contract lives in
+        # kvcache/base.py). These buffers feed the indexer -- qsa_index_norm_rope and
+        # qsa_mqa_paged -- whose tl.dot has no fp8 path, so an e4m3 q_index does not
+        # fail here, it fails at CUDA-graph capture with "Unsupported rhs dtype
+        # fp8e4nv". --kv-cache-dtype fp8 quantizes only the KV tiers; the index tiers
+        # stay 16-bit by design (kvcache/qsa_pool.py).
         self.dtype = self.kvcache.dtype
+        assert self.dtype.itemsize == 2, (
+            f"QSA block selection needs a 16-bit compute dtype, got {self.dtype} -- "
+            "the KV pool must report its compute dtype, not e4m3 codes"
+        )
         self.index_head_dim = self.kvcache.index_head_dim
         self.ratio = self.kvcache.index_ratio
         self.ring_capacity = self.kvcache.ring_capacity
@@ -299,6 +309,8 @@ def qsa_forward(
 
         self._update_index_cache(index, md, slot)
         indices = self._select(index, md, slot)
+        # K/V scale tensors are independent of the BF16 index tier, so selection is
+        # quantization-agnostic; only sparse K/V attention reconstructs the codes.
         return qsa_sparse_paged_attention(
             q,
             self.kvcache.k_cache(layer_id),
@@ -307,6 +319,11 @@ def qsa_forward(
             md.block_table,
             md.token_to_req,
             torch.empty_like(q),
+            k_scale=self.kvcache.k_scale(layer_id),
+            v_scale=self.kvcache.v_scale(layer_id),
+            kv_quant=self.kvcache.kv_quant,
+            k_block_scale=self.kvcache.k_block_scale(layer_id),
+            v_block_scale=self.kvcache.v_block_scale(layer_id),
         )
 
     def _plan_index_writes(self, md: QSASparseMetadata, batch: Batch) -> None:
diff --git a/python/freetoken/attention/triton.py b/python/freetoken/attention/triton.py
index 8bf9792d5..2b0137cb4 100644
--- a/python/freetoken/attention/triton.py
+++ b/python/freetoken/attention/triton.py
@@ -151,10 +151,19 @@ def forward(
 
         k_raw = self.kvcache.k_cache(layer_id)
         v_raw = self.kvcache.v_cache(layer_id)
-        kv_heads, head_dim = k_raw.shape[-2], k_raw.shape[-1]
-        assert head_dim == q.shape[-1]
-        k_cache = k_raw.view(-1, kv_heads, head_dim)
-        v_cache = v_raw.view(-1, kv_heads, head_dim)
+        kv_heads, stored_dim = k_raw.shape[-2], k_raw.shape[-1]
+        head_dim = q.shape[-1]
+        kv_quant = getattr(self.kvcache, "kv_quant", "none")
+        assert stored_dim == (head_dim // 2 if kv_quant == "nvfp4" else head_dim)
+        k_cache = k_raw.view(-1, kv_heads, stored_dim)
+        v_cache = v_raw.view(-1, kv_heads, stored_dim)
+        k_block_scale = self.kvcache.k_block_scale(layer_id) if kv_quant == "nvfp4" else None
+        v_block_scale = self.kvcache.v_block_scale(layer_id) if kv_quant == "nvfp4" else None
+        # An fp8 KV pool hands us its per-(token, head) scales; a 16-bit pool returns
+        # None and every kernel below keeps its original (scale-free) code path.
+        k_scale = self.kvcache.k_scale(layer_id)
+        v_scale = self.kvcache.v_scale(layer_id)
+        assert (k_scale is None) == (v_scale is None), "K and V scales come as a pair"
 
         spec = attn_spec or AttentionSpec()
         block_ends = batch.mm_block_ends if spec.bidirectional_mm_blocks else None
@@ -182,6 +191,11 @@ def forward(
                 sm_scale=scale,
                 sliding_window=spec.sliding_window,
                 sinks=spec.sinks,
+                k_scale=k_scale,
+                v_scale=v_scale,
+                kv_quant=kv_quant,
+                k_block_scale=k_block_scale,
+                v_block_scale=v_block_scale,
             )
         if (
             (not metadata.is_decode)
@@ -202,6 +216,11 @@ def forward(
                 sinks=spec.sinks,
                 k_extend=k.view(q.shape[0], kv_heads, head_dim),
                 v_extend=v.view(q.shape[0], kv_heads, head_dim),
+                k_scale=k_scale,
+                v_scale=v_scale,
+                kv_quant=kv_quant,
+                k_block_scale=k_block_scale,
+                v_block_scale=v_block_scale,
                 block_ends=block_ends,
             )
         if block_ends is not None:
@@ -217,6 +236,11 @@ def forward(
             sm_scale=scale,
             sliding_window=spec.sliding_window,
             sinks=spec.sinks,
+            k_scale=k_scale,
+            v_scale=v_scale,
+            kv_quant=kv_quant,
+            k_block_scale=k_block_scale,
+            v_block_scale=v_block_scale,
         )
 
     def prepare_metadata(self, batch: Batch) -> None:
diff --git a/python/freetoken/checkpoint/convert.py b/python/freetoken/checkpoint/convert.py
index a875d77f3..af321408f 100644
--- a/python/freetoken/checkpoint/convert.py
+++ b/python/freetoken/checkpoint/convert.py
@@ -8,6 +8,8 @@
 * offload experts = exactly what ``load_expert_banks(parallel=True)`` produces (post
   backend-repack pinned banks + alpha scale vectors) -> ``kind="experts_bank"`` (alphas are
   told apart at load by their reserved names, so they need no separate kind).
+* V4.1 Engram tables = original compressed FP8 or FP4 bytes in separate mapped files,
+  with their dtype and scale layout recorded in ``engram_tables.json``.
 
 The output directory is a self-contained checkpoint (config + tokenizer copied), so you can
 point ``--model`` straight at it; the load path auto-detects the FTW and reads it (FTW).
@@ -296,6 +298,10 @@ def convert_checkpoint(
 
     _progress("finalize")  # writing shard index + copying config/tokenizer
     copied = _copy_metadata(model_path, out_dir)
+    if getattr(mc, "dsv41_args", None) is not None and mc.dsv41_args.engram_layer_ids:
+        from freetoken.models.deepseek_v41.engram import export_engram_tables
+
+        copied.extend(export_engram_tables(model_path, out_dir, mc.dsv41_args))
 
     try:
         fingerprint = _source_fingerprint(model_path, mc, device=dev)
diff --git a/python/freetoken/checkpoint/ftw.py b/python/freetoken/checkpoint/ftw.py
index e2bd0de24..bd52b0aef 100644
--- a/python/freetoken/checkpoint/ftw.py
+++ b/python/freetoken/checkpoint/ftw.py
@@ -616,7 +616,8 @@ def _read_layer(job):
     # the file names the banks the legacy way; the quant_format tag names the (kind, kernel) they were packed for
     sources = {canonical_role(name): views for name, views in sources.items()}
     quant_format = reader.meta("quant_format")
-    kind, kernel = kind_kernel_for(quant_format) if quant_format is not None else (None, None)
+    # Q4_0 still uses its GGUF provider rather than a quant method.
+    kind, kernel = kind_kernel_for(quant_format) if quant_format not in (None, "q4_0") else (None, None)
 
     # a failed mlock leaves a LOCKED layer pageable; the log and labels report what the banks actually settled at
     applied = list(residency)
diff --git a/python/freetoken/core.py b/python/freetoken/core.py
index e07e442ca..494076c69 100644
--- a/python/freetoken/core.py
+++ b/python/freetoken/core.py
@@ -39,6 +39,10 @@ class Req:
     uid: int
     sampling_params: SamplingParams
     cache_handle: BaseCacheHandle
+    # Optional precomputed multimodal soft-token embeddings (GPU, [num_image_tokens,
+    # hidden]) scattered at image-token positions during this request's prefill.
+    mm_embeds: torch.Tensor | None = None
+    media: list[dict] | None = None
     # per-item processor outputs and the tokenizer's precomputed mrope rows and delta
     mm_items: list | None = None
     mrope_positions_full: torch.Tensor | None = None  # [3, prompt_len] int32, CPU
diff --git a/python/freetoken/engine/config.py b/python/freetoken/engine/config.py
index 44c5a38e7..5dd497165 100644
--- a/python/freetoken/engine/config.py
+++ b/python/freetoken/engine/config.py
@@ -66,6 +66,12 @@ class EngineConfig:
     cuda_graph_bs: List[int] | None = None
     cuda_graph_max_bs: int | None = None
     page_size: int = 1
+    # KV-cache storage quantization: "none" stores the compute dtype, "fp8" stores e4m3
+    # codes plus one fp32 scale per (token, slab, layer, kv head) -- about 2x the tokens
+    # per GiB, at a small accuracy cost. --kv-cache-dtype; resolved from "auto" by
+    # _adjust_config, which also refuses it on a pool family or attention backend that
+    # cannot read the scales.
+    kv_quant: str = "none"
     memory_ratio: float = 0.9
     # Hybrid GDN models default to the HybridRadixCache (cross-request GDN-state prefix reuse);
     # `--cache-type naive` opts out. linear_state_cache_ratio sizes the GDN snapshot cache as
diff --git a/python/freetoken/engine/engine.py b/python/freetoken/engine/engine.py
index 06a74ee5f..99d2f23bc 100644
--- a/python/freetoken/engine/engine.py
+++ b/python/freetoken/engine/engine.py
@@ -89,6 +89,8 @@ def _required_attn_types(model_config) -> frozenset[AttnType]:
     dsv4_args marks DSV4 (the real config declares a DSV4 attention group)."""
     specs_fn = getattr(model_config, "kv_cache_group_specs", None)
     if specs_fn is None:
+        if getattr(model_config, "dsv41_args", None) is not None:
+            return frozenset({AttnType.DSV41})
         if getattr(model_config, "dsv4_args", None) is not None:
             return frozenset({AttnType.DSV4})
         return frozenset({AttnType.FULL})
@@ -118,12 +120,48 @@ def _backend_requirements_met(name: str) -> bool:
     return True
 
 
-def _resolve_auto_attention_backend(required: frozenset[AttnType]) -> str:
+# --kv-cache-dtype spellings -> the stored EngineConfig.kv_quant value.
+KV_QUANT_ALIASES = {
+    "auto": "none", "bf16": "none", "none": "none", "fp8": "fp8",
+    "nvfp4": "nvfp4", "fp8-fp4": "fp8-fp4",
+}
+
+
+def _resolve_kv_quant(value: str | None) -> str:
+    """Normalize a --kv-cache-dtype spelling to EngineConfig.kv_quant."""
+    key = (value or "auto").strip().lower()
+    if key not in KV_QUANT_ALIASES:
+        raise ValueError(
+            f"unknown --kv-cache-dtype {value!r}; expected one of "
+            f"{', '.join(sorted(KV_QUANT_ALIASES))}"
+        )
+    return KV_QUANT_ALIASES[key]
+
+
+def _backend_supports_kv_quant(name: str, kv_quant: str) -> bool:
+    """Whether every comma part of an attention-backend string can read a quantized
+    KV pool (an unquantized pool needs nothing from the backend)."""
+    if kv_quant == "none":
+        return True
+    if kv_quant not in ("fp8", "nvfp4", "fp8-fp4"):
+        return False
+    return all(
+        getattr(attention_backend_info(part.strip()), f"supports_{kv_quant.replace('-', '_')}_kv")
+        for part in name.split(",")
+    )
+
+
+def _resolve_auto_attention_backend(
+    required: frozenset[AttnType], *, kv_quant: str = "none"
+) -> str:
     """First candidate (in per-type priority order) whose arch condition holds,
-    whose packages are installed, and whose every comma part serves ALL required
-    types. Reproduces the historical hardware tree for FULL-only models:
+    whose packages are installed, whose every comma part serves ALL required
+    types, and which can decode a quantized KV cache when one is configured.
+    Reproduces the historical hardware tree for FULL-only models:
     sm_100 -> trtllm, sm_90+sgl_kernel -> "fa,fi", flashinfer -> fi, else triton."""
     candidates: list[tuple[str, bool]] = []
+    if AttnType.DSV41 in required:
+        candidates.append(("dsv41_sparse", True))
     if AttnType.DSV4 in required:
         candidates.append(("dsv4_sparse", True))
     if required & {AttnType.MLA, AttnType.DSA}:
@@ -148,10 +186,18 @@ def _resolve_auto_attention_backend(required: frozenset[AttnType]) -> str:
             continue
         if not _backend_requirements_met(name):
             continue
+        if not _backend_supports_kv_quant(name, kv_quant):
+            continue
         return name
     raise RuntimeError(
         "No attention backend can serve attention types "
-        f"{sorted(t.value for t in required)} on this machine."
+        f"{sorted(t.value for t in required)} on this machine"
+        + (
+            f" with a {kv_quant} KV cache"
+            if kv_quant != "none"
+            else ""
+        )
+        + "."
     )
 
 
@@ -177,7 +223,7 @@ def _validate_attention_backend_choice(config, override, required: frozenset[Att
             valid = [
                 name
                 for name in (
-                    "fa", "fi", "trtllm", "triton", "dsa", "dsv4_sparse", "m3_sparse",
+                    "fa", "fi", "trtllm", "triton", "dsa", "dsv4_sparse", "dsv41_sparse", "m3_sparse",
                     "qsa_sparse",
                 )
                 if required <= attention_backend_info(name).supported_types
@@ -196,6 +242,23 @@ def _validate_attention_backend_choice(config, override, required: frozenset[Att
                 f"SWA models require, got {config.attention_backend!r}."
             )
 
+    # A quantized KV pool is only readable by a backend that applies its per-(token,
+    # head) scales; one that hands the cache to an external kernel would silently
+    # attend to raw e4m3 codes. Rejected here, before any weight is resident.
+    kv_quant = getattr(config, "kv_quant", "none")
+    if not _backend_supports_kv_quant(config.attention_backend, kv_quant):
+        fp8_backends = [
+            name
+            for name in ("trtllm", "fi", "fa", "triton")
+            if required <= attention_backend_info(name).supported_types
+            and _backend_supports_kv_quant(name, kv_quant)
+        ]
+        raise ValueError(
+            f"--kv-cache-dtype {kv_quant} needs an attention backend that decodes the KV "
+            f"scales; {config.attention_backend!r} does not. Valid for this model: "
+            f"{', '.join(fp8_backends) or 'none'} (or use --kv-cache-dtype bf16)."
+        )
+
     # An explicitly-selected backend may require a package that isn't installed. Auto
     # never resolves to one of these when its package is missing, so this only fires for
     # explicit --attention-backend choices.
@@ -258,6 +321,8 @@ def _make_dummy_weight_state_dict(
             t = torch.empty(param.shape, dtype=param.dtype, device=device)
             t.view(torch.uint8).random_(0, 16)
             state_dict[key] = t
+        elif param.dtype == torch.float8_e8m0fnu:
+            state_dict[key] = torch.full(param.shape, 127, dtype=torch.uint8, device=device).view(param.dtype)
         elif param.dtype.is_floating_point or param.dtype.is_complex:
             state_dict[key] = torch.randn(param.shape, dtype=param.dtype, device=device)
         elif param.dtype == torch.uint8 and key.endswith("weight_scale_inv"):
@@ -342,14 +407,6 @@ def __init__(self, config: EngineConfig):
                 )
             # before the residency snapshot, so streamed blocks are not charged as resident weights
             self.model.place_encoder_weights(config.mm.encoder_weights)
-        post_weights_free = self._sync_get_memory()[0]
-        self._weights_bytes = self._baseline_free - post_weights_free
-        # Pool-budget baseline for the desktop cache sliders: free VRAM after the weights are
-        # resident but before ANY runtime cache pool (MoE expert cache below, KV pages, GDN
-        # state) is allocated. This is the stable "if all free VRAM went to one pool" budget —
-        # unlike a query-time mem_get_info it doesn't drift with allocator caching, CUDA
-        # graphs, or other processes. Cross-rank MIN, deterministic across ranks.
-        self._post_weights_free = post_weights_free
         self.moe_offload_cache = None
         self.cpu_moe_executor = None
         # Host-side auxiliary stores (qwen4_exp's pinned PLE table): after the weights so a
@@ -358,6 +415,11 @@ def __init__(self, config: EngineConfig):
         self._host_tables_bytes = 0
         if hasattr(self.model, "load_host_tables"):
             self._host_tables_bytes = int(self.model.load_host_tables(config) or 0)
+        # Auxiliary host tables can own persistent GPU staging (V4.1 Engram). Measure it
+        # with the resident weights before dividing the remaining budget between pools.
+        post_weights_free = self._sync_get_memory()[0]
+        self._weights_bytes = self._baseline_free - post_weights_free
+        self._post_weights_free = post_weights_free
         if is_offload_moe_strategy(config.moe_strategy):
             self._init_offload_moe_cache(config)
         if hasattr(self.model, "prepare_for_runtime"):
@@ -395,6 +457,12 @@ def __init__(self, config: EngineConfig):
         self.ctx.kv_cache = self.kv_cache = create_kv_pool(
             config, self.num_pages, device=self.device, dtype=self.dtype
         )
+        if config.kv_quant == "fp8-fp4":
+            logger.info_rank0(
+                "KV storage fp8-fp4: window=FP8 E4M3/UE8M0 block32, "
+                "compressed=FP4 E2M1/E4M3 block16, index=MXFP4 E2M1/UE8M0 block32; "
+                f"pool={self.kv_cache.total_bytes() / (1024 ** 2):.2f} MiB"
+            )
 
         # ======================= Linear (GatedDeltaNet) state initialization ========================
         linear_group = config.model_config.linear_attention_group()
@@ -552,6 +620,9 @@ def _run_mm_encoder(self, batch: Batch) -> None:
             parts.append(cache.get_slice(item_hash, lo, hi, self.device))
             cache.consume(item_hash, uid, hi - lo)
         if parts:
+            # The scheduler appends legacy precomputed rows after the MMItem gather rows.
+            if batch.mm_embeds is not None:
+                parts.append(batch.mm_embeds)
             batch.mm_embeds = torch.cat([p.to(self.dtype) for p in parts], dim=0)
 
     def _resolve_auto_moe_cache_size(self, config: EngineConfig, banks, method=None) -> tuple[int, int, bool]:
@@ -1104,7 +1175,7 @@ def _is_unified_memory_gpu(index: "int | None" = None) -> bool:
 
 def _fused_resident_ok(model_config) -> bool:
     """Whether the resident ('fused') MoE path can hold this model's experts.
-   
+
        FIXME: auto resolves to fused only for bf16 and fp8_block experts; drop this gate once the other quant formats support fused.
        """
     expert_quant = getattr(model_config, "expert_quant", "none")
@@ -1149,21 +1220,25 @@ def _resolve_cache_type(has_linear_attention: bool, requested: str) -> str:
 
 
 def _adjust_dsv4_config(config: EngineConfig, override) -> None:
-    """DSV4 engine-config reconciliation at config-resolution time (before the pool exists).
-    Syncs the resolved runtime config into the opaque ``dsv4_args`` payload, sets
-    page_size to the window page P, forces single-chunk prefill, and clamps cuda_graph_bs/max_bs to
-    the DSV4 decode batch size.
-    """
+    """Resolve the V4/V4.1 paged-window geometry before allocating the model."""
     model_config = config.model_config
-    model_config.dsv4_args.max_seq_len = config.max_seq_len
-    model_config.dsv4_args.max_batch_size = config.max_running_req + 1  # +1 dummy
+    is_v41 = getattr(model_config, "dsv41_args", None) is not None
+    args = model_config.dsv41_args if is_v41 else model_config.dsv4_args
+    args.max_seq_len = config.max_seq_len
+    args.max_batch_size = config.max_running_req + 1  # +1 dummy
     # config.swa_full_tokens_ratio is the DSV4 window/full ratio directly (default sizing);
     # a runtime rebuild pins an absolute window via swa_num_pages_override instead.
     # DSV4's KV page IS the P-token window page (window == radix reuse granularity == lcm of
     # the compress ratios), so max_num_tokens = num_pages * page_size holds like every model.
-    P = model_config.dsv4_args.window_size
+    P = args.window_size
     override("page_size", P)
-    logger.info_rank0(f"DSV4 KV pages are {P}-token window pages; page_size set to {P}")
+    model_name = "DSV4.1" if is_v41 else "DSV4"
+    logger.info_rank0(f"{model_name} KV pages are {P}-token window pages; page_size set to {P}")
+    if is_v41:
+        if not 0 < config.swa_full_tokens_ratio <= 1:
+            raise ValueError("DeepSeek-V4.1 swa_full_tokens_ratio must be in (0, 1]")
+        if getattr(config, "max_extend_tokens", config.max_seq_len) < min(P, config.max_seq_len):
+            raise ValueError(f"DeepSeek-V4.1 --max-prefill-length must fit one {P}-token page")
     # The generic CacheManager materializes DSV4 'radix' as the shared SWARadixCache (is_swa);
     # 'naive' stays naive with the pool's swa currency riding swa_paged.
     if getattr(config, "cache_type", "radix") != "naive":
@@ -1172,7 +1247,7 @@ def _adjust_dsv4_config(config: EngineConfig, override) -> None:
     # honored, as is an explicit 'naive'. Don't let max_extend_tokens force a second chunk within
     # one prompt (the pool's prefill_chunk_budget still chunks prompts larger than the window
     # pool); prefill batches ragged (bs>=1), each segment resuming from its own cached_len.
-    if getattr(config, "max_extend_tokens", 0) < config.max_seq_len:
+    if not is_v41 and getattr(config, "max_extend_tokens", 0) < config.max_seq_len:
         override("max_extend_tokens", config.max_seq_len)
 
     # DSV4 decode batches at most max_running_req rows; its full-loc snapshot is sized to that,
@@ -1423,7 +1498,15 @@ def override(attr: str, value: Any):  # this is dangerous, use with caution
 
     model_config = config.model_config
     single_stream_only = getattr(model_config, "single_stream_only", False)
-    is_dsv4 = getattr(model_config, "dsv4_args", None) is not None
+    if getattr(model_config, "dsv41_args", None) is not None:
+        if config.dtype != torch.bfloat16:
+            raise ValueError("DeepSeek-V4.1 requires --dtype bfloat16 for its quantized attention path")
+        if (getattr(config, "cuda_graph_max_bs", None) or 0) > 0 or getattr(config, "cuda_graph_bs", None):
+            raise ValueError("DeepSeek-V4.1 currently uses eager execution; set --cuda-graph-max-bs 0")
+        override("cuda_graph_max_bs", 0)
+        override("cuda_graph_bs", [])
+    is_dsv4 = (getattr(model_config, "dsv4_args", None) is not None
+               or getattr(model_config, "dsv41_args", None) is not None)
     has_swa_attention = getattr(model_config, "has_swa_attention", False)
     has_linear_attention = getattr(model_config, "has_linear_attention", False)
     is_moe = getattr(model_config, "is_moe", False)
@@ -1494,6 +1577,40 @@ def override(attr: str, value: Any):  # this is dangerous, use with caution
     # lists, then validate whatever is now selected (explicit or auto) -- every
     # comma part must serve every required type, with packages/arch available.
     required_attn_types = _required_attn_types(model_config)
+    # Resolve KV quantization BEFORE the backend tree: a quantized pool narrows both
+    # which pool families are usable and which backend auto may pick.
+    kv_quant = _resolve_kv_quant(getattr(config, "kv_quant", "none"))
+    override("kv_quant", kv_quant)
+    if kv_quant == "nvfp4":
+        if required_attn_types - {AttnType.FULL, AttnType.SWA, AttnType.QSA, AttnType.MLA, AttnType.DSA}:
+            raise ValueError(
+                "--kv-cache-dtype nvfp4 requires a paged FULL, hybrid-SWA, QSA, or MLA/DSA KV pool"
+            )
+        for spec in model_config.kv_cache_group_specs():
+            if spec.head_dim % 16:
+                raise ValueError("--kv-cache-dtype nvfp4 requires head_dim divisible by 16")
+    if kv_quant == "fp8-fp4":
+        if required_attn_types != {AttnType.DSV41}:
+            raise ValueError("--kv-cache-dtype fp8-fp4 requires DeepSeek-V4.1 attention")
+        from freetoken.kvcache.dsv41_layout import dsv41_row_bytes
+
+        dsv41_row_bytes(model_config.dsv41_args, kv_quant)
+    elif kv_quant != "none":
+        # Quantized codes are wired through the pools that hand their rows to a Triton
+        # kernel: plain paged and hybrid-SWA, QSA sparse, and DSA/MLA. QSA's index
+        # tier and DSA's index-key/tail tiers stay bf16; the DSA kernel dequantizes
+        # selected latent rows with their per-token scale. Other sparse families have
+        # no scale-read path and remain rejected before weights load.
+        quant_unsupported = required_attn_types - {
+            AttnType.FULL, AttnType.SWA, AttnType.QSA, AttnType.MLA, AttnType.DSA,
+        }
+        if quant_unsupported:
+            raise ValueError(
+                f"--kv-cache-dtype {kv_quant} is implemented for the plain paged, "
+                "hybrid-SWA, QSA sparse and DSA/MLA KV pools; this model also needs "
+                f"{', '.join(sorted(t.value for t in quant_unsupported))} attention "
+                "(use --kv-cache-dtype bf16)."
+            )
     _dtype = getattr(config, "dtype", None)  # duck-typed test configs omit it
     if (
         required_attn_types & {AttnType.BSA, AttnType.QSA}
@@ -1521,7 +1638,7 @@ def override(attr: str, value: Any):  # this is dangerous, use with caution
     if config.attention_backend == "auto":
         override(
             "attention_backend",
-            _resolve_auto_attention_backend(required_attn_types),
+            _resolve_auto_attention_backend(required_attn_types, kv_quant=kv_quant),
         )
         logger.info_rank0(f"Auto-selected attention backend: {config.attention_backend}")
     _validate_attention_backend_choice(config, override, required_attn_types)
diff --git a/python/freetoken/kernel/aot_models.py b/python/freetoken/kernel/aot_models.py
index a8ef85d43..142c7d7a1 100644
--- a/python/freetoken/kernel/aot_models.py
+++ b/python/freetoken/kernel/aot_models.py
@@ -11,12 +11,14 @@
 (a drifted derivation misses the prebuilt cache by spec name and falls back to
 JIT, which needs nvcc):
 
-- store: ``element_size = num_kv_heads * head_dim * 2`` (bf16 KV row), one per
-  paged-KV attention group (kvcache/mha_pool.py, kvcache/hybrid_swa_pool.py).
-  DSV4 writes its MLA latent via torch scatter and contributes nothing.
+- store: ``element_size = num_kv_heads * head_dim * dtype_bytes``, one per
+  paged-KV attention group (kvcache/mha_pool.py, kvcache/hybrid_swa_pool.py) and
+  per KV width: 2 for the 16-bit cache, 1 for an fp8 one (``--kv-cache-dtype
+  fp8``, kvcache/mha_pool.py). DSV4/V4.1 write MLA latents via torch scatter and
+  contribute nothing.
 - index: ``element_size = hidden_size * 2`` (bf16 embedding row) paired with
   the runtime ``num_splits_for`` rule (layers/embedding.py -> kernel/index.py).
-  DSV4 (plain nn.Embedding) and GGUF embeddings (GGUFEmbedding) bypass it.
+  DSV4/V4.1 (plain nn.Embedding) and GGUF embeddings (GGUFEmbedding) bypass it.
 
 The whole table targets the shipped serving configuration: TP=1 (TP>1 shards
 kv heads, shrinking the store row) and a 2-byte compute dtype (``--dtype
@@ -33,7 +35,8 @@
 
 from .index import num_splits_for
 
-KV_CACHE_DTYPE_BYTES = 2  # every current model allocates bf16 paged KV
+KV_CACHE_DTYPE_BYTES = 2  # the default paged KV is the 16-bit compute dtype
+FP8_KV_CACHE_DTYPE_BYTES = 1  # --kv-cache-dtype fp8 stores one e4m3 code per element
 EMBED_DTYPE_BYTES = 2  # embedding weights stay bf16 on the indexing() path
 
 
@@ -332,6 +335,16 @@ def expert_bank_row_bytes(fmt: str, hidden_size: int, moe_intermediate_size: int
         expert_formats=("ds_fp4",),
         embed_indexing=False,  # plain nn.Embedding
     ),
+    AotModel(
+        name="s-zaizen/DeepSeek-V4.1-Flash-NVFP4",
+        architecture="DeepseekV41ForCausalLM",
+        hidden_size=5120,
+        kv_groups=(),
+        top_k=6,
+        moe_intermediate_size=2304,
+        expert_formats=("nvfp4",),
+        embed_indexing=False,
+    ),
     # ---- dense checkpoints (store/index only, no expert banks) ----
     AotModel(
         name="Qwen/Qwen3.6-27B",
@@ -395,8 +408,9 @@ def expert_bank_row_bytes(fmt: str, hidden_size: int, moe_intermediate_size: int
 )
 
 
-def store_element_sizes(model: AotModel) -> set[int]:
-    return {kv * hd * KV_CACHE_DTYPE_BYTES for kv, hd in model.kv_groups}
+def store_element_sizes(model: AotModel, dtype_bytes: int = KV_CACHE_DTYPE_BYTES) -> set[int]:
+    """Store-kernel row sizes for one model's paged-KV groups at a given bytes/elem."""
+    return {kv * hd * dtype_bytes for kv, hd in model.kv_groups}
 
 
 def index_variants(model: AotModel) -> set[tuple[int, int]]:
@@ -418,9 +432,17 @@ def fast_index_copy_feature_sizes(model: AotModel) -> set[int]:
 
 
 def aggregate_store_element_sizes() -> tuple[int, ...]:
+    """Every store row size the runtime can ask for.
+
+    Both KV widths ship: the 16-bit default and the fp8 (``--kv-cache-dtype fp8``)
+    code buffer, whose rows are exactly half as wide. A missing size is not a
+    correctness bug -- it is a kernel-cache miss that falls back to JIT and fails
+    the ``FREETOKEN_DISABLE_JIT=1`` release gate.
+    """
     sizes: set[int] = set()
     for model in SUPPORTED_MODELS:
-        sizes.update(store_element_sizes(model))
+        for dtype_bytes in (KV_CACHE_DTYPE_BYTES, FP8_KV_CACHE_DTYPE_BYTES):
+            sizes.update(store_element_sizes(model, dtype_bytes))
     return tuple(sorted(sizes))
 
 
diff --git a/python/freetoken/kernel/triton/attention.py b/python/freetoken/kernel/triton/attention.py
index 3f53c5fc0..0515df16e 100644
--- a/python/freetoken/kernel/triton/attention.py
+++ b/python/freetoken/kernel/triton/attention.py
@@ -6,6 +6,33 @@
 import triton
 import triton.language as tl
 
+from freetoken.kernel.triton.e4m3_compat import KV_TILE_SCALE
+from freetoken.kernel.triton.e4m3_compat import kv_load_e4m3_tile_f32 as _kv_load_f32
+from freetoken.kernel.triton.e4m3_compat import (
+    kv_load_e4m3_tile_scaled16 as _kv_load_s16,
+)
+from freetoken.kernel.triton.kv_nvfp4 import load_nvfp4
+
+
+@triton.jit
+def _kv_dequant_scale(scale_ptr, slots, stride, kv_head, mask):
+    """Per-(token, kv_head) dequant scale for an fp8 KV tile, pre-multiplied by the
+    2**8 that :func:`kv_load_e4m3_tile_scaled16` leaves on the tile it returns.
+
+    Applied to a dot's OUTPUT rather than to K or V. The scale is constant down the
+    reduction dim, so ``scores[m,n] = (sum_d q[m,d]*k[d,n]) * s_k[n]`` and
+    ``p @ (diag(s_v) @ v) = (p * s_v[None,:]) @ v`` both hold: scaling the
+    ``BLOCK_M x BLOCK_N`` result costs less than scaling the ``BLOCK_D x BLOCK_N`` K
+    tile or the ``BLOCK_N x BLOCK_DV`` V tile, by a factor of head_dim / BLOCK_M.
+
+    It is also the more accurate order, which is why the 16-bit tile is safe: every
+    value the loader returns carries at most the code's own 3 mantissa bits and lies
+    in +-1.75, so narrowing it to the compute dtype is lossless, whereas multiplying
+    by a general scale first and narrowing after rounds the product.
+    """
+    s = tl.load(scale_ptr + slots * stride + kv_head, mask=mask, other=0.0)
+    return s * KV_TILE_SCALE
+
 
 _MAX_KV_SPLITS = 8
 _MIN_BLOCK_KV = 32
@@ -18,25 +45,34 @@ def _optin_smem_bytes(device_index: int) -> int:
     return int(getattr(props, "shared_memory_per_block_optin", 0))
 
 
-def _select_extend_tile(head_dim: int, block_d: int, smem_optin: int) -> tuple[int, int]:
+def _select_extend_tile(
+    head_dim: int, block_d: int, smem_optin: int, kv_bytes: int = 2
+) -> tuple[int, int]:
     """Pick ``(BLOCK_M, BLOCK_N)`` for the extend/prefill kernel, shared-memory aware.
 
-    Larger tiles run materially faster (~2x for head_dim 512 on H100) but their bf16
-    q/k/v tiles need about ``(BLOCK_M + 2 * BLOCK_N) * BLOCK_D * 2`` bytes of shared
-    memory, which overflows consumer GPUs (sm_89 ~99KB opt-in) once head_dim >= 256.
-    Keep the fast tiles where the device's opt-in shared memory fits them (datacenter
-    A100/H100); shrink only where it does not. ``smem_optin == 0`` (unknown) conservatively
-    selects the small tiles, i.e. the prior consumer-safe behavior.
+    Larger tiles run materially faster (~2x for head_dim 512 on H100) but their q/k/v
+    tiles need shared memory, which overflows consumer GPUs (sm_89 ~99KB opt-in) once
+    head_dim >= 256. Keep the fast tiles where the device's opt-in shared memory fits
+    them (datacenter A100/H100); shrink only where it does not. ``smem_optin == 0``
+    (unknown) conservatively selects the small tiles, i.e. the prior consumer-safe
+    behavior.
+
+    ``kv_bytes`` is the KV cache's element size: q is always 2 bytes/element but K and
+    V follow the cache, so a 1-byte fp8 cache fits a tile a 16-bit one cannot. Passing
+    2 reproduces the previous budget exactly, so the 16-bit ladder is unchanged.
     """
     budget = smem_optin * 0.8  # headroom for scores/acc/alignment/triton scratch
 
     def fits(block_m: int, block_n: int) -> bool:
-        return (block_m + 2 * block_n) * block_d * 2 <= budget
+        return (block_m * 2 + 2 * block_n * kv_bytes) * block_d <= budget
 
     if head_dim <= 128:
         return 128, 64
     if head_dim <= 256:
-        return (128, 64) if fits(128, 64) else (64, 32)
+        if fits(128, 64):
+            return 128, 64
+        # Reachable by an fp8 cache where a 16-bit one falls through to 64x32.
+        return (64, 64) if fits(64, 64) else (64, 32)
     if head_dim <= 384:
         return (32, 64) if fits(32, 64) else (32, 32)
     return (32, 64) if fits(32, 64) else (16, 16)
@@ -47,6 +83,10 @@ def _paged_attention_kernel(
     q_ptr,
     k_ptr,
     v_ptr,
+    k_scale_ptr,
+    v_scale_ptr,
+    k_block_ptr,
+    v_block_ptr,
     o_ptr,
     indptr_ptr,
     indices_ptr,
@@ -60,6 +100,8 @@ def _paged_attention_kernel(
     stride_kh,
     stride_vs,
     stride_vh,
+    stride_kss,
+    stride_vss,
     stride_ot,
     stride_oh,
     GROUP: tl.constexpr,
@@ -68,6 +110,8 @@ def _paged_attention_kernel(
     BLOCK_N: tl.constexpr,
     SLIDING_WINDOW: tl.constexpr,
     HAS_SINKS: tl.constexpr,
+    HAS_KV_SCALE: tl.constexpr,
+    KV_NVFP4: tl.constexpr,
 ):
     q_tok = tl.program_id(0)
     q_head = tl.program_id(1)
@@ -107,14 +151,36 @@ def _paged_attention_kernel(
         skip_tile = tl.max(mask_n.to(tl.int32), axis=0) == 0
         if not skip_tile:
             slots = tl.load(indices_ptr + kv_start + offs_n, mask=offs_n < kv_len, other=0)
-            k = tl.load(
-                k_ptr
-                + slots[:, None] * stride_ks
-                + kv_head * stride_kh
-                + offs_d[None, :],
-                mask=(offs_n[:, None] < kv_len) & mask_d[None, :],
-                other=0.0,
-            ).to(tl.float32)
+            if KV_NVFP4:
+                k = load_nvfp4(
+                    k_ptr, k_block_ptr, k_scale_ptr, slots[:, None], kv_head,
+                    offs_d[None, :],
+                    offs_n[:, None] < kv_len, stride_ks, stride_kh, stride_kss, D,
+                )
+            elif HAS_KV_SCALE:
+                # fp8 KV: the codes carry magnitude, the per-(token, head) fp32 scale
+                # restores it. Index math stays int32 like the 16-bit path below.
+                s_k = tl.load(
+                    k_scale_ptr + slots * stride_kss + kv_head,
+                    mask=offs_n < kv_len,
+                    other=0.0,
+                )
+                k = _kv_load_f32(
+                    k_ptr
+                    + slots[:, None] * stride_ks
+                    + kv_head * stride_kh
+                    + offs_d[None, :],
+                    (offs_n[:, None] < kv_len) & mask_d[None, :],
+                ) * s_k[:, None]
+            else:
+                k = tl.load(
+                    k_ptr
+                    + slots[:, None] * stride_ks
+                    + kv_head * stride_kh
+                    + offs_d[None, :],
+                    mask=(offs_n[:, None] < kv_len) & mask_d[None, :],
+                    other=0.0,
+                ).to(tl.float32)
             scores = tl.sum(q[None, :] * k, axis=1) * sm_scale
             scores = tl.where(mask_n, scores, -float("inf"))
 
@@ -124,14 +190,34 @@ def _paged_attention_kernel(
             alpha = tl.exp(m_i - m_new)
             p = tl.exp(scores - m_new)
 
-            v = tl.load(
-                v_ptr
-                + slots[:, None] * stride_vs
-                + kv_head * stride_vh
-                + offs_d[None, :],
-                mask=(offs_n[:, None] < kv_len) & mask_d[None, :],
-                other=0.0,
-            ).to(tl.float32)
+            if KV_NVFP4:
+                v = load_nvfp4(
+                    v_ptr, v_block_ptr, v_scale_ptr, slots[:, None], kv_head,
+                    offs_d[None, :],
+                    offs_n[:, None] < kv_len, stride_vs, stride_vh, stride_vss, D,
+                )
+            elif HAS_KV_SCALE:
+                s_v = tl.load(
+                    v_scale_ptr + slots * stride_vss + kv_head,
+                    mask=offs_n < kv_len,
+                    other=0.0,
+                )
+                v = _kv_load_f32(
+                    v_ptr
+                    + slots[:, None] * stride_vs
+                    + kv_head * stride_vh
+                    + offs_d[None, :],
+                    (offs_n[:, None] < kv_len) & mask_d[None, :],
+                ) * s_v[:, None]
+            else:
+                v = tl.load(
+                    v_ptr
+                    + slots[:, None] * stride_vs
+                    + kv_head * stride_vh
+                    + offs_d[None, :],
+                    mask=(offs_n[:, None] < kv_len) & mask_d[None, :],
+                    other=0.0,
+                ).to(tl.float32)
             acc = acc * alpha + tl.sum(p[:, None] * v, axis=0)
             l_i = l_i * alpha + tl.sum(p, axis=0)
             m_i = m_new
@@ -149,6 +235,10 @@ def _decode_grouped_stage1_kernel(
     q_ptr,
     k_ptr,
     v_ptr,
+    k_scale_ptr,
+    v_scale_ptr,
+    k_block_ptr,
+    v_block_ptr,
     sm_scale,
     indptr_ptr,
     indices_ptr,
@@ -162,6 +252,8 @@ def _decode_grouped_stage1_kernel(
     stride_kh,
     stride_vs,
     stride_vh,
+    stride_kss,
+    stride_vss,
     stride_mid_ob,
     stride_mid_oh,
     stride_mid_os,
@@ -179,6 +271,8 @@ def _decode_grouped_stage1_kernel(
     D: tl.constexpr,
     DV: tl.constexpr,
     SLIDING_WINDOW: tl.constexpr,
+    HAS_KV_SCALE: tl.constexpr,
+    KV_NVFP4: tl.constexpr,
 ):
     batch_id = tl.program_id(0)
     head_block_id = tl.program_id(1)
@@ -224,7 +318,10 @@ def _decode_grouped_stage1_kernel(
 
     if split_end > split_start:
         q = tl.load(q_ptr + q_offsets, mask=mask_h[:, None] & mask_d[None, :], other=0.0)
-        q = q.to(k_ptr.dtype.element_ty)
+        if not HAS_KV_SCALE:
+            # A 16-bit cache feeds tl.dot as-is; an fp8 cache is decoded up to q's own
+            # compute dtype below, so q must NOT be narrowed to the (1-byte) cache type.
+            q = q.to(k_ptr.dtype.element_ty)
 
         for rel_start in tl.range(split_start, split_end, BLOCK_N):
             rel_offs = rel_start + tl.arange(0, BLOCK_N)
@@ -232,24 +329,56 @@ def _decode_grouped_stage1_kernel(
             logical_offs = effective_start + rel_offs
             slots = tl.load(indices_ptr + kv_start + logical_offs, mask=mask_n, other=0)
 
-            k = tl.load(
-                k_ptr + slots[None, :] * stride_ks + k_base_offsets,
-                mask=mask_n[None, :] & mask_d[:, None],
-                other=0.0,
-            )
+            if KV_NVFP4:
+                k = load_nvfp4(
+                    k_ptr, k_block_ptr, k_scale_ptr, slots[None, :], kv_head,
+                    offs_d[:, None],
+                    mask_n[None, :], stride_ks, stride_kh, stride_kss, D,
+                ).to(q.dtype)
+            elif HAS_KV_SCALE:
+                s_k = _kv_dequant_scale(k_scale_ptr, slots, stride_kss, kv_head, mask_n)
+                k = _kv_load_s16(
+                    k_ptr + slots[None, :] * stride_ks + k_base_offsets,
+                    mask_n[None, :] & mask_d[:, None],
+                ).to(q.dtype)
+            else:
+                k = tl.load(
+                    k_ptr + slots[None, :] * stride_ks + k_base_offsets,
+                    mask=mask_n[None, :] & mask_d[:, None],
+                    other=0.0,
+                )
             scores = tl.dot(q, k) * sm_scale
+            # NVFP4's loader already applies both row and block scales.
+            if HAS_KV_SCALE and not KV_NVFP4:
+                scores = scores * s_k[None, :]
             scores = tl.where(mask_h[:, None] & mask_n[None, :], scores, -float("inf"))
 
-            v = tl.load(
-                v_ptr + slots[:, None] * stride_vs + v_base_offsets,
-                mask=mask_n[:, None] & mask_dv[None, :],
-                other=0.0,
-            )
+            if KV_NVFP4:
+                v = load_nvfp4(
+                    v_ptr, v_block_ptr, v_scale_ptr, slots[:, None], kv_head,
+                    offs_dv[None, :],
+                    mask_n[:, None], stride_vs, stride_vh, stride_vss, D,
+                ).to(q.dtype)
+            elif HAS_KV_SCALE:
+                s_v = _kv_dequant_scale(v_scale_ptr, slots, stride_vss, kv_head, mask_n)
+                v = _kv_load_s16(
+                    v_ptr + slots[:, None] * stride_vs + v_base_offsets,
+                    mask_n[:, None] & mask_dv[None, :],
+                ).to(q.dtype)
+            else:
+                v = tl.load(
+                    v_ptr + slots[:, None] * stride_vs + v_base_offsets,
+                    mask=mask_n[:, None] & mask_dv[None, :],
+                    other=0.0,
+                )
 
             m_new = tl.maximum(tl.max(scores, axis=1), m_i)
             alpha = tl.exp(m_i - m_new)
             p = tl.exp(scores - m_new[:, None])
-            acc = acc * alpha[:, None] + tl.dot(p.to(v.dtype), v)
+            # p itself stays unscaled: l_i is the softmax denominator and knows
+            # nothing about V's quantization.
+            pv = (p * s_v[None, :]) if HAS_KV_SCALE and not KV_NVFP4 else p
+            acc = acc * alpha[:, None] + tl.dot(pv.to(v.dtype), v)
             l_i = l_i * alpha + tl.sum(p, axis=1)
             m_i = m_new
 
@@ -349,6 +478,27 @@ def _decode_stage2_kernel(
     )
 
 
+def _validate_kv_format(q, k, v, kr, vr, quant, kb, vb):
+    quant = quant if quant is not None else ("fp8" if kr is not None else "none")
+    if quant not in ("none", "fp8", "nvfp4"):
+        raise ValueError(f"unknown kv_quant {quant!r}")
+    d = q.shape[-1]
+    assert k.shape == v.shape
+    assert k.stride(-1) == v.stride(-1) == 1
+    assert (kr is not None) == (vr is not None) == (quant != "none")
+    if quant == "nvfp4":
+        assert d % 16 == 0 and k.shape[-1] == d // 2
+        for codes, row, block in ((k, kr, kb), (v, vr, vb)):
+            assert codes.dtype == torch.uint8
+            assert row.dtype == torch.float32 and row.is_contiguous()
+            assert block is not None and block.dtype == torch.uint8
+            assert block.shape == (*codes.shape[:2], d // 16) and block.is_contiguous()
+            assert codes.device == row.device == block.device == q.device
+    else:
+        assert k.shape[-1] == d and kb is None and vb is None
+    return quant == "nvfp4"
+
+
 def decode_paged_attention(
     q: torch.Tensor,
     k_cache: torch.Tensor,
@@ -364,16 +514,31 @@ def decode_paged_attention(
     sliding_window: int | None = None,
     sinks: torch.Tensor | None = None,
     out: torch.Tensor | None = None,
+    k_scale: torch.Tensor | None = None,
+    v_scale: torch.Tensor | None = None,
+    kv_quant: str | None = None,
+    k_block_scale: torch.Tensor | None = None,
+    v_block_scale: torch.Tensor | None = None,
 ) -> torch.Tensor:
-    """SGLang-style split-k grouped decode attention for one query per request."""
+    """SGLang-style split-k grouped decode attention for one query per request.
+
+    ``kv_quant="nvfp4"`` requires packed half-width codes, FP32 row scales and
+    uint8 E4M3 block scales. Omitted ``kv_quant`` retains FP8 scale inference.
+
+    ``k_scale`` / ``v_scale`` (``[num_slots, num_kv_heads]`` fp32) turn the cache into
+    an fp8 KV cache: every code row is multiplied by its own token/head scale. Both
+    must be given together; ``None`` keeps the 16-bit path byte-identical.
+    """
 
     assert q.is_cuda and k_cache.is_cuda and v_cache.is_cuda
+    assert (k_scale is None) == (v_scale is None), "K and V scales come as a pair"
     assert q.dim() == 3 and k_cache.dim() == 3 and v_cache.dim() == 3
     batch, num_q_heads, head_dim = q.shape
     num_kv_heads = k_cache.shape[1]
     assert batch == indptr.numel() - 1
     assert v_cache.shape[1] == num_kv_heads
-    assert k_cache.shape[-1] == head_dim and v_cache.shape[-1] == head_dim
+    nvfp4 = _validate_kv_format(q, k_cache, v_cache, k_scale, v_scale,
+                                kv_quant, k_block_scale, v_block_scale)
     assert num_q_heads % num_kv_heads == 0
     assert attn_logits.shape[0] >= batch
     assert attn_logits.shape[1] >= num_q_heads
@@ -391,6 +556,17 @@ def decode_paged_attention(
     o = out if out is not None else torch.empty_like(q)
     sinks_arg = sinks if sinks is not None else q
     group = num_q_heads // num_kv_heads
+    # Unused pointer args still need a real tensor (same convention as sinks_arg).
+    has_kv_scale = k_scale is not None
+    k_scale_arg = k_scale if has_kv_scale else k_cache
+    v_scale_arg = v_scale if has_kv_scale else v_cache
+    k_block_arg = k_block_scale if nvfp4 else k_cache
+    v_block_arg = v_block_scale if nvfp4 else v_cache
+    if has_kv_scale:
+        assert k_scale.shape == v_scale.shape == (k_cache.shape[0], num_kv_heads), (
+            tuple(k_scale.shape),
+            tuple(k_cache.shape),
+        )
     # valid_block_h = heads computed per program (drives the grid + head indexing); block_h =
     # power-of-two tile size for tl.arange. They differ only for non-power-of-two GQA groups
     # (e.g. 6), where block_h rounds up and the kernel masks the extra lanes.
@@ -405,6 +581,10 @@ def decode_paged_attention(
         q,
         k_cache,
         v_cache,
+        k_scale_arg,
+        v_scale_arg,
+        k_block_arg,
+        v_block_arg,
         sm_scale,
         indptr,
         indices,
@@ -418,6 +598,8 @@ def decode_paged_attention(
         k_cache.stride(1),
         v_cache.stride(0),
         v_cache.stride(1),
+        k_scale_arg.stride(0),
+        v_scale_arg.stride(0),
         attn_logits.stride(0),
         attn_logits.stride(1),
         attn_logits.stride(2),
@@ -435,6 +617,8 @@ def decode_paged_attention(
         D=head_dim,
         DV=head_dim,
         SLIDING_WINDOW=sliding_window or 0,
+        HAS_KV_SCALE=has_kv_scale,
+        KV_NVFP4=nvfp4,
         num_warps=4,
         num_stages=2,
     )
@@ -471,6 +655,10 @@ def _extend_attention_kernel(
     q_ptr,
     k_ptr,
     v_ptr,
+    k_scale_ptr,
+    v_scale_ptr,
+    k_block_ptr,
+    v_block_ptr,
     o_ptr,
     qo_indptr_ptr,
     kv_indptr_ptr,
@@ -485,6 +673,8 @@ def _extend_attention_kernel(
     stride_kh,
     stride_vs,
     stride_vh,
+    stride_kss,
+    stride_vss,
     stride_ot,
     stride_oh,
     GROUP: tl.constexpr,
@@ -495,6 +685,8 @@ def _extend_attention_kernel(
     BLOCK_N: tl.constexpr,
     SLIDING_WINDOW: tl.constexpr,
     HAS_SINKS: tl.constexpr,
+    HAS_KV_SCALE: tl.constexpr,
+    KV_NVFP4: tl.constexpr,
     HAS_BLOCKS: tl.constexpr,
 ):
     seq_id = tl.program_id(0)
@@ -555,15 +747,33 @@ def _extend_attention_kernel(
         skip_tile = tl.max(tl.max(final_mask.to(tl.int32), axis=1), axis=0) == 0
         if not skip_tile:
             slots = tl.load(kv_indices_ptr + kv_start + kv_offsets, mask=mask_n, other=0)
-            k = tl.load(
-                k_ptr
-                + slots[None, :] * stride_ks
-                + kv_head * stride_kh
-                + offs_d[:, None],
-                mask=mask_n[None, :] & mask_d[:, None],
-                other=0.0,
-            )
+            if KV_NVFP4:
+                k = load_nvfp4(
+                    k_ptr, k_block_ptr, k_scale_ptr, slots[None, :], kv_head,
+                    offs_d[:, None],
+                    mask_n[None, :], stride_ks, stride_kh, stride_kss, D,
+                ).to(q.dtype)
+            elif HAS_KV_SCALE:
+                s_k = _kv_dequant_scale(k_scale_ptr, slots, stride_kss, kv_head, mask_n)
+                k = _kv_load_s16(
+                    k_ptr
+                    + slots[None, :] * stride_ks
+                    + kv_head * stride_kh
+                    + offs_d[:, None],
+                    mask_n[None, :] & mask_d[:, None],
+                ).to(q.dtype)
+            else:
+                k = tl.load(
+                    k_ptr
+                    + slots[None, :] * stride_ks
+                    + kv_head * stride_kh
+                    + offs_d[:, None],
+                    mask=mask_n[None, :] & mask_d[:, None],
+                    other=0.0,
+                )
             scores = tl.dot(q.to(k.dtype), k) * sm_scale
+            if HAS_KV_SCALE and not KV_NVFP4:
+                scores = scores * s_k[None, :]
             scores = tl.where(final_mask, scores, -float("inf"))
 
             row_max = tl.max(scores, axis=1)
@@ -572,15 +782,33 @@ def _extend_attention_kernel(
             alpha = tl.exp(m_i - m_new)
             p = tl.exp(scores - m_new[:, None])
 
-            v = tl.load(
-                v_ptr
-                + slots[:, None] * stride_vs
-                + kv_head * stride_vh
-                + offs_dv[None, :],
-                mask=mask_n[:, None] & mask_dv[None, :],
-                other=0.0,
-            )
-            acc = acc * alpha[:, None] + tl.dot(p.to(v.dtype), v)
+            if KV_NVFP4:
+                v = load_nvfp4(
+                    v_ptr, v_block_ptr, v_scale_ptr, slots[:, None], kv_head,
+                    offs_dv[None, :],
+                    mask_n[:, None], stride_vs, stride_vh, stride_vss, D,
+                ).to(q.dtype)
+            elif HAS_KV_SCALE:
+                s_v = _kv_dequant_scale(v_scale_ptr, slots, stride_vss, kv_head, mask_n)
+                v = _kv_load_s16(
+                    v_ptr
+                    + slots[:, None] * stride_vs
+                    + kv_head * stride_vh
+                    + offs_dv[None, :],
+                    mask_n[:, None] & mask_dv[None, :],
+                ).to(q.dtype)
+            else:
+                v = tl.load(
+                    v_ptr
+                    + slots[:, None] * stride_vs
+                    + kv_head * stride_vh
+                    + offs_dv[None, :],
+                    mask=mask_n[:, None] & mask_dv[None, :],
+                    other=0.0,
+                )
+            # p stays unscaled for the l_i denominator below.
+            pv = (p * s_v[None, :]) if HAS_KV_SCALE and not KV_NVFP4 else p
+            acc = acc * alpha[:, None] + tl.dot(pv.to(v.dtype), v)
             l_i = l_i * alpha + tl.sum(p, axis=1)
             m_i = m_new
 
@@ -602,6 +830,10 @@ def _extend_attention_split_kernel(
     v_extend_ptr,
     k_cache_ptr,
     v_cache_ptr,
+    k_scale_ptr,
+    v_scale_ptr,
+    k_block_ptr,
+    v_block_ptr,
     o_ptr,
     qo_indptr_ptr,
     kv_indptr_ptr,
@@ -620,6 +852,8 @@ def _extend_attention_split_kernel(
     stride_kch,
     stride_vcs,
     stride_vch,
+    stride_kss,
+    stride_vss,
     stride_ot,
     stride_oh,
     GROUP: tl.constexpr,
@@ -630,6 +864,8 @@ def _extend_attention_split_kernel(
     BLOCK_N: tl.constexpr,
     SLIDING_WINDOW: tl.constexpr,
     HAS_SINKS: tl.constexpr,
+    HAS_KV_SCALE: tl.constexpr,
+    KV_NVFP4: tl.constexpr,
     HAS_BLOCKS: tl.constexpr,
 ):
     seq_id = tl.program_id(0)
@@ -684,15 +920,33 @@ def _extend_attention_split_kernel(
 
         if not skip_tile:
             slots = tl.load(kv_indices_ptr + kv_start + kv_offsets, mask=mask_n, other=0)
-            k = tl.load(
-                k_cache_ptr
-                + slots[None, :] * stride_kcs
-                + kv_head * stride_kch
-                + offs_d[:, None],
-                mask=mask_n[None, :] & mask_d[:, None],
-                other=0.0,
-            )
+            if KV_NVFP4:
+                k = load_nvfp4(
+                    k_cache_ptr, k_block_ptr, k_scale_ptr, slots[None, :], kv_head,
+                    offs_d[:, None],
+                    mask_n[None, :], stride_kcs, stride_kch, stride_kss, D,
+                ).to(q.dtype)
+            elif HAS_KV_SCALE:
+                s_k = _kv_dequant_scale(k_scale_ptr, slots, stride_kss, kv_head, mask_n)
+                k = _kv_load_s16(
+                    k_cache_ptr
+                    + slots[None, :] * stride_kcs
+                    + kv_head * stride_kch
+                    + offs_d[:, None],
+                    mask_n[None, :] & mask_d[:, None],
+                ).to(q.dtype)
+            else:
+                k = tl.load(
+                    k_cache_ptr
+                    + slots[None, :] * stride_kcs
+                    + kv_head * stride_kch
+                    + offs_d[:, None],
+                    mask=mask_n[None, :] & mask_d[:, None],
+                    other=0.0,
+                )
             scores = tl.dot(q.to(k.dtype), k) * sm_scale
+            if HAS_KV_SCALE and not KV_NVFP4:
+                scores = scores * s_k[None, :]
             scores = tl.where(final_mask, scores, -float("inf"))
 
             row_max = tl.max(scores, axis=1)
@@ -701,15 +955,33 @@ def _extend_attention_split_kernel(
             alpha = tl.exp(m_i - m_new)
             p = tl.exp(scores - m_new[:, None])
 
-            v = tl.load(
-                v_cache_ptr
-                + slots[:, None] * stride_vcs
-                + kv_head * stride_vch
-                + offs_dv[None, :],
-                mask=mask_n[:, None] & mask_dv[None, :],
-                other=0.0,
-            )
-            acc = acc * alpha[:, None] + tl.dot(p.to(v.dtype), v)
+            if KV_NVFP4:
+                v = load_nvfp4(
+                    v_cache_ptr, v_block_ptr, v_scale_ptr, slots[:, None], kv_head,
+                    offs_dv[None, :],
+                    mask_n[:, None], stride_vcs, stride_vch, stride_vss, D,
+                ).to(q.dtype)
+            elif HAS_KV_SCALE:
+                s_v = _kv_dequant_scale(v_scale_ptr, slots, stride_vss, kv_head, mask_n)
+                v = _kv_load_s16(
+                    v_cache_ptr
+                    + slots[:, None] * stride_vcs
+                    + kv_head * stride_vch
+                    + offs_dv[None, :],
+                    mask_n[:, None] & mask_dv[None, :],
+                ).to(q.dtype)
+            else:
+                v = tl.load(
+                    v_cache_ptr
+                    + slots[:, None] * stride_vcs
+                    + kv_head * stride_vch
+                    + offs_dv[None, :],
+                    mask=mask_n[:, None] & mask_dv[None, :],
+                    other=0.0,
+                )
+            # p stays unscaled for the l_i denominator below.
+            pv = (p * s_v[None, :]) if HAS_KV_SCALE and not KV_NVFP4 else p
+            acc = acc * alpha[:, None] + tl.dot(pv.to(v.dtype), v)
             l_i = l_i * alpha + tl.sum(p, axis=1)
             m_i = m_new
 
@@ -793,18 +1065,32 @@ def extend_paged_attention(
     out: torch.Tensor | None = None,
     k_extend: torch.Tensor | None = None,
     v_extend: torch.Tensor | None = None,
+    k_scale: torch.Tensor | None = None,
+    v_scale: torch.Tensor | None = None,
+    kv_quant: str | None = None,
+    k_block_scale: torch.Tensor | None = None,
+    v_block_scale: torch.Tensor | None = None,
     block_ends: torch.Tensor | None = None,
 ) -> torch.Tensor:
-    """Block-tiled causal prefill/extend attention over paged KV cache; block_ends holds per query token the end of the multimodal span it sits in (0 for none), whose later keys the row also attends."""
+    """Block-tiled causal prefill/extend attention over paged KV cache.
+
+    ``k_scale`` / ``v_scale`` mark an fp8 KV cache (see ``decode_paged_attention``).
+    The ``k_extend`` / ``v_extend`` rows are the current request's own K/V and stay in
+    the compute dtype either way, so only the cached prefix is decoded.
+    ``block_ends`` holds the exclusive end of each query's multimodal span (0 for
+    text), allowing image rows to attend to later keys within that span.
+    """
 
     assert q.is_cuda and k_cache.is_cuda and v_cache.is_cuda
+    assert (k_scale is None) == (v_scale is None), "K and V scales come as a pair"
     assert q.dim() == 3 and k_cache.dim() == 3 and v_cache.dim() == 3
     num_q_tokens, num_q_heads, head_dim = q.shape
     num_kv_heads = k_cache.shape[1]
     assert qo_indptr.numel() == kv_indptr.numel()
     assert prefix_lens.numel() == qo_indptr.numel() - 1
     assert v_cache.shape[1] == num_kv_heads
-    assert k_cache.shape[-1] == head_dim and v_cache.shape[-1] == head_dim
+    nvfp4 = _validate_kv_format(q, k_cache, v_cache, k_scale, v_scale,
+                                kv_quant, k_block_scale, v_block_scale)
     assert num_q_heads % num_kv_heads == 0
     if sinks is not None:
         assert sinks.is_cuda
@@ -816,14 +1102,27 @@ def extend_paged_attention(
         assert block_ends.is_cuda and block_ends.dtype == torch.int32 and block_ends.numel() == num_q_tokens
     o = out if out is not None else torch.empty_like(q)
     sinks_arg = sinks if sinks is not None else q
+    has_kv_scale = k_scale is not None
+    # Unused pointer args still need a real tensor (same convention as sinks_arg).
+    k_scale_arg = k_scale if has_kv_scale else k_cache
+    v_scale_arg = v_scale if has_kv_scale else v_cache
+    k_block_arg = k_block_scale if nvfp4 else k_cache
+    v_block_arg = v_block_scale if nvfp4 else v_cache
+    if has_kv_scale:
+        assert k_scale.shape == v_scale.shape == (k_cache.shape[0], num_kv_heads), (
+            tuple(k_scale.shape),
+            tuple(k_cache.shape),
+        )
     block_ends_arg = block_ends if block_ends is not None else qo_indptr
     block_d = triton.next_power_of_2(head_dim)
     block_dv = triton.next_power_of_2(head_dim)
     # Tile size is shared-memory bound: keep the fast (large) tiles on GPUs whose opt-in
     # shared memory fits them, shrink on consumer GPUs (sm_89 ~99KB) where the default
     # 128x64 overflows once head_dim >= 256 (e.g. gemma4: SWA 256, full-attention 512).
+    # NVFP4 restores full-width tiles with block scales; keep its 16-bit budget.
+    kv_tile_bytes = 2 if nvfp4 else k_cache.element_size()
     block_m, block_n = _select_extend_tile(
-        head_dim, block_d, _optin_smem_bytes(q.device.index)
+        head_dim, block_d, _optin_smem_bytes(q.device.index), kv_tile_bytes
     )
     grid = (qo_indptr.numel() - 1, num_q_heads, triton.cdiv(max_q_len, block_m))
     if k_extend is not None or v_extend is not None:
@@ -839,6 +1138,10 @@ def extend_paged_attention(
             v_extend,
             k_cache,
             v_cache,
+            k_scale_arg,
+            v_scale_arg,
+            k_block_arg,
+            v_block_arg,
             o,
             qo_indptr,
             kv_indptr,
@@ -857,6 +1160,8 @@ def extend_paged_attention(
             k_cache.stride(1),
             v_cache.stride(0),
             v_cache.stride(1),
+            k_scale_arg.stride(0),
+            v_scale_arg.stride(0),
             o.stride(0),
             o.stride(1),
             GROUP=num_q_heads // num_kv_heads,
@@ -867,6 +1172,8 @@ def extend_paged_attention(
             BLOCK_N=block_n,
             SLIDING_WINDOW=sliding_window or 0,
             HAS_SINKS=sinks is not None,
+            HAS_KV_SCALE=has_kv_scale,
+            KV_NVFP4=nvfp4,
             HAS_BLOCKS=block_ends is not None,
             num_warps=8,
             num_stages=1,
@@ -877,6 +1184,10 @@ def extend_paged_attention(
         q,
         k_cache,
         v_cache,
+        k_scale_arg,
+        v_scale_arg,
+        k_block_arg,
+        v_block_arg,
         o,
         qo_indptr,
         kv_indptr,
@@ -891,6 +1202,8 @@ def extend_paged_attention(
         k_cache.stride(1),
         v_cache.stride(0),
         v_cache.stride(1),
+        k_scale_arg.stride(0),
+        v_scale_arg.stride(0),
         o.stride(0),
         o.stride(1),
         GROUP=num_q_heads // num_kv_heads,
@@ -901,6 +1214,8 @@ def extend_paged_attention(
         BLOCK_N=block_n,
         SLIDING_WINDOW=sliding_window or 0,
         HAS_SINKS=sinks is not None,
+        HAS_KV_SCALE=has_kv_scale,
+        KV_NVFP4=nvfp4,
         HAS_BLOCKS=block_ends is not None,
         num_warps=8,
         num_stages=1,
@@ -921,20 +1236,28 @@ def paged_attention(
     sinks: torch.Tensor | None = None,
     out: torch.Tensor | None = None,
     block_n: int = 32,
+    k_scale: torch.Tensor | None = None,
+    v_scale: torch.Tensor | None = None,
+    kv_quant: str | None = None,
+    k_block_scale: torch.Tensor | None = None,
+    v_block_scale: torch.Tensor | None = None,
 ) -> torch.Tensor:
     """Paged causal attention for one layer.
 
     ``q`` is ``[num_query_tokens, num_q_heads, head_dim]``. KV cache tensors are
     flattened to ``[num_slots, num_kv_heads, head_dim]``. ``indptr`` and
-    ``indices`` describe each request's logical KV slots in order.
+    ``indices`` describe each request's logical KV slots in order. ``k_scale`` /
+    ``v_scale`` (``[num_slots, num_kv_heads]`` fp32) mark an fp8 codes cache.
     """
 
     assert q.is_cuda and k_cache.is_cuda and v_cache.is_cuda
+    assert (k_scale is None) == (v_scale is None), "K and V scales come as a pair"
     assert q.dim() == 3 and k_cache.dim() == 3 and v_cache.dim() == 3
     num_tokens, num_q_heads, head_dim = q.shape
     num_kv_heads = k_cache.shape[1]
     assert v_cache.shape[1] == num_kv_heads
-    assert k_cache.shape[-1] == head_dim and v_cache.shape[-1] == head_dim
+    nvfp4 = _validate_kv_format(q, k_cache, v_cache, k_scale, v_scale,
+                                kv_quant, k_block_scale, v_block_scale)
     assert num_q_heads % num_kv_heads == 0
     if sinks is not None:
         assert sinks.is_cuda
@@ -944,12 +1267,26 @@ def paged_attention(
 
     o = out if out is not None else torch.empty_like(q)
     sinks_arg = sinks if sinks is not None else q
+    has_kv_scale = k_scale is not None
+    k_scale_arg = k_scale if has_kv_scale else k_cache
+    v_scale_arg = v_scale if has_kv_scale else v_cache
+    k_block_arg = k_block_scale if nvfp4 else k_cache
+    v_block_arg = v_block_scale if nvfp4 else v_cache
+    if has_kv_scale:
+        assert k_scale.shape == v_scale.shape == (k_cache.shape[0], num_kv_heads), (
+            tuple(k_scale.shape),
+            tuple(k_cache.shape),
+        )
     block_d = triton.next_power_of_2(head_dim)
     grid = (num_tokens, num_q_heads)
     _paged_attention_kernel[grid](
         q,
         k_cache,
         v_cache,
+        k_scale_arg,
+        v_scale_arg,
+        k_block_arg,
+        v_block_arg,
         o,
         indptr,
         indices,
@@ -963,6 +1300,8 @@ def paged_attention(
         k_cache.stride(1),
         v_cache.stride(0),
         v_cache.stride(1),
+        k_scale_arg.stride(0),
+        v_scale_arg.stride(0),
         o.stride(0),
         o.stride(1),
         GROUP=num_q_heads // num_kv_heads,
@@ -971,6 +1310,8 @@ def paged_attention(
         BLOCK_N=block_n,
         SLIDING_WINDOW=sliding_window or 0,
         HAS_SINKS=sinks is not None,
+        HAS_KV_SCALE=has_kv_scale,
+        KV_NVFP4=nvfp4,
         num_warps=8 if head_dim >= 256 else 4,
         num_stages=2,
     )
diff --git a/python/freetoken/kernel/triton/dsv41/indexer.py b/python/freetoken/kernel/triton/dsv41/indexer.py
new file mode 100644
index 000000000..272817fec
--- /dev/null
+++ b/python/freetoken/kernel/triton/dsv41/indexer.py
@@ -0,0 +1,134 @@
+"""Paged V4.1 index scores without materializing per-head score tensors."""
+
+from __future__ import annotations
+
+import torch
+import triton
+import triton.language as tl
+
+from freetoken.kernel.triton.kv_nvfp4 import _decode_e2m1
+from .quant import _ue8m0_to_f32
+
+
+@triton.jit
+def _scores(q, weights, pool, table, ids, valid, out,
+            qs, qh, qd, ws, wh, ps, pd, ts, td, ids_s, ids_k, os, ok,
+            K: tl.constexpr, RATIO: tl.constexpr, H: tl.constexpr, D: tl.constexpr,
+            BH: tl.constexpr, BK: tl.constexpr, PACKED: tl.constexpr):
+    row = tl.program_id(0)
+    j = tl.program_id(1) * BK + tl.arange(0, BK)
+    h, d = tl.arange(0, BH), tl.arange(0, D)
+    logical = tl.load(ids + row * ids_s + j * ids_k, j < K, other=-1)
+    limit = tl.load(valid + row)
+    live = (j < K) & (logical >= 0) & (logical < limit)
+    loc = tl.load(table + row * ts + tl.maximum(logical, 0) * RATIO * td, live, other=-1)
+    live = live & (loc >= 0)
+    key_rows = pool + (tl.maximum(loc, 0) // RATIO)[:, None].to(tl.int64) * ps
+    if PACKED:
+        raw = tl.load(key_rows + d[None, :] // 2, live[:, None], other=0).to(tl.int32)
+        code = tl.where((d[None, :] & 1) == 0, raw & 15, raw >> 4)
+        scale = tl.load(key_rows + D // 2 + d[None, :] // 32, live[:, None], other=0)
+        k = (_decode_e2m1(code) * _ue8m0_to_f32(scale)).to(tl.bfloat16).to(q.dtype.element_ty)
+    else:
+        k = tl.load(key_rows + d[None, :] * pd, live[:, None], other=0)
+    query = tl.load(q + row * qs + h[:, None] * qh + d[None, :] * qd, h[:, None] < H, other=0)
+    w = tl.load(weights + row * ws + h * wh, h < H, other=0)
+    # The reference einsum and elementwise product each round to the query dtype.
+    dot = tl.dot(query, tl.trans(k)).to(query.dtype).to(tl.float32)
+    weighted = (tl.maximum(dot, 0.0) * w[:, None].to(tl.float32)).to(query.dtype).to(tl.float32)
+    score = tl.sum(weighted, 0).to(query.dtype)
+    tl.store(out + row * os + j * ok, tl.where(live, score, -float("inf")), j < K)
+
+
+def index_scores(q, weights, pool, table, ids, ratio, valid):
+    """Score logical key IDs [queries, keys] through each query's page-table row."""
+    if ids.shape[0] != q.shape[0] or table.shape[0] != q.shape[0]:
+        raise ValueError("Index query, ID, and page-table row counts differ")
+    count, heads, dim = q.shape
+    packed = pool.dtype == torch.uint8
+    if packed and (dim % 32 or pool.ndim != 2 or pool.shape[1] != dim // 2 + dim // 32
+                   or not pool.is_contiguous()):
+        raise ValueError("Packed index rows need E2M1 bytes followed by one UE8M0 scale per 32 values")
+    width = ids.shape[1]
+    out = torch.empty((count, width), device=q.device, dtype=q.dtype)
+    if not width:
+        return out
+    if not q.is_cuda:
+        for row in range(count):
+            live = (ids[row] >= 0) & (ids[row] < valid[row])
+            positions = (ids[row].clamp_min(0) * ratio).clamp_max(table.shape[1] - 1)
+            loc = table[row, positions].long()
+            live = live & (loc >= 0)
+            keys = pool[loc.clamp_min(0) // ratio]
+            if packed:
+                from .quant import unpack_fp4
+                keys = unpack_fp4(keys, block_size=32, scale_format="e8m0", dtype=torch.bfloat16)
+            dot = (q[row].float() @ keys.float().T).to(q.dtype)
+            score = (dot.relu() * weights[row, :, None]).sum(0)
+            out[row] = score.masked_fill(~live, -torch.inf)
+        return out
+    _scores[(count, triton.cdiv(width, 64))](
+        q, weights, pool, table, ids, valid, out,
+        *q.stride(), *weights.stride(), *pool.stride(), *table.stride(), *ids.stride(), *out.stride(),
+        K=width, RATIO=ratio, H=heads, D=dim, BH=max(16, triton.next_power_of_2(heads)), BK=64,
+        PACKED=packed, num_warps=4, num_stages=2,
+    )
+    return out
+
+
+def _merge_topk(values, indices, next_values, next_indices, count):
+    values = torch.cat((values, next_values), -1)
+    indices = torch.cat((indices, next_indices), -1)
+    # Equal quantized scores prefer earlier positions, independent of tile/chunk widths.
+    by_position = indices.argsort(dim=-1, stable=True)
+    by_score = values.gather(-1, by_position).argsort(dim=-1, descending=True, stable=True)
+    selected = by_position.gather(-1, by_score[..., :count])
+    return values.gather(-1, selected), indices.gather(-1, selected)
+
+
+def select_indices(q, weights, pool, table, valid, width, ratio, topk,
+                   *, candidates=None, candidate_topk=0, block_size=8,
+                   query_tile=32, key_tile=4096):
+    """Exact tiled top-k, optionally publishing or consuming candidate block IDs.
+
+    Working score memory is bounded by query_tile * key_tile. The source retains
+    candidate IDs instead of a query-by-context boolean mask.
+    """
+    if key_tile % block_size:
+        raise ValueError("Index key tile must contain whole candidate blocks")
+    result, block_result = [], []
+    final_k = min(topk, width)
+    final_c = min(candidate_topk, triton.cdiv(width, block_size))
+    for first in range(0, q.shape[0], query_tile):
+        last = min(first + query_tile, q.shape[0])
+        query, gate, mapping, limits = q[first:last], weights[first:last], table[first:last], valid[first:last]
+        empty = q.new_empty((last - first, 0))
+        best, best_ids = empty, torch.empty_like(empty, dtype=torch.int64)
+        cbest, cbest_ids = empty, torch.empty_like(empty, dtype=torch.int64)
+        search_width = width if candidates is None else candidates.shape[-1] * block_size
+        for offset in range(0, search_width, key_tile):
+            end = min(offset + key_tile, search_width)
+            if candidates is None:
+                ids = torch.arange(offset, end, device=q.device).expand(last - first, -1)
+            else:
+                blocks = candidates[first:last, offset // block_size:triton.cdiv(end, block_size)]
+                ids = blocks[:, :, None] * block_size + torch.arange(block_size, device=q.device)
+                ids = torch.where(blocks[:, :, None] >= 0, ids, -1).flatten(1)[:, :end-offset]
+            scores = index_scores(query, gate, pool, mapping, ids, ratio, limits)
+            best, best_ids = _merge_topk(best, best_ids, scores, ids, final_k)
+            if candidate_topk:
+                padded = torch.nn.functional.pad(scores, (0, -scores.shape[-1] % block_size), value=-torch.inf)
+                block_scores = padded.unflatten(-1, (-1, block_size)).amax(-1)
+                block_ids = torch.arange(offset // block_size, triton.cdiv(end, block_size), device=q.device)
+                block_ids = block_ids.expand(last - first, -1)
+                newest = (limits - 1) // block_size
+                block_scores = block_scores.masked_fill((block_ids == newest[:, None]) & (limits[:, None] > 0), torch.inf)
+                cbest, cbest_ids = _merge_topk(cbest, cbest_ids, block_scores, block_ids, final_c)
+        ids = torch.where(torch.isfinite(best), best_ids, width)
+        ids = ids.sort(-1).values
+        result.append(torch.where(ids < width, ids, -1))
+        if candidate_topk:
+            block_result.append(torch.where(cbest > -torch.inf, cbest_ids, -1))
+    selected = torch.cat(result, 0) if result else torch.empty((0, final_k), device=q.device, dtype=torch.int64)
+    blocks = torch.cat(block_result, 0) if block_result else None
+    return selected, blocks
diff --git a/python/freetoken/kernel/triton/dsv41/quant.py b/python/freetoken/kernel/triton/dsv41/quant.py
new file mode 100644
index 000000000..ab5bba50c
--- /dev/null
+++ b/python/freetoken/kernel/triton/dsv41/quant.py
@@ -0,0 +1,345 @@
+"""DeepSeek V4.1 block-32 FP8 linears and the two CSA2 FP4 scale formats."""
+
+from __future__ import annotations
+
+import torch
+import triton
+import triton.language as tl
+
+from freetoken.kernel.triton.dsv4.fp8_linear import (
+    _log2_ceil,
+    _round_fp4,
+    act_quant_fp8,
+    act_quant_fp8_roundtrip,
+)
+from freetoken.kernel.triton.e4m3_compat import (
+    e4m3_f32_to_u8,
+    e4m3_kernel_view,
+    e4m3_native_cx,
+    e4m3_u8_to_f32,
+    round_e4m3,
+)
+from freetoken.kernel.triton.kv_nvfp4 import _decode_e2m1, _encode_e2m1
+
+
+def _check_block(x: torch.Tensor, block_size: int) -> None:
+    if block_size not in (16, 32, 64, 128) or x.ndim < 1 or x.shape[-1] % block_size:
+        raise ValueError(f"shape {tuple(x.shape)} is incompatible with block size {block_size}")
+
+
+def fp8_roundtrip(x: torch.Tensor, block_size: int = 32) -> torch.Tensor:
+    _check_block(x, block_size)
+    if x.numel() == 0:
+        return x.clone()
+    if x.is_cuda:
+        return act_quant_fp8_roundtrip(x, block_size)
+    groups = x.float().reshape(*x.shape[:-1], -1, block_size)
+    scale = torch.exp2(torch.ceil(torch.log2(groups.abs().amax(-1, keepdim=True).clamp_min(1e-4) / 448)))
+    values = (groups / scale).clamp(-448, 448).to(torch.float8_e4m3fn).float()
+    return (values * scale).reshape_as(x).to(x.dtype)
+
+
+@triton.jit
+def _fp4_roundtrip_kernel(X, Y, M: tl.constexpr, K: tl.constexpr,
+                          GROUP: tl.constexpr, E4M3: tl.constexpr, ROWS: tl.constexpr):
+    rows = tl.program_id(0) * ROWS + tl.arange(0, ROWS)
+    cols = tl.program_id(1) * GROUP + tl.arange(0, GROUP)
+    offsets = rows[:, None] * K + cols[None, :]
+    x = tl.load(X + offsets, rows[:, None] < M, other=0).to(tl.float32)
+    amax = tl.max(tl.abs(x), 1)
+    if E4M3:
+        scale = round_e4m3(tl.clamp(amax / 6.0, 2.0 ** -9, 448.0))
+    else:
+        exponent = _log2_ceil(tl.maximum(amax, 6.0 * (2.0 ** -126)) * (1.0 / 6.0))
+        scale = tl.exp2(exponent.to(tl.float32))
+    value = _round_fp4(tl.clamp(tl.div_rn(x, scale[:, None]), -6, 6)) * scale[:, None]
+    tl.store(Y + offsets, value, rows[:, None] < M)
+
+
+def fp4_roundtrip(x: torch.Tensor, block_size: int = 16, scale_format: str = "e4m3") -> torch.Tensor:
+    _check_block(x, block_size)
+    if scale_format not in ("e4m3", "e8m0"):
+        raise ValueError(f"unsupported FP4 scale format: {scale_format}")
+    if x.numel() == 0:
+        return x.clone()
+    if x.is_cuda:
+        source = x.contiguous()
+        out = torch.empty_like(source)
+        k = x.shape[-1]
+        m = x.numel() // k
+        _fp4_roundtrip_kernel[(triton.cdiv(m, 32), k // block_size)](
+            source, out, m, k, block_size, scale_format == "e4m3", 32,
+            num_warps=4, enable_fp_fusion=False,
+        )
+        return out
+    groups = x.float().reshape(*x.shape[:-1], -1, block_size)
+    amax = groups.abs().amax(-1, keepdim=True)
+    if scale_format == "e4m3":
+        scale = (amax / 6).clamp(2.0**-9, 448).to(torch.float8_e4m3fn).float()
+    else:
+        scale = torch.exp2(torch.ceil(torch.log2(amax.clamp_min(6 * 2.0**-126) / 6)))
+    values = (groups / scale).clamp(-6, 6)
+    magnitudes = torch.tensor([0, .5, 1, 1.5, 2, 3, 4, 6], device=x.device)
+    distance = (values.abs().unsqueeze(-1) - magnitudes).abs()
+    # Ties go to an even FP4 code, rather than always to the smaller magnitude.
+    even = torch.tensor([0, 2, 4, 6, 1, 3, 5, 7], device=x.device)
+    codes = even[distance[..., even].argmin(-1)]
+    result = magnitudes[codes] * values.sign() * scale
+    return result.reshape_as(x).to(x.dtype)
+
+
+@triton.jit
+def _e4m3_to_f32(code):
+    value = e4m3_u8_to_f32(code)
+    return tl.where((code.to(tl.int32) & 127) == 127, float("nan"), value)
+
+
+@triton.jit
+def _ue8m0_to_f32(code):
+    bits = code.to(tl.uint32) << 23
+    bits = tl.where(code == 0, 0x00400000, bits)
+    bits = tl.where(code == 255, 0x7FC00000, bits).to(tl.uint32)
+    return bits.to(tl.float32, bitcast=True)
+
+
+@triton.jit
+def _pack_kernel(X, Y, M: tl.constexpr, K: tl.constexpr, WIDTH: tl.constexpr,
+                 GROUP: tl.constexpr, FP4: tl.constexpr, E4M3: tl.constexpr,
+                 ROWS: tl.constexpr):
+    rows = tl.program_id(0) * ROWS + tl.arange(0, ROWS)
+    group = tl.program_id(1)
+    cols = group * GROUP + tl.arange(0, GROUP)
+    x = tl.load(X + rows[:, None] * K + cols[None, :], rows[:, None] < M, other=0).to(tl.float32)
+    amax = tl.max(tl.abs(x), 1)
+    if FP4:
+        if E4M3:
+            scale = round_e4m3(tl.clamp(amax / 6.0, 2.0 ** -9, 448.0))
+            scale_code = e4m3_f32_to_u8(scale)
+        else:
+            exponent = _log2_ceil(tl.maximum(amax, 6.0 * (2.0 ** -126)) * (1.0 / 6.0))
+            scale = tl.exp2(exponent.to(tl.float32))
+            scale_code = (exponent + 127).to(tl.uint8)
+        normalized = tl.clamp(tl.div_rn(x, scale[:, None]), -6, 6)
+        codes = _encode_e2m1(normalized).reshape(ROWS, GROUP // 2, 2)
+        low, high = tl.split(codes)
+        packed = low | (high << 4)
+        out_cols = group * (GROUP // 2) + tl.arange(0, GROUP // 2)
+        tl.store(Y + rows[:, None] * WIDTH + out_cols[None, :], packed, rows[:, None] < M)
+        scale_offset = K // 2
+    else:
+        exponent = _log2_ceil(tl.maximum(amax, 1e-4) * (1.0 / 448.0))
+        scale = tl.exp2(exponent.to(tl.float32))
+        normalized = tl.clamp(x / scale[:, None], -448.0, 448.0)
+        if e4m3_native_cx():
+            codes = normalized.to(tl.float8e4nv).to(tl.uint8, bitcast=True)
+        else:
+            codes = e4m3_f32_to_u8(round_e4m3(normalized))
+        tl.store(Y + rows[:, None] * WIDTH + cols[None, :], codes, rows[:, None] < M)
+        scale_code = (exponent + 127).to(tl.uint8)
+        scale_offset = K
+    tl.store(Y + rows * WIDTH + scale_offset + group, scale_code, rows < M)
+
+
+def _pack(x, block_size, *, fp4, scale_format):
+    _check_block(x, block_size)
+    if x.dtype not in (torch.bfloat16, torch.float16, torch.float32):
+        raise ValueError("packed V4.1 KV input must be BF16, FP16, or FP32")
+    if scale_format not in ("e4m3", "e8m0"):
+        raise ValueError(f"unsupported FP4 scale format: {scale_format}")
+    k = x.shape[-1]
+    code_width = k // 2 if fp4 else k
+    out = torch.empty((*x.shape[:-1], code_width + k // block_size), dtype=torch.uint8, device=x.device)
+    if not x.numel():
+        return out
+    if x.is_cuda:
+        source = x.contiguous()
+        _pack_kernel[(triton.cdiv(x.numel() // k, 32), k // block_size)](
+            source, out, x.numel() // k, k, out.shape[-1], block_size,
+            fp4, scale_format == "e4m3", 32, num_warps=4, enable_fp_fusion=False,
+        )
+        return out
+    groups = x.float().reshape(*x.shape[:-1], -1, block_size)
+    amax = groups.abs().amax(-1, keepdim=True)
+    if fp4 and scale_format == "e4m3":
+        scales = (amax / 6).clamp(2.0**-9, 448).to(torch.float8_e4m3fn)
+        scale = scales.float()
+        scale_codes = scales.view(torch.uint8)
+    else:
+        maximum, floor = (6, 6 * 2.0**-126) if fp4 else (448, 1e-4)
+        exponent = torch.ceil(torch.log2(amax.clamp_min(floor) / maximum))
+        scale = torch.exp2(exponent)
+        scale_codes = (exponent + 127).to(torch.uint8)
+    if fp4:
+        values = (groups / scale).clamp(-6, 6)
+        grid = torch.tensor([0, .5, 1, 1.5, 2, 3, 4, 6], dtype=torch.float32, device=x.device)
+        even = torch.tensor([0, 2, 4, 6, 1, 3, 5, 7], device=x.device)
+        distance = (values.abs().unsqueeze(-1) - grid).abs()
+        codes = even[distance[..., even].argmin(-1)].to(torch.uint8) | ((values < 0).to(torch.uint8) << 3)
+        pairs = codes.reshape(*x.shape[:-1], k // 2, 2)
+        out[..., :code_width] = pairs[..., 0] | (pairs[..., 1] << 4)
+    else:
+        codes = (groups / scale).clamp(-448, 448).to(torch.float8_e4m3fn).view(torch.uint8)
+        out[..., :code_width] = codes.reshape_as(x)
+    out[..., code_width:] = scale_codes.squeeze(-1)
+    return out
+
+
+def pack_fp8(x: torch.Tensor, block_size: int = 32) -> torch.Tensor:
+    """Keep E4M3 codes followed by one UE8M0 byte per block in each packed row."""
+    return _pack(x, block_size, fp4=False, scale_format="e8m0")
+
+
+def pack_fp4(x: torch.Tensor, block_size: int = 16, scale_format: str = "e4m3") -> torch.Tensor:
+    """Keep low-first E2M1 nibbles followed by the original block-scale bytes."""
+    return _pack(x, block_size, fp4=True, scale_format=scale_format)
+
+
+@triton.jit
+def _unpack_kernel(X, Y, M: tl.constexpr, K: tl.constexpr, WIDTH: tl.constexpr,
+                   GROUP: tl.constexpr, FP4: tl.constexpr, E4M3: tl.constexpr,
+                   ROWS: tl.constexpr):
+    rows = tl.program_id(0) * ROWS + tl.arange(0, ROWS)
+    group = tl.program_id(1)
+    cols = group * GROUP + tl.arange(0, GROUP)
+    if FP4:
+        packed = tl.load(X + rows[:, None] * WIDTH + cols[None, :] // 2,
+                         rows[:, None] < M, other=0).to(tl.int32)
+        value = _decode_e2m1(tl.where((cols[None, :] & 1) == 0, packed & 15, packed >> 4))
+        scale_offset = K // 2
+    else:
+        code = tl.load(X + rows[:, None] * WIDTH + cols[None, :], rows[:, None] < M, other=0)
+        value = _e4m3_to_f32(code)
+        scale_offset = K
+    scale_code = tl.load(X + rows * WIDTH + scale_offset + group, rows < M, other=0)
+    if E4M3:
+        scale = _e4m3_to_f32(scale_code)
+    else:
+        scale = _ue8m0_to_f32(scale_code)
+    tl.store(Y + rows[:, None] * K + cols[None, :], value * scale[:, None], rows[:, None] < M)
+
+
+def _unpack(packed, block_size, dtype, *, fp4, scale_format):
+    if packed.ndim < 1 or packed.dtype != torch.uint8 or block_size not in (16, 32, 64, 128):
+        raise ValueError("packed V4.1 KV requires uint8 rows and a supported block size")
+    unit = (block_size // 2 if fp4 else block_size) + 1
+    if packed.shape[-1] % unit or scale_format not in ("e4m3", "e8m0"):
+        raise ValueError("packed V4.1 KV row width or scale format is invalid")
+    if dtype not in (torch.bfloat16, torch.float16, torch.float32):
+        raise ValueError("unpacked V4.1 KV dtype must be BF16, FP16, or FP32")
+    k = packed.shape[-1] // unit * block_size
+    out = torch.empty((*packed.shape[:-1], k), device=packed.device, dtype=dtype)
+    if not out.numel():
+        return out
+    if packed.is_cuda:
+        source = packed.contiguous()
+        _unpack_kernel[(triton.cdiv(out.numel() // k, 32), k // block_size)](
+            source, out, out.numel() // k, k, packed.shape[-1], block_size,
+            fp4, scale_format == "e4m3", 32, num_warps=4, enable_fp_fusion=False,
+        )
+        return out
+    code_width = k // 2 if fp4 else k
+    raw = packed[..., :code_width].contiguous()
+    if fp4:
+        codes = torch.stack((raw & 15, raw >> 4), -1).flatten(-2).long()
+        grid = torch.tensor([0., .5, 1., 1.5, 2., 3., 4., 6., -0., -.5, -1., -1.5, -2., -3., -4., -6.],
+                            dtype=torch.float32, device=packed.device)
+        values = grid[codes]
+    else:
+        values = raw.view(torch.float8_e4m3fn).float()
+    scale_dtype = torch.float8_e4m3fn if scale_format == "e4m3" else torch.float8_e8m0fnu
+    scale = packed[..., code_width:].contiguous().view(scale_dtype).float()
+    out.copy_((values.reshape(*packed.shape[:-1], -1, block_size) * scale[..., None]).flatten(-2))
+    return out
+
+
+def unpack_fp8(packed: torch.Tensor, block_size: int = 32, dtype=torch.bfloat16) -> torch.Tensor:
+    return _unpack(packed, block_size, dtype, fp4=False, scale_format="e8m0")
+
+
+def unpack_fp4(packed: torch.Tensor, block_size: int = 16, scale_format: str = "e4m3",
+               dtype=torch.bfloat16) -> torch.Tensor:
+    return _unpack(packed, block_size, dtype, fp4=True, scale_format=scale_format)
+
+
+@triton.jit
+def _gemm(A, W, SA, SW, Y, M: tl.constexpr, N: tl.constexpr, K: tl.constexpr,
+          GROUP: tl.constexpr, BM: tl.constexpr):
+    rows = tl.program_id(0) * BM + tl.arange(0, BM)
+    cols = tl.program_id(1) * GROUP + tl.arange(0, GROUP)
+    ki = tl.arange(0, GROUP)
+    acc = tl.zeros((BM, GROUP), tl.float32)
+    for kb in range(K // GROUP):
+        a = tl.load(A + rows[:, None] * K + kb * GROUP + ki[None, :], rows[:, None] < M, other=0.0)
+        w = tl.load(W + cols[:, None] * K + kb * GROUP + ki[None, :], cols[:, None] < N, other=0.0)
+        if e4m3_native_cx():
+            product = tl.dot(a, tl.trans(w), out_dtype=tl.float32)
+        else:
+            product = tl.dot(a, tl.trans(e4m3_u8_to_f32(w).to(tl.bfloat16)), out_dtype=tl.float32)
+        sa = tl.load(SA + rows * (K // GROUP) + kb, rows < M, other=127)
+        sw = tl.load(SW + tl.program_id(1) * (K // GROUP) + kb)
+        scale = tl.exp2(sa.to(tl.float32) - 127) * tl.exp2(sw.to(tl.float32) - 127)
+        acc += product * scale[:, None]
+    tl.store(Y + rows[:, None] * N + cols[None, :], acc, (rows[:, None] < M) & (cols[None, :] < N))
+
+
+@triton.jit
+def _gemv(A, W, SA, SW, PART, N: tl.constexpr, K: tl.constexpr,
+          GROUP: tl.constexpr, BN: tl.constexpr, SPLITS: tl.constexpr):
+    rows = tl.program_id(0) * BN + tl.arange(0, BN)
+    ki = tl.arange(0, GROUP)
+    split = tl.program_id(1)
+    acc = tl.zeros((BN,), tl.float32)
+    for kb in range(split, K // GROUP, SPLITS):
+        a = tl.load(A + kb * GROUP + ki).to(tl.float32)
+        raw_w = tl.load(W + rows[:, None] * K + kb * GROUP + ki[None, :], rows[:, None] < N, other=0.0)
+        if e4m3_native_cx():
+            w = raw_w.to(tl.float32)
+        else:
+            w = e4m3_u8_to_f32(raw_w)
+        sw = tl.load(SW + (rows // GROUP) * (K // GROUP) + kb, rows < N, other=127)
+        sa = tl.load(SA + kb)
+        acc += tl.sum(w * a[None, :], 1) * tl.exp2(sw.to(tl.float32) - 127) * tl.exp2(sa.to(tl.float32) - 127)
+    tl.store(PART + split * N + rows, acc, rows < N)
+
+
+def block_fp8_linear(x: torch.Tensor, weight: torch.Tensor, scale: torch.Tensor,
+                     bias: torch.Tensor | None = None, block_size: int = 32) -> torch.Tensor:
+    _check_block(x, block_size)
+    if weight.ndim != 2 or weight.shape[1] != x.shape[-1] or weight.shape[0] % block_size:
+        raise ValueError("FP8 linear weight dimensions do not match the input/block size")
+    n, k = weight.shape
+    if weight.dtype != torch.float8_e4m3fn or tuple(scale.shape) != (n // block_size, k // block_size):
+        raise ValueError("expected E4M3 weights and a per-block E8M0 scale matrix")
+    if scale.dtype not in (torch.float8_e8m0fnu, torch.uint8):
+        raise ValueError("FP8 linear scales must be E8M0 values or unsigned E8M0 codes")
+    if x.device != weight.device or x.device != scale.device:
+        raise ValueError("FP8 linear operands must be on the same device")
+    shape = (*x.shape[:-1], n)
+    if x.numel() == 0:
+        return x.new_empty(shape)
+    if x.is_cuda:
+        activation, act_scale = act_quant_fp8(x, block_size)
+        w = e4m3_kernel_view(weight.contiguous())
+        sw = scale.view(torch.uint8).contiguous()
+        m = activation.shape[0]
+        if m == 1:
+            splits = min(16, k // block_size)
+            partial = torch.empty((splits, n), device=x.device, dtype=torch.float32)
+            _gemv[(triton.cdiv(n, 16), splits)](
+                activation, w, act_scale, sw, partial, n, k, block_size, 16, splits,
+                num_warps=4, enable_fp_fusion=False,
+            )
+            out = partial.sum(0).to(x.dtype).reshape(shape)
+        else:
+            out = x.new_empty((m, n))
+            _gemm[(triton.cdiv(m, 32), n // block_size)](
+                activation, w, act_scale, sw, out, m, n, k, block_size, 32,
+                num_warps=4, enable_fp_fusion=False,
+            )
+            out = out.reshape(shape)
+    else:
+        activation = fp8_roundtrip(x, block_size).float().reshape(-1, k)
+        codes = scale.view(torch.uint8).float()
+        factors = torch.exp2(codes - 127).repeat_interleave(block_size, 0).repeat_interleave(block_size, 1)
+        out = (activation @ (weight.float() * factors).T).to(x.dtype).reshape(shape)
+    return out if bias is None else out + bias.to(out.dtype)
diff --git a/python/freetoken/kernel/triton/dsv41/sparse_attn.py b/python/freetoken/kernel/triton/dsv41/sparse_attn.py
new file mode 100644
index 000000000..293aad709
--- /dev/null
+++ b/python/freetoken/kernel/triton/dsv41/sparse_attn.py
@@ -0,0 +1,156 @@
+"""Gather native V4.1 FP8-window/FP4-compressed KV directly into attention tiles."""
+
+from __future__ import annotations
+
+import torch
+import triton
+import triton.language as tl
+
+from freetoken.kernel.triton.dsv4.sparse_attn import (
+    BLOCK_H, BLOCK_T, _sparse_attn_splitk_merge_kernel, split_count,
+)
+from freetoken.kernel.triton.kv_nvfp4 import _decode_e2m1
+from .quant import _e4m3_to_f32, _ue8m0_to_f32
+
+
+@triton.jit
+def _load_mixed(win, cmp, slots, is_window, live, dim,
+                win_stride, cmp_stride, D: tl.constexpr):
+    base = tl.where(is_window, win, cmp)
+    stride = tl.where(is_window, win_stride, cmp_stride)
+    row = base[:, None] + slots[:, None].to(tl.int64) * stride[:, None]
+    # One selected data load avoids staging two full KV tiles on consumer GPUs.
+    data = tl.load(row + dim[None, :], live[:, None] & (is_window[:, None] | (dim[None, :] < D // 2)),
+                   other=0).to(tl.int32)
+    packed = tl.gather(data, tl.broadcast_to(dim[None, :] // 2, (slots.shape[0], D)), 1)
+    raw = tl.where(is_window[:, None], data, packed)
+    group = tl.arange(0, D // 16)
+    scale_offset = tl.where(is_window[:, None], D + group[None, :] // 2, D // 2 + group[None, :])
+    scale_code = tl.load(row + scale_offset, live[:, None], other=0)
+    nibble = tl.where((dim[None, :] & 1) == 0, raw & 15, raw >> 4)
+    fp8_value = _e4m3_to_f32(raw)
+    value = tl.where(is_window[:, None], fp8_value, _decode_e2m1(nibble))
+    fp4_scale = _e4m3_to_f32(scale_code)
+    scale = tl.where(is_window[:, None], _ue8m0_to_f32(scale_code),
+                     fp4_scale)
+    # The previous pool held the roundtrip's BF16 result, not its FP32 product.
+    restored = value.reshape(slots.shape[0], D // 16, 16) * scale[:, :, None]
+    return restored.reshape(slots.shape[0], D).to(tl.bfloat16)
+
+
+@triton.jit
+def _attend(q, win, cmp, out, lse, sink, indices, counts,
+            scale, H, TOPK, N_WINDOW,
+            qb, qm, qh, qd, win_stride, cmp_stride,
+            ob, om, oh, os, od, lb, lm, lh, ls, ib, im, it, cb, cm,
+            D: tl.constexpr, BH: tl.constexpr, BT: tl.constexpr,
+            HAS_COUNTS: tl.constexpr, SPLITS: tl.constexpr):
+    query_split, batch, head_block = tl.program_id(0), tl.program_id(1), tl.program_id(2)
+    query, split = query_split // SPLITS, query_split % SPLITS
+    h = head_block * BH + tl.arange(0, BH)
+    d = tl.arange(0, D)
+    hm = h < H
+    active = TOPK
+    if HAS_COUNTS:
+        active = N_WINDOW + tl.load(counts + batch * cb + query * cm)
+    first = 0
+    end = active
+    if SPLITS > 1:
+        width = tl.cdiv(tl.cdiv(active, SPLITS), BT) * BT
+        first = split * width
+        end = tl.minimum(first + width, active)
+    maximum = tl.full((BH,), -float("inf"), tl.float32)
+    denominator = tl.zeros((BH,), tl.float32)
+    acc = tl.zeros((BH, D), tl.float32)
+    if end > first:
+        query_values = tl.load(q + batch * qb + query * qm + h[:, None] * qh + d[None, :] * qd,
+                               hm[:, None], other=0).to(tl.float32)
+        for start in range(first, end, BT):
+            columns = start + tl.arange(0, BT)
+            slots = tl.load(indices + batch * ib + query * im + columns * it,
+                            columns < end, other=-1)
+            live = slots >= 0
+            kv = _load_mixed(win, cmp, slots, columns < N_WINDOW, live, d,
+                             win_stride, cmp_stride, D)
+            # Promote after choosing each dot's layout so shared KV tiles stay BF16.
+            # Reusing one FP32 tile here exceeds the RTX 5090 shared-memory limit.
+            scores = tl.dot(query_values, tl.trans(kv).to(tl.float32)) * scale
+            scores = tl.where(live[None, :], scores, -float("inf"))
+            next_max = tl.maximum(maximum, tl.max(scores, 1))
+            alpha = tl.where(next_max == -float("inf"), 1.0, tl.exp(maximum - next_max))
+            p = tl.where(live[None, :], tl.exp(scores - next_max[:, None]), 0.0)
+            denominator = denominator * alpha + tl.sum(p, 1)
+            acc = acc * alpha[:, None] + tl.dot(p, kv.to(tl.float32))
+            maximum = next_max
+    if SPLITS == 1:
+        sink_value = tl.load(sink + h, hm, other=0).to(tl.float32)
+        denominator += tl.exp(sink_value - maximum)
+        result = acc / denominator[:, None]
+    else:
+        result = tl.where(denominator[:, None] == 0, 0.0, acc / denominator[:, None])
+        logsum = tl.where(denominator == 0, -float("inf"), maximum + tl.log(denominator))
+        tl.store(lse + batch * lb + query * lm + h * lh + split * ls, logsum, hm)
+    tl.store(out + batch * ob + query * om + h[:, None] * oh + split * os + d[None, :] * od,
+             result.to(out.dtype.element_ty), hm[:, None])
+
+
+def sparse_attn_paged(q, window_pool, cmp_pool, attn_sink, topk_idxs, n_window,
+                      softmax_scale, cmp_counts=None, *, force_splits=None):
+    """Attend to packed selected rows; the full cache never expands to BF16."""
+    if q.ndim != 4 or not q.is_cuda:
+        raise ValueError("Packed sparse attention needs CUDA queries shaped [batch, queries, heads, dim]")
+    b, m, h, d = q.shape
+    topk = topk_idxs.shape[-1]
+    if d % 32 or d & (d - 1) or not 0 <= n_window <= topk:
+        raise ValueError("Packed V4.1 attention needs a power-of-two dimension divisible by 32 and a valid window width")
+    for tier, (pool, width) in enumerate(((window_pool, d + d // 32), (cmp_pool, d // 2 + d // 16))):
+        if pool.dtype != torch.uint8 or pool.ndim != 2 or not pool.is_contiguous():
+            raise ValueError("Packed KV rows must be contiguous uint8 tensors")
+        if pool.device != q.device:
+            raise ValueError("Packed KV rows and queries must share a device")
+        # Window-only layers alias the window pool; no compressed column is read.
+        if pool.shape[1] != width and not (tier == 1 and n_window == topk):
+            raise ValueError("Packed KV row width does not match the query dimension")
+    if tuple(topk_idxs.shape) != (b, m, topk) or attn_sink.numel() != h:
+        raise ValueError("Sparse indices or attention sink do not match the queries")
+    if any(t.device != q.device for t in (topk_idxs, attn_sink)):
+        raise ValueError("Sparse indices, attention sink, and queries must share a device")
+    if cmp_counts is not None and cmp_counts.device != q.device:
+        raise ValueError("Compressed counts and queries must share a device")
+    q = q.contiguous()
+    idx = topk_idxs.contiguous().to(torch.int32)
+    sink = attn_sink.contiguous().to(torch.float32)
+    out = torch.empty_like(q)
+    if not b or not m:
+        return out
+    has_counts = cmp_counts is not None
+    if has_counts:
+        counts = cmp_counts.contiguous().to(torch.int32).view(b, m)
+        cb, cm = counts.stride()
+    else:
+        counts, cb, cm = idx, 0, 0
+    splits = split_count(b, m, h, topk, q.device) if force_splits is None else force_splits
+    if not isinstance(splits, int) or splits < 0:
+        raise ValueError("force_splits must be a nonnegative integer")
+    splits = max(1, splits)
+    if splits > 1:
+        partial = torch.empty((b, m, h, splits, d), device=q.device, dtype=torch.float32)
+        logsum = torch.empty((b, m, h, splits), device=q.device, dtype=torch.float32)
+        lstrides = logsum.stride()
+    else:
+        partial, logsum = out.unsqueeze(3), out
+        lstrides = (0, 0, 0, 0)
+    _attend[(m * splits, b, triton.cdiv(h, BLOCK_H))](
+        q, window_pool, cmp_pool, partial, logsum, sink, idx, counts,
+        float(softmax_scale), h, topk, n_window,
+        *q.stride(), window_pool.stride(0), cmp_pool.stride(0),
+        *partial.stride(), *lstrides, *idx.stride(), cb, cm,
+        D=d, BH=BLOCK_H, BT=BLOCK_T, HAS_COUNTS=has_counts, SPLITS=splits,
+        num_warps=8, num_stages=1,
+    )
+    if splits > 1:
+        _sparse_attn_splitk_merge_kernel[(m, b, h)](
+            partial, logsum, out, sink, *partial.stride(), *logsum.stride(), *out.stride(),
+            D=d, NUM_SPLITS=splits, num_warps=4,
+        )
+    return out
diff --git a/python/freetoken/kernel/triton/e4m3_compat.py b/python/freetoken/kernel/triton/e4m3_compat.py
index 1d9f744ce..a23b7cde1 100644
--- a/python/freetoken/kernel/triton/e4m3_compat.py
+++ b/python/freetoken/kernel/triton/e4m3_compat.py
@@ -5,8 +5,18 @@
 Affected kernels branch on :func:`e4m3_native_cx` (a compile-time constexpr): the
 native branch stays byte-identical on sm_89+, the emulated branch is dead-code
 eliminated there. When the emulated branch is active, wrappers must pass e4m3
-tensors as ``.view(torch.uint8)`` and allocate act-quant outputs as bf16 -- use
-the host-side twin :func:`e4m3_native` for those decisions.
+tensors as ``.view(torch.uint8)`` and allocate act-quant outputs as bf16.
+
+TWO independent probes answer "is fp8e4nv native here": :func:`e4m3_native` (host,
+torch's device capability -- decides what a buffer is ALLOCATED as) and
+:func:`e4m3_native_cx` (compile-time, triton's target -- decides which arm a kernel
+compiles). They cannot be merged, because a constexpr function that referenced the
+host one would not survive triton's cache-key AST walk, so on a box where the probes
+disagree the host holds real fp8 tensors while kernels take the emulated arm.
+:func:`warn_if_probes_disagree` says so once at startup. Code handed a buffer whose
+dtype the host already settled -- the KV pool -- must not re-ask at all and picks its
+decode from the pointer: :func:`kv_load_e4m3_tile_f32`. Trusting the probe over the
+tensor is what produced "cannot cast int32 to fp8e4nv" at CUDA graph capture.
 
 ``FREETOKEN_FORCE_E4M3_EMU=1`` (or true/yes/on) forces the emulated path on any
 GPU (for A/B validation against the native fp8 unit). The flag is read ONCE at
@@ -47,8 +57,9 @@ def _env_force() -> bool:
 
 
 def e4m3_native() -> bool:
-    """Host-side twin of :func:`e4m3_native_cx`: True when kernels take fp8e4nv
-    tensors directly. False: pass ``.view(torch.uint8)`` and bf16 act buffers."""
+    """Host-side probe: does THIS device take fp8e4nv tensors directly? True: kernels
+    get fp8 tensors, False: pass ``.view(torch.uint8)`` and bf16 act buffers. NOT
+    necessarily the answer :func:`e4m3_native_cx` gives -- see this module's header."""
     global _native
     if _env_force() != FORCE_EMU:
         raise RuntimeError(
@@ -64,9 +75,47 @@ def e4m3_native() -> bool:
 
             # one process runs on one GPU, so its convention is that GPU's; None (-> the current device) only before the process binds
             _native = torch.cuda.get_device_capability(assigned_visible_gpu()) >= (8, 9)
+        warn_if_probes_disagree()
     return _native
 
 
+_warned_disagree = False
+
+
+def warn_if_probes_disagree() -> None:
+    """Log once when the two native-fp8e4nv probes answer differently.
+
+    :func:`e4m3_native` decides what buffers the host ALLOCATES while
+    :func:`e4m3_native_cx` decides which arm a kernel compiles, and the two cannot be
+    unified (that function's docstring explains triton's cache-key walk). A box where
+    triton's probe under-reports therefore runs every e4m3 kernel through the software
+    decode -- bit-exact per this module's header, but slower -- and any kernel that
+    trusts the probe over the tensor it was handed stops compiling outright. Reads the
+    latched ``_native`` rather than calling back into :func:`e4m3_native`."""
+    global _warned_disagree
+    if _warned_disagree:
+        return
+    _warned_disagree = True
+    if FORCE_EMU:  # emulating by request is not a disagreement
+        return
+    try:
+        triton_native = target_info.cuda_capability_geq(8, 9)
+        if triton_native == bool(_native):
+            return
+        major, minor = torch.cuda.get_device_capability()
+    except Exception:  # noqa: BLE001 -- no driver/no target yet: nothing to compare
+        return
+    from freetoken.utils import init_logger
+
+    init_logger(__name__).warning(
+        "native fp8e4nv disagreement: torch reports sm_%d%d for this device but "
+        "triton's target probe says %s, so e4m3 kernels compile the software-decode "
+        "branch (bit-exact, slower). The fp8 KV cache is unaffected -- it follows the "
+        "buffer it was given.",
+        major, minor, "supported" if triton_native else "unsupported",
+    )
+
+
 def e4m3_kernel_view(t: torch.Tensor) -> torch.Tensor:
     """An e4m3 tensor as the branched kernels expect it: unchanged when native,
     the uint8 view otherwise (the fp8 pointer type is illegal pre-sm_89)."""
@@ -84,7 +133,17 @@ def e4m3_native_cx():
     """Compile-time: does the compilation target have native fp8e4nv (sm_89+)?
     Delegates to ``target_info`` (reads the active driver's target, so
     cross-compilation tests that patch ``driver.active.get_current_target``
-    resolve consistently)."""
+    resolve consistently).
+
+    It CANNOT defer to :func:`e4m3_native`, however much one verdict per process is
+    what we want: triton hashes a constexpr function by walking its AST
+    (runtime/jit.py: cache_key -> record_reference), and a bare reference to a plain
+    python function raises "Unsupported function referenced: <function e4m3_native>"
+    -- trying that once disabled every e4m3 kernel at once, PLE gather included.
+    Module attributes (``target_info.whatever``) survive the walk, plain functions do
+    not. The probes therefore stay separate, :func:`warn_if_probes_disagree` reports
+    when they disagree, and code handed a buffer the host already typed -- the KV
+    pool -- ignores this function and follows the pointer: kv_load_e4m3_tile_f32."""
     return not FORCE_EMU and target_info.cuda_capability_geq(8, 9)
 
 
@@ -122,3 +181,84 @@ def round_e4m3(x):
     y_norm = ((b + 524287 + lsb) & 0xFFF00000).to(tl.float32, bitcast=True)
     y_sub = (x + 24576.0) - 24576.0
     return tl.where(tl.abs(x) >= 0.015625, y_norm, y_sub)
+
+
+@jit
+def e4m3_f32_to_u8(x):
+    """Encode an fp32 value that ALREADY lies on the e4m3 grid -- the output of
+    :func:`round_e4m3`, clamped to +-448 -- into its e4m3 byte code. This is the
+    encoder the pre-sm_89 emulated path needs to STORE fp8-sized data (the fp8
+    type itself is unavailable there, so the bytes live in a uint8 buffer that
+    :func:`e4m3_u8_to_f32` decodes back).
+
+    Normal range: read the (unbiased) exponent and the now-zero-padded fp32
+    mantissa back out of the fp32 header. Subnormal range (|x| < 2^-6, grid step
+    2^-9): the value is an exact multiple of 2^-9, so ``|x| * 512`` IS the mantissa
+    field -- the sign bit has to be carried in by hand, since that branch never
+    looks at the header. The same 0.015625 boundary as :func:`round_e4m3` keeps the
+    two consistent: ``e4m3_u8_to_f32(e4m3_f32_to_u8(round_e4m3(x)))`` is x's
+    single-rounded value for every input, and no code it emits is a NaN pattern (the
+    caller's +-448 clamp caps the code at 0x7E). ``-0.0`` encodes as 0x00 after
+    round_e4m3 (which documents returning +0.0 for it).
+    """
+    u = x.to(tl.uint32, bitcast=True)
+    sign = ((u >> 31) & 1).to(tl.int32)
+    exp = ((u >> 23) & 0xFF).to(tl.int32) - 127
+    mant = ((u >> 20) & 7).to(tl.int32)
+    normal = (sign << 7) | ((exp + 7) << 3) | mant
+    sub = (sign << 7) | (tl.abs(x) * 512.0).to(tl.int32)
+    return tl.where(tl.abs(x) >= 0.015625, normal, sub).to(tl.uint8)
+
+
+@jit
+def kv_load_e4m3_tile_f32(ptrs, mask):
+    """Load a tile of KV e4m3 codes and widen it to fp32.
+
+    Straight-line on purpose: no probe, no dtype test, so there is no arm left to
+    prune. The pools keep their codes in a plain byte buffer on EVERY architecture
+    (kv_quant.kv_codes_dtype), so the fp8e4nv type never reaches Triton through here.
+    Both ways of choosing an arm were tried on real hardware and each broke the run:
+    the compile-time fp8-native answer is a second, independent verdict that can
+    disagree with the host that allocated the buffer, and a comparison against the
+    pointer's element type is NOT statically pruned -- Triton type-checks the arm that
+    should have been dead, and an int mask fill against an fp8 pointer is rejected
+    ("cannot cast int32 to fp8e4nv", raised at CUDA graph capture on sm_100).
+
+    What remains is the decode that already runs wherever the fp8 type is unavailable,
+    bit-exact per this module's header: the same load-with-int-fill and software
+    widening as kernel/triton/ple.py and nvfp4_linear.py. Callers use this only in
+    their quantized branch -- the 16-bit path keeps its own tl.load, so bf16 attention
+    is untouched instruction for instruction -- and the dense paged kernels and the
+    QSA sparse one read the same pool, hence this helper lives here.
+    """
+    return e4m3_u8_to_f32(tl.load(ptrs, mask=mask, other=0))
+
+
+# What kv_load_e4m3_tile_scaled16 leaves on its tile for the caller to repay.
+# tl.constexpr, not a plain float: a @jit function that references a module-level
+# python float fails triton's cache-key AST walk with a CompilationError, the same
+# way this module's header describes for @constexpr_function and host functions.
+# Host-side callers (tests) want ``KV_TILE_SCALE.value``.
+KV_TILE_SCALE = tl.constexpr(256.0)
+
+
+@jit
+def kv_load_e4m3_tile_scaled16(ptrs, mask):
+    """Load a tile of KV e4m3 codes as fp16 holding the value times 1/KV_TILE_SCALE.
+
+    :func:`kv_load_e4m3_tile_f32` without its last two steps: the widen to fp32 and
+    the ``* 256.0`` that puts the tile back on the true e4m3 scale exist only so the
+    result is directly usable, and they are what force the tile to 32 bits. A caller
+    that must apply a per-(token, kv_head) dequant scale anyway can fold ``2**8`` into
+    that scale instead -- exactly, since it is a power of two -- and keep the tile
+    16-bit. Callers owe that fold; see :data:`KV_TILE_SCALE`.
+
+    Same straight-line load-with-int-fill as that function, for the same reason (see
+    its docstring), and the same bit placement, folded: for ``v = 128s + r`` it builds
+    ``32768s + 128r`` with two masks, two widens, two shifts and an or, while
+    ``(v + (v & 0x80)) << 7 = (256s + r) << 7`` is the same number in one widen, one
+    mask, one add and one shift. Verified identical on all 256 codes, NaN patterns
+    (0x7F/0xFF -> +-480 after the caller's fold) included.
+    """
+    w = tl.load(ptrs, mask=mask, other=0).to(tl.uint16)
+    return ((w + (w & 0x80)) << 7).to(tl.float16, bitcast=True)
diff --git a/python/freetoken/kernel/triton/glm_dsa_sparse.py b/python/freetoken/kernel/triton/glm_dsa_sparse.py
index 8aaeeb9ba..d73ece8cf 100644
--- a/python/freetoken/kernel/triton/glm_dsa_sparse.py
+++ b/python/freetoken/kernel/triton/glm_dsa_sparse.py
@@ -28,6 +28,9 @@
 import triton
 import triton.language as tl
 
+from freetoken.kernel.triton.e4m3_compat import kv_load_e4m3_tile_f32
+from freetoken.kernel.triton.kv_nvfp4 import load_nvfp4
+
 BLOCK_H = 16
 BLOCK_T = 32
 MAX_SPLITS = 32
@@ -36,11 +39,12 @@
 
 @triton.jit
 def _glm_dsa_sparse_kernel(
-    q_ptr, pool_ptr, o_ptr, idx_ptr, cnt_ptr,
+    q_ptr, pool_ptr, pool_scale_ptr, pool_block_ptr, o_ptr, idx_ptr, cnt_ptr,
     scale,
     H, TOPK,
     stride_qb, stride_qm, stride_qh, stride_qd,
     stride_pn, stride_pd,
+    stride_ps,
     stride_ob, stride_om, stride_oh, stride_od,
     stride_ib, stride_im, stride_it,
     stride_nb, stride_nm,
@@ -50,6 +54,8 @@ def _glm_dsa_sparse_kernel(
     BLOCK_T: tl.constexpr,
     HAS_COUNTS: tl.constexpr,
     HAS_ROPE: tl.constexpr,
+    HAS_FP8: tl.constexpr,
+    HAS_NVFP4: tl.constexpr,
 ):
     pid_m = tl.program_id(0)
     pid_b = tl.program_id(1)
@@ -61,11 +67,15 @@ def _glm_dsa_sparse_kernel(
 
     q_base = q_ptr + pid_b * stride_qb + pid_m * stride_qm + offs_h[:, None] * stride_qh
     q_v = tl.load(q_base + offs_v[None, :] * stride_qd, mask=h_mask[:, None], other=0.0).to(tl.float32)
+    if HAS_NVFP4:
+        q_v = q_v.to(q_ptr.dtype.element_ty)
     if HAS_ROPE:
         # NoPE checkpoints (glm5_next) have D_R == 0: tl.arange needs a non-empty
         # span, so the whole rope half is compiled out on the constexpr.
         offs_r = tl.arange(0, D_R)
         q_r = tl.load(q_base + (D_V + offs_r[None, :]) * stride_qd, mask=h_mask[:, None], other=0.0).to(tl.float32)
+        if HAS_NVFP4:
+            q_r = q_r.to(q_ptr.dtype.element_ty)
 
     m_i = tl.full((BLOCK_H,), -float("inf"), dtype=tl.float32)
     l_i = tl.zeros((BLOCK_H,), dtype=tl.float32)
@@ -82,11 +92,34 @@ def _glm_dsa_sparse_kernel(
         idxs = tl.load(idx_base + offs_t * stride_it, mask=t_mask, other=-1)
         valid = idxs >= 0
         kv_base = pool_ptr + idxs[:, None] * stride_pn
-        kv_v = tl.load(kv_base + offs_v[None, :] * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
+        if HAS_NVFP4:
+            kv_v = load_nvfp4(
+                pool_ptr, pool_block_ptr, pool_scale_ptr,
+                idxs[:, None], 0, offs_v[None, :], valid[:, None],
+                stride_pn, 0, stride_ps, D_V + D_R,
+            ).to(q_ptr.dtype.element_ty)
+        elif HAS_FP8:
+            row_scale = tl.load(pool_scale_ptr + idxs * stride_ps, mask=valid, other=0.0)
+            kv_v = kv_load_e4m3_tile_f32(
+                kv_base + offs_v[None, :] * stride_pd, valid[:, None]
+            ) * row_scale[:, None]
+        else:
+            kv_v = tl.load(kv_base + offs_v[None, :] * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
 
         scores = tl.dot(q_v, tl.trans(kv_v))
         if HAS_ROPE:
-            kv_r = tl.load(kv_base + (D_V + offs_r[None, :]) * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
+            if HAS_NVFP4:
+                kv_r = load_nvfp4(
+                    pool_ptr, pool_block_ptr, pool_scale_ptr,
+                    idxs[:, None], 0, offs_r[None, :], valid[:, None],
+                    stride_pn, 0, stride_ps, D_V + D_R, DIM_OFFSET=D_V,
+                ).to(q_ptr.dtype.element_ty)
+            elif HAS_FP8:
+                kv_r = kv_load_e4m3_tile_f32(
+                    kv_base + (D_V + offs_r[None, :]) * stride_pd, valid[:, None]
+                ) * row_scale[:, None]
+            else:
+                kv_r = tl.load(kv_base + (D_V + offs_r[None, :]) * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
             scores += tl.dot(q_r, tl.trans(kv_r))
         scores = scores * scale
         scores = tl.where(valid[None, :], scores, -float("inf"))
@@ -99,7 +132,7 @@ def _glm_dsa_sparse_kernel(
         acc = acc * alpha[:, None] + tl.dot(p.to(kv_v.dtype), kv_v)
         m_i = m_new
 
-    o = acc / l_i[:, None]
+    o = tl.where(l_i[:, None] > 0, acc / l_i[:, None], 0.0)
     o_ptrs = o_ptr + pid_b * stride_ob + pid_m * stride_om + offs_h[:, None] * stride_oh + offs_v[None, :] * stride_od
     tl.store(o_ptrs, o.to(o_ptr.dtype.element_ty), mask=h_mask[:, None])
 
@@ -200,11 +233,12 @@ def glm_dsa_decode_logits(
 
 @triton.jit
 def _glm_dsa_splitk_kernel(
-    q_ptr, pool_ptr, mid_o_ptr, mid_lse_ptr, idx_ptr, cnt_ptr,
+    q_ptr, pool_ptr, pool_scale_ptr, pool_block_ptr, mid_o_ptr, mid_lse_ptr, idx_ptr, cnt_ptr,
     scale,
     H, TOPK,
     stride_qb, stride_qm, stride_qh, stride_qd,
     stride_pn, stride_pd,
+    stride_ps,
     stride_mb, stride_mm, stride_mh, stride_ms, stride_md,
     stride_lb, stride_lm, stride_lh, stride_ls,
     stride_ib, stride_im, stride_it,
@@ -215,6 +249,8 @@ def _glm_dsa_splitk_kernel(
     BLOCK_T: tl.constexpr,
     HAS_COUNTS: tl.constexpr,
     HAS_ROPE: tl.constexpr,
+    HAS_FP8: tl.constexpr,
+    HAS_NVFP4: tl.constexpr,
     NUM_SPLITS: tl.constexpr,
 ):
     """Stage 1 (decode flash-decoding): each program reduces one BLOCK_T-aligned slice of
@@ -245,9 +281,13 @@ def _glm_dsa_splitk_kernel(
     if split_end > split_start:
         q_base = q_ptr + pid_b * stride_qb + pid_m * stride_qm + offs_h[:, None] * stride_qh
         q_v = tl.load(q_base + offs_v[None, :] * stride_qd, mask=h_mask[:, None], other=0.0).to(tl.float32)
+        if HAS_NVFP4:
+            q_v = q_v.to(q_ptr.dtype.element_ty)
         if HAS_ROPE:
             offs_r = tl.arange(0, D_R)
             q_r = tl.load(q_base + (D_V + offs_r[None, :]) * stride_qd, mask=h_mask[:, None], other=0.0).to(tl.float32)
+            if HAS_NVFP4:
+                q_r = q_r.to(q_ptr.dtype.element_ty)
         idx_base = idx_ptr + pid_b * stride_ib + pid_m * stride_im
 
         for start in range(split_start, split_end, BLOCK_T):
@@ -256,11 +296,34 @@ def _glm_dsa_splitk_kernel(
             idxs = tl.load(idx_base + offs_t * stride_it, mask=t_mask, other=-1)
             valid = idxs >= 0
             kv_base = pool_ptr + idxs[:, None] * stride_pn
-            kv_v = tl.load(kv_base + offs_v[None, :] * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
+            if HAS_NVFP4:
+                kv_v = load_nvfp4(
+                    pool_ptr, pool_block_ptr, pool_scale_ptr,
+                    idxs[:, None], 0, offs_v[None, :], valid[:, None],
+                    stride_pn, 0, stride_ps, D_V + D_R,
+                ).to(q_ptr.dtype.element_ty)
+            elif HAS_FP8:
+                row_scale = tl.load(pool_scale_ptr + idxs * stride_ps, mask=valid, other=0.0)
+                kv_v = kv_load_e4m3_tile_f32(
+                    kv_base + offs_v[None, :] * stride_pd, valid[:, None]
+                ) * row_scale[:, None]
+            else:
+                kv_v = tl.load(kv_base + offs_v[None, :] * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
 
             scores = tl.dot(q_v, tl.trans(kv_v))
             if HAS_ROPE:
-                kv_r = tl.load(kv_base + (D_V + offs_r[None, :]) * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
+                if HAS_NVFP4:
+                    kv_r = load_nvfp4(
+                        pool_ptr, pool_block_ptr, pool_scale_ptr,
+                        idxs[:, None], 0, offs_r[None, :], valid[:, None],
+                        stride_pn, 0, stride_ps, D_V + D_R, DIM_OFFSET=D_V,
+                    ).to(q_ptr.dtype.element_ty)
+                elif HAS_FP8:
+                    kv_r = kv_load_e4m3_tile_f32(
+                        kv_base + (D_V + offs_r[None, :]) * stride_pd, valid[:, None]
+                    ) * row_scale[:, None]
+                else:
+                    kv_r = tl.load(kv_base + (D_V + offs_r[None, :]) * stride_pd, mask=valid[:, None], other=0.0).to(tl.float32)
                 scores += tl.dot(q_r, tl.trans(kv_r))
             scores = scores * scale
             scores = tl.where(valid[None, :], scores, -float("inf"))
@@ -323,7 +386,7 @@ def _glm_dsa_merge_kernel(
         l_i = l_i * alpha + beta
         m_i = m_new
 
-    o = acc / l_i
+    o = tl.where(l_i > 0, acc / l_i, 0.0)
     o_ptrs = (
         o_ptr + pid_b * stride_ob + pid_m * stride_om + pid_h * stride_oh
         + offs_v * stride_od
@@ -351,7 +414,10 @@ def glm_dsa_sparse_attn(
     softmax_scale: float,
     counts: torch.Tensor | None = None,  # [b, m] int32 live columns per query (device-read)
     d_v: int = 512,
+    pool_scale: torch.Tensor | None = None,  # [rows] fp32, one scale per quantized latent row
     force_splits: int | None = None,     # tests only: 0 = single-program, N = split-k N
+    kv_quant: str | None = None,
+    pool_block_scale: torch.Tensor | None = None,
 ) -> torch.Tensor:
     """Sparse MLA attention over gathered latent rows; returns ``[b, m, h, d_v]``.
 
@@ -359,13 +425,45 @@ def glm_dsa_sparse_attn(
     list across all queries with STRIDE 0 -- the identity-selection dense path: every
     query reads the same position-ordered row list, causally bounded by its own
     ``counts[q] = position + 1``, with zero per-query index materialization.
+
+    NVFP4 restores tiles to Q's compute dtype before the dot products, with FP32
+    accumulators. FP32 dot operands exceed consumer GPU shared memory at width 576.
     """
     b, m, h, d = q.shape
     d_r = d - d_v
     topk = topk_idxs.shape[-1]
-    assert pool.shape[-1] == d, (pool.shape, d)
+    if kv_quant is None:
+        kv_quant = "fp8" if pool_scale is not None else "none"
+    assert kv_quant in ("none", "fp8", "nvfp4"), kv_quant
+    has_fp8 = kv_quant == "fp8"
+    has_nvfp4 = kv_quant == "nvfp4"
+    stored_dim = d // 2 if has_nvfp4 else d
+    assert pool.ndim == 2 and pool.shape[-1] == stored_dim, (pool.shape, d)
+    if has_fp8 or has_nvfp4:
+        assert pool_scale is not None
+        assert pool.dtype == torch.uint8, pool.dtype
+        assert pool_scale.shape == (pool.shape[0],), (pool_scale.shape, pool.shape)
+        assert pool_scale.dtype is torch.float32, pool_scale.dtype
+        scale_pool = pool_scale.contiguous()
+    else:
+        assert pool_scale is None
+        assert pool.is_floating_point()
+        scale_pool = pool
+    if has_nvfp4:
+        assert d % 16 == 0
+        assert pool_block_scale is not None
+        assert pool_block_scale.shape == (pool.shape[0], d // 16)
+        assert pool_block_scale.dtype == torch.uint8
+        assert pool_block_scale.device == pool.device
+        block_pool = pool_block_scale.contiguous()
+    else:
+        assert pool_block_scale is None
+        block_pool = pool
+    assert pool.device == q.device
+    if has_fp8 or has_nvfp4:
+        assert pool_scale.device == pool.device
     q = q.contiguous()
-    pool_2d = pool.reshape(-1, d)
+    pool_2d = pool.reshape(-1, stored_dim)
     assert pool_2d.stride(-1) == 1
     idx = topk_idxs.contiguous().to(torch.int32)
     broadcast_m = idx.shape[1] == 1 and m > 1
@@ -379,25 +477,30 @@ def glm_dsa_sparse_attn(
     else:
         cnt, stride_nb, stride_nm = idx, 0, 0
 
+    # Packed gathers need additional layout conversions at the 512-wide latent size.
+    block_t = 16 if has_nvfp4 else BLOCK_T
     n_splits = _split_count(b, m, h, topk, q.device) if force_splits is None else force_splits
     if n_splits:
         mid_o = q.new_empty(b, m, h, n_splits, d_v, dtype=torch.float32)
         mid_lse = q.new_empty(b, m, h, n_splits, dtype=torch.float32)
         grid1 = (m * n_splits, b, triton.cdiv(h, BLOCK_H))
         _glm_dsa_splitk_kernel[grid1](
-            q, pool_2d, mid_o, mid_lse, idx, cnt,
+            q, pool_2d, scale_pool, block_pool, mid_o, mid_lse, idx, cnt,
             float(softmax_scale),
             h, topk,
             q.stride(0), q.stride(1), q.stride(2), q.stride(3),
             pool_2d.stride(0), pool_2d.stride(1),
+            scale_pool.stride(0) if has_fp8 or has_nvfp4 else 0,
             mid_o.stride(0), mid_o.stride(1), mid_o.stride(2), mid_o.stride(3), mid_o.stride(4),
             mid_lse.stride(0), mid_lse.stride(1), mid_lse.stride(2), mid_lse.stride(3),
             idx.stride(0), 0 if broadcast_m else idx.stride(1), idx.stride(2),
             stride_nb, stride_nm,
             D_V=d_v, D_R=d_r,
-            BLOCK_H=BLOCK_H, BLOCK_T=BLOCK_T,
-            HAS_COUNTS=has_counts, HAS_ROPE=d_r > 0, NUM_SPLITS=n_splits,
-            num_warps=4, num_stages=2,
+            BLOCK_H=BLOCK_H, BLOCK_T=block_t,
+            HAS_COUNTS=has_counts, HAS_ROPE=d_r > 0, HAS_FP8=has_fp8, HAS_NVFP4=has_nvfp4, NUM_SPLITS=n_splits,
+            # The 512-wide latent accumulator exceeds the 99 KiB shared-memory
+            # limit of consumer Blackwell GPUs with a two-stage pipeline.
+            num_warps=4, num_stages=1,
         )
         grid2 = (m, b, h)
         _glm_dsa_merge_kernel[grid2](
@@ -412,18 +515,19 @@ def glm_dsa_sparse_attn(
 
     grid = (m, b, triton.cdiv(h, BLOCK_H))
     _glm_dsa_sparse_kernel[grid](
-        q, pool_2d, o, idx, cnt,
+        q, pool_2d, scale_pool, block_pool, o, idx, cnt,
         float(softmax_scale),
         h, topk,
         q.stride(0), q.stride(1), q.stride(2), q.stride(3),
         pool_2d.stride(0), pool_2d.stride(1),
+        scale_pool.stride(0) if has_fp8 or has_nvfp4 else 0,
         o.stride(0), o.stride(1), o.stride(2), o.stride(3),
         idx.stride(0), 0 if broadcast_m else idx.stride(1), idx.stride(2),
         stride_nb, stride_nm,
         D_V=d_v, D_R=d_r,
-        BLOCK_H=BLOCK_H, BLOCK_T=BLOCK_T,
-        HAS_COUNTS=has_counts, HAS_ROPE=d_r > 0,
-        num_warps=4, num_stages=2,
+        BLOCK_H=BLOCK_H, BLOCK_T=block_t,
+        HAS_COUNTS=has_counts, HAS_ROPE=d_r > 0, HAS_FP8=has_fp8, HAS_NVFP4=has_nvfp4,
+        num_warps=4, num_stages=1,
     )
     return o
 
diff --git a/python/freetoken/kernel/triton/kv_nvfp4.py b/python/freetoken/kernel/triton/kv_nvfp4.py
new file mode 100644
index 000000000..04727deec
--- /dev/null
+++ b/python/freetoken/kernel/triton/kv_nvfp4.py
@@ -0,0 +1,154 @@
+"""Row-scaled NVFP4 KV storage, with low nibble first and 16-wide blocks.
+
+Each (token, head) row has an FP32 scale; each block has an E4M3 scale
+stored as uint8. Reconstruction is E2M1 * block_scale * row_scale.
+Row-local second-level scales keep appends independent of the cached prefix.
+"""
+
+import torch
+import triton
+import triton.language as tl
+
+from .e4m3_compat import e4m3_f32_to_u8, e4m3_u8_to_f32, round_e4m3
+
+
+@triton.jit
+def _encode_e2m1(x):
+    a = tl.abs(x)
+    code = tl.full(x.shape, 0, tl.int32)
+    code = tl.where(a > 0.25, 1, code)
+    code = tl.where(a >= 0.75, 2, code)
+    code = tl.where(a > 1.25, 3, code)
+    code = tl.where(a >= 1.75, 4, code)
+    code = tl.where(a > 2.5, 5, code)
+    code = tl.where(a >= 3.5, 6, code)
+    code = tl.where(a > 5.0, 7, code)
+    return code | tl.where(x < 0, 8, 0)
+
+
+@triton.jit
+def _decode_e2m1(code):
+    # Place E2M1 bits in FP16, then compensate the exponent-bias difference (15 - 1).
+    bits = ((code & 8).to(tl.uint16) << 12) | ((code & 7).to(tl.uint16) << 9)
+    return bits.to(tl.float16, bitcast=True).to(tl.float32) * 16384.0
+
+
+@triton.jit
+def load_nvfp4(ptr, block_ptr, row_ptr, slots, head, dims, slot_mask,
+               stride_slot, stride_head, stride_row, D: tl.constexpr,
+               DIM_OFFSET: tl.constexpr = 0):
+    TRANSPOSE: tl.constexpr = dims.shape[0] != 1
+    WIDTH: tl.constexpr = dims.shape[0] if TRANSPOSE else dims.shape[1]
+    TOKENS: tl.constexpr = slots.shape[0] * slots.shape[1]
+    slot = slots.reshape(TOKENS).to(tl.int64)
+    valid = slot_mask.reshape(TOKENS)
+    dim = DIM_OFFSET + tl.arange(0, WIDTH)
+    packed = tl.load(
+        ptr + slot[:, None] * stride_slot + head * stride_head + (dim[None, :] // 2),
+        valid[:, None] & (dim[None, :] < D), other=0,
+    ).to(tl.int32)
+    codes = tl.where((dim[None, :] & 1) == 0, packed & 15, packed >> 4)
+    block_dim = dim // 16
+    block = tl.load(
+        block_ptr + (slot[:, None] * stride_row + head) * (D // 16) + block_dim[None, :],
+        valid[:, None] & (dim[None, :] < D), other=0,
+    )
+    row = tl.load(row_ptr + slot * stride_row + head, valid, other=0)
+    value = _decode_e2m1(codes) * e4m3_u8_to_f32(block) * row[:, None]
+    if TRANSPOSE:
+        return tl.trans(value)
+    else:
+        return value
+
+
+@triton.jit
+def _quantize_row(src, dst, block_ptr, row_ptr, t, h, slot, stride_src,
+                  HEADS: tl.constexpr, D: tl.constexpr, BLOCKS: tl.constexpr):
+    blocks = tl.arange(0, BLOCKS)
+    dims = blocks[:, None] * 16 + tl.arange(0, 16)[None, :]
+    x = tl.load(src + t * stride_src + h * D + dims, dims < D, other=0).to(tl.float32)
+    amax = tl.max(tl.abs(x), 1)
+    row_scale = tl.maximum(tl.max(amax, 0), 1e-10) / (6.0 * 448.0)
+    block_scale = round_e4m3(tl.minimum(tl.div_rn(amax, 6.0 * row_scale), 448.0))
+    # Quantize against the scale actually stored, including E4M3 rounding/underflow.
+    denom = tl.where(block_scale > 0, block_scale * row_scale, 1.0)
+    normalized = tl.where(block_scale[:, None] > 0, tl.div_rn(x, denom[:, None]), 0.0)
+    codes = _encode_e2m1(normalized).reshape(BLOCKS, 8, 2)
+    lo, hi = tl.split(codes)
+    packed = lo | (hi << 4)
+    byte_dims = blocks[:, None] * 8 + tl.arange(0, 8)[None, :]
+    tl.store(dst + (slot * HEADS + h) * (D // 2) + byte_dims, packed, byte_dims < D // 2)
+    tl.store(block_ptr + (slot * HEADS + h) * (D // 16) + blocks,
+             e4m3_f32_to_u8(block_scale), blocks < D // 16)
+    tl.store(row_ptr + slot * HEADS + h, row_scale)
+
+
+@triton.jit
+def _scatter_rows(src, dst, block, row, indices, stride_src,
+                  D: tl.constexpr, BLOCKS: tl.constexpr):
+    t = tl.program_id(0)
+    slot = tl.load(indices + t).to(tl.int64)
+    _quantize_row(src, dst, block, row, t, 0, slot, stride_src, 1, D, BLOCKS)
+
+
+def quantize_nvfp4_rows_to_cache(rows, out_loc, cache, scales, block_scales) -> None:
+    """Quantize a single MLA latent slab, with one second-level scale per token."""
+    tokens, dim = rows.shape
+    slots = cache.shape[0]
+    assert dim % 16 == 0 and rows.stride(1) == 1
+    assert rows.dtype in (torch.float16, torch.bfloat16, torch.float32)
+    assert cache.shape == (slots, dim // 2) and cache.dtype == torch.uint8
+    assert scales.shape == (slots,) and scales.dtype == torch.float32
+    assert block_scales.shape == (slots, dim // 16) and block_scales.dtype == torch.uint8
+    assert out_loc.shape == (tokens,) and out_loc.dtype in (torch.int32, torch.int64)
+    for tensor in (cache, scales, block_scales, out_loc):
+        assert tensor.is_contiguous() and tensor.device == rows.device
+    assert rows.is_cuda
+    if tokens:
+        _scatter_rows[(tokens,)](
+            rows, cache, block_scales, scales, out_loc, rows.stride(0),
+            D=dim, BLOCKS=triton.next_power_of_2(dim // 16),
+            num_warps=4, enable_fp_fusion=False,
+        )
+
+
+@triton.jit
+def _scatter(k, v, kc, vc, kb, vb, kr, vr, indices, stride_k, stride_v,
+             HEADS: tl.constexpr, D: tl.constexpr, BLOCKS: tl.constexpr):
+    t, h = tl.program_id(0), tl.program_id(1)
+    slot = tl.load(indices + t).to(tl.int64)
+    _quantize_row(k, kc, kb, kr, t, h, slot, stride_k, HEADS, D, BLOCKS)
+    _quantize_row(v, vc, vb, vr, t, h, slot, stride_v, HEADS, D, BLOCKS)
+
+
+def quantize_nvfp4_to_cache(
+    k: torch.Tensor,
+    v: torch.Tensor,
+    out_loc: torch.Tensor,
+    k_cache: torch.Tensor,
+    v_cache: torch.Tensor,
+    k_scale: torch.Tensor,
+    v_scale: torch.Tensor,
+    k_block_scale: torch.Tensor,
+    v_block_scale: torch.Tensor,
+) -> None:
+    tokens, width = k.shape
+    slots, heads, packed_dim = k_cache.shape
+    dim = packed_dim * 2
+    assert dim % 16 == 0 and width == heads * dim
+    assert v.shape == k.shape and k.stride(1) == v.stride(1) == 1
+    assert k.dtype in (torch.float16, torch.bfloat16, torch.float32) and v.dtype == k.dtype
+    assert out_loc.shape == (tokens,) and out_loc.dtype in (torch.int32, torch.int64)
+    assert out_loc.is_contiguous()
+    for codes, row, block in ((k_cache, k_scale, k_block_scale), (v_cache, v_scale, v_block_scale)):
+        assert codes.shape == (slots, heads, packed_dim) and codes.dtype == torch.uint8
+        assert row.shape == (slots, heads) and row.dtype == torch.float32
+        assert block.shape == (slots, heads, dim // 16) and block.dtype == torch.uint8
+        assert codes.is_contiguous() and row.is_contiguous() and block.is_contiguous()
+        assert codes.device == row.device == block.device == k.device
+    assert k.is_cuda and v.device == out_loc.device == k.device
+    if tokens:
+        _scatter[(tokens, heads)](k, v, k_cache, v_cache, k_block_scale, v_block_scale,
+                                 k_scale, v_scale, out_loc, k.stride(0), v.stride(0),
+                                 HEADS=heads, D=dim, BLOCKS=triton.next_power_of_2(dim // 16),
+                                 num_warps=4, enable_fp_fusion=False)
diff --git a/python/freetoken/kernel/triton/kv_quant.py b/python/freetoken/kernel/triton/kv_quant.py
new file mode 100644
index 000000000..c10ad9c4b
--- /dev/null
+++ b/python/freetoken/kernel/triton/kv_quant.py
@@ -0,0 +1,249 @@
+"""FP8 (e4m3) KV-cache storage: per-token, per-head symmetric quantization.
+
+One KV row is one ``(token, kv_head)`` slice of ``head_dim`` elements. It is stored
+as ``head_dim`` e4m3 bytes plus ONE fp32 scale shared by the whole row:
+
+    scale = max(amax(row) / 448, eps)          # 448 == e4m3 finite max
+    code  = round_e4m3(clamp(row / scale))     # RNE onto the e4m3 grid
+    read  = code.to(f32) * scale               # in the attention kernels
+
+Granularity rationale: an fp32 scale per (token, head) costs ``4 / head_dim`` bytes
+per element (3% at head_dim 128, 6% at 64) while tracking each key's own magnitude,
+which is what keeps a quantized KV from collapsing on outlier heads. A coarser
+per-tensor scale needs no storage at all but has no headroom for them; a finer
+per-element "scale" is the format itself.
+
+Architectures below sm_89 have no fp8e4nv type in Triton (see
+:mod:`freetoken.kernel.triton.e4m3_compat`), so the codes live in a plain ``uint8``
+buffer on EVERY architecture and are decoded by :func:`e4m3_u8_to_f32`. That keeps one
+set of bytes and one set of numbers across GPUs, and it is why :func:`kv_codes_dtype`
+is a constant rather than a question: the fp8 type never appears in a kernel
+signature, so nothing here can disagree with the host that allocated the buffer.
+Choosing the encode/decode per target -- by an arch probe, or by testing the pointer's
+element type -- is what broke this feature twice on real hardware (see
+:func:`kv_load_e4m3_tile_f32`).
+
+The write path replaces ``kernel.store_cache`` for a quantized pool: the plain store
+kernel is a raw byte copy that requires the source and the cache to share a dtype,
+and quantization is exactly the step where the two diverge. Folding the scatter into
+the quantization kernel keeps that to a single launch (and a single HBM round trip)
+under CUDA-graph capture, where the slot ids arrive as a device tensor.
+"""
+
+from __future__ import annotations
+
+import torch
+import triton
+import triton.language as tl
+
+from freetoken.kernel.triton.e4m3_compat import (
+    e4m3_f32_to_u8,
+    round_e4m3,
+)
+
+FP8 = torch.float8_e4m3fn
+KV_SCALE_DTYPE = torch.float32
+KV_QUANT_FP8 = "fp8"
+
+
+def kv_codes_dtype() -> torch.dtype:
+    """Storage dtype of one quantized KV element: e4m3 bytes in a uint8 buffer.
+
+    A constant, deliberately. Every attempt to answer this per target -- triton's
+    compile-time probe, then testing the pointer's element type at the load -- ended up
+    picking an arm that did not match the buffer the host had just allocated (see the
+    module header). torch still reads these bytes as fp8 whenever real numbers are
+    wanted: :func:`codes_to_f32`.
+    """
+    return torch.uint8
+
+
+def alloc_codes(shape: tuple[int, ...], device: torch.device) -> torch.Tensor:
+    """A zero-filled code buffer of :func:`kv_codes_dtype` -- bytes, on every arch.
+
+    Zero-filling matters because the pools read slots that were never written: a stale
+    byte decodes to a real number (0x7F/0xFF even to NaN once reinterpreted as fp8),
+    while 0x00 is exactly 0.0 (same reasoning as kvcache/bsa_pool.py).
+    """
+    return torch.zeros(shape, dtype=kv_codes_dtype(), device=device)
+
+
+def codes_to_f32(codes: torch.Tensor) -> torch.Tensor:
+    """Decode a code buffer to fp32 ON THE HOST (torch's own e4m3 cast).
+
+    Works on either storage dtype -- uint8 bytes are reinterpreted as fp8 first -- so
+    a test or a debugging tool reads the same numbers on sm_86 as on sm_90, and does
+    it through torch rather than through the software decoder it is checking.
+    """
+    if codes.dtype is not FP8:
+        codes = codes.view(FP8)
+    return codes.to(torch.float32)
+
+
+@triton.jit
+def _kv_quant_scatter_kernel(
+    k_src,
+    v_src,
+    k_dst,
+    v_dst,
+    k_scale,
+    v_scale,
+    idx_ptr,
+    stride_xs,  # K source row pitch, in elements (the qkv slice is wider than one row)
+    stride_vx,  # V source row pitch. K and V need not share one: a .clamp() on one side
+                # leaves it densely packed, so reusing K's pitch reads V off its rows.
+    stride_kd,  # K cache row pitch, in elements (== HEADS * D)
+    stride_vd,
+    stride_ks,  # scale row pitch, in elements (== HEADS)
+    stride_vs,
+    D: tl.constexpr,
+    BLOCK_D: tl.constexpr,
+):
+    """One program per (token, kv_head): quantize the row and write it to slot
+    ``idx_ptr[token]`` (the same ``out_loc`` the bf16 store scatters through)."""
+    t = tl.program_id(0)
+    h = tl.program_id(1)
+    pos = tl.load(idx_ptr + t).to(tl.int64)  # int64: slots * row can pass 2**31
+    d = tl.arange(0, BLOCK_D)
+    mask = d < D
+
+    off = h * D + d
+    xk = tl.load(k_src + t * stride_xs + off, mask=mask, other=0.0).to(tl.float32)
+    xv = tl.load(v_src + t * stride_vx + off, mask=mask, other=0.0).to(tl.float32)
+
+    # 448 == e4m3 finite max; 1e-10 is the amax floor of the activation quant in
+    # kernel/triton/fp8_block_linear.py (literals keep the kernel self-contained).
+    sk = tl.maximum(tl.max(tl.abs(xk), axis=0), 1e-10) / 448.0
+    sv = tl.maximum(tl.max(tl.abs(xv), axis=0), 1e-10) / 448.0
+    qk = tl.clamp(xk / sk, -448.0, 448.0)
+    qv = tl.clamp(xv / sv, -448.0, 448.0)
+
+    # Straight-line, like the reader: round onto the e4m3 grid in ONE step (RNE) and
+    # pack the bits into the byte buffer. No fp8 type on either side -- the reason this
+    # feature stopped compiling twice is explained in kv_codes_dtype's docstring.
+    out_k = e4m3_f32_to_u8(round_e4m3(qk))
+    out_v = e4m3_f32_to_u8(round_e4m3(qv))
+
+    tl.store(k_dst + pos * stride_kd + h * D + d, out_k, mask=mask)
+    tl.store(v_dst + pos * stride_vd + h * D + d, out_v, mask=mask)
+    tl.store(k_scale + pos * stride_ks + h, sk)
+    tl.store(v_scale + pos * stride_vs + h, sv)
+
+
+def quantize_kv_to_cache(
+    k: torch.Tensor,
+    v: torch.Tensor,
+    out_loc: torch.Tensor,
+    k_cache: torch.Tensor,
+    v_cache: torch.Tensor,
+    k_scale: torch.Tensor,
+    v_scale: torch.Tensor,
+) -> None:
+    """Quantize fresh K/V rows into an fp8 KV pool.
+
+    ``k``/``v``  : ``[T, num_kv_heads * head_dim]`` compute-dtype rows -- exactly what
+        the attention backends hand to ``store_kv`` (a slice of the qkv projection,
+        so the row pitch may be wider than the row itself).
+    ``out_loc``  : ``[T]`` device slot index per row (int32 or int64).
+    ``*_cache``  : ``[num_slots, num_kv_heads, head_dim]`` of :func:`kv_codes_dtype`.
+    ``*_scale``  : ``[num_slots, num_kv_heads]`` fp32, indexed by the SAME slot.
+    """
+    tokens = k.shape[0]
+    if tokens == 0:
+        return
+    assert k.dim() == 2 and v.shape == k.shape, (k.shape, v.shape)
+    assert k.stride(1) == 1 and v.stride(1) == 1, "K/V rows must be contiguous"
+    heads, dim = k_cache.shape[1], k_cache.shape[2]
+    assert k.shape[1] == heads * dim, (tuple(k.shape), tuple(k_cache.shape))
+    assert k_cache.shape == v_cache.shape, (tuple(k_cache.shape), tuple(v_cache.shape))
+    assert k_scale.shape == v_scale.shape == (k_cache.shape[0], heads), (
+        tuple(k_scale.shape),
+        tuple(k_cache.shape),
+    )
+    assert k_cache.dtype == v_cache.dtype == kv_codes_dtype(), (
+        k_cache.dtype,
+        kv_codes_dtype(),
+    )
+    assert k_scale.dtype == KV_SCALE_DTYPE, k_scale.dtype
+    _kv_quant_scatter_kernel[(tokens, heads)](
+        k,
+        v,
+        k_cache,
+        v_cache,
+        k_scale,
+        v_scale,
+        out_loc,
+        k.stride(0),
+        v.stride(0),
+        k_cache.stride(0),
+        v_cache.stride(0),
+        k_scale.stride(0),
+        v_scale.stride(0),
+        D=dim,
+        BLOCK_D=triton.next_power_of_2(dim),
+        num_warps=1,
+    )
+
+
+@triton.jit
+def _kv_quant_rows_scatter_kernel(
+    src, dst, scale_ptr, idx_ptr,
+    stride_src, stride_dst, stride_scale,
+    D: tl.constexpr,
+    BLOCK_D: tl.constexpr,
+):
+    """Quantize one whole latent row per program and scatter it by ``out_loc``."""
+    t = tl.program_id(0)
+    pos = tl.load(idx_ptr + t).to(tl.int64)
+    d = tl.arange(0, BLOCK_D)
+    mask = d < D
+    x = tl.load(src + t * stride_src + d, mask=mask, other=0.0).to(tl.float32)
+    s = tl.maximum(tl.max(tl.abs(x), axis=0), 1e-10) / 448.0
+    q = tl.clamp(x / s, -448.0, 448.0)
+    tl.store(dst + pos * stride_dst + d, e4m3_f32_to_u8(round_e4m3(q)), mask=mask)
+    tl.store(scale_ptr + pos * stride_scale, s)
+
+
+def quantize_rows_to_cache(
+    rows: torch.Tensor,
+    out_loc: torch.Tensor,
+    cache: torch.Tensor,
+    scales: torch.Tensor,
+) -> None:
+    """Quantize and scatter one scale-bearing FP8 row per token.
+
+    MLA stores its absorbed K/V state as one latent row, rather than separate K and
+    V heads. Its scale granularity is therefore one FP32 value per ``(token,
+    layer)`` row, which is the single-head MLA case already priced by
+    ``kv_scale_bytes_per_token``.
+    """
+    tokens = rows.shape[0]
+    if tokens == 0:
+        return
+    assert rows.dim() == 2 and rows.stride(1) == 1, tuple(rows.shape)
+    assert cache.dim() == 2 and cache.shape[1] == rows.shape[1], (
+        tuple(cache.shape), tuple(rows.shape),
+    )
+    assert scales.shape == (cache.shape[0],), (tuple(scales.shape), tuple(cache.shape))
+    assert cache.dtype == kv_codes_dtype(), cache.dtype
+    assert scales.dtype == KV_SCALE_DTYPE, scales.dtype
+    dim = rows.shape[1]
+    _kv_quant_rows_scatter_kernel[(tokens,)](
+        rows, cache, scales, out_loc,
+        rows.stride(0), cache.stride(0), scales.stride(0),
+        D=dim, BLOCK_D=triton.next_power_of_2(dim),
+        num_warps=4 if dim > 256 else 1,
+    )
+
+
+__all__ = [
+    "FP8",
+    "KV_QUANT_FP8",
+    "KV_SCALE_DTYPE",
+    "alloc_codes",
+    "codes_to_f32",
+    "kv_codes_dtype",
+    "quantize_kv_to_cache",
+    "quantize_rows_to_cache",
+]
+
diff --git a/python/freetoken/kernel/triton/qsa/attend.py b/python/freetoken/kernel/triton/qsa/attend.py
index 541e27c68..f6e223385 100644
--- a/python/freetoken/kernel/triton/qsa/attend.py
+++ b/python/freetoken/kernel/triton/qsa/attend.py
@@ -9,12 +9,19 @@
 import triton
 import triton.language as tl
 
+from freetoken.kernel.triton.e4m3_compat import kv_load_e4m3_tile_f32
+from freetoken.kernel.triton.kv_nvfp4 import load_nvfp4
+
 
 @triton.jit
 def _qsa_sparse_paged_gqa_splitk_kernel(
     q_ptr,
     k_cache_ptr,
     v_cache_ptr,
+    k_scale_ptr,
+    v_scale_ptr,
+    k_block_scale_ptr,
+    v_block_scale_ptr,
     indices_ptr,
     block_table_ptr,
     token_to_req_ptr,
@@ -29,6 +36,8 @@ def _qsa_sparse_paged_gqa_splitk_kernel(
     stride_v_block,
     stride_v_token,
     stride_v_head,
+    stride_kss,
+    stride_vss,
     stride_indices_row,
     stride_table_req,
     stride_output_row,
@@ -46,6 +55,11 @@ def _qsa_sparse_paged_gqa_splitk_kernel(
     NUM_TILES: tl.constexpr,
     BLOCK_M: tl.constexpr,
     BLOCK_N: tl.constexpr,
+    # e4m3 KV pool (kvcache/mha_pool.py): read codes + per-token row scales instead of
+    # 16-bit values. The bf16 branch below stays exactly as it was, instruction for
+    # instruction, for the unquantized default.
+    HAS_KV_SCALE: tl.constexpr,
+    KV_NVFP4: tl.constexpr,
 ) -> None:
     # row * stride can overflow int32 for large row counts.
     row = tl.program_id(0).to(tl.int64)
@@ -101,24 +115,94 @@ def _qsa_sparse_paged_gqa_splitk_kernel(
         valid &= (physical_page >= 0) & (physical_page < num_cache_blocks)
         # physical_page * block stride can overflow int32 for large caches.
         safe_page = tl.maximum(physical_page, 0).to(tl.int64)
-        keys = tl.load(
-            k_cache_ptr
-            + safe_page[None, :] * stride_k_block
-            + page_offset[None, :] * stride_k_token
-            + kv_head * stride_k_head
-            + dim_offsets[:, None],
-            mask=valid[None, :],
-            other=0.0,
-        )
-        values = tl.load(
-            v_cache_ptr
-            + safe_page[:, None] * stride_v_block
-            + page_offset[:, None] * stride_v_token
-            + kv_head * stride_v_head
-            + dim_offsets[None, :],
-            mask=valid[:, None],
-            other=0.0,
-        )
+        if KV_NVFP4:
+            scale_slot = safe_page * PAGE_SIZE + page_offset
+            keys = load_nvfp4(
+                k_cache_ptr,
+                k_block_scale_ptr,
+                k_scale_ptr,
+                scale_slot[None, :],
+                kv_head,
+                dim_offsets[:, None],
+                valid[None, :],
+                stride_k_token,
+                stride_k_head,
+                stride_kss,
+                HEAD_DIM,
+            ).to(query.dtype)
+            values = load_nvfp4(
+                v_cache_ptr,
+                v_block_scale_ptr,
+                v_scale_ptr,
+                scale_slot[:, None],
+                kv_head,
+                dim_offsets[None, :],
+                valid[:, None],
+                stride_v_token,
+                stride_v_head,
+                stride_vss,
+                HEAD_DIM,
+            ).to(query.dtype)
+        elif HAS_KV_SCALE:
+            # The scale row is the slot the code lives in: QSA pins page_size to this
+            # kernel's PAGE_SIZE (attention/__init__.py registers page_sizes=(64,)), so
+            # slot = page * PAGE_SIZE + offset addresses k_scale/v_scale exactly.
+            scale_slot = safe_page * PAGE_SIZE + page_offset  # safe_page is int64
+            # Invalid columns mask the codes to 0.0 and the scale to 1.0, and slots that
+            # were never written read back 0.0 * 0.0 -- both buffers are zero-filled.
+            # Either way the operand stays finite, so the -inf row mask below is what
+            # decides such a column's fate rather than a NaN poisoning the row.
+            s_k = tl.load(
+                k_scale_ptr + scale_slot[None, :] * stride_kss + kv_head,
+                mask=valid[None, :],
+                other=1.0,
+            )
+            s_v = tl.load(
+                v_scale_ptr + scale_slot[:, None] * stride_vss + kv_head,
+                mask=valid[:, None],
+                other=1.0,
+            )
+            keys = (
+                kv_load_e4m3_tile_f32(
+                    k_cache_ptr
+                    + safe_page[None, :] * stride_k_block
+                    + page_offset[None, :] * stride_k_token
+                    + kv_head * stride_k_head
+                    + dim_offsets[:, None],
+                    valid[None, :],
+                )
+                * s_k
+            ).to(query.dtype)
+            values = (
+                kv_load_e4m3_tile_f32(
+                    v_cache_ptr
+                    + safe_page[:, None] * stride_v_block
+                    + page_offset[:, None] * stride_v_token
+                    + kv_head * stride_v_head
+                    + dim_offsets[None, :],
+                    valid[:, None],
+                )
+                * s_v
+            ).to(query.dtype)
+        else:
+            keys = tl.load(
+                k_cache_ptr
+                + safe_page[None, :] * stride_k_block
+                + page_offset[None, :] * stride_k_token
+                + kv_head * stride_k_head
+                + dim_offsets[:, None],
+                mask=valid[None, :],
+                other=0.0,
+            )
+            values = tl.load(
+                v_cache_ptr
+                + safe_page[:, None] * stride_v_block
+                + page_offset[:, None] * stride_v_token
+                + kv_head * stride_v_head
+                + dim_offsets[None, :],
+                mask=valid[:, None],
+                other=0.0,
+            )
         scores = tl.dot(query, keys)
         # Scaling scores avoids re-quantizing a scaled query to BF16.
         scores *= softmax_scale_log2
@@ -232,8 +316,13 @@ def qsa_sparse_paged_attention(
     block_table: torch.Tensor,
     token_to_req: torch.Tensor,
     out: torch.Tensor | None = None,
+    k_scale: torch.Tensor | None = None,
+    v_scale: torch.Tensor | None = None,
+    kv_quant: str | None = None,
+    k_block_scale: torch.Tensor | None = None,
+    v_block_scale: torch.Tensor | None = None,
 ) -> torch.Tensor:
-    """Run sparse GQA directly over paged BF16 K/V caches."""
+    """Run sparse GQA over bf16, FP8, or packed NVFP4 K/V caches."""
 
     if q.ndim != 3 or k_cache.ndim != 4 or v_cache.shape != k_cache.shape:
         raise ValueError("QSA sparse attention received invalid Q/K/V shapes")
@@ -243,11 +332,52 @@ def qsa_sparse_paged_attention(
         raise ValueError("QSA sparse attention metadata has invalid shapes")
     if logical_indices.shape[1] <= 0:
         raise ValueError("QSA sparse attention requires a positive selection width")
-    if q.shape[2] != k_cache.shape[3] or q.shape[1] % k_cache.shape[2]:
+    kv_quant = kv_quant if kv_quant is not None else ("fp8" if k_scale is not None else "none")
+    if kv_quant not in ("none", "fp8", "nvfp4"):
+        raise ValueError(f"unknown QSA kv_quant {kv_quant!r}")
+    stored_dim = q.shape[2] // 2 if kv_quant == "nvfp4" else q.shape[2]
+    if (
+        (kv_quant == "nvfp4" and q.shape[2] % 16)
+        or k_cache.shape[3] != stored_dim
+        or q.shape[1] % k_cache.shape[2]
+    ):
         raise ValueError("QSA sparse attention requires valid grouped-query heads")
     head_dim = q.shape[2]
     assert head_dim >= 16 and (head_dim & (head_dim - 1)) == 0
-    assert q.dtype == k_cache.dtype == v_cache.dtype
+    if (k_scale is None) != (v_scale is None) or (k_block_scale is None) != (v_block_scale is None):
+        raise ValueError("QSA sparse attention requires both KV scale tensors")
+    nvfp4 = kv_quant == "nvfp4"
+    if nvfp4 and (k_scale is None or k_block_scale is None):
+        raise ValueError("QSA NVFP4 requires row and block scale tensors")
+    if not nvfp4 and (k_block_scale is not None or v_block_scale is not None):
+        raise ValueError("QSA block scales require kv_quant='nvfp4'")
+    if k_scale is not None:
+        # The pool hands out 1-byte e4m3 codes plus one fp32 row scale per
+        # (slot, kv_head); the kernel rebuilds that slot as
+        # page * PAGE_SIZE + page_offset, which is exact only because QSA pins
+        # page_size to PAGE_SIZE (page_sizes=(64,) in attention/__init__.py).
+        if k_cache.element_size() != 1 or v_cache.element_size() != 1:
+            raise ValueError("QSA KV scales require 1-byte e4m3 code caches")
+        if k_scale.dtype is not torch.float32 or v_scale.dtype is not torch.float32:
+            raise ValueError("QSA KV scale tensors must be float32")
+        want = (k_cache.shape[0] * k_cache.shape[1], k_cache.shape[2])
+        if k_scale.shape != want or v_scale.shape != want:
+            raise ValueError(f"QSA KV scale tensors must have shape {want}")
+        assert k_scale.stride(1) == v_scale.stride(1) == 1
+        if nvfp4:
+            block_want = (*want, head_dim // 16)
+            if (
+                k_cache.dtype is not torch.uint8
+                or v_cache.dtype is not torch.uint8
+                or k_block_scale.dtype is not torch.uint8
+                or v_block_scale.dtype is not torch.uint8
+                or k_block_scale.shape != block_want
+                or v_block_scale.shape != block_want
+            ):
+                raise ValueError(f"QSA NVFP4 block scales must have shape {block_want}")
+            assert k_block_scale.stride(2) == v_block_scale.stride(2) == 1
+    else:
+        assert q.dtype == k_cache.dtype == v_cache.dtype
     assert logical_indices.dtype == block_table.dtype == torch.int32
     assert token_to_req.dtype == torch.int32
     assert q.stride(2) == k_cache.stride(3) == v_cache.stride(3) == 1
@@ -303,6 +433,12 @@ def qsa_sparse_paged_attention(
         q,
         k_cache,
         v_cache,
+        # Never dereferenced while HAS_KV_SCALE is False -- pass the caches so the
+        # launch stays type-valid without a second None-handling path.
+        k_cache if k_scale is None else k_scale,
+        v_cache if v_scale is None else v_scale,
+        k_cache if k_block_scale is None else k_block_scale,
+        v_cache if v_block_scale is None else v_block_scale,
         logical_indices,
         block_table,
         token_to_req,
@@ -317,6 +453,8 @@ def qsa_sparse_paged_attention(
         v_cache.stride(0),
         v_cache.stride(1),
         v_cache.stride(2),
+        0 if k_scale is None else k_scale.stride(0),
+        0 if v_scale is None else v_scale.stride(0),
         logical_indices.stride(0),
         block_table.stride(0),
         out.stride(0),
@@ -334,6 +472,8 @@ def qsa_sparse_paged_attention(
         NUM_TILES=num_tiles,
         BLOCK_M=block_m,
         BLOCK_N=block_n,
+        HAS_KV_SCALE=k_scale is not None,
+        KV_NVFP4=nvfp4,
         num_warps=partial_warps,
         num_stages=2,
     )
diff --git a/python/freetoken/kernel/triton/qsa/score.py b/python/freetoken/kernel/triton/qsa/score.py
index 49d702082..7b65da8d6 100644
--- a/python/freetoken/kernel/triton/qsa/score.py
+++ b/python/freetoken/kernel/triton/qsa/score.py
@@ -139,6 +139,14 @@ def qsa_mqa_paged(
         raise ValueError("QSA request mapping and positions must match query rows")
     if sequence_lengths.shape != (page_table.shape[0],):
         raise ValueError("QSA sequence lengths must match page-table requests")
+    if q.dtype not in (torch.bfloat16, torch.float16) or k_cache.dtype is not q.dtype:
+        # The dot below is a plain 16-bit matmul: an fp8 operand is not a slow path, it
+        # is a triton compile error that surfaces mid-CUDA-graph-capture. The KV cache
+        # may well be e4m3 codes (--kv-cache-dtype fp8) -- what must never reach here
+        # are those codes; the compressed index keys are their own 16-bit tier.
+        raise ValueError(
+            f"QSA scoring is 16-bit only, got query={q.dtype} keys={k_cache.dtype}"
+        )
     score_divisor = math.sqrt(q.shape[2]) if score_scale is None else score_scale
     columns = logits.shape[1]
     if not q.shape[0] or not columns:
diff --git a/python/freetoken/kvcache/__init__.py b/python/freetoken/kvcache/__init__.py
index 1792da2b3..3cc670d60 100644
--- a/python/freetoken/kvcache/__init__.py
+++ b/python/freetoken/kvcache/__init__.py
@@ -31,6 +31,11 @@ def resolve_pool_class(model_config: ModelConfig) -> type[BaseKVCachePool]:
     cover duck-typed test configs that don't implement it."""
     from freetoken.attention import AttnType
 
+    if getattr(model_config, "dsv41_args", None) is not None:
+        from .dsv41_paged_pool import DSV41PagedKVCache
+
+        return DSV41PagedKVCache
+
     specs_fn = getattr(model_config, "kv_cache_group_specs", None)
     if specs_fn is None:
         if getattr(model_config, "dsv4_args", None) is not None:
@@ -42,6 +47,10 @@ def resolve_pool_class(model_config: ModelConfig) -> type[BaseKVCachePool]:
         return MHAKVCache
     specs = list(specs_fn())
     types = {spec.attn_type for spec in specs}
+    if AttnType.DSV41 in types:
+        from .dsv41_paged_pool import DSV41PagedKVCache
+
+        return DSV41PagedKVCache
     if AttnType.DSV4 in types:
         from .dsv4_paged_pool import DSV4PagedKVCache
 
@@ -76,6 +85,17 @@ def resolve_pool_class(model_config: ModelConfig) -> type[BaseKVCachePool]:
     return MHAKVCache
 
 
+def _reject_unsupported_quant(pool: str, kv_quant: str) -> None:
+    """A pool family that has no fp8 store/scale-read path must say so at startup,
+    not silently serve a 16-bit cache the budget priced for an fp8 one."""
+    if kv_quant != "none":
+        raise ValueError(
+            f"--kv-cache-dtype {kv_quant} is not implemented for the {pool} KV pool "
+            "(only the plain paged / hybrid-SWA pools, served by the triton attention "
+            "backend); use --kv-cache-dtype bf16."
+        )
+
+
 def create_kv_pool(config, num_pages: int, device: torch.device, dtype: torch.dtype):
     """Build the engine's KV pool for ``num_pages`` USABLE pages (the dummy page and every
     secondary tier -- window pool, index slab, state rings -- are derived here or inside
@@ -85,10 +105,26 @@ def create_kv_pool(config, num_pages: int, device: torch.device, dtype: torch.dt
     from .dsv4_paged_pool import DSV4PagedKVCache
 
     model_config = config.model_config
+    kv_quant = getattr(config, "kv_quant", "none")
+    if getattr(model_config, "dsv41_args", None) is not None:
+        from .dsv41_cost_model import _dsv41_pool_sizes
+        from .dsv41_paged_pool import DSV41PagedKVCache
+
+        if kv_quant not in ("none", "fp8-fp4"):
+            _reject_unsupported_quant("DeepSeek-V4.1 paged", kv_quant)
+        pool = DSV41PagedKVCache(
+            sizes=_dsv41_pool_sizes(config, num_pages + 1), args=model_config.dsv41_args,
+            device=device, dtype=dtype, P=model_config.dsv41_args.window_size,
+            n_scratch=config.max_running_req + 1,
+            kv_quant=kv_quant,
+        )
+        pool._init_paged_state(config.max_running_req, config.cache_type != "naive")
+        return pool
     if resolve_pool_class(model_config) is DSV4PagedKVCache:
         # DSV4 is driven by the generic CacheManager over the shared page table; the pool is
         # the only DSV4-specific piece (the swa_pool plug-in: window tier + cmp/idx/state
         # shadows). Sizing reads dsv4_args, never the group spec.
+        _reject_unsupported_quant("DSV4 paged", kv_quant)
         pool = DSV4PagedKVCache(
             sizes=_dsv4_pool_sizes(config, num_pages + 1),  # +1 for dummy page
             args=model_config.dsv4_args,
@@ -117,6 +153,7 @@ def create_kv_pool(config, num_pages: int, device: torch.device, dtype: torch.dt
         device=device,
         dtype=dtype,
         num_req_slots=config.max_running_req + 1,  # + 1 for the dummy request row
+        kv_quant=kv_quant,
     )
 
 
@@ -128,7 +165,22 @@ def create_kvcache_pool(
     device: torch.device,
     num_swa_tokens: int | None = None,
     num_req_slots: int | None = None,
+    kv_quant: str = "none",
 ) -> BaseKVCachePool:
+    if kv_quant == "fp8-fp4":
+        raise ValueError("--kv-cache-dtype fp8-fp4 requires the DeepSeek-V4.1 paged pool")
+    if kv_quant == "nvfp4":
+        from freetoken.attention import AttnType
+
+        if any(
+            spec.attn_type not in (AttnType.FULL, AttnType.SWA, AttnType.QSA, AttnType.MLA, AttnType.DSA)
+            or spec.head_dim % 16
+            for spec in model_config.kv_cache_group_specs()
+        ):
+            raise ValueError(
+                "--kv-cache-dtype nvfp4 requires paged FULL, hybrid-SWA, QSA, or MLA/DSA groups "
+                "with head_dim divisible by 16"
+            )
     if model_config.has_swa_attention:
         from .hybrid_swa_pool import HybridSWAKVCache
 
@@ -140,6 +192,7 @@ def create_kvcache_pool(
             num_swa_tokens=num_swa_tokens,
             device=device,
             dtype=dtype,
+            kv_quant=kv_quant,
         )
 
     from .mha_pool import MHAKVCache
@@ -167,6 +220,7 @@ def create_kvcache_pool(
     if len(kv_specs) == 1 and kv_specs[0].attn_type == _AttnType.BSA:
         from .bsa_pool import BSAKVCache
 
+        _reject_unsupported_quant("block-sparse (BSA)", kv_quant)
         spec = kv_specs[0]
         assert layer_ids is None, "hybrid-linear x BSA has no pool support yet"
         return BSAKVCache(
@@ -204,6 +258,7 @@ def create_kvcache_pool(
             index_ratio=spec.index_ratio,
             num_req_slots=num_req_slots,
             layer_ids=spec.layer_ids,
+            kv_quant=kv_quant,
             mrope=model_config.model_is_mrope,
         )
 
@@ -225,6 +280,7 @@ def create_kvcache_pool(
                 index_head_dim=spec.index_head_dim,
                 num_index_layers=spec.num_index_layers,
                 layer_ids=layer_ids,
+                kv_quant=kv_quant,
             )
             if spec.index_ratio > 1:
                 # kpool tail rings are keyed by Req.table_idx; + 1 covers the dummy request row.
@@ -244,6 +300,7 @@ def create_kvcache_pool(
             dtype=dtype,
             device=device,
             layer_ids=layer_ids,
+            kv_quant=kv_quant,
         )
 
     spec = kv_specs[0] if len(kv_specs) == 1 else None
@@ -256,6 +313,7 @@ def create_kvcache_pool(
         device=device,
         dtype=dtype,
         layer_ids=layer_ids,
+        kv_quant=kv_quant,
     )
 
 
diff --git a/python/freetoken/kvcache/base.py b/python/freetoken/kvcache/base.py
index 95669e8c8..276f53524 100644
--- a/python/freetoken/kvcache/base.py
+++ b/python/freetoken/kvcache/base.py
@@ -9,6 +9,9 @@
 
 logger = init_logger(__name__)
 
+# One fp32 scale per (token, slab, layer, kv head) rides alongside fp8 KV codes.
+FP8_KV_SCALE_BYTES = 4
+
 
 class CacheRebuildRejected(Exception):
     """A runtime cache rebuild was rejected BEFORE any destructive free (e.g. the
@@ -16,23 +19,63 @@ class CacheRebuildRejected(Exception):
     this is recoverable, unlike a failure after the free."""
 
 
+def kv_storage_bytes_per_elem(config) -> int:
+    """Storage bytes of ONE cached KV element under the configured quantization.
+
+    ``kv_quant == "fp8"`` stores e4m3 codes (1 byte) instead of the 16-bit compute
+    dtype; anything else is the compute dtype itself. Single source for the pool's
+    allocation, the budget math below, and the AOT kernel-shape table.
+    """
+    quant = getattr(config, "kv_quant", "none")
+    if quant == "fp8":
+        return 1
+    if quant != "none":
+        raise ValueError(f"unknown kv_quant {quant!r}")
+    return config.dtype.itemsize
+
+
+def kv_scale_bytes_per_token(spec, config) -> int:
+    """Sidecar bytes per token: FP32 row scales plus NVFP4 E4M3 block scales.
+    The unquantized pool has no scales.
+
+    Priced here rather than inside the pool so ``kv_cost`` and the pool's own
+    allocation can never disagree -- the same rule the 16-bit path follows."""
+    if getattr(config, "kv_quant", "none") not in ("fp8", "nvfp4"):
+        return 0
+    heads = div_even(spec.num_kv_heads, config.tp_info.size, allow_replicate=True)
+    scale_bytes = FP8_KV_SCALE_BYTES
+    if getattr(config, "kv_quant", "none") == "nvfp4":
+        scale_bytes += spec.head_dim // 16
+    return (1 if spec.mla else 2) * spec.num_layers * heads * scale_bytes
+
+
 def spec_kv_bytes_per_token(spec, config) -> int:
     """One paged-KV group's bytes per token: (1|2 slabs) x head_dim x local kv heads x dtype
-    x layers, plus the bf16 DSA index-key slab when the spec carries indexer dims. Pure
-    per-spec arithmetic -- pool families compose it over THEIR OWN groups; no family
-    branching here. (2 bytes/elem == the torch.bfloat16 dsa_pool.DSAKVCache._alloc
-    hardcodes; keep the two in lockstep if the slab dtype ever changes.)
+    x layers, plus the fp8 scale sidecar when the cache is quantized, plus the bf16 DSA
+    index-key slab when the spec carries indexer dims. Pure per-spec arithmetic -- pool
+    families compose it over THEIR OWN groups; no family branching here. (2 bytes/elem ==
+    the torch.bfloat16 dsa_pool.DSAKVCache._alloc hardcodes; keep the two in lockstep if the
+    slab dtype ever changes.)
 
     ``index_ratio`` > 1 (QSA) stores one index key per token group, not per token; that slab's
     ring and scratch rows are fixed-size and priced in QSAKVCache.kv_cost instead."""
+    if getattr(config, "kv_quant", "none") == "nvfp4":
+        if spec.head_dim % 16:
+            raise ValueError("NVFP4 KV requires head_dim divisible by 16")
+        row_bytes = spec.head_dim // 2
+    else:
+        row_bytes = spec.head_dim * kv_storage_bytes_per_elem(config)
     per_token = (
         (1 if spec.mla else 2)  # MLA latent groups store one slab (V aliases K)
-        * spec.head_dim
+        * row_bytes
         * div_even(spec.num_kv_heads, config.tp_info.size, allow_replicate=True)
-        * config.dtype.itemsize
         * spec.num_layers
     )
-    return per_token + spec.index_head_dim * spec.num_index_layers * 2 // spec.index_ratio
+    return (
+        per_token
+        + kv_scale_bytes_per_token(spec, config)
+        + spec.index_head_dim * spec.num_index_layers * 2 // spec.index_ratio
+    )
 
 
 class BaseKVCachePool(ABC):
@@ -45,6 +88,12 @@ class BaseKVCachePool(ABC):
     # model re-bound after a rebuild; the engine asks before it resizes.
     needs_rebind_on_rebuild: ClassVar[bool] = False
 
+    # KV storage quantization: "none" keeps the compute dtype, "fp8" keeps e4m3 codes plus
+    # one fp32 scale per (token, slab, layer, kv head). A pool that does not implement the
+    # quantized allocation/store/scale-view trio stays at "none"; create_kv_pool rejects a
+    # quantization the family does not implement, so nothing here is ever silently ignored.
+    kv_quant: str = "none"
+
     # ---- sizing/cost classmethods: run BEFORE the pool exists (startup budget solve,
     # --moe-cache-auto). The engine measures memory and passes bytes in; each pool family
     # implements kv_cost for ITS OWN buffers only (the engine sums families, e.g. adds
@@ -154,13 +203,44 @@ def store_kv(
         layer_id: int,
     ) -> None: ...
 
+    def k_scale(self, index: int) -> torch.Tensor | None:
+        """fp32 ``[num_slots, local_kv_heads]`` scale of ``k_cache(index)``, indexed by the
+        same slot; None when the pool is not quantized."""
+        return None
+
+    def v_scale(self, index: int) -> torch.Tensor | None:
+        """fp32 ``[num_slots, local_kv_heads]`` scale of ``v_cache(index)``; see k_scale."""
+        return None
+
+    def k_block_scale(self, index: int) -> torch.Tensor | None:
+        """NVFP4 E4M3 scales for ``k_cache(index)``, or ``None`` for other layouts."""
+        return None
+
+    def v_block_scale(self, index: int) -> torch.Tensor | None:
+        """NVFP4 E4M3 scales for ``v_cache(index)``, or ``None`` for other layouts."""
+        return None
+
     @property
     @abstractmethod
     def device(self) -> torch.device: ...
 
     @property
     @abstractmethod
-    def dtype(self) -> torch.dtype: ...
+    def dtype(self) -> torch.dtype:
+        """The pool's COMPUTE dtype: the dtype of the K/V rows a backend hands to
+        ``store_kv``, and what backends size their scratch with. A quantized (e4m3)
+        pool still answers 16-bit here -- code bytes handed to a scratch buffer end up
+        as the rhs of a ``tl.dot`` that has no fp8 path (QSA's indexer died this way at
+        graph capture). See :attr:`store_dtype` for what the buffer holds."""
+        ...
+
+    @property
+    def store_dtype(self) -> torch.dtype:
+        """Element type of the KV buffer: e4m3/uint8 codes on a quantized pool, else
+        :attr:`dtype`. Only code that touches the buffer itself needs this; attention
+        backends that cannot apply the row scales are refused a quantized pool up
+        front (``BackendInfo.supports_fp8_kv``)."""
+        return self.dtype
 
     @property
     @abstractmethod
diff --git a/python/freetoken/kvcache/cache_status.py b/python/freetoken/kvcache/cache_status.py
index 10169d651..6318089d2 100644
--- a/python/freetoken/kvcache/cache_status.py
+++ b/python/freetoken/kvcache/cache_status.py
@@ -7,7 +7,7 @@ def _supports_swa_ratio(config) -> bool:
     """Whether ``swa_full_tokens_ratio`` sizes a separate window pool for this model -- DSV4
     (always) or a radix-SWA model (Gemma). Gates the ratio in telemetry and rebuild."""
     mc = config.model_config
-    if mc.dsv4_args is not None:
+    if mc.dsv4_args is not None or getattr(mc, "dsv41_args", None) is not None:
         return True
     return mc.has_swa_attention and config.cache_type == "swa_radix"
 
@@ -129,8 +129,8 @@ def _swa() -> int:
         from .hybrid_swa_pool import _swa_pool_floor
 
         mc = config.model_config
-        if mc.dsv4_args is not None:
-            P = mc.dsv4_args.window_size
+        if mc.dsv4_args is not None or getattr(mc, "dsv41_args", None) is not None:
+            P = (getattr(mc, "dsv41_args", None) or mc.dsv4_args).window_size
             return int(_dsv4_window_floor_pages(config, P) * P)
         if not (mc.has_swa_attention and config.cache_type == "swa_radix"):
             return 0
@@ -165,8 +165,8 @@ def compute_cache_pools(engine: "Engine") -> Dict[str, int]:
             # (num_swa_pages, usable count). Same source as the scheduler's _current_cache_geometry.
             # Both 0 for models without a window pool. Lets a client denominate the swa control.
             mc = config.model_config
-            if mc.dsv4_args is not None:
-                pools["swa_page_size"] = int(mc.dsv4_args.window_size or 0)
+            if mc.dsv4_args is not None or getattr(mc, "dsv41_args", None) is not None:
+                pools["swa_page_size"] = int((getattr(mc, "dsv41_args", None) or mc.dsv4_args).window_size or 0)
                 sizes = getattr(engine.kv_cache, "sizes", None)  # usable = physical minus dummy
                 if sizes is not None:
                     pools["num_swa_pages"] = max(0, int(sizes.n_win_pages) - 1)
diff --git a/python/freetoken/kvcache/dsa_pool.py b/python/freetoken/kvcache/dsa_pool.py
index 13225cec0..6d9e2d4f3 100644
--- a/python/freetoken/kvcache/dsa_pool.py
+++ b/python/freetoken/kvcache/dsa_pool.py
@@ -43,8 +43,12 @@ def __init__(
         dtype: torch.dtype,
         device: torch.device,
         layer_ids: "tuple[int, ...] | None" = None,
+        kv_quant: str = "none",
     ) -> None:
         self._latent_dim = latent_dim
+        if kv_quant == "nvfp4" and latent_dim % 16:
+            raise ValueError("NVFP4 KV requires latent_dim divisible by 16")
+        self._stored_dim = latent_dim // 2 if kv_quant == "nvfp4" else latent_dim
         if layer_ids is None:
             self._num_layers = num_layers
             self._layer_index: dict[int, int] | None = None
@@ -54,6 +58,7 @@ def __init__(
         self._page_size = page_size
         self._dtype = dtype
         self._device = device
+        self.kv_quant = kv_quant
         self._alloc(num_pages)
 
     def _local_layer(self, layer_id: int) -> int:
@@ -61,15 +66,29 @@ def _local_layer(self, layer_id: int) -> int:
 
     def _alloc(self, num_pages: int) -> None:
         self._num_pages = num_pages
-        self._kv_buffer = torch.empty(
-            (1, self._num_layers, num_pages, self._page_size, 1, self._latent_dim),
-            device=self._device,
-            dtype=self._dtype,
-        )
+        shape = (1, self._num_layers, num_pages, self._page_size, 1, self._stored_dim)
+        self._block_scale_buffer = None
+        if self.kv_quant in ("fp8", "nvfp4"):
+            self._kv_buffer = torch.zeros(shape, device=self._device, dtype=torch.uint8)
+            self._scale_buffer = torch.zeros(
+                (self._num_layers, num_pages * self._page_size),
+                device=self._device,
+                dtype=torch.float32,
+            )
+            if self.kv_quant == "nvfp4":
+                self._block_scale_buffer = torch.zeros(
+                    (self._num_layers, num_pages * self._page_size, self._latent_dim // 16),
+                    device=self._device, dtype=torch.uint8,
+                )
+        elif self.kv_quant == "none":
+            self._kv_buffer = torch.empty(shape, device=self._device, dtype=self._dtype)
+            self._scale_buffer = None
+        else:
+            raise ValueError(f"unknown kv_quant {self.kv_quant!r}")
 
     # -- views (addressed by GLOBAL layer id; remapped when layer_ids was given) --
     def k_cache(self, layer_id: int) -> torch.Tensor:
-        """Paged latent view ``[num_pages, page_size, latent_dim]``."""
+        """Paged latent view; NVFP4 stores ``latent_dim // 2`` bytes per row."""
         return self._kv_buffer[0, self._local_layer(layer_id)].view(
             self._num_pages, self._page_size, -1
         )
@@ -79,8 +98,20 @@ def v_cache(self, layer_id: int) -> torch.Tensor:
         return self.k_cache(layer_id)
 
     def latent_rows(self, layer_id: int) -> torch.Tensor:
-        """Row-flat latent view ``[num_pages * page_size, latent_dim]``."""
-        return self._kv_buffer[0, self._local_layer(layer_id)].view(-1, self._latent_dim)
+        """Row-flat latent view ``[num_pages * page_size, stored_dim]``."""
+        return self._kv_buffer[0, self._local_layer(layer_id)].view(-1, self._stored_dim)
+
+    def latent_scale(self, layer_id: int) -> torch.Tensor | None:
+        """One FP32 scale per latent row, or ``None`` for the compute-dtype pool."""
+        if self._scale_buffer is None:
+            return None
+        return self._scale_buffer[self._local_layer(layer_id)]
+
+    def latent_block_scale(self, layer_id: int) -> torch.Tensor | None:
+        """E4M3 bytes per 16 latent elements, present only for NVFP4."""
+        if self._block_scale_buffer is None:
+            return None
+        return self._block_scale_buffer[self._local_layer(layer_id)]
 
     # -- writes -----------------------------------------------------------------
     def store_kv(
@@ -97,7 +128,24 @@ def store_kv(
         store.cu (two-width store).
         """
         rows = self.latent_rows(layer_id)
-        split = rows.shape[1] - k_rope.shape[-1]
+        split = self._latent_dim - k_rope.shape[-1]
+        assert c_kv.shape == (out_loc.numel(), split)
+        assert k_rope.shape[0] == c_kv.shape[0]
+        if self.kv_quant == "nvfp4":
+            from freetoken.kernel.triton.kv_nvfp4 import quantize_nvfp4_rows_to_cache
+
+            latent = torch.cat((c_kv, k_rope), dim=-1) if k_rope.shape[-1] else c_kv
+            quantize_nvfp4_rows_to_cache(
+                latent, out_loc, rows, self.latent_scale(layer_id),
+                self.latent_block_scale(layer_id),
+            )
+            return
+        if self.kv_quant == "fp8":
+            from freetoken.kernel.triton.kv_quant import quantize_rows_to_cache
+
+            latent = torch.cat((c_kv, k_rope), dim=-1).contiguous()
+            quantize_rows_to_cache(latent, out_loc, rows, self.latent_scale(layer_id))
+            return
         rows[out_loc, :split] = c_kv
         rows[out_loc, split:] = k_rope
 
@@ -105,6 +153,8 @@ def rebuild(self, num_pages: int) -> None:
         """In-place resize (frees the old slab first; object identity preserved --
         callers re-derive views per forward, same contract as MHAKVCache.rebuild)."""
         self._kv_buffer = None
+        self._scale_buffer = None
+        self._block_scale_buffer = None
         if self._device.type == "cuda":
             torch.cuda.synchronize(self._device)
             torch.cuda.empty_cache()
@@ -127,7 +177,13 @@ def rebuild_from_config(
 
     def unit_bytes(self) -> tuple[int, int]:
         buf = self._kv_buffer
-        return int(buf.numel() * buf.element_size()) // (self._num_pages * self._page_size), 0
+        tokens = self._num_pages * self._page_size
+        kv = int(buf.numel() * buf.element_size()) // tokens
+        if self._scale_buffer is not None:
+            kv += int(self._scale_buffer.numel() * self._scale_buffer.element_size()) // tokens
+        if self._block_scale_buffer is not None:
+            kv += self._block_scale_buffer.numel() // tokens
+        return kv, 0
 
     # -- pool properties ----------------------------------------------------------
     @property
@@ -138,6 +194,10 @@ def device(self) -> torch.device:
     def dtype(self) -> torch.dtype:
         return self._dtype
 
+    @property
+    def store_dtype(self) -> torch.dtype:
+        return self._kv_buffer.dtype
+
     @property
     def num_layers(self) -> int:
         return self._num_layers
@@ -158,12 +218,13 @@ def __init__(
         index_head_dim: int,
         num_index_layers: int,
         layer_ids: "tuple[int, ...] | None" = None,
+        kv_quant: str = "none",
     ) -> None:
         self._index_head_dim = index_head_dim
         self._num_index_layers = num_index_layers
         super().__init__(
             latent_dim, num_layers, num_pages, page_size, dtype, device,
-            layer_ids=layer_ids,
+            layer_ids=layer_ids, kv_quant=kv_quant,
         )
 
     def _index_rows(self, num_pages: int) -> int:
@@ -231,6 +292,21 @@ def _index_rows(self, num_pages: int) -> int:
         # 1/ratio shadow of every token slot + one scratch row per request slot.
         return num_pages * self._page_size // self._index_ratio + self._num_req_slots
 
+    @classmethod
+    def kv_cost(cls, config) -> tuple[int, int, int, int]:
+        per_page, fixed, page_size, reserve = super().kv_cost(config)
+        for spec in config.model_config.kv_cache_group_specs():
+            if spec.mla and spec.index_ratio > 1:
+                row_bytes = spec.index_head_dim * spec.num_index_layers * 2
+                fixed += (config.max_running_req + 1) * row_bytes * (2 * spec.index_ratio + 1)
+        return per_page, fixed, page_size, reserve
+
+    def unit_bytes(self) -> tuple[int, int]:
+        # Scratch rows and the two tail rings are fixed costs, not token capacity.
+        kv, swa = MLAKVCache.unit_bytes(self)
+        index_bytes = self._num_index_layers * self._index_head_dim * 2 // self._index_ratio
+        return kv + index_bytes, swa
+
     @property
     def cmp_scratch_base(self) -> int:
         """First scratch row (== shadow row count); request ``table_idx`` offsets it."""
diff --git a/python/freetoken/kvcache/dsv41_cost_model.py b/python/freetoken/kvcache/dsv41_cost_model.py
new file mode 100644
index 000000000..8e70f9b5b
--- /dev/null
+++ b/python/freetoken/kvcache/dsv41_cost_model.py
@@ -0,0 +1,133 @@
+"""Byte accounting for V4.1's shared compressed KV and per-layer windows."""
+
+from __future__ import annotations
+
+import math
+
+from .dsv4_cost_model import DSV4PoolSizes, _dsv4_window_floor_pages
+from .dsv41_layout import dsv41_row_bytes
+
+
+def source_layers(args) -> tuple[tuple[int | None, ...], tuple[int | None, ...]]:
+    ratios = tuple(args.compress_ratios)[:args.n_layers]
+    if len(ratios) != args.n_layers or any(r not in (0, 1, 2) for r in ratios):
+        raise ValueError("DeepSeek-V4.1 requires one compression ratio (0, 1, or 2) per layer")
+    kv_ids, index_ids = set(args.kv_source_layers), set(args.index_source_layers)
+    if not kv_ids <= index_ids or any(i < 0 or i >= args.n_layers for i in index_ids):
+        raise ValueError("DeepSeek-V4.1 KV sources must be valid index source layers")
+    kv_map, index_map = [], []
+    kv = index = None
+    for layer, ratio in enumerate(ratios):
+        if layer in kv_ids:
+            kv = layer
+        if layer in index_ids:
+            index = layer
+        if ratio:
+            if kv is None or index is None or ratios[kv] != ratio or ratios[index] != ratio:
+                raise ValueError(f"DeepSeek-V4.1 layer {layer} has no matching KV/index source")
+            kv_map.append(kv)
+            index_map.append(index)
+        else:
+            if layer in kv_ids or layer in index_ids:
+                raise ValueError("Sliding-window-only layers cannot publish compressed KV/indexes")
+            kv_map.append(None)
+            index_map.append(None)
+    return tuple(kv_map), tuple(index_map)
+
+
+def dsv41_pool_sizes(num_pages, args, swa_ratio, P=128, n_win_pages=None):
+    source_layers(args)
+    if num_pages < 1 or P <= 0 or P % 2:
+        raise ValueError("DeepSeek-V4.1 requires positive pages with an even window size")
+    if n_win_pages is None:
+        n_win_pages = math.ceil(swa_ratio * num_pages)
+    n_win_pages = min(num_pages, max(1, n_win_pages))
+    sizes = DSV4PoolSizes(P, swa_ratio, num_pages * P, n_win_pages * P, n_win_pages)
+    owners = set(args.kv_source_layers)
+    for layer, ratio in enumerate(args.compress_ratios[:args.n_layers]):
+        owns = layer in owners
+        sizes.cmp_blocks.append(sizes.full_token // ratio if owns else None)
+        sizes.idx_blocks.append(sizes.full_token // ratio if owns else None)
+        sizes.state_slots.append(n_win_pages * 2 if owns and ratio == 2 else None)
+        sizes.ring_sizes.append(2 if owns and ratio == 2 else None)
+        sizes.idx_state_slots.append(None)
+    return sizes
+
+
+def dsv41_pool_bytes(sizes, args, n_scratch=1, kv_quant="none"):
+    window_bytes, cmp_bytes, idx_bytes = dsv41_row_bytes(args, kv_quant)
+    total = sizes.n_win_slots * args.n_layers * window_bytes
+    total += (sizes.full_token + 1) * 8
+    for layer in args.kv_source_layers:
+        total += (sizes.cmp_blocks[layer] + n_scratch) * cmp_bytes
+        total += (sizes.idx_blocks[layer] + n_scratch) * idx_bytes
+        if sizes.state_slots[layer] is not None:
+            total += (sizes.state_slots[layer] + 1) * args.head_dim * 8
+    return int(total)
+
+
+def dsv41_unit_bytes(args, P=128, kv_quant="none"):
+    window_bytes, cmp_bytes, idx_bytes = dsv41_row_bytes(args, kv_quant)
+    full = 8.0
+    window = args.n_layers * window_bytes
+    for layer in args.kv_source_layers:
+        ratio = args.compress_ratios[layer]
+        full += (cmp_bytes + idx_bytes) / ratio
+        if ratio == 2:
+            window += 2 * args.head_dim * 8 / P
+    return int(math.ceil(full)), int(math.ceil(window))
+
+
+def _dsv41_pool_sizes(config, num_pages, num_swa_pages=None):
+    args = config.model_config.dsv41_args
+    P = args.window_size
+    floor = _dsv4_window_floor_pages(config, P)
+    target = num_swa_pages if num_swa_pages is not None else config.swa_num_pages_override
+    win = max(floor, int(target) + 1 if target is not None else math.ceil(config.swa_full_tokens_ratio * num_pages))
+    return dsv41_pool_sizes(num_pages, args, config.swa_full_tokens_ratio, P, win)
+
+
+def dsv41_auto_cost_model(config):
+    args = config.model_config.dsv41_args
+    window_bytes, cmp_bytes, idx_bytes = dsv41_row_bytes(args, getattr(config, "kv_quant", "none"))
+    P = args.window_size
+    floor = _dsv4_window_floor_pages(config, P)
+    full = P * 8
+    window = P * args.n_layers * window_bytes
+    fixed = 8
+    for layer in args.kv_source_layers:
+        ratio = args.compress_ratios[layer]
+        full += P // ratio * (cmp_bytes + idx_bytes)
+        fixed += (config.max_running_req + 1) * (cmp_bytes + idx_bytes)
+        if ratio == 2:
+            window += 2 * args.head_dim * 8
+            fixed += args.head_dim * 8
+    fixed += full  # The planner counts usable pages; the pool also owns one dummy page.
+    if config.swa_num_pages_override is not None:
+        fixed += max(floor, config.swa_num_pages_override + 1) * window
+        per_page = full
+    else:
+        ratio = config.swa_full_tokens_ratio
+        per_page = math.ceil(full + ratio * window)
+        # Bound both the working-set floor and ceil(ratio * physical_pages) rounding.
+        fixed += math.ceil(max(floor * (1 - ratio), ratio + 1) * window)
+    return per_page, fixed, P, floor * P
+
+
+def dsv41_solve_num_pages(config, available_bytes):
+    args = config.model_config.dsv41_args
+    kv_quant = getattr(config, "kv_quant", "none")
+    floor = _dsv4_window_floor_pages(config, args.window_size) + 1
+    def cost(pages):
+        return dsv41_pool_bytes(_dsv41_pool_sizes(config, pages), args, config.max_running_req + 1, kv_quant)
+    if cost(floor) > available_bytes:
+        raise ValueError("KV budget cannot fit the DeepSeek-V4.1 window working set")
+    full_bytes, _ = dsv41_unit_bytes(args, args.window_size, kv_quant)
+    lo, hi = floor, max(floor + 1, available_bytes // max(1, full_bytes * args.window_size) + 1)
+    while lo + 1 < hi:
+        mid = (lo + hi) // 2
+        if cost(mid) <= available_bytes:
+            lo = mid
+        else:
+            hi = mid
+    return lo - 1
diff --git a/python/freetoken/kvcache/dsv41_layout.py b/python/freetoken/kvcache/dsv41_layout.py
new file mode 100644
index 000000000..82bf1bed6
--- /dev/null
+++ b/python/freetoken/kvcache/dsv41_layout.py
@@ -0,0 +1,18 @@
+"""Storage row sizes shared by the V4.1 pool and its budget planner."""
+
+from __future__ import annotations
+
+
+def dsv41_row_bytes(args, kv_quant="none") -> tuple[int, int, int]:
+    if kv_quant not in ("none", "fp8-fp4"):
+        raise ValueError(f"Unsupported DeepSeek-V4.1 KV storage format: {kv_quant}")
+    head_dim, index_dim = args.head_dim, args.index_head_dim
+    if head_dim <= 0 or index_dim <= 0:
+        raise ValueError("DeepSeek-V4.1 KV dimensions must be positive")
+    if kv_quant == "none":
+        return head_dim * 2, head_dim * 2, index_dim * 2
+    if head_dim % 32 or index_dim % 32:
+        raise ValueError("DeepSeek-V4.1 fp8-fp4 KV dimensions must be divisible by 32")
+    return (head_dim + head_dim // 32,
+            head_dim // 2 + head_dim // 16,
+            index_dim // 2 + index_dim // 32)
diff --git a/python/freetoken/kvcache/dsv41_paged_pool.py b/python/freetoken/kvcache/dsv41_paged_pool.py
new file mode 100644
index 000000000..c56d9859a
--- /dev/null
+++ b/python/freetoken/kvcache/dsv41_paged_pool.py
@@ -0,0 +1,138 @@
+"""V4.1 source-owned compressed/index keys on the shared page-table allocator."""
+
+from __future__ import annotations
+
+import torch
+
+from .dsv4_paged_pool import CompressStateRing, DSV4PagedKVCache
+from .dsv4_cost_model import _dsv4_window_floor_pages
+from .dsv41_cost_model import (
+    _dsv41_pool_sizes, dsv41_auto_cost_model, dsv41_pool_bytes,
+    dsv41_solve_num_pages, dsv41_unit_bytes, source_layers,
+)
+from .dsv41_layout import dsv41_row_bytes
+
+
+class DSV41PagedKVCache(DSV4PagedKVCache):
+    def __init__(self, sizes, args, device, dtype=torch.bfloat16, P=128, n_scratch=1,
+                 *, kv_quant="none"):
+        dsv41_row_bytes(args, kv_quant)
+        if dtype != torch.bfloat16:
+            raise ValueError("DeepSeek-V4.1 KV pools require logical dtype bfloat16")
+        self._kv_quant = kv_quant
+        super().__init__(sizes, args, device, dtype, P, n_scratch)
+
+    @property
+    def kv_quant(self):
+        return self._kv_quant
+
+    def _alloc_buffers(self):
+        sizes, device, dtype = self.sizes, self._device, self._dtype
+        row_bytes = dsv41_row_bytes(self.args, self.kv_quant)
+        if self.kv_quant == "fp8-fp4":
+            dtype = torch.uint8
+            window_width, cmp_width, idx_width = row_bytes
+        else:
+            window_width, cmp_width, idx_width = (width // 2 for width in row_bytes)
+        self.kv_sources, self.index_sources = source_layers(self.args)
+        self.full_to_window = torch.full((sizes.full_token + 1,), -1, dtype=torch.int64, device=device)
+        if not hasattr(self, "full_loc_map"):
+            self.full_loc_map = None
+        self.window_pool = [torch.zeros(sizes.n_win_slots, window_width, device=device, dtype=dtype)
+                            for _ in range(self._n_layers)]
+        self.cmp_pool = [None] * self._n_layers
+        self.idx_pool = [None] * self._n_layers
+        self.state_ring = [None] * self._n_layers
+        self.indexer_state_ring = [None] * self._n_layers
+        self.cmp_scratch_base = [None] * self._n_layers
+        self.idx_scratch_base = [None] * self._n_layers
+        for layer in self.args.kv_source_layers:
+            count = sizes.cmp_blocks[layer]
+            self.cmp_scratch_base[layer] = count
+            self.idx_scratch_base[layer] = count
+            self.cmp_pool[layer] = torch.zeros(count + self.n_scratch, cmp_width, device=device, dtype=dtype)
+            self.idx_pool[layer] = torch.zeros(count + self.n_scratch, idx_width, device=device, dtype=dtype)
+            if self.compress_ratios[layer] == 2:
+                self.state_ring[layer] = CompressStateRing(sizes.state_slots[layer], 2, False,
+                                                          self.head_dim, device)
+
+    def _store_rows(self, pool, rows, values):
+        if pool is None:
+            raise ValueError("Only a DeepSeek-V4.1 source layer owns compressed KV rows")
+        if (values.ndim != 2 or rows.ndim != 1 or values.shape[0] != rows.numel()
+                or values.shape[1] != pool.shape[1]):
+            raise ValueError("DeepSeek-V4.1 KV row shape does not match its storage layout")
+        if self.kv_quant == "fp8-fp4":
+            if values.dtype != torch.uint8:
+                raise ValueError("DeepSeek-V4.1 fp8-fp4 KV writes require packed uint8 rows")
+        else:
+            if not values.is_floating_point():
+                raise ValueError("DeepSeek-V4.1 unquantized KV writes require floating-point rows")
+            values = values.to(self._dtype)
+        pool.index_copy_(0, rows, values)
+
+    def store_window(self, k, layer_id, window_slot):
+        self._store_rows(self.window_pool[layer_id], window_slot, k)
+
+    def store_compressed(self, kv, layer_id, cmp_slot):
+        self._store_rows(self.cmp_pool[layer_id], cmp_slot, kv)
+
+    def store_indexer(self, k, layer_id, idx_slot):
+        self._store_rows(self.idx_pool[layer_id], idx_slot, k)
+
+    def ring_size(self, layer_id):
+        ring = self.state_ring[layer_id]
+        if ring is None:
+            raise ValueError(f"Layer {layer_id} has no pending compression state")
+        return ring.ring_size
+
+    @classmethod
+    def kv_cost(cls, config):
+        return dsv41_auto_cost_model(config)
+
+    @classmethod
+    def solve_num_pages(cls, config, available_memory):
+        dsv41_row_bytes(config.model_config.dsv41_args, getattr(config, "kv_quant", "none"))
+        if config.num_page_override is None:
+            return dsv41_solve_num_pages(config, available_memory)
+        pages = config.num_page_override
+        if pages < cls.min_kv_tokens(config) // config.page_size:
+            raise ValueError("--num-pages is below the DeepSeek-V4.1 window working-set floor")
+        return pages
+
+    @classmethod
+    def min_kv_tokens(cls, config):
+        P = config.model_config.dsv41_args.window_size
+        return _dsv4_window_floor_pages(config, P) * P
+
+    def unit_bytes(self):
+        return dsv41_unit_bytes(self.args, self.P, self.kv_quant)
+
+    def _validate_rebuild_format(self, config):
+        from .base import CacheRebuildRejected
+
+        if getattr(config, "kv_quant", "none") != self.kv_quant:
+            raise CacheRebuildRejected("Changing DeepSeek-V4.1 KV storage format requires a model restart")
+
+    def rebuild_from_config(self, config, num_pages, *, num_swa_pages=None):
+        self._validate_rebuild_format(config)
+        self.rebuild(_dsv41_pool_sizes(config, num_pages + 1, num_swa_pages))
+
+    def validate_rebuild(self, config, *, num_pages, target_moe, per_expert_bytes,
+                         baseline_free, weights_bytes, current_num_pages,
+                         extra_fixed_bytes=0, extra_note="", num_swa_pages=None, **targets):
+        from freetoken.engine.cache_budget import net_cache_budget_bytes
+        from .base import CacheRebuildRejected
+
+        self._validate_rebuild_format(config)
+        if num_pages is not None and num_pages * self.P < self.min_kv_tokens(config):
+            raise CacheRebuildRejected("DeepSeek-V4.1 KV pool is below its window working-set floor")
+        sizes = self.sizes
+        if num_pages is not None or num_swa_pages is not None:
+            sizes = _dsv41_pool_sizes(config, (num_pages if num_pages is not None else current_num_pages) + 1,
+                                     num_swa_pages)
+        budget = net_cache_budget_bytes(config.memory_ratio, baseline_free, weights_bytes, extra_fixed_bytes)
+        need = target_moe * per_expert_bytes + dsv41_pool_bytes(
+            sizes, self.args, config.max_running_req + 1, self.kv_quant)
+        if need > budget:
+            raise CacheRebuildRejected(f"Requested V4.1 cache needs {need} bytes, exceeding budget {budget}")
diff --git a/python/freetoken/kvcache/hybrid_swa_pool.py b/python/freetoken/kvcache/hybrid_swa_pool.py
index 41c3880e3..bd69b2884 100644
--- a/python/freetoken/kvcache/hybrid_swa_pool.py
+++ b/python/freetoken/kvcache/hybrid_swa_pool.py
@@ -23,6 +23,62 @@ class _KVGroupStorage:
     k_buffer: torch.Tensor
     v_buffer: torch.Tensor
     storage_shape: tuple[int, int, int]
+    logical_head_dim: int
+    # (2, num_layers, num_slots, local_kv_heads) fp32, or None for an unquantized group.
+    scale_buffer: torch.Tensor | None = None
+    # (2, num_layers, num_slots, local_kv_heads, head_dim // 16) uint8, or None.
+    block_scale_buffer: torch.Tensor | None = None
+
+
+def _alloc_group_storage(
+    *,
+    num_layers: int,
+    local_kv_heads: int,
+    head_dim: int,
+    device: torch.device,
+    store_dtype: torch.dtype,
+    outer_size: int,
+    inner_size: int,
+    kv_quant: str,
+) -> _KVGroupStorage:
+    """One group's code buffer (+ fp8 scale buffer), shared by the initial allocation
+    and the in-place rebuild so the two can never drift.
+
+    A quantized buffer is zero-filled: a stale e4m3 code decodes to a real number, so
+    an unwritten slot (the dummy page, a padded request row) would poison attention,
+    while code 0x00 is exactly 0.0. One memset per allocation, same as
+    kvcache/bsa_pool.py. The 16-bit buffer keeps torch.empty.
+    """
+    stored_dim = head_dim // 2 if kv_quant == "nvfp4" else head_dim
+    shape = (2, num_layers, outer_size, inner_size, local_kv_heads, stored_dim)
+    if kv_quant != "none":
+        from freetoken.kernel.triton.kv_quant import alloc_codes
+
+        buffer = alloc_codes(shape, device)
+        scale = torch.zeros(
+            (2, num_layers, outer_size * inner_size, local_kv_heads),
+            device=device,
+            dtype=torch.float32,
+        )
+    else:
+        buffer = torch.empty(shape, device=device, dtype=store_dtype)
+        scale = None
+    block_scale = None
+    if kv_quant == "nvfp4":
+        block_scale = torch.zeros(
+            (2, num_layers, outer_size * inner_size, local_kv_heads, head_dim // 16),
+            device=device,
+            dtype=torch.uint8,
+        )
+    return _KVGroupStorage(
+        buffer=buffer,
+        k_buffer=buffer[0],
+        v_buffer=buffer[1],
+        storage_shape=(outer_size * inner_size, local_kv_heads, stored_dim),
+        logical_head_dim=head_dim,
+        scale_buffer=scale,
+        block_scale_buffer=block_scale,
+    )
 
 
 class HybridSWAKVCache(BaseKVCachePool):
@@ -37,14 +93,26 @@ def __init__(
         dtype: torch.dtype,
         device: torch.device,
         num_swa_tokens: int | None = None,
+        kv_quant: str = "none",
     ) -> None:
+        if kv_quant not in ("none", "fp8", "nvfp4"):
+            raise ValueError(f"unsupported hybrid-SWA kv_quant {kv_quant!r}")
         specs = {group.name: group for group in groups if group.num_layers > 0}
         if set(specs) != {"full", "swa"}:
             raise ValueError(f"HybridSWAKVCache requires full and swa groups, got {sorted(specs)}")
+        if kv_quant == "nvfp4" and any(spec.head_dim % 16 for spec in specs.values()):
+            raise ValueError("NVFP4 KV requires head_dim divisible by 16")
+
+        from .mha_pool import _kv_store_dtype
 
         self._num_layers = num_layers
         self._device = device
-        self._dtype = dtype
+        self.kv_quant = kv_quant
+        self._compute_dtype = dtype
+        # What the BUFFER holds -- fp8/uint8 codes when quantized -- reported as
+        # store_dtype. The dtype property keeps answering the compute dtype, same
+        # contract as MHAKVCache (kvcache/base.py): backends size their scratch with it.
+        self._store_dtype = _kv_store_dtype(dtype, kv_quant)
         self._full_num_tokens = num_full_pages * page_size
         self._swa_num_tokens = num_swa_tokens if num_swa_tokens is not None else self._full_num_tokens
         self._page_size = page_size
@@ -63,6 +131,7 @@ def __init__(
             inner_size=page_size,
             dtype=dtype,
             device=device,
+            kv_quant=kv_quant,
         )
         self.swa_kv_pool = self._allocate_group(
             specs["swa"],
@@ -71,6 +140,7 @@ def __init__(
             inner_size=1,
             dtype=dtype,
             device=device,
+            kv_quant=kv_quant,
         )
         self._storages = {
             "full": self.full_kv_pool,
@@ -88,18 +158,20 @@ def _allocate_group(
         inner_size: int,
         dtype: torch.dtype,
         device: torch.device,
+        kv_quant: str = "none",
     ) -> _KVGroupStorage:
+        from .mha_pool import _kv_store_dtype
+
         local_kv_heads = div_even(spec.num_kv_heads, tp_size, allow_replicate=True)
-        buffer = torch.empty(
-            (2, spec.num_layers, outer_size, inner_size, local_kv_heads, spec.head_dim),
+        return _alloc_group_storage(
+            num_layers=spec.num_layers,
+            local_kv_heads=local_kv_heads,
+            head_dim=spec.head_dim,
             device=device,
-            dtype=dtype,
-        )
-        return _KVGroupStorage(
-            buffer=buffer,
-            k_buffer=buffer[0],
-            v_buffer=buffer[1],
-            storage_shape=(outer_size * inner_size, local_kv_heads, spec.head_dim),
+            store_dtype=_kv_store_dtype(dtype, kv_quant),
+            outer_size=outer_size,
+            inner_size=inner_size,
+            kv_quant=kv_quant,
         )
 
     @staticmethod
@@ -200,6 +272,26 @@ def v_cache(self, index: int) -> torch.Tensor:
         ref = self.layers_mapping[index]
         return self._storages[ref.group].v_buffer[ref.index]
 
+    def k_scale(self, index: int) -> torch.Tensor | None:
+        ref = self.layers_mapping[index]
+        scale = self._storages[ref.group].scale_buffer
+        return None if scale is None else scale[0][ref.index]
+
+    def v_scale(self, index: int) -> torch.Tensor | None:
+        ref = self.layers_mapping[index]
+        scale = self._storages[ref.group].scale_buffer
+        return None if scale is None else scale[1][ref.index]
+
+    def k_block_scale(self, index: int) -> torch.Tensor | None:
+        ref = self.layers_mapping[index]
+        scale = self._storages[ref.group].block_scale_buffer
+        return None if scale is None else scale[0][ref.index]
+
+    def v_block_scale(self, index: int) -> torch.Tensor | None:
+        ref = self.layers_mapping[index]
+        scale = self._storages[ref.group].block_scale_buffer
+        return None if scale is None else scale[1][ref.index]
+
     def store_kv(
         self,
         k: torch.Tensor,
@@ -207,13 +299,41 @@ def store_kv(
         out_loc: torch.Tensor,
         layer_id: int,
     ) -> None:
-        from freetoken.kernel import store_cache
-
         ref = self.layers_mapping[layer_id]
         storage = self._storages[ref.group]
         indices = out_loc
         if ref.group == "swa":
             indices = self.translate_loc_from_full_to_swa(out_loc)
+        if self.kv_quant == "nvfp4":
+            from freetoken.kernel.triton.kv_nvfp4 import quantize_nvfp4_to_cache
+
+            quantize_nvfp4_to_cache(
+                k=k,
+                v=v,
+                out_loc=indices,
+                k_cache=storage.k_buffer[ref.index].view(storage.storage_shape),
+                v_cache=storage.v_buffer[ref.index].view(storage.storage_shape),
+                k_scale=storage.scale_buffer[0][ref.index],
+                v_scale=storage.scale_buffer[1][ref.index],
+                k_block_scale=storage.block_scale_buffer[0][ref.index],
+                v_block_scale=storage.block_scale_buffer[1][ref.index],
+            )
+            return
+        if self.kv_quant == "fp8":
+            from freetoken.kernel.triton.kv_quant import quantize_kv_to_cache
+
+            quantize_kv_to_cache(
+                k=k,
+                v=v,
+                out_loc=indices,
+                k_cache=storage.k_buffer[ref.index].view(storage.storage_shape),
+                v_cache=storage.v_buffer[ref.index].view(storage.storage_shape),
+                k_scale=storage.scale_buffer[0][ref.index],
+                v_scale=storage.scale_buffer[1][ref.index],
+            )
+            return
+        from freetoken.kernel import store_cache
+
         store_cache(
             k_cache=storage.k_buffer[ref.index].view(storage.storage_shape),
             v_cache=storage.v_buffer[ref.index].view(storage.storage_shape),
@@ -236,7 +356,12 @@ def device(self) -> torch.device:
 
     @property
     def dtype(self) -> torch.dtype:
-        return self._dtype
+        return self._compute_dtype
+
+    @property
+    def store_dtype(self) -> torch.dtype:
+        """Element type of the KV buffers: e4m3/uint8 codes when quantized."""
+        return self._store_dtype
 
     @property
     def num_layers(self) -> int:
@@ -245,24 +370,31 @@ def num_layers(self) -> int:
     @staticmethod
     def _group_geometry(group: _KVGroupStorage) -> tuple:
         # Everything the realloc needs that does NOT pin the old buffer alive: layer count,
-        # kv heads, head_dim, device, dtype. (Plain ints + device/dtype handles, no tensor.)
-        _, num_layers, _old_outer, _old_inner, local_kv_heads, head_dim = group.buffer.shape
-        return (num_layers, local_kv_heads, head_dim, group.buffer.device, group.buffer.dtype)
+        # kv heads, logical head_dim, device, storage dtype, and quantized sidecars.
+        # (Plain ints + device/dtype handles, no tensor.)
+        _, num_layers, _old_outer, _old_inner, local_kv_heads, _stored_dim = group.buffer.shape
+        return (
+            num_layers,
+            local_kv_heads,
+            group.logical_head_dim,
+            group.buffer.device,
+            group.buffer.dtype,
+            "nvfp4" if group.block_scale_buffer is not None else ("fp8" if group.scale_buffer is not None else "none"),
+        )
 
     @staticmethod
     def _alloc_group(geom: tuple, outer_size: int, inner_size: int) -> _KVGroupStorage:
         # Only the outer (page/token) dimension changes; the rest comes from ``geom``.
-        num_layers, local_kv_heads, head_dim, device, dtype = geom
-        buffer = torch.empty(
-            (2, num_layers, outer_size, inner_size, local_kv_heads, head_dim),
+        num_layers, local_kv_heads, head_dim, device, store_dtype, kv_quant = geom
+        return _alloc_group_storage(
+            num_layers=num_layers,
+            local_kv_heads=local_kv_heads,
+            head_dim=head_dim,
             device=device,
-            dtype=dtype,
-        )
-        return _KVGroupStorage(
-            buffer=buffer,
-            k_buffer=buffer[0],
-            v_buffer=buffer[1],
-            storage_shape=(outer_size * inner_size, local_kv_heads, head_dim),
+            store_dtype=store_dtype,
+            outer_size=outer_size,
+            inner_size=inner_size,
+            kv_quant=kv_quant,
         )
 
     def rebuild(self, num_full_pages: int, num_swa_tokens: int | None = None) -> None:
@@ -346,10 +478,22 @@ def unit_bytes(self) -> tuple[int, int]:
         full = self.full_kv_pool.buffer
         swa = self.swa_kv_pool.buffer
         full_tokens = int(full.shape[2]) * int(full.shape[3])
-        return (
-            int(full.numel() * full.element_size()) // full_tokens,
-            int(swa.numel() * swa.element_size()) // self._swa_num_tokens,
-        )
+        kv = int(full.numel() * full.element_size()) // full_tokens
+        swa_b = int(swa.numel() * swa.element_size()) // self._swa_num_tokens
+        # fp8 codes are priced with their scale sidecar, matching kv_cost exactly.
+        if self.full_kv_pool.scale_buffer is not None:
+            fs = self.full_kv_pool.scale_buffer
+            kv += int(fs.numel() * fs.element_size()) // full_tokens
+        if self.swa_kv_pool.scale_buffer is not None:
+            ss = self.swa_kv_pool.scale_buffer
+            swa_b += int(ss.numel() * ss.element_size()) // self._swa_num_tokens
+        if self.full_kv_pool.block_scale_buffer is not None:
+            fs = self.full_kv_pool.block_scale_buffer
+            kv += int(fs.numel() * fs.element_size()) // full_tokens
+        if self.swa_kv_pool.block_scale_buffer is not None:
+            ss = self.swa_kv_pool.block_scale_buffer
+            swa_b += int(ss.numel() * ss.element_size()) // self._swa_num_tokens
+        return kv, swa_b
 
 
 # ---- SWA pool sizing (pure arithmetic; the pool family's geometry formulas) ----
diff --git a/python/freetoken/kvcache/mha_pool.py b/python/freetoken/kvcache/mha_pool.py
index 8ed280b96..f198887a7 100644
--- a/python/freetoken/kvcache/mha_pool.py
+++ b/python/freetoken/kvcache/mha_pool.py
@@ -9,6 +9,19 @@
 from .base import BaseKVCachePool
 
 
+def _kv_store_dtype(dtype: torch.dtype, kv_quant: str) -> torch.dtype:
+    """Storage dtype of the KV buffer for a quantization mode."""
+    if kv_quant == "none":
+        return dtype
+    if kv_quant == "nvfp4":
+        return torch.uint8
+    if kv_quant == "fp8":
+        from freetoken.kernel.triton.kv_quant import kv_codes_dtype
+
+        return kv_codes_dtype()
+    raise ValueError(f"unknown kv_quant {kv_quant!r}")
+
+
 class MHAKVCache(BaseKVCachePool):
     """
     Base class for key-value caches.
@@ -20,6 +33,16 @@ class MHAKVCache(BaseKVCachePool):
     that hold no paged KV; passing the full-attention layer ids here allocates one
     storage slab per KV layer (not per model layer) and remaps the global id to its
     dense slot, avoiding a multiple-x over-allocation of unused slabs.
+
+    ``kv_quant="fp8"`` halves the cache: rows become e4m3 codes and every
+    ``(token, slab, layer, kv head)`` row carries one fp32 scale (see
+    :mod:`freetoken.kernel.triton.kv_quant`). The codes buffer keeps the exact same
+    shape as the 16-bit one, so ``k_cache``/``v_cache`` and every index into them are
+    unchanged -- only the element type, and ``store_kv``'s write path, differ.
+
+    ``kv_quant="nvfp4"`` packs two E2M1 values per byte and adds one E4M3 scale
+    per 16 values, alongside the FP32 row scale. Logical head_dim is retained
+    separately so rebuild never mistakes the packed width for the model width.
     """
 
     def __init__(
@@ -32,10 +55,17 @@ def __init__(
         dtype: torch.dtype,
         device: torch.device,
         layer_ids: Sequence[int] | None = None,
+        kv_quant: str = "none",
     ) -> None:
         tp_info = get_tp_info()
         local_kv_heads = div_even(num_kv_heads, tp_info.size, allow_replicate=True)
+        _kv_store_dtype(dtype, kv_quant)
+        if kv_quant == "nvfp4" and head_dim % 16:
+            raise ValueError("NVFP4 KV requires head_dim divisible by 16")
+        self._head_dim = head_dim
         self._num_layers = num_layers
+        self.kv_quant = kv_quant
+        self._compute_dtype = dtype
         if layer_ids is None:
             num_storage_layers = num_layers
             self._layer_map: list[int] | None = None
@@ -47,15 +77,49 @@ def __init__(
                     raise ValueError(f"KV layer id {global_id} outside [0, {num_layers})")
                 layer_map[global_id] = dense
             self._layer_map = layer_map
-        self._kv_buffer = torch.empty(
-            (2, num_storage_layers, num_pages, page_size, local_kv_heads, head_dim),
-            device=device,
-            dtype=dtype,
-        )
+        self._device = device
+        self._alloc(num_pages, page_size, num_storage_layers, local_kv_heads, head_dim)
+
+    def _alloc(
+        self,
+        num_pages: int,
+        page_size: int,
+        num_storage_layers: int,
+        local_kv_heads: int,
+        head_dim: int,
+    ) -> None:
+        """Allocate the code buffer (and, when quantized, the scale buffer).
+
+        A quantized buffer is zero-filled -- e4m3 has NaN bit patterns, so an
+        unwritten slot (the dummy page, a padded request's row) must not read back as
+        one. The 16-bit buffer keeps ``torch.empty``: it is bytes-sized, never
+        interpreted, and the memset would cost real startup time on a large cache.
+        """
+        stored_dim = head_dim // 2 if self.kv_quant == "nvfp4" else head_dim
+        shape = (2, num_storage_layers, num_pages, page_size, local_kv_heads, stored_dim)
+        self._block_scale_buffer = None
+        if self.kv_quant in ("fp8", "nvfp4"):
+            from freetoken.kernel.triton.kv_quant import alloc_codes
+
+            self._kv_buffer = alloc_codes(shape, self._device)
+            self._scale_buffer = torch.zeros(
+                (2, num_storage_layers, num_pages * page_size, local_kv_heads),
+                device=self._device,
+                dtype=torch.float32,
+            )
+        else:
+            self._kv_buffer = torch.empty(
+                shape, device=self._device, dtype=self._compute_dtype
+            )
+            self._scale_buffer = None
+        if self.kv_quant == "nvfp4":
+            self._block_scale_buffer = torch.zeros(
+                (2, num_storage_layers, num_pages * page_size, local_kv_heads, head_dim // 16),
+                device=self._device, dtype=torch.uint8,
+            )
         self._k_buffer = self._kv_buffer[0]
         self._v_buffer = self._kv_buffer[1]
-        self._device = device
-        self._storage_shape = (num_pages * page_size, local_kv_heads, head_dim)
+        self._storage_shape = (num_pages * page_size, local_kv_heads, stored_dim)
 
     def rebuild(self, num_pages: int) -> None:
         """Reallocate the KV buffer for ``num_pages`` pages IN PLACE.
@@ -64,23 +128,18 @@ def rebuild(self, num_pages: int) -> None:
         existing buffer; only the page count changes. Views and ``_storage_shape`` are
         refreshed. Object identity is preserved so cached backend references stay valid.
         """
-        _, num_storage_layers, _old_pages, page_size, local_kv_heads, head_dim = self._kv_buffer.shape
-        dtype = self._kv_buffer.dtype
+        _, num_storage_layers, _old_pages, page_size, local_kv_heads, _stored_dim = self._kv_buffer.shape
+        head_dim = self._head_dim
         device = self._device
+        self._block_scale_buffer = None
         self._k_buffer = None
         self._v_buffer = None
         self._kv_buffer = None
+        self._scale_buffer = None
         if device.type == "cuda":
             torch.cuda.synchronize(device)
             torch.cuda.empty_cache()
-        self._kv_buffer = torch.empty(
-            (2, num_storage_layers, num_pages, page_size, local_kv_heads, head_dim),
-            device=device,
-            dtype=dtype,
-        )
-        self._k_buffer = self._kv_buffer[0]
-        self._v_buffer = self._kv_buffer[1]
-        self._storage_shape = (num_pages * page_size, local_kv_heads, head_dim)
+        self._alloc(num_pages, page_size, num_storage_layers, local_kv_heads, head_dim)
 
     @classmethod
     def kv_cost(cls, config) -> tuple[int, int, int, int]:
@@ -101,7 +160,13 @@ def rebuild_from_config(
     def unit_bytes(self) -> tuple[int, int]:
         buf = self._kv_buffer
         tokens = int(buf.shape[2]) * int(buf.shape[3])
-        return int(buf.numel() * buf.element_size()) // tokens, 0
+        kv = int(buf.numel() * buf.element_size()) // tokens
+        if self._scale_buffer is not None:
+            sc = self._scale_buffer
+            kv += int(sc.numel() * sc.element_size()) // tokens
+        if self._block_scale_buffer is not None:
+            kv += self._block_scale_buffer.numel() // tokens
+        return kv, 0
 
     def _dense(self, layer_id: int) -> int:
         if self._layer_map is None:
@@ -117,6 +182,26 @@ def k_cache(self, index: int) -> torch.Tensor:
     def v_cache(self, index: int) -> torch.Tensor:
         return self._v_buffer[self._dense(index)]
 
+    def k_scale(self, index: int) -> torch.Tensor | None:
+        if self._scale_buffer is None:
+            return None
+        return self._scale_buffer[0][self._dense(index)]
+
+    def v_scale(self, index: int) -> torch.Tensor | None:
+        if self._scale_buffer is None:
+            return None
+        return self._scale_buffer[1][self._dense(index)]
+
+    def k_block_scale(self, index: int) -> torch.Tensor | None:
+        if self._block_scale_buffer is None:
+            return None
+        return self._block_scale_buffer[0][self._dense(index)]
+
+    def v_block_scale(self, index: int) -> torch.Tensor | None:
+        if self._block_scale_buffer is None:
+            return None
+        return self._block_scale_buffer[1][self._dense(index)]
+
     def store_kv(
         self,
         k: torch.Tensor,
@@ -124,9 +209,33 @@ def store_kv(
         out_loc: torch.Tensor,
         layer_id: int,
     ) -> None:
+        dense = self._dense(layer_id)
+        if self.kv_quant == "nvfp4":
+            from freetoken.kernel.triton.kv_nvfp4 import quantize_nvfp4_to_cache
+
+            quantize_nvfp4_to_cache(
+                k, v, out_loc,
+                self._k_buffer[dense].view(self._storage_shape),
+                self._v_buffer[dense].view(self._storage_shape),
+                self.k_scale(layer_id), self.v_scale(layer_id),
+                self.k_block_scale(layer_id), self.v_block_scale(layer_id),
+            )
+            return
+        if self.kv_quant == "fp8":
+            from freetoken.kernel.triton.kv_quant import quantize_kv_to_cache
+
+            quantize_kv_to_cache(
+                k=k,
+                v=v,
+                out_loc=out_loc,
+                k_cache=self._k_buffer[dense].view(self._storage_shape),
+                v_cache=self._v_buffer[dense].view(self._storage_shape),
+                k_scale=self._scale_buffer[0][dense],
+                v_scale=self._scale_buffer[1][dense],
+            )
+            return
         from freetoken.kernel import store_cache
 
-        dense = self._dense(layer_id)
         store_cache(
             k_cache=self._k_buffer[dense].view(self._storage_shape),
             v_cache=self._v_buffer[dense].view(self._storage_shape),
@@ -141,6 +250,15 @@ def device(self) -> torch.device:
 
     @property
     def dtype(self) -> torch.dtype:
+        """The COMPUTE dtype (kvcache/base.py): what ``store_kv`` receives and what
+        backends size their scratch with -- still 16-bit on an fp8 pool."""
+        return self._compute_dtype
+
+    @property
+    def store_dtype(self) -> torch.dtype:
+        """Element type of ``_kv_buffer``: fp8/uint8 codes when quantized, else the
+        compute dtype. Reading the buffer itself needs THIS, and a backend that cannot
+        apply the row scales never gets a quantized pool (supports_fp8_kv gate)."""
         return self._kv_buffer.dtype
 
     @property
diff --git a/python/freetoken/kvcache/qsa_pool.py b/python/freetoken/kvcache/qsa_pool.py
index 7de70977d..530eefe7b 100644
--- a/python/freetoken/kvcache/qsa_pool.py
+++ b/python/freetoken/kvcache/qsa_pool.py
@@ -65,8 +65,11 @@ def __init__(
         num_req_slots: int,
         ring_capacity: int | None = None,
         layer_ids: Sequence[int] | None = None,
+        kv_quant: str = "none",
         mrope: bool = False,
     ) -> None:
+        if kv_quant not in ("none", "fp8", "nvfp4"):
+            raise ValueError(f"unsupported QSA kv_quant {kv_quant!r}")
         if index_ratio < 1 or page_size % index_ratio != 0:
             # slot // index_ratio only names one group when a group never straddles a page.
             raise ValueError(
@@ -102,6 +105,14 @@ def __init__(
             dtype=dtype,
             device=device,
             layer_ids=layer_ids,
+            # "fp8" quantizes ONLY the KV tiers: codes + one scale per (slot, kv_head)
+            # replace the bf16 K/V, and store_kv's fused writer replaces the
+            # separate qsa_store_rows. The three index tiers below stay bf16 no matter
+            # what is asked for -- block selection reads a different tensor
+            # (models/qwen3_8_flash_next.py builds index_k in the engine dtype), so
+            # --kv-cache-dtype fp8 leaves retrieval quality, and the score kernel's
+            # dtype asserts, untouched.
+            kv_quant=kv_quant,
         )
         self._zero_kv_slabs()
         self._alloc_index_tiers(num_pages)
@@ -109,8 +120,11 @@ def __init__(
     def _zero_kv_slabs(self) -> None:
         # Defense-in-depth: the attend kernels pos-mask every K/V load (the real fix for
         # torch.empty's recycled NaN/Inf bit patterns), but a zeroed slab keeps any future
-        # unmasked read finite instead of model-poisoning. One memset per (re)allocation.
-        self._kv_buffer.zero_()
+        # unmasked read finite instead of model-poisoning. One memset per (re)allocation,
+        # through the byte view so it works whether the slab holds bf16 values or e4m3
+        # codes -- a memset is the one op both representations support, and the codes
+        # live in bytes on every architecture (kv_quant.kv_codes_dtype).
+        self._kv_buffer.view(torch.uint8).zero_()
 
     def _alloc_index_tiers(self, num_pages: int) -> None:
         # ZERO-initialized: the score kernel reads whole rows of blocks unmasked and relies on
@@ -156,6 +170,10 @@ def rebuild(self, num_pages: int) -> None:
             self._kv_buffer = None
             self._k_buffer = None
             self._v_buffer = None
+            # Same reason as above on an fp8 pool: a grown K/V slab whose scales are gone
+            # would serve quantized rows at the wrong scale rather than fail.
+            self._scale_buffer = None
+            self._block_scale_buffer = None
             raise
 
     @classmethod
diff --git a/python/freetoken/layers/rotary.py b/python/freetoken/layers/rotary.py
index 0d04645d9..899ec50c1 100644
--- a/python/freetoken/layers/rotary.py
+++ b/python/freetoken/layers/rotary.py
@@ -160,6 +160,8 @@ def __init__(
     ) -> None:
         super().__init__(*args, **kwargs)
         assert self.is_neox, "mrope is defined on the NeoX half-rotation layout"
+        if self._cos_sin_cache.shape[1] != self.rotary_dim:
+            raise ValueError("mrope does not support partial proportional rotary embeddings")
         half = self.rotary_dim // 2
         assert sum(mrope_section) == half, (mrope_section, half)
         self._section_table = build_section_table(tuple(mrope_section), layout)
@@ -200,16 +202,18 @@ def _get_rope(
     base: float,
     rope_scaling: Dict[str, Any] | None = None,
     is_neox: bool = True,
+    *,
+    rotary_cls: Callable[..., RotaryEmbedding] = RotaryEmbedding,
 ) -> RotaryEmbedding:
     if rope_scaling is None:
-        return RotaryEmbedding(head_dim, rotary_dim, max_position, base, is_neox=is_neox)
+        return rotary_cls(head_dim, rotary_dim, max_position, base, is_neox=is_neox)
     # need to test some cases:
     match rope_scaling["rope_type"]:
         case "default":
-            return RotaryEmbedding(head_dim, rotary_dim, max_position, base, is_neox=is_neox)
+            return rotary_cls(head_dim, rotary_dim, max_position, base, is_neox=is_neox)
 
         case "proportional":
-            return RotaryEmbedding(
+            return rotary_cls(
                 head_dim,
                 rotary_dim,
                 max_position,
@@ -240,7 +244,7 @@ def post_process(inv_freq: torch.Tensor) -> torch.Tensor:
                 factor = (1 - smooth) / scaling_factor + smooth
                 return factor * inv_freq
 
-            return RotaryEmbedding(
+            return rotary_cls(
                 head_dim, rotary_dim, max_position, base, post_process, is_neox=is_neox
             )
 
@@ -296,7 +300,7 @@ def post_process(inv_freq: torch.Tensor) -> torch.Tensor:
                 )
                 return (inv_freq / factor) * ramp + inv_freq * (1 - ramp)
 
-            return RotaryEmbedding(
+            return rotary_cls(
                 head_dim,
                 rotary_dim,
                 max_position,
@@ -331,14 +335,13 @@ def get_rope(
     rope_map = dict(rope_scaling) if rope_scaling is not None else None
 
     def build() -> RotaryEmbedding:
+        rotary_cls = RotaryEmbedding
         if mrope_section is not None:
-            assert rope_map is None or rope_map.get("rope_type", "default") == "default"
-            return MRotaryEmbedding(
-                head_dim, rotary_dim, max_position, base,
-                is_neox=is_neox, mrope_section=tuple(mrope_section),
-                layout=mrope_layout,
+            rotary_cls = functools.partial(
+                MRotaryEmbedding, mrope_section=tuple(mrope_section), layout=mrope_layout,
             )
-        return _get_rope(head_dim, rotary_dim, max_position, base, rope_map, is_neox)
+        return _get_rope(head_dim, rotary_dim, max_position, base, rope_map, is_neox,
+                         rotary_cls=rotary_cls)
 
     t = torch.tensor([])
     if t.device == torch.device("meta"):
diff --git a/python/freetoken/message/backend.py b/python/freetoken/message/backend.py
index ccc1d7921..97527e8fd 100644
--- a/python/freetoken/message/backend.py
+++ b/python/freetoken/message/backend.py
@@ -74,6 +74,10 @@ class UserMsg(BaseBackendMsg):
     uid: int
     input_ids: torch.Tensor  # CPU 1D int32 tensor
     sampling_params: SamplingParams
+    # Optional precomputed multimodal soft-token embeddings (GPU tensor). Only used by
+    # the in-process offline path; remains None for the (serialized) online path.
+    mm_embeds: torch.Tensor | None = None
+    media: list[dict] | None = None
     # per-image processor outputs, in prompt order
     mm_items: List[MMItem] | None = None
     # precomputed [3, len(input_ids)] mrope positions and decode delta; None for text-only requests and 1-D rope models
diff --git a/python/freetoken/message/tokenizer.py b/python/freetoken/message/tokenizer.py
index 9442ee059..1549b0fc8 100644
--- a/python/freetoken/message/tokenizer.py
+++ b/python/freetoken/message/tokenizer.py
@@ -72,6 +72,7 @@ class TokenizeMsg(BaseTokenizerMsg):
     sampling_params: SamplingParams
     chat_template_kwargs: Dict[str, Any] | None = None
     tools: List[Dict[str, Any]] | None = None
+    media: List[Dict[str, Any]] | None = None
     images: List[bytes] | None = None
 
 
diff --git a/python/freetoken/message/utils.py b/python/freetoken/message/utils.py
index bf2c260e9..b84879c66 100644
--- a/python/freetoken/message/utils.py
+++ b/python/freetoken/message/utils.py
@@ -1,8 +1,8 @@
 from __future__ import annotations
 
 from typing import Any, Dict, Type
+import math
 
-import numpy as np
 import torch
 
 
@@ -11,6 +11,10 @@
 # may legitimately use our tag key as a field name. Wrapping such a dict keeps the decoder from
 # reading it as a serialized class -- without this, a request could crash the tokenizer worker.
 _RAW_DICT_KEY = "__raw_dict__"
+_TENSOR_DTYPES = {str(dtype): dtype for dtype in (
+    torch.bool, torch.uint8, torch.int8, torch.int16, torch.int32, torch.int64,
+    torch.float16, torch.bfloat16, torch.float32, torch.float64,
+)}
 
 def _serialize_any(value: Any) -> Any:
     if isinstance(value, dict):
@@ -31,16 +35,14 @@ def serialize_type(self) -> Dict:
     serialized = {}
 
     if isinstance(self, torch.Tensor):
-        assert not self.is_cuda, "wire tensors must live on CPU"
-        t = self.contiguous()
+        if self.device.type != "cpu" or str(self.dtype) not in _TENSOR_DTYPES:
+            raise ValueError("message tensors must use a supported CPU dtype")
         serialized["__type__"] = "Tensor"
-        serialized["dtype"] = str(t.dtype)
-        # 1-D tensors omit the shape so the payload matches the legacy wire format.
-        if t.dim() != 1:
-            serialized["shape"] = list(t.shape)
-        if t.dtype == torch.bfloat16:
-            t = t.view(torch.uint16)  # numpy has no bf16; ship the raw bytes
-        serialized["buffer"] = t.numpy().tobytes()
+        serialized["buffer"] = self.detach().contiguous().reshape(-1).view(torch.uint8).numpy().tobytes()
+        serialized["dtype"] = str(self.dtype)
+        # Keep the original 1-D wire format readable by older workers.
+        if self.dim() != 1:
+            serialized["shape"] = list(self.shape)
         return serialized
 
     # normal type
@@ -71,15 +73,18 @@ def deserialize_type(cls_map: Dict[str, Type], data: Dict) -> Any:
     type_name = data["__type__"]
     if type_name == "Tensor":
         buffer = data["buffer"]
-        dtype_str = data["dtype"].replace("torch.", "")
-        assert isinstance(buffer, bytes)
-        is_bf16 = dtype_str == "bfloat16"
-        np_tensor = np.frombuffer(buffer, dtype=getattr(np, "uint16" if is_bf16 else dtype_str))
-        tensor = torch.from_numpy(np_tensor.copy())
-        if is_bf16:
-            tensor = tensor.view(torch.bfloat16)
-        shape = data.get("shape")
-        return tensor if shape is None else tensor.view(shape)
+        dtype = _TENSOR_DTYPES.get(data["dtype"])
+        if dtype is None or not isinstance(buffer, bytes):
+            raise ValueError("invalid serialized tensor dtype or data")
+        itemsize = torch.empty((), dtype=dtype).element_size()
+        shape = data.get("shape", [len(buffer) // itemsize])
+        if (not isinstance(shape, (list, tuple)) or len(shape) > 8
+                or any(type(n) is not int or n < 0 for n in shape)
+                or math.prod(shape) * itemsize != len(buffer)):
+            raise ValueError("serialized tensor shape does not match its data")
+        if not buffer:
+            return torch.empty(shape, dtype=dtype)
+        return torch.frombuffer(bytearray(buffer), dtype=dtype).reshape(shape)
 
     cls = cls_map.get(type_name)
     if cls is None:
diff --git a/python/freetoken/mm/processor.py b/python/freetoken/mm/processor.py
index 48211a5b8..acc5e3087 100644
--- a/python/freetoken/mm/processor.py
+++ b/python/freetoken/mm/processor.py
@@ -200,7 +200,8 @@ def get_mm_processor(model_path: str, mm: MultimodalConfig | None = None) -> MMP
     served = [e for e in spec.encoders if getattr(config, e.config_key, None) is not None and e.kind not in mm.disabled_encoders]
     if spec.mm_processor is None or not served:
         return None
-    check_mm_pad_shift(config.text_config.vocab_size)
+    text_config = getattr(config, "text_config", None) or config
+    check_mm_pad_shift(text_config.vocab_size)
     module, _, cls = spec.mm_processor.partition(":")
     return getattr(importlib.import_module(module), cls)(config, model_path, mm)
 
diff --git a/python/freetoken/models/config.py b/python/freetoken/models/config.py
index d7b504a76..f4722dfda 100644
--- a/python/freetoken/models/config.py
+++ b/python/freetoken/models/config.py
@@ -5,7 +5,10 @@
 from freetoken.attention.base import AttnType
 
 # State-dict key prefixes of the vision stack; load_weight drops them when the engine serves text-only.
-VISION_KEY_PREFIXES = ("vision_tower.", "embed_vision.", "vision_embedder.", "visual.")
+VISION_KEY_PREFIXES = (
+    "vision_tower.", "embed_vision.", "vision_embedder.", "visual.",
+    "vision.", "aligner.", "image_",
+)
 
 
 def detect_expert_quant(hf_config: Any) -> str:
@@ -208,11 +211,22 @@ class on purpose: subclassing SWAAttentionGroupConfig would flip has_swa_attenti
     sliding_window: int  # the P-token window page
 
 
+@dataclass(frozen=True)
+class DSV41AttentionGroupConfig(BaseAttentionGroupConfig):
+    kind: ClassVar[Literal["dsv41"]] = "dsv41"
+    cache_kind: ClassVar[Literal["dsv41_paged"]] = "dsv41_paged"
+
+    num_kv_heads: int
+    head_dim: int
+    sliding_window: int
+
+
 AttentionGroupConfig: TypeAlias = (
     FullAttentionGroupConfig
     | SWAAttentionGroupConfig
     | LinearGatedDeltaGroupConfig
     | DSV4AttentionGroupConfig
+    | DSV41AttentionGroupConfig
 )
 
 
@@ -320,6 +334,7 @@ class ModelConfig:
     # CSA/HCA compressors, Lightning Indexer, manifold-constrained Hyper-Connections,
     # hash routing). Opaque to model-agnostic engine code; None for non-DSV4 models.
     dsv4_args: Any | None = None
+    dsv41_args: Any | None = None
     # GLM-5.2 (glm_moe_dsa) MLA/DSA payload (GlmMoeDsaArgs): the MLA low-rank dims and the
     # DSA indexer geometry the model module needs. Opaque to model-agnostic engine code;
     # None for every other model.
@@ -442,6 +457,8 @@ def attn_type_for_layer(self, layer_id: int) -> AttnType:
             return AttnType.SWA
         if isinstance(group, DSV4AttentionGroupConfig):
             return AttnType.DSV4
+        if isinstance(group, DSV41AttentionGroupConfig):
+            return AttnType.DSV41
         return _full_group_attn_type(group)
 
     def kv_cache_group_specs(self) -> Tuple[KVCacheGroupSpec, ...]:
@@ -484,7 +501,7 @@ def kv_cache_group_specs(self) -> Tuple[KVCacheGroupSpec, ...]:
                         attn_type=AttnType.SWA,
                     )
                 )
-            elif isinstance(group, DSV4AttentionGroupConfig):
+            elif isinstance(group, (DSV4AttentionGroupConfig, DSV41AttentionGroupConfig)):
                 # Matrix/taxonomy entry only: DSV4 sizing never reads this spec
                 # (the pool prices itself from dsv4_args), and is_swa/mla stay
                 # False so no generic spec walker treats it as SWA or MLA.
@@ -495,7 +512,7 @@ def kv_cache_group_specs(self) -> Tuple[KVCacheGroupSpec, ...]:
                         num_kv_heads=group.num_kv_heads,
                         head_dim=group.head_dim,
                         sliding_window=group.sliding_window,
-                        attn_type=AttnType.DSV4,
+                        attn_type=(AttnType.DSV41 if isinstance(group, DSV41AttentionGroupConfig) else AttnType.DSV4),
                     )
                 )
         return tuple(specs)
diff --git a/python/freetoken/models/deepseek_v41/NOTICE b/python/freetoken/models/deepseek_v41/NOTICE
new file mode 100644
index 000000000..1dc9ace5c
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/NOTICE
@@ -0,0 +1,32 @@
+DeepSeek-V4.1 model support adapts code from:
+https://huggingface.co/s-zaizen/DeepSeek-V4.1-Flash-NVFP4
+
+The prompt encoder, image processor, vision tower, Engram normalization and
+hashing, and model/attention/quantization operations adapt the checkpoint's
+encoding/encoding.py and inference reference implementation. FreeToken changes
+add native NVFP4 expert loading, demand-paged Engram tables, paged CSA2 attention,
+API transport, bounded image loading, chunked prefill and device placement.
+
+The upstream license is reproduced below.
+
+MIT License
+
+Copyright (c) 2023 DeepSeek
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.
diff --git a/python/freetoken/models/deepseek_v41/__init__.py b/python/freetoken/models/deepseek_v41/__init__.py
new file mode 100644
index 000000000..c04a0d4ea
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/__init__.py
@@ -0,0 +1,22 @@
+"""DeepSeek-V4.1 model support; media/tokenizer workers need no GPU imports."""
+
+from importlib import import_module
+
+_EXPORTS = {
+    "DeepseekV41Args": "args", "load_args": "args", "parse_config": "config",
+    "checkpoint_quant_config": "config", "DeepseekV41QuantConfig": "config",
+    "DeepseekV41ForCausalLM": "model", "iter_weights": "weight",
+    "iter_expert_pieces": "weight", "iter_vision_weights": "weight",
+    "is_expert_tensor": "weight",
+}
+
+
+def __getattr__(name):
+    if name not in _EXPORTS:
+        raise AttributeError(name)
+    value = getattr(import_module(f"{__name__}.{_EXPORTS[name]}"), name)
+    globals()[name] = value
+    return value
+
+
+__all__ = list(_EXPORTS)
diff --git a/python/freetoken/models/deepseek_v41/args.py b/python/freetoken/models/deepseek_v41/args.py
new file mode 100644
index 000000000..62642f097
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/args.py
@@ -0,0 +1,205 @@
+"""DeepSeek-V4.1 shapes, including the shared CSA2 sources and vision tower."""
+
+from __future__ import annotations
+
+import json
+import os
+from dataclasses import dataclass, fields
+from typing import Any
+
+
+def as_dict(value: Any) -> dict:
+    if value is None:
+        return {}
+    if isinstance(value, dict):
+        return value
+    if hasattr(value, "to_dict"):
+        return value.to_dict()
+    return vars(value)
+
+
+@dataclass
+class DeepseekV41Args:
+    max_batch_size: int = 1
+    max_seq_len: int = 4096
+    dtype: str = "fp8"
+    vocab_size: int = 129280
+    dim: int = 5120
+    moe_inter_dim: int = 2304
+    n_layers: int = 40
+    n_mtp_layers: int = 3
+    n_heads: int = 64
+    n_routed_experts: int = 384
+    n_shared_experts: int = 1
+    n_activated_experts: int = 6
+    score_func: str = "sqrtsoftplus"
+    gate_temp: float = 1.0
+    norm_topk_prob: bool = True
+    route_scale: float = 1.5
+    swiglu_limit: float = 10.0
+    q_lora_rank: int = 1280
+    head_dim: int = 512
+    rope_head_dim: int = 64
+    norm_eps: float = 1e-20
+    o_groups: int = 8
+    o_lora_rank: int = 1024
+    window_size: int = 128
+    compress_ratios: tuple[int, ...] = (0, 0) + (2,) * 18 + (1,) * 20 + (0,) * 3
+    kv_source_layers: tuple[int, ...] = (2, 8, 14, 20)
+    index_source_layers: tuple[int, ...] = (2, 8, 14, 20, 24, 28, 32, 36)
+    compress_rope_theta: float = 160000.0
+    original_seq_len: int = 65536
+    rope_theta: float = 10000.0
+    rope_factor: float = 16.0
+    beta_fast: int = 32
+    beta_slow: int = 1
+    index_n_heads: int = 32
+    index_head_dim: int = 128
+    index_topk: int = 512
+    candidate_source_layer: int = 20
+    candidate_topk_blocks: int = 2048
+    candidate_block_size: int = 8
+    hc_mult: int = 4
+    hc_sinkhorn_iters: int = 20
+    hc_eps: float = 1e-6
+    engram_layer_ids: tuple[int, ...] = (1, 14)
+    engram_num_embeddings: tuple[int, ...] = (384006168, 384016682)
+    engram_max_ngram_size: int = 4
+    engram_vocab_size: int = 16000000
+    engram_n_heads: int = 8
+    engram_head_dim: int = 256
+    engram_pad_id: int = 2
+    engram_compressed_vocab_size: int = 99092
+    engram_dtype: str = "fp8"
+    engram_block_size: int = 32
+    engram_scale_fmt: str = "ue8m0"
+    vision_n_layers: int = 0
+    vision_dim: int = 1024
+    vision_n_heads: int = 16
+    vision_inter_dim: int = 2816
+    vision_patch_size: int = 14
+    vision_rope_theta: float = 10000.0
+    vision_downsample_ratio: int = 3
+    vision_max_n_token: int = 1024
+    vision_min_pixels: int = 295936
+    vision_max_wh_ratio: int | None = None
+    image_token_id: int = 129264
+
+    def __post_init__(self):
+        for name in ("compress_ratios", "kv_source_layers", "index_source_layers",
+                     "engram_layer_ids", "engram_num_embeddings"):
+            setattr(self, name, tuple(getattr(self, name)))
+        if self.n_layers <= 0 or len(self.compress_ratios) < self.n_layers:
+            raise ValueError("DeepSeek-V4.1 needs a compression ratio for every backbone layer")
+        if self.n_shared_experts != 1:
+            raise ValueError("DeepSeek-V4.1 requires exactly one shared expert")
+        if not 0 < self.n_activated_experts <= self.n_routed_experts or self.gate_temp <= 0:
+            raise ValueError("Invalid DeepSeek-V4.1 expert routing configuration")
+        if self.score_func not in {"sqrtsoftplus", "softmax", "sigmoid"}:
+            raise ValueError(f"Unsupported DeepSeek-V4.1 routing score: {self.score_func}")
+        if self.hc_mult <= 0 or self.hc_mult & (self.hc_mult - 1) or self.hc_sinkhorn_iters < 1:
+            raise ValueError("DeepSeek-V4.1 requires power-of-two hc_mult and positive Sinkhorn iterations")
+        if self.n_heads % self.o_groups or not 0 <= self.rope_head_dim <= self.head_dim:
+            raise ValueError("Invalid DeepSeek-V4.1 attention geometry")
+        for name in ("kv_source_layers", "index_source_layers"):
+            sources = getattr(self, name)
+            if sources != tuple(sorted(set(sources))):
+                raise ValueError(f"{name} must contain unique increasing layer IDs")
+            if any(i < 0 or i >= self.n_layers or not self.compress_ratios[i] for i in sources):
+                raise ValueError(f"{name} must name compressed backbone layers")
+            for layer, ratio in enumerate(self.compress_ratios[:self.n_layers]):
+                if ratio not in (0, 1, 2):
+                    raise ValueError(f"Unsupported CSA2 compression ratio {ratio}")
+                source = max((i for i in sources if i <= layer), default=-1)
+                if ratio and (source < 0 or self.compress_ratios[source] != ratio):
+                    raise ValueError(f"Layer {layer} has no compatible {name} source")
+        if not set(self.kv_source_layers).issubset(self.index_source_layers):
+            raise ValueError("Every KV source must also own an indexer")
+        if self.candidate_source_layer >= 0:
+            if self.candidate_source_layer not in self.index_source_layers:
+                raise ValueError("Candidate source must own an indexer")
+            if self.candidate_block_size <= 0 or self.candidate_topk_blocks <= 0:
+                raise ValueError("Candidate selection sizes must be positive")
+        if len(self.engram_layer_ids) != len(self.engram_num_embeddings):
+            raise ValueError("Engram layers and table sizes must have equal lengths")
+        if any(i < 0 or i >= self.n_layers for i in self.engram_layer_ids):
+            raise ValueError("Engram layer must belong to the backbone")
+        if self.engram_dtype not in {"fp8", "fp4"}:
+            raise ValueError("DeepSeek-V4.1 Engram tables require FP8 or FP4 weights")
+        if self.engram_block_size != 32 or self.engram_scale_fmt != "ue8m0":
+            raise ValueError("DeepSeek-V4.1 Engram tables require block_size=32 and UE8M0 scales")
+        if self.engram_layer_ids and (self.engram_head_dim <= 0 or self.engram_head_dim % 32):
+            raise ValueError("DeepSeek-V4.1 Engram head dimension must be a positive multiple of 32")
+
+    @property
+    def nope_head_dim(self) -> int:
+        return self.head_dim - self.rope_head_dim
+
+    @property
+    def vision_enabled(self) -> bool:
+        return self.vision_n_layers > 0
+
+
+_TEXT_NAMES = {
+    "hidden_size": "dim", "moe_intermediate_size": "moe_inter_dim",
+    "num_hidden_layers": "n_layers", "num_nextn_predict_layers": "n_mtp_layers",
+    "num_attention_heads": "n_heads", "num_experts_per_tok": "n_activated_experts",
+    "scoring_func": "score_func", "routed_scaling_factor": "route_scale",
+    "qk_rope_head_dim": "rope_head_dim", "rms_norm_eps": "norm_eps",
+    "sliding_window": "window_size", "kv_source_layer_ids": "kv_source_layers",
+    "index_source_layer_ids": "index_source_layers",
+    "candidate_source_layer_id": "candidate_source_layer", "engram_pad_token_id": "engram_pad_id",
+}
+_VISION_NAMES = {
+    "num_hidden_layers": "vision_n_layers", "hidden_size": "vision_dim",
+    "num_attention_heads": "vision_n_heads", "intermediate_size": "vision_inter_dim",
+    "patch_size": "vision_patch_size", "rope_theta": "vision_rope_theta",
+    "downsample_ratio": "vision_downsample_ratio", "max_image_tokens": "vision_max_n_token",
+    "min_pixels": "vision_min_pixels", "max_wh_ratio": "vision_max_wh_ratio",
+}
+
+
+def load_args(config_or_path: Any, **overrides) -> DeepseekV41Args:
+    """Read either native inference fields or the HF nested configuration without model code."""
+    if isinstance(config_or_path, (str, os.PathLike)):
+        path = os.fspath(config_or_path)
+        if os.path.isdir(path):
+            path = os.path.join(path, "config.json")
+        if os.path.isfile(path):
+            with open(path, encoding="utf-8") as f:
+                raw = json.load(f)
+        else:
+            from freetoken.utils import cached_load_hf_config
+
+            raw = as_dict(cached_load_hf_config(path))
+    else:
+        raw = as_dict(config_or_path)
+    text = as_dict(raw.get("text_config", raw))
+    valid = {field.name for field in fields(DeepseekV41Args)}
+    kwargs = {_TEXT_NAMES.get(k, k): v for k, v in text.items()
+              if _TEXT_NAMES.get(k, k) in valid}
+    scaling = as_dict(text.get("rope_scaling"))
+    for src, dst in (("factor", "rope_factor"), ("beta_fast", "beta_fast"),
+                     ("beta_slow", "beta_slow"),
+                     ("original_max_position_embeddings", "original_seq_len")):
+        if src in scaling:
+            kwargs[dst] = scaling[src]
+    for src, dst in _VISION_NAMES.items():
+        if src in as_dict(raw.get("vision_config")):
+            kwargs[dst] = as_dict(raw["vision_config"])[src]
+    if "vision_config" in raw and raw["vision_config"] is None:
+        kwargs["vision_n_layers"] = 0
+    if "image_token_id" in raw:
+        kwargs["image_token_id"] = raw["image_token_id"]
+    quant = as_dict(raw.get("quantization_config"))
+    for name in ("engram_dtype", "engram_block_size", "engram_scale_fmt"):
+        if name in quant:
+            kwargs[name] = quant[name]
+    # HF dtype describes activations; resident FP8 storage comes from quantization_config.
+    if "text_config" in raw:
+        kwargs["dtype"] = "fp8"
+    kwargs.update(overrides)
+    return DeepseekV41Args(**kwargs)
+
+
+__all__ = ["DeepseekV41Args", "load_args"]
diff --git a/python/freetoken/models/deepseek_v41/attention.py b/python/freetoken/models/deepseek_v41/attention.py
new file mode 100644
index 000000000..ee054a253
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/attention.py
@@ -0,0 +1,255 @@
+"""CSA2 attention, compressed-source sharing, and two-stage index selection."""
+
+from __future__ import annotations
+
+import math
+
+import torch
+import torch.nn.functional as F
+from torch import nn
+
+from freetoken.core import get_global_ctx
+from freetoken.kernel.triton.dsv41.indexer import select_indices
+from freetoken.kernel.triton.dsv41.quant import fp4_roundtrip, fp8_roundtrip, pack_fp4, pack_fp8
+from .layers import Linear, RMSNorm
+
+
+def rotary_frequencies(args, compressed, device):
+    dim = args.rope_head_dim
+    base = args.compress_rope_theta if compressed else args.rope_theta
+    freq = 1.0 / base ** (torch.arange(0, dim, 2, device=device, dtype=torch.float32) / dim)
+    if compressed and args.original_seq_len > 0:
+        def correction(rotations):
+            return dim * math.log(args.original_seq_len / (rotations * 2 * math.pi)) / (2 * math.log(base))
+        low, high = max(math.floor(correction(args.beta_fast)), 0), min(math.ceil(correction(args.beta_slow)), dim - 1)
+        if low == high:
+            high += .001
+        smooth = 1 - ((torch.arange(dim // 2, device=device) - low) / (high - low)).clamp(0, 1)
+        freq = freq / args.rope_factor * (1 - smooth) + freq * smooth
+    return freq
+
+
+def apply_rope(x, positions, inv_freq, inverse=False):
+    """Rotate only the tail; derive frequencies for active tokens instead of a 1M table."""
+    dim = inv_freq.numel() * 2
+    tail = torch.view_as_complex(x[..., -dim:].float().unflatten(-1, (-1, 2)))
+    phase = positions.float()[:, None] * inv_freq[None, :]
+    frequencies = torch.polar(torch.ones_like(phase), -phase if inverse else phase)
+    frequencies = frequencies.view(positions.numel(), *([1] * (tail.ndim - 2)), -1)
+    x[..., -dim:] = torch.view_as_real(tail * frequencies).flatten(-2).to(x.dtype)
+    return x
+
+
+class Compressor(nn.Module):
+    def __init__(self, args, layer_id):
+        super().__init__()
+        self.ratio = args.compress_ratios[layer_id]
+        self.norm = RMSNorm(args.head_dim, args.norm_eps)
+        self.wkv = Linear(args.dim, args.head_dim, kind="fp32" if self.ratio == 2 else "bf16")
+        if self.ratio == 2:
+            self.wgate = Linear(args.dim, args.head_dim, kind="fp32")
+
+    def project(self, x):
+        if self.ratio == 1:
+            return self.norm(self.wkv(x)), None
+        return self.wkv(x.float()), self.wgate(x.float())
+
+    def pool(self, kv, scores, dtype):
+        return self.norm((kv * scores.softmax(-2)).sum(-2).to(dtype))
+
+
+class Indexer(nn.Module):
+    def __init__(self, args, layer_id):
+        super().__init__()
+        self.n_heads, self.head_dim = args.index_n_heads, args.index_head_dim
+        self.wq_b = Linear(args.q_lora_rank, self.n_heads * self.head_dim, kind="fp8")
+        self.weights_proj = Linear(args.dim, self.n_heads, kind="bf16")
+        self.scale = (self.head_dim * self.n_heads) ** -.5
+        if layer_id in args.kv_source_layers:
+            self.wk = Linear(args.head_dim, self.head_dim, kind="bf16")
+            self.k_norm = RMSNorm(self.head_dim, args.norm_eps)
+
+    def query(self, x, qr, positions, freq):
+        q = self.wq_b(qr).unflatten(-1, (self.n_heads, self.head_dim))
+        q = fp4_roundtrip(apply_rope(q, positions, freq), block_size=32, scale_format="e8m0")
+        return q, self.weights_proj(x) * self.scale
+
+    def keys(self, latent, positions, freq, *, packed=False):
+        k = self.k_norm(self.wk(latent))
+        quantize = pack_fp4 if packed else fp4_roundtrip
+        return quantize(apply_rope(k, positions, freq), block_size=32, scale_format="e8m0")
+
+
+class Attention(nn.Module):
+    def __init__(self, layer_id, args):
+        super().__init__()
+        self.layer_id, self.args = layer_id, args
+        self.dim, self.n_heads, self.head_dim = args.dim, args.n_heads, args.head_dim
+        self.n_groups, self.o_lora_rank = args.o_groups, args.o_lora_rank
+        self.ratio = self.compress_ratio = args.compress_ratios[layer_id]
+        self.window_size = args.window_size
+        self.is_kv_source = layer_id in args.kv_source_layers
+        self.is_index_source = layer_id in args.index_source_layers
+        self.wq_a = Linear(args.dim, args.q_lora_rank, kind="fp8")
+        self.q_norm = RMSNorm(args.q_lora_rank, args.norm_eps)
+        self.wq_b = Linear(args.q_lora_rank, args.n_heads * args.head_dim, kind="fp8")
+        self.wkv = Linear(args.dim, args.head_dim, kind="fp8")
+        self.kv_norm = RMSNorm(args.head_dim, args.norm_eps)
+        self.wo_a = nn.Parameter(torch.empty(args.o_groups * args.o_lora_rank,
+                                            args.n_heads * args.head_dim // args.o_groups,
+                                            dtype=torch.bfloat16), requires_grad=False)
+        self.wo_b = Linear(args.o_groups * args.o_lora_rank, args.dim, kind="fp8")
+        self.attn_sink = nn.Parameter(torch.empty(args.n_heads, dtype=torch.float32), requires_grad=False)
+        self.softmax_scale = args.head_dim ** -.5
+        self.compressor = Compressor(args, layer_id) if self.is_kv_source else None
+        self.indexer = Indexer(args, layer_id) if self.is_index_source else None
+        self.inv_freq = None
+        self.fp8_fp4 = False
+
+    @property
+    def attn(self):
+        return get_global_ctx().attn_backend
+
+    def bind(self, pool, device):
+        self.inv_freq = rotary_frequencies(self.args, bool(self.ratio), device)
+        self.fp8_fp4 = getattr(pool, "kv_quant", "none") == "fp8-fp4"
+
+    def reset(self):
+        pass
+
+    def _project(self, x, positions):
+        qr = self.q_norm(self.wq_a(x))
+        q = apply_rope(self.wq_b(qr).unflatten(-1, (self.n_heads, self.head_dim)), positions, self.inv_freq)
+        kv = apply_rope(self.kv_norm(self.wkv(x)), positions, self.inv_freq)
+        quantize = pack_fp8 if self.fp8_fp4 else fp8_roundtrip
+        return qr, q, quantize(kv, block_size=32)
+
+    def _output(self, o, positions):
+        o = apply_rope(o, positions, self.inv_freq, inverse=True).reshape(-1, self.n_groups,
+                                                                        self.n_heads * self.head_dim // self.n_groups)
+        projected = torch.einsum("tgd,grd->tgr", o, self.wo_a.view(self.n_groups, self.o_lora_rank, -1))
+        return self.wo_b(projected.flatten(1))
+
+    def _prefill_compress(self, x, ti, start, slots):
+        compressor, pool, ratio = self.compressor, self.attn.pool, self.ratio
+        kv, score = compressor.project(x)
+        end = start + x.shape[0]
+        if ratio == 1:
+            latent = kv
+        else:
+            if start % ratio:
+                ws = self.attn.window_slots_of(ti, start - 1, start)
+                state = pool.get_state(self.layer_id, pool.state_loc(ws, 2, pool.P))
+                old_kv, old_score = state.split(self.head_dim, -1)
+                kv, score = torch.cat((old_kv, kv), 0), torch.cat((old_score, score), 0)
+            complete = kv.shape[0] // ratio * ratio
+            latent = (compressor.pool(kv[:complete].unflatten(0, (-1, ratio)),
+                                      score[:complete].unflatten(0, (-1, ratio)), x.dtype)
+                      if complete else x.new_empty((0, self.head_dim)))
+            if end % ratio:
+                state_loc = pool.state_loc(slots[-1:], 2, pool.P)
+                pool.set_state(self.layer_id, state_loc, torch.cat((kv[-1:], score[-1:]), -1))
+        group_positions = torch.arange(start // ratio * ratio, end // ratio * ratio, ratio, device=x.device)
+        rows = self.attn.compress_rows_of(ti, group_positions, ratio)
+        self._publish(latent, group_positions, rows, rows)
+
+    def _publish(self, latent, positions, cmp_rows, idx_rows):
+        if not latent.shape[0]:
+            return
+        keys = self.indexer.keys(latent, positions, self.inv_freq, packed=self.fp8_fp4)
+        self.attn.scatter_compressed(self.layer_id, "idx", idx_rows, keys)
+        quantize = pack_fp4 if self.fp8_fp4 else fp4_roundtrip
+        compressed = quantize(apply_rope(latent, positions, self.inv_freq),
+                              block_size=16, scale_format="e4m3")
+        self.attn.scatter_compressed(self.layer_id, "attn", cmp_rows, compressed)
+
+    def _select(self, x, qr, positions, mapping, key, width):
+        if not self.is_index_source:
+            return self.attn.shared_indices[key]
+        args, backend = self.args, self.attn
+        q, weights = self.indexer.query(x, qr, positions, self.inv_freq)
+        source = backend.pool.kv_sources[self.layer_id]
+        candidates = backend.shared_candidates[key] if self.layer_id > args.candidate_source_layer >= 0 else None
+        publish = self.layer_id == args.candidate_source_layer
+        indices, blocks = select_indices(
+            q, weights, backend.pool.idx_pool[source], mapping,
+            (positions + 1) // self.ratio, width, self.ratio, args.index_topk,
+            candidates=candidates, candidate_topk=args.candidate_topk_blocks if publish else 0,
+            block_size=args.candidate_block_size or 8,
+        )
+        backend.shared_indices[key] = indices
+        if publish:
+            backend.shared_candidates[key] = blocks
+        return indices
+
+    def forward_ragged(self, x, segments, flat_positions):
+        shape = x.shape
+        x = x.reshape(-1, self.dim)
+        backend = self.attn
+        if self.layer_id == 0:
+            backend.begin_forward()
+        qr, q, kv = self._project(x, flat_positions)
+        windows, compressed = [], []
+        for offset, length, ti, start in segments:
+            end = start + length
+            pos = flat_positions[offset:offset + length]
+            slots = backend.window_slots_of(ti, start, end)
+            backend.store_window(kv[offset:offset + length], self.layer_id, slots)
+            win_positions = pos[:, None] - self.window_size + 1 + torch.arange(self.window_size, device=x.device)
+            locs = backend.pool.full_loc_map[ti, win_positions.clamp_min(0)]
+            win = backend.pool.translate_full_to_window(locs)
+            windows.append(torch.where(win_positions >= 0, win, -1))
+            if self.ratio:
+                sx, sqr = x[offset:offset + length], qr[offset:offset + length]
+                if self.is_kv_source:
+                    self._prefill_compress(sx, ti, start, slots)
+                key = (offset, length, ti, start)
+                mapping = backend.pool.full_loc_map[ti:ti + 1].expand(length, -1)
+                picks = self._select(sx, sqr, pos, mapping, key, end // self.ratio)
+                compressed.append(backend.blocks_to_global(picks[None], self.ratio, ti=ti)[0])
+        win = torch.cat(windows, 0)
+        if self.ratio:
+            width = max(part.shape[-1] for part in compressed)
+            cmp = torch.cat([F.pad(part, (0, width - part.shape[-1]), value=-1) for part in compressed], 0)
+            picks = torch.cat((win, cmp), -1)
+        else:
+            picks = win
+        o = backend.attend(q[None], self.layer_id, picks[None], self.window_size,
+                           self.attn_sink, self.softmax_scale, has_compression=bool(self.ratio))[0]
+        return self._output(o, flat_positions).view(shape)
+
+    def decode_step(self, x, pos, rows, cmp_stage_cap, wctx=None):
+        shape = x.shape
+        x = x.reshape(-1, self.dim)
+        backend, pool = self.attn, self.attn.pool
+        if self.layer_id == 0:
+            backend.begin_forward()
+        if wctx is None:
+            wctx = get_global_ctx().batch.attn_metadata.window_ctx(pos, rows)
+        slots, previous_slots, window_picks = wctx
+        qr, q, kv = self._project(x, pos)
+        backend.store_window(kv, self.layer_id, slots)
+        picks = window_picks
+        if self.ratio:
+            if self.is_kv_source:
+                latent, score = self.compressor.project(x)
+                if self.ratio == 2:
+                    old = pool.get_state(self.layer_id, pool.state_loc(previous_slots, 2, pool.P))
+                    previous_kv, previous_score = old.split(self.head_dim, -1)
+                    pooled = self.compressor.pool(torch.stack((previous_kv, latent), 1),
+                                                  torch.stack((previous_score, score), 1), x.dtype)
+                    pool.set_state(self.layer_id, pool.state_loc(slots, 2, pool.P), torch.cat((latent, score), -1))
+                    latent = pooled
+                complete = (pos + 1) % self.ratio == 0
+                cmp_rows = backend.decode_compress_rows(rows, pos, self.ratio, self.layer_id, "attn", complete)
+                idx_rows = backend.decode_compress_rows(rows, pos, self.ratio, self.layer_id, "idx", complete)
+                self._publish(latent, pos // self.ratio * self.ratio, cmp_rows, idx_rows)
+            indices = self._select(x, qr, pos, backend.snapshot()[rows], "decode", (cmp_stage_cap + 1) // self.ratio)
+            cmp = backend.blocks_to_global(indices[:, None], self.ratio, rows=rows)
+            picks = torch.cat((window_picks, cmp), -1)
+        o = backend.attend(q[:, None], self.layer_id, picks, self.window_size, self.attn_sink,
+                           self.softmax_scale, has_compression=bool(self.ratio))[:, 0]
+        return self._output(o, pos).view(shape)
+
+
+DeepseekV41Attention = Attention
diff --git a/python/freetoken/models/deepseek_v41/config.py b/python/freetoken/models/deepseek_v41/config.py
new file mode 100644
index 000000000..77813ec23
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/config.py
@@ -0,0 +1,113 @@
+"""Engine configuration for the mixed FP8/NVFP4 DeepSeek-V4.1 checkpoint."""
+
+from __future__ import annotations
+
+from freetoken.layers.quantization import QuantConfig, QuantKind, QuantScheme, WeightDesc
+from freetoken.layers.quantization.scheme import nvfp4_scheme
+from freetoken.models.config import DSV41AttentionGroupConfig, ModelConfig, RotaryConfig
+
+from .args import as_dict, load_args
+
+
+class DeepseekV41QuantConfig(QuantConfig):
+    """Native block-32 projections and W4A16 NVFP4 routed experts."""
+
+    dialect = "deepseek_v41"
+    STORAGE = {
+        QuantKind.FP8_BLOCK: {"weight": "weight", "weight_scale_inv": "scale"},
+        QuantKind.NVFP4: {"weight": "weight", "weight_scale": "weight_scale",
+                          "weight_global": "weight_scale_2"},
+    }
+
+    def __init__(self, args):
+        super().__init__()
+        self._fp8 = QuantScheme(QuantKind.FP8_BLOCK, WeightDesc("e4m3", (32, 32), "e8m0"),
+                                {"weight", "weight_scale_inv"})
+        self._nvfp4 = nvfp4_scheme(input_scale=False)
+        self._experts = {f"layers.{i}.ffn.experts" for i in range(args.n_layers)}
+        self._projections = set()
+        for i in range(args.n_layers):
+            prefix = f"layers.{i}"
+            self._projections.update(f"{prefix}.attn.{name}" for name in ("wq_a", "wq_b", "wkv", "wo_b"))
+            self._projections.update(f"{prefix}.ffn.shared_experts.{name}" for name in ("w1", "w2", "w3"))
+            if i in args.index_source_layers:
+                self._projections.add(f"{prefix}.attn.indexer.wq_b")
+            if i in args.engram_layer_ids:
+                self._projections.add(f"{prefix}.engram.wkv")
+
+    def scheme_for_name(self, name):
+        prefix, _, suffix = name.partition(".experts")
+        if prefix + ".experts" in self._experts:
+            if not suffix or (len(parts := suffix.split(".")) == 3
+                              and parts[1].isdigit() and parts[2] in {"w1", "w2", "w3"}):
+                return self._nvfp4
+        return self._fp8 if name in self._projections else None
+
+
+def checkpoint_quant_config(hf_config):
+    return DeepseekV41QuantConfig(parse_config(hf_config).dsv41_args)
+
+
+def parse_config(hf_config) -> ModelConfig:
+    raw = as_dict(hf_config)
+    args = load_args(hf_config)
+    quant = as_dict(raw.get("quantization_config"))
+    layers = as_dict(quant.get("quantized_layers"))
+    algo = str(quant.get("moe_quant_algo", "")).lower()
+    if not algo and layers:
+        algos = {str(as_dict(v).get("quant_algo", "")).lower() for k, v in layers.items()
+                 if k.startswith("layers.") and k.endswith(".ffn.experts")}
+        algo = "nvfp4" if algos == {"nvfp4"} else ""
+    if not algo and str(quant.get("expert_dtype", "")).lower() == "nvfp4":
+        algo = "nvfp4"
+    if (algo != "nvfp4" or int(quant.get("group_size", 16)) != 16
+            or int(quant.get("expert_block_size", 16)) != 16):
+        raise ValueError("DeepSeek-V4.1 requires NVFP4 routed experts with group_size=16")
+    # ModelOpt leaves expert_dtype='fp4' in its converted NVFP4 config.
+    if str(quant.get("expert_dtype", "nvfp4")).lower() not in {"nvfp4", "fp4"}:
+        raise ValueError("DeepSeek-V4.1 requires NVFP4 routed experts")
+    if str(quant.get("expert_scale_fmt", "e4m3")).lower() != "e4m3":
+        raise ValueError("DeepSeek-V4.1 NVFP4 routed experts require E4M3 block scales")
+    if quant.get("expert_global_scale", True) is not True:
+        raise ValueError("DeepSeek-V4.1 NVFP4 routed experts require a global scale")
+    for layer in range(args.n_layers):
+        entry = layers.get(f"layers.{layer}.ffn.experts")
+        if layers and (entry is None or str(as_dict(entry).get("quant_algo", "")).lower() != "nvfp4"
+                       or int(as_dict(entry).get("group_size", 16)) != 16):
+            raise ValueError(f"Missing or incompatible NVFP4 quantization for backbone layer {layer}")
+    if tuple(quant.get("weight_block_size", (32, 32))) != (32, 32):
+        raise ValueError("DeepSeek-V4.1 resident FP8 projections require 32x32 weight blocks")
+    if quant.get("scale_fmt", "ue8m0") != "ue8m0":
+        raise ValueError("DeepSeek-V4.1 resident FP8 projections require UE8M0 scales")
+    text = as_dict(raw.get("text_config", raw))
+    max_position = int(text.get("max_position_embeddings", args.original_seq_len * args.rope_factor))
+    return ModelConfig(
+        num_layers=args.n_layers, num_qo_heads=args.n_heads, num_kv_heads=1,
+        head_dim=args.head_dim, hidden_size=args.dim, vocab_size=args.vocab_size,
+        intermediate_size=args.moe_inter_dim, rms_norm_eps=args.norm_eps,
+        rotary_config=RotaryConfig(args.head_dim, args.rope_head_dim, max_position,
+                                  args.rope_theta, {
+                                      "rope_type": "yarn", "factor": args.rope_factor,
+                                      "beta_fast": args.beta_fast, "beta_slow": args.beta_slow,
+                                      "original_max_position_embeddings": args.original_seq_len,
+                                  }),
+        hidden_act="swiglu_clamp", tie_word_embeddings=False,
+        num_experts=args.n_routed_experts, num_experts_per_tok=args.n_activated_experts,
+        moe_intermediate_size=args.moe_inter_dim, norm_topk_prob=args.norm_topk_prob,
+        model_type="deepseek_v41", architectures=["DeepseekV41ForCausalLM"],
+        moe_strategy="offload", moe_enabled=True, expert_quant="nvfp4",
+        weight_block_size=(32, 32), n_shared_experts=args.n_shared_experts,
+        shared_expert_intermediate_size=args.moe_inter_dim,
+        routed_scaling_factor=args.route_scale, swiglu_limit=args.swiglu_limit,
+        hidden_act_alpha=1.0, attn_sm_scale=args.head_dim ** -0.5,
+        vision_config=raw.get("vision_config") if args.vision_enabled else None,
+        image_token_id=args.image_token_id if args.vision_enabled else None,
+        dsv41_args=args,
+        attention_groups=(DSV41AttentionGroupConfig(
+            name="dsv41", layer_ids=tuple(range(args.n_layers)), num_kv_heads=1,
+            head_dim=args.head_dim, sliding_window=args.window_size,
+        ),),
+    )
+
+
+__all__ = ["parse_config", "checkpoint_quant_config", "DeepseekV41QuantConfig"]
diff --git a/python/freetoken/models/deepseek_v41/encoding.py b/python/freetoken/models/deepseek_v41/encoding.py
new file mode 100644
index 000000000..d673cae65
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/encoding.py
@@ -0,0 +1,689 @@
+"""
+DeepSeek-V4.1 Text and Vision Encoding (upstream MIT attribution: NOTICE).
+
+A fully self-contained implementation for encoding/decoding DeepSeek-V4.1 chat
+messages with tool calling, thinking mode, quick instruction tasks, and image
+content blocks. No dependency on encoding_dsv4.
+
+V4.1 changes relative to V4:
+
+1. DSML tag names: tool calls are wrapped in "<｜DSML｜ calls>" blocks with
+   "<｜DSML｜ invoke>" / "<｜DSML｜ parameter>" tags (leading-space tag names).
+2. Numeric reasoning effort: "Reasoning Effort: {budget} (range 1-100, ...)".
+   Accepts an int in [1, 100] or one of "low"/"high"/"xhigh"/"max"
+   (mapped to 25/50/75/100). Defaults to "high". Only rendered in thinking mode.
+3. Mid-conversation system messages are supported via the "<｜System｜>" token.
+   A mid-conversation system message behaves like a user message for the purpose
+   of appending the assistant generation header.
+"""
+
+from typing import Any, Dict, List, Union, Optional, Tuple
+import copy
+import json
+import re
+
+IS_DSV41 = True
+
+# Special Tokens
+
+bos_token: str = "<｜begin▁of▁sentence｜>"
+eos_token: str = "<｜end▁of▁sentence｜>"
+thinking_start_token: str = "<think>"
+thinking_end_token: str = "</think>"
+dsml_token: str = "｜DSML｜"
+
+USER_SP_TOKEN = "<｜User｜>"
+ASSISTANT_SP_TOKEN = "<｜Assistant｜>"
+LATEST_REMINDER_SP_TOKEN = "<｜latest_reminder｜>"
+
+IMAGE_PLACEHOLDER = "<｜deepseek_image｜>"
+
+
+# Task special tokens for internal classification tasks
+DS_TASK_SP_TOKENS = {
+    "action": "<｜action｜>",
+    "query": "<｜query｜>",
+    "authority": "<｜authority｜>",
+    "domain": "<｜domain｜>",
+    "title": "<｜title｜>",
+    "read_url": "<｜read_url｜>",
+}
+VALID_TASKS = set(DS_TASK_SP_TOKENS.keys())
+
+# Templates
+
+system_msg_template: str = "{content}"
+user_msg_template: str = "{content}"
+latest_reminder_msg_template: str = "{content}"
+assistant_msg_template: str = "{reasoning}{content}{tool_calls}" + eos_token
+assistant_msg_wo_eos_template: str = "{reasoning}{content}{tool_calls}"
+thinking_template: str = "{reasoning_content}"
+
+response_format_template: str = (
+    "## Response Format:\n\nYou MUST strictly adhere to the following schema to reply:\n{schema}"
+)
+
+tool_output_template: str = (
+    "<tool_result>{content}</tool_result>"
+)
+
+# Utility Functions
+
+def to_json(value: Any) -> str:
+    """Serialize a value to JSON string."""
+    try:
+        return json.dumps(value, ensure_ascii=False)
+    except:
+        return json.dumps(value, ensure_ascii=True)
+
+
+def tools_from_openai_format(tools):
+    """Extract function definitions from OpenAI-format tool list."""
+    return [tool["function"] for tool in tools]
+
+
+def tool_calls_from_openai_format(tool_calls):
+    """Convert OpenAI-format tool calls to internal format."""
+    return [
+        {
+            "name": tool_call["function"]["name"],
+            "arguments": tool_call["function"]["arguments"],
+        }
+        for tool_call in tool_calls
+    ]
+
+
+def merge_tool_messages(messages: List[Dict[str, Any]]) -> List[Dict[str, Any]]:
+    """
+    Merge tool messages into the preceding user message using content_blocks format.
+
+    DeepSeek-V4.1 does not have a standalone "tool" role; instead, tool results
+    are encoded as <tool_result> blocks within user messages.
+    """
+    merged: List[Dict[str, Any]] = []
+
+    for msg in messages:
+        msg = copy.deepcopy(msg)
+        role = msg.get("role")
+
+        if role == "tool":
+            # Convert tool message to a user message with tool_result block
+            tool_block = {
+                "type": "tool_result",
+                "tool_use_id": msg.get("tool_call_id", ""),
+                "content": msg.get("content", ""),
+            }
+            # Merge into previous message if it's already a user (merged tool)
+            if merged and merged[-1].get("role") == "user" and "content_blocks" in merged[-1]:
+                merged[-1]["content_blocks"].append(tool_block)
+            else:
+                merged.append({
+                    "role": "user",
+                    "content_blocks": [tool_block],
+                })
+        elif role == "user":
+            content_blocks = msg.get("content_blocks")
+            if content_blocks is None:
+                content_blocks = [{"type": "text", "text": msg.get("content", "")}]
+            if merged and merged[-1].get("role") == "user" and "content_blocks" in merged[-1] and merged[-1].get("task") is None:
+                merged[-1]["content_blocks"].extend(content_blocks)
+            else:
+                # Preserve structured content and all message-level metadata.
+                new_msg = msg
+                new_msg["content_blocks"] = content_blocks
+                merged.append(new_msg)
+        else:
+            merged.append(msg)
+
+    return merged
+
+
+def sort_tool_results_by_call_order(messages: List[Dict[str, Any]]) -> List[Dict[str, Any]]:
+    """
+    Sort tool_result blocks within user messages by the order of tool_calls
+    in the preceding assistant message.
+    """
+    last_tool_call_order: Dict[str, int] = {}
+
+    for msg in messages:
+        role = msg.get("role")
+        if role == "assistant" and msg.get("tool_calls"):
+            last_tool_call_order = {}
+            for idx, tc in enumerate(msg["tool_calls"]):
+                tc_id = tc.get("id") or tc.get("function", {}).get("id", "")
+                if tc_id:
+                    last_tool_call_order[tc_id] = idx
+
+        elif role == "user" and msg.get("content_blocks"):
+            tool_blocks = [b for b in msg["content_blocks"] if b.get("type") == "tool_result"]
+            if len(tool_blocks) > 1 and last_tool_call_order:
+                sorted_blocks = sorted(
+                    tool_blocks,
+                    key=lambda b: last_tool_call_order.get(b.get("tool_use_id", ""), 0)
+                )
+                sorted_idx = 0
+                new_blocks = []
+                for block in msg["content_blocks"]:
+                    if block.get("type") == "tool_result":
+                        new_blocks.append(sorted_blocks[sorted_idx])
+                        sorted_idx += 1
+                    else:
+                        new_blocks.append(block)
+                msg["content_blocks"] = new_blocks
+
+    return messages
+
+
+# Vision Message Preprocessing
+
+def _is_image_block(block: Dict[str, Any]) -> bool:
+    """Return whether a content block is an OpenAI/Anthropic/internal image."""
+    return isinstance(block, dict) and block.get("type") in ("image", "image_url")
+
+
+def _extract_image(block: Dict[str, Any]) -> Dict[str, Any]:
+    """Normalize a supported image block into an internal image record."""
+    record: Dict[str, Any] = {"type": "image"}
+    if block.get("type") == "image_url":
+        image_url = block.get("image_url")
+        if isinstance(image_url, str):
+            record["url"] = image_url
+        else:
+            record["url"] = (image_url or {}).get("url", "")
+    else:
+        for key in ("source", "url", "data"):
+            if key in block:
+                record[key] = block[key]
+    if not any(record.get(key) for key in ("source", "url", "data")):
+        raise ValueError("Image block does not contain a valid source")
+    return record
+
+
+def _process_image_blocks(
+    blocks: List[Any], image_placeholder: str = IMAGE_PLACEHOLDER
+) -> Tuple[List[Any], List[Dict[str, Any]]]:
+    """Replace image blocks and collect their records in one ordered traversal."""
+    new_blocks: List[Any] = []
+    images: List[Dict[str, Any]] = []
+    for block in blocks:
+        if not isinstance(block, dict):
+            new_blocks.append(block)
+            continue
+        if _is_image_block(block):
+            new_blocks.append({"type": "text", "text": image_placeholder})
+            images.append(_extract_image(block))
+        elif block.get("type") == "tool_result" and isinstance(block.get("content"), list):
+            block = copy.copy(block)
+            block["content"], nested_images = _process_image_blocks(
+                block["content"], image_placeholder)
+            new_blocks.append(block)
+            images.extend(nested_images)
+        elif block.get("type") == "text":
+            text = block.get("text") or ""
+            if IMAGE_PLACEHOLDER in text:
+                raise ValueError(
+                    f"Text block contains image placeholder '{IMAGE_PLACEHOLDER}': "
+                    f"'{text[:100]}'. Images should be separate content blocks."
+                )
+            new_blocks.append(block)
+        else:
+            new_blocks.append(block)
+    return new_blocks, images
+
+
+def _validate_no_image_sp_tokens(msg: Dict[str, Any]) -> None:
+    """Reject user-supplied image placeholder tokens in textual fields."""
+    content = msg.get("content")
+    if isinstance(content, str) and IMAGE_PLACEHOLDER in content:
+        raise ValueError(
+            f"Message content contains image special token '{IMAGE_PLACEHOLDER}'. "
+            "Images should be provided as image content blocks."
+        )
+    reasoning_content = msg.get("reasoning_content")
+    if isinstance(reasoning_content, str) and IMAGE_PLACEHOLDER in reasoning_content:
+        raise ValueError(
+            f"reasoning_content contains image special token '{IMAGE_PLACEHOLDER}'"
+        )
+
+
+def process_image_messages(
+    messages: List[Dict[str, Any]],
+) -> Tuple[List[Dict[str, Any]], List[Dict[str, Any]]]:
+    """Normalize image blocks and return their records in prompt order."""
+    processed: List[Dict[str, Any]] = []
+    images: List[Dict[str, Any]] = []
+    for msg in messages:
+        msg = copy.deepcopy(msg)
+        _validate_no_image_sp_tokens(msg)
+
+        if isinstance(msg.get("content"), list) and "content_blocks" not in msg:
+            msg["content_blocks"] = msg.pop("content")
+
+        if msg.get("content_blocks"):
+            msg["content_blocks"], message_images = _process_image_blocks(
+                msg["content_blocks"])
+            images.extend(message_images)
+            if not isinstance(msg.get("content"), str):
+                texts = [
+                    block.get("text", "")
+                    for block in msg["content_blocks"]
+                    if isinstance(block, dict) and block.get("type") == "text"
+                ]
+                msg["content"] = "\n\n".join(texts)
+
+        processed.append(msg)
+    return processed, images
+
+
+SYSTEM_SP_TOKEN = "<｜System｜>"
+
+tool_calls_block_name: str = " calls"
+tool_call_tag_name: str = " invoke"
+tool_parameter_tag_name: str = " parameter"
+
+tool_call_template: str = (
+    "<{dsml_token}{tool_call_tag_name} name=\"{name}\">\n{arguments}\n</{dsml_token}{tool_call_tag_name}>"
+)
+tool_calls_template = (
+    "<{dsml_token}{tc_block_name}>\n{tool_calls}\n</{dsml_token}{tc_block_name}>"
+)
+
+# Reasoning Effort (numeric budget)
+
+REASONING_EFFORT_TEMPLATE = (
+    "Reasoning Effort: {budget} "
+    "(range 1-100, the higher the value, the more thorough the reasoning)\n\n"
+)
+
+REASONING_EFFORT_MAPPINGS: Dict[str, int] = {
+    "low": 25,
+    "high": 50,
+    "xhigh": 75,
+    "max": 100,
+}
+DEFAULT_REASONING_EFFORT = "high"
+
+
+def render_reasoning_effort(
+    index: int,
+    thinking_mode: str,
+    effort: Union[str, int, None],
+) -> str:
+    """Render the V4.1 numeric reasoning effort prefix (thinking mode, index 0 only)."""
+    if effort is None:
+        effort = DEFAULT_REASONING_EFFORT
+    assert (
+        type(effort) is int and 1 <= effort <= 100
+    ) or effort in REASONING_EFFORT_MAPPINGS, (
+        "Invalid reasoning effort for deepseek_v41: "
+        f"{effort}, should be int within [1,100] or {list(REASONING_EFFORT_MAPPINGS)}"
+    )
+    if type(effort) is str:
+        effort = REASONING_EFFORT_MAPPINGS[effort]
+    if index == 0 and thinking_mode == "thinking":
+        return REASONING_EFFORT_TEMPLATE.format(budget=effort)
+    return ""
+
+
+# Tools rendering
+
+TOOLS_TEMPLATE = """## Tools
+
+You have access to a set of tools to help answer the user's question. You can invoke tools by writing a "<{dsml_token}{tc_block_name}>" block like the following:
+
+<{dsml_token}{tc_block_name}>
+<{dsml_token}{tool_call_tag_name} name="$TOOL_NAME">
+<{dsml_token}{tool_parameter_tag_name} name="$PARAMETER_NAME" string="true|false">$PARAMETER_VALUE</{dsml_token}{tool_parameter_tag_name}>
+...
+</{dsml_token}{tool_call_tag_name}>
+<{dsml_token}{tool_call_tag_name} name="$TOOL_NAME2">
+...
+</{dsml_token}{tool_call_tag_name}>
+</{dsml_token}{tc_block_name}>
+
+String parameters should be specified as is and set `string="true"`. For all other types (numbers, booleans, arrays, objects), pass the value in JSON format and set `string="false"`.
+
+If thinking_mode is enabled (triggered by {thinking_start_token}), you MUST output your complete reasoning inside {thinking_start_token}...{thinking_end_token} BEFORE any tool calls or final response.
+
+Otherwise, output directly after {thinking_end_token} with tool calls or final response.
+
+### Available Tool Schemas
+
+{tool_schemas}
+
+You MUST strictly follow the above defined tool name and parameter schemas to invoke tool calls.
+"""
+
+
+def render_tools(tools: List[Dict[str, Union[str, Dict[str, Any]]]]) -> str:
+    """Render tool schemas into the V4.1 system prompt format."""
+    tools_json = [to_json(t) for t in tools]
+
+    return TOOLS_TEMPLATE.format(
+        tool_schemas="\n".join(tools_json),
+        dsml_token=dsml_token,
+        tc_block_name=tool_calls_block_name,
+        tool_call_tag_name=tool_call_tag_name,
+        tool_parameter_tag_name=tool_parameter_tag_name,
+        thinking_start_token=thinking_start_token,
+        thinking_end_token=thinking_end_token,
+    )
+
+
+def encode_arguments_to_dsml(tool_call: Dict[str, Any]) -> str:
+    """Encode tool call arguments into V4.1 DSML parameter format."""
+    p_dsml_template = (
+        '<{dsml_token}{tool_parameter_tag_name} name="{key}" string="{is_str}">'
+        '{value}</{dsml_token}{tool_parameter_tag_name}>'
+    )
+    P_dsml_strs = []
+
+    arguments = tool_call["arguments"]
+    if not isinstance(arguments, dict):
+        # Tolerate JSON strings, including double-encoded ones.
+        for _ in range(2):
+            if isinstance(arguments, str):
+                try:
+                    arguments = json.loads(arguments)
+                except Exception:
+                    break
+            else:
+                break
+        if not isinstance(arguments, dict):
+            arguments = {"arguments": tool_call["arguments"]}
+
+    for k, v in arguments.items():
+        P_dsml_strs.append(p_dsml_template.format(
+            dsml_token=dsml_token,
+            tool_parameter_tag_name=tool_parameter_tag_name,
+            key=k,
+            is_str="true" if isinstance(v, str) else "false",
+            value=v if isinstance(v, str) else to_json(v),
+        ))
+
+    return "\n".join(P_dsml_strs)
+
+
+# Message Rendering
+
+def find_last_user_index(messages: List[Dict[str, Any]]) -> int:
+    """
+    Find the index of the last user message.
+
+    V4.1 supports mid-conversation system messages, which count as user
+    messages for the purposes of the assistant generation header.
+    """
+    last_user_index = -1
+    for idx in range(len(messages) - 1, -1, -1):
+        role = messages[idx].get("role")
+        if role == "user" or (role == "system" and idx > 0):
+            last_user_index = idx
+            break
+    return last_user_index
+
+
+def render_message(
+    index: int,
+    messages: List[Dict[str, Any]],
+    thinking_mode: str,
+    drop_thinking: bool = True,
+    reasoning_effort: Union[str, int, None] = None,
+) -> str:
+    """
+    Render a single message at the given index into its V4.1 encoded string form.
+    """
+    assert 0 <= index < len(messages)
+    assert thinking_mode in ["chat", "thinking"], f"Invalid thinking_mode `{thinking_mode}`"
+
+    msg = messages[index]
+    last_user_idx = find_last_user_index(messages)
+
+    role = msg.get("role")
+    content = msg.get("content")
+    tools = msg.get("tools")
+    response_format = msg.get("response_format")
+    tool_calls = msg.get("tool_calls")
+    reasoning_content = msg.get("reasoning_content")
+    wo_eos = msg.get("wo_eos", False)
+
+    if tools:
+        tools = tools_from_openai_format(tools)
+    if tool_calls:
+        tool_calls = tool_calls_from_openai_format(tool_calls)
+
+    # Reasoning effort prefix (thinking mode, index 0 only)
+    reasoning_effort_prompt = render_reasoning_effort(index, thinking_mode, reasoning_effort)
+    # System token leads the conversation when there is a reasoning effort prompt
+    # or the first message is a system message.
+    prompt = SYSTEM_SP_TOKEN if index == 0 and (reasoning_effort_prompt or role == "system") else ""
+    prompt += reasoning_effort_prompt
+
+    if role == "system":
+        if index > 0:
+            # Mid-conversation system message
+            prompt += SYSTEM_SP_TOKEN
+        prompt += system_msg_template.format(content=content or "")
+        if tools:
+            prompt += "\n\n" + render_tools(tools)
+        if response_format:
+            prompt += "\n\n" + response_format_template.format(schema=to_json(response_format))
+
+    elif role == "user":
+        prompt += USER_SP_TOKEN
+
+        # Handle content blocks (tool results mixed with text)
+        content_blocks = msg.get("content_blocks")
+        if content_blocks:
+            parts = []
+            for block in content_blocks:
+                block_type = block.get("type")
+                if block_type == "text":
+                    parts.append(block.get("text", ""))
+                elif block_type == "tool_result":
+                    tool_content = block.get("content", "")
+                    if isinstance(tool_content, list):
+                        text_parts = []
+                        for b in tool_content:
+                            if b.get("type") == "text":
+                                text_parts.append(b.get("text", ""))
+                            else:
+                                text_parts.append(f"[Unsupported {b.get('type')}]")
+                        tool_content = "\n\n".join(text_parts)
+                    parts.append(tool_output_template.format(content=tool_content))
+                else:
+                    parts.append(f"[Unsupported {block_type}]")
+            prompt += "\n\n".join(parts)
+        else:
+            prompt += content or ""
+
+    elif role == "latest_reminder":
+        prompt += LATEST_REMINDER_SP_TOKEN + latest_reminder_msg_template.format(content=content)
+
+    elif role == "tool":
+        raise NotImplementedError("deepseek_v41 merges tool messages into user; please preprocess with merge_tool_messages()")
+
+    elif role == "assistant":
+        thinking_part = ""
+        tc_content = ""
+
+        if tool_calls:
+            tc_list = [
+                tool_call_template.format(
+                    dsml_token=dsml_token,
+                    tool_call_tag_name=tool_call_tag_name,
+                    name=tc.get("name"),
+                    arguments=encode_arguments_to_dsml(tc)
+                )
+                for tc in tool_calls
+            ]
+            tc_content += '\n\n' + tool_calls_template.format(
+                dsml_token=dsml_token,
+                tool_calls="\n".join(tc_list),
+                tc_block_name=tool_calls_block_name,
+            )
+
+        summary_content = content or ""
+        rc = reasoning_content or ""
+
+        # Check if previous message has a task - if so, this is a task output (no thinking)
+        prev_has_task = index - 1 >= 0 and messages[index - 1].get("task") is not None
+
+        if thinking_mode == "thinking" and not prev_has_task:
+            if not drop_thinking or index > last_user_idx:
+                thinking_part = thinking_template.format(reasoning_content=rc) + thinking_end_token
+            else:
+                thinking_part = ""
+
+        if wo_eos:
+            prompt += assistant_msg_wo_eos_template.format(
+                reasoning=thinking_part,
+                content=summary_content,
+                tool_calls=tc_content,
+            )
+        else:
+            prompt += assistant_msg_template.format(
+                reasoning=thinking_part,
+                content=summary_content,
+                tool_calls=tc_content,
+            )
+    else:
+        raise NotImplementedError(f"Unknown role: {role}")
+
+    # Append transition tokens based on what follows
+    if index + 1 < len(messages) and messages[index + 1].get("role") not in ["assistant", "latest_reminder"]:
+        return prompt
+
+    task = messages[index].get("task")
+    if task is not None:
+        # Task special token for internal classification tasks
+        assert task in VALID_TASKS, f"Invalid task: '{task}'. Valid tasks are: {list(VALID_TASKS)}"
+        task_sp_token = DS_TASK_SP_TOKENS[task]
+
+        if task != "action":
+            # Non-action tasks: append task sp token directly after the message
+            prompt += task_sp_token
+        else:
+            # Action task: append Assistant + thinking token + action sp token
+            prompt += ASSISTANT_SP_TOKEN
+            prompt += thinking_end_token if thinking_mode != "thinking" else thinking_start_token
+            prompt += task_sp_token
+
+    elif role == "user" or (role == "system" and index > 0):
+        # Normal generation: append Assistant + thinking token
+        # (mid-conversation system messages also trigger the assistant header)
+        prompt += ASSISTANT_SP_TOKEN
+        if not drop_thinking and thinking_mode == "thinking":
+            prompt += thinking_start_token
+        elif drop_thinking and thinking_mode == "thinking" and index >= last_user_idx:
+            prompt += thinking_start_token
+        else:
+            prompt += thinking_end_token
+
+    return prompt
+
+
+# Main Encoding Function
+
+def _drop_thinking_messages(messages: List[Dict[str, Any]]) -> List[Dict[str, Any]]:
+    """
+    Drop reasoning_content and non-essential messages before the last user message.
+    Same as V4, but uses the V4.1 last-user definition (mid systems count).
+    """
+    last_user_idx = find_last_user_index(messages)
+    result = []
+    keep_roles = {"user", "system", "tool", "latest_reminder", "direct_search_results"}
+
+    for idx, msg in enumerate(messages):
+        role = msg.get("role")
+        if role in keep_roles or idx >= last_user_idx:
+            result.append(msg)
+        elif role == "assistant":
+            msg = copy.copy(msg)
+            msg.pop("reasoning_content", None)
+            result.append(msg)
+
+    return result
+
+
+def _encode_messages_text(
+    messages: List[Dict[str, Any]],
+    thinking_mode: str,
+    context: Optional[List[Dict[str, Any]]] = None,
+    drop_thinking: bool = True,
+    add_default_bos_token: bool = True,
+    reasoning_effort: Union[str, int, None] = None,
+) -> str:
+    """Encode preprocessed (text-only) messages into the V4.1 prompt format."""
+    context = context if context else []
+
+    # Preprocess: merge tool messages and sort tool results
+    messages = merge_tool_messages(messages)
+    messages = sort_tool_results_by_call_order(context + messages)[len(context):]
+    if context:
+        context = merge_tool_messages(context)
+        context = sort_tool_results_by_call_order(context)
+
+    full_messages = context + messages
+
+    prompt = bos_token if add_default_bos_token and len(context) == 0 else ""
+
+    # Resolve drop_thinking: if any message has tools defined, don't drop thinking
+    effective_drop_thinking = drop_thinking
+    if any(m.get("tools") for m in full_messages):
+        effective_drop_thinking = False
+
+    if thinking_mode == "thinking" and effective_drop_thinking:
+        full_messages = _drop_thinking_messages(full_messages)
+        num_to_render = len(full_messages) - len(_drop_thinking_messages(context))
+        context_len = len(full_messages) - num_to_render
+    else:
+        num_to_render = len(messages)
+        context_len = len(context)
+
+    for idx in range(num_to_render):
+        prompt += render_message(
+            idx + context_len,
+            full_messages,
+            thinking_mode=thinking_mode,
+            drop_thinking=effective_drop_thinking,
+            reasoning_effort=reasoning_effort,
+        )
+
+    return prompt
+
+
+def encode_messages(
+    messages: List[Dict[str, Any]],
+    thinking_mode: str,
+    context: Optional[List[Dict[str, Any]]] = None,
+    drop_thinking: bool = True,
+    add_default_bos_token: bool = True,
+    reasoning_effort: Union[str, int, None] = None,
+    return_multi_modal_data: bool = False,
+) -> Any:
+    """Encode text or multimodal messages into the DeepSeek-V4.1 prompt format.
+
+    Text-only calls return the prompt string. When return_multi_modal_data is
+    true, the result is ``(prompt, media_data)``.
+    """
+    context = context or []
+    processed_context, _ = process_image_messages(context) if context else ([], [])
+    normalized = copy.deepcopy(messages)
+    for message in normalized:
+        if message.get("role") == "user" and isinstance(message.get("content"), list):
+            message.setdefault("content_blocks", message.pop("content"))
+    ordered = merge_tool_messages(normalized)
+    ordered = sort_tool_results_by_call_order(copy.deepcopy(context) + ordered)[len(context):]
+    # Sort tool results before collecting images so payload order matches the rendered spans.
+    processed_messages, images = process_image_messages(ordered)
+    prompt = _encode_messages_text(
+        processed_messages,
+        thinking_mode=thinking_mode,
+        context=processed_context if processed_context else None,
+        drop_thinking=drop_thinking,
+        add_default_bos_token=add_default_bos_token,
+        reasoning_effort=reasoning_effort,
+    )
+    if return_multi_modal_data:
+        return prompt, {"images": images}
+    return prompt
diff --git a/python/freetoken/models/deepseek_v41/engram.py b/python/freetoken/models/deepseek_v41/engram.py
new file mode 100644
index 000000000..cecde1e06
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/engram.py
@@ -0,0 +1,391 @@
+"""V4.1 conditional memory, with demand-paged FP8/FP4 tables and request-local hashes."""
+
+from __future__ import annotations
+
+import json
+import math
+import os
+import struct
+from contextlib import contextmanager
+from dataclasses import dataclass
+
+import numpy as np
+import torch
+from torch import nn
+
+
+def compressed_token_map(tokenizer) -> np.ndarray:
+    from tokenizers import Regex, normalizers
+
+    sentinel = "\ue000"
+    normalize = normalizers.Sequence([
+        normalizers.NFKC(), normalizers.NFD(), normalizers.StripAccents(),
+        normalizers.Lowercase(), normalizers.Replace(Regex(r"[ \t\r\n]+"), " "),
+        normalizers.Replace(Regex(r"^ $"), sentinel), normalizers.Strip(),
+        normalizers.Replace(sentinel, " "),
+    ])
+    keys = {}
+    result = np.empty(len(tokenizer), dtype=np.int64)
+    backend = tokenizer.backend_tokenizer
+    for token_id in range(len(tokenizer)):
+        text = backend.decode([token_id], skip_special_tokens=False)
+        key = backend.id_to_token(token_id) if "\ufffd" in text else normalize.normalize_str(text) or text
+        if key not in keys:
+            keys[key] = len(keys)
+        result[token_id] = keys[key]
+    return result
+
+
+@dataclass(frozen=True)
+class HashLayout:
+    layer_ids: tuple[int, ...]
+    max_ngram_size: int
+    heads: int
+    primes: np.ndarray
+    offsets: np.ndarray
+    multipliers: np.ndarray
+
+    @classmethod
+    def from_args(cls, args):
+        from sympy import nextprime
+
+        layers = tuple(args.engram_layer_ids)
+        max_ngram = args.engram_max_ngram_size
+        heads = args.engram_n_heads
+        if not layers or max_ngram < 2 or heads < 1 or args.engram_compressed_vocab_size < 1:
+            raise ValueError("invalid Engram hash geometry")
+        primes, used = [], set()
+        for _ in layers:
+            per_layer = []
+            for _ in range(max_ngram - 1):
+                current = args.engram_vocab_size - 1
+                for _ in range(heads):
+                    current = int(nextprime(current))
+                    while current in used:
+                        current = int(nextprime(current))
+                    used.add(current)
+                    per_layer.append(current)
+            primes.append(per_layer)
+        primes = np.asarray(primes, dtype=np.int64)
+        if tuple(primes.sum(1)) != tuple(args.engram_num_embeddings):
+            raise ValueError("Engram hash buckets do not match the checkpoint table sizes")
+        offsets = np.cumsum(np.concatenate((np.zeros((len(layers), 1), dtype=np.int64), primes[:, :-1]), 1), 1)
+        bound = max(1, (np.iinfo(np.int64).max // args.engram_compressed_vocab_size) // 2)
+        multipliers = np.stack([
+            np.random.default_rng(10007 * layer).integers(0, bound, size=max_ngram, dtype=np.int64) * 2 + 1
+            for layer in layers
+        ])
+        return cls(layers, max_ngram, heads, primes, offsets, multipliers)
+
+
+def hash_token_run(ids: np.ndarray, image_mask: np.ndarray, token_map: np.ndarray,
+                   layout: HashLayout, pad_id: int, prefix_length: int = 0) -> np.ndarray:
+    ids = np.asarray(ids, dtype=np.int64)
+    image_mask = np.asarray(image_mask, dtype=bool)
+    if ids.ndim != 1 or image_mask.shape != ids.shape or not 0 <= prefix_length <= len(ids):
+        raise ValueError("invalid Engram token run")
+    ids = np.where(image_mask, pad_id, ids)
+    if np.any(ids < 0) or np.any(ids >= len(token_map)):
+        raise ValueError("Engram input token is outside the tokenizer vocabulary")
+    compressed = token_map[ids]
+    positions = np.arange(prefix_length, len(ids))
+    blocked = np.zeros(len(positions), dtype=bool)
+    lookbacks = []
+    for shift in range(layout.max_ngram_size):
+        source = positions - shift
+        safe = source.clip(0)
+        blocked |= (source < 0) | image_mask[safe]
+        lookbacks.append(np.where(blocked, token_map[pad_id], compressed[safe]))
+    tokens = np.stack(lookbacks, -1)
+    products = tokens[:, None, :] * layout.multipliers[None, :, :]
+    rolling = products[..., 0]
+    hashes = []
+    for shift in range(1, layout.max_ngram_size):
+        rolling = np.bitwise_xor(rolling, products[..., shift])
+        start = (shift - 1) * layout.heads
+        hashes.append(rolling[..., None] % layout.primes[None, :, start:start + layout.heads])
+    return np.concatenate(hashes, -1) + layout.offsets[None, :, :]
+
+
+def _header(path: str):
+    with open(path, "rb") as stream:
+        prefix = stream.read(8)
+        if len(prefix) != 8:
+            raise ValueError(f"truncated safetensors header: {path}")
+        size = struct.unpack("<Q", prefix)[0]
+        if size > 64 << 20:
+            raise ValueError(f"oversized safetensors header: {path}")
+        raw = stream.read(size)
+    if len(raw) != size:
+        raise ValueError(f"truncated safetensors header: {path}")
+    return json.loads(raw), 8 + size
+
+
+class DiskEngramTable:
+    def __init__(self, weights: np.ndarray, scales: np.ndarray, *, dtype="fp8",
+                 block_size=32, scale_fmt="ue8m0"):
+        if dtype not in ("fp8", "fp4") or block_size != 32 or scale_fmt != "ue8m0":
+            raise ValueError("Engram requires FP8 or packed FP4 with block-32 UE8M0 scales")
+        if weights.dtype != np.uint8 or scales.dtype != np.uint8:
+            raise ValueError("Engram mapped weights and scales must use raw uint8 storage")
+        head_dim = weights.shape[1] * (2 if dtype == "fp4" else 1) if weights.ndim == 2 else 0
+        if (weights.ndim != 2 or head_dim == 0 or head_dim % 32
+                or scales.shape != (weights.shape[0], head_dim // 32)):
+            raise ValueError("Engram table requires one E8M0 scale per 32 weight channels")
+        self.weights = weights
+        self.scales = scales
+        self.num_rows, self.head_dim = weights.shape[0], head_dim
+        self.dtype, self.block_size, self.scale_fmt = dtype, block_size, scale_fmt
+
+    @classmethod
+    def from_checkpoint(cls, folder: str, layer_id: int, *, args=None):
+        manifest_path = os.path.join(folder, "engram_tables.json")
+        if os.path.isfile(manifest_path):
+            with open(manifest_path, encoding="utf-8") as stream:
+                manifest = json.load(stream)
+            if manifest.get("version") not in (1, 2):
+                raise ValueError("Unsupported Engram table manifest version")
+            info = manifest["layers"][str(layer_id)]
+            format_info = ({"dtype": "fp8", "block_size": 32, "scale_fmt": "ue8m0"}
+                           if manifest["version"] == 1 else
+                           {name: info[name] for name in ("dtype", "block_size", "scale_fmt")})
+            arrays = []
+            for kind in ("weight", "scale"):
+                entry = info[kind]
+                path = os.path.abspath(os.path.join(folder, entry["file"]))
+                if os.path.commonpath((os.path.abspath(folder), path)) != os.path.abspath(folder):
+                    raise ValueError("Engram table path escapes the checkpoint")
+                arrays.append(cls._map(path, entry["shape"], entry["offset"]))
+            table = cls(*arrays, **format_info)
+            table._validate_config(args)
+            return table
+        with open(os.path.join(folder, "model.safetensors.index.json"), encoding="utf-8") as stream:
+            index = json.load(stream)["weight_map"]
+        arrays = []
+        dtype = None
+        for kind in ("weight", "scale"):
+            key = f"layers.{layer_id}.engram.embed.{kind}"
+            if key not in index:
+                raise ValueError(f"missing Engram tensor {key}")
+            path = os.path.abspath(os.path.join(folder, index[key]))
+            if os.path.commonpath((os.path.abspath(folder), path)) != os.path.abspath(folder):
+                raise ValueError("Engram tensor path escapes the checkpoint")
+            header, base = _header(path)
+            entry = header[key]
+            if kind == "weight":
+                dtype = {"F8_E4M3": "fp8", "U8": "fp4"}.get(entry["dtype"])
+                if dtype is None:
+                    raise ValueError(f"{key} must have dtype F8_E4M3 or U8")
+            elif entry["dtype"] != "F8_E8M0":
+                raise ValueError(f"{key} must have dtype F8_E8M0")
+            begin, end = entry["data_offsets"]
+            if end - begin != math.prod(entry["shape"]):
+                raise ValueError(f"invalid byte extent for {key}")
+            arrays.append(cls._map(path, entry["shape"], base + begin))
+        table = cls(*arrays, dtype=dtype)
+        table._validate_config(args)
+        return table
+
+    def _validate_config(self, args):
+        if args is not None and (self.dtype, self.block_size, self.scale_fmt) != (
+                getattr(args, "engram_dtype", "fp8"), getattr(args, "engram_block_size", 32),
+                getattr(args, "engram_scale_fmt", "ue8m0")):
+            raise ValueError("Engram table format disagrees with model configuration")
+
+    @staticmethod
+    def _map(path, shape, offset):
+        if len(shape) != 2 or min(shape) < 1 or offset < 0 or offset + math.prod(shape) > os.path.getsize(path):
+            raise ValueError(f"invalid Engram table extent: {path}")
+        return np.memmap(path, mode="r", dtype=np.uint8, offset=offset, shape=tuple(shape))
+
+    def lookup(self, row_ids: np.ndarray) -> torch.Tensor:
+        ids = np.asarray(row_ids, dtype=np.int64)
+        if np.any(ids < 0) or np.any(ids >= self.num_rows):
+            raise ValueError("Engram row index outside table")
+        unique, inverse = np.unique(ids, return_inverse=True)
+        weights = torch.from_numpy(np.array(self.weights[unique], copy=True))
+        if self.dtype == "fp4":
+            codes = torch.stack((weights & 15, weights >> 4), -1).flatten(-2).long()
+            values = torch.tensor([0., .5, 1., 1.5, 2., 3., 4., 6.,
+                                   -0., -.5, -1., -1.5, -2., -3., -4., -6.],
+                                  dtype=torch.float32, device="cpu")
+            weights = values[codes]
+        else:
+            weights = weights.view(torch.float8_e4m3fn).float()
+        scales = torch.from_numpy(np.array(self.scales[unique], copy=True)).view(torch.float8_e8m0fnu).float()
+        rows = (weights.reshape(-1, self.head_dim // 32, 32) * scales[..., None]).flatten(-2).to(torch.bfloat16)
+        return rows[torch.from_numpy(inverse)].reshape(*ids.shape, self.head_dim)
+
+
+class Engram(nn.Module):
+    def __init__(self, args, layer_id: int):
+        super().__init__()
+        from .layers import Linear
+
+        self.args = args
+        self.layer_id = layer_id
+        self.layer_hash_index = tuple(args.engram_layer_ids).index(layer_id)
+        self.dim, self.hc_mult, self.eps = args.dim, args.hc_mult, args.norm_eps
+        self.hash_cols = (args.engram_max_ngram_size - 1) * args.engram_n_heads
+        self.wkv = Linear(self.hash_cols * args.engram_head_dim, args.dim * (args.hc_mult + 1))
+        self.q_weight = nn.Parameter(torch.ones(args.hc_mult, args.dim, dtype=torch.bfloat16), requires_grad=False)
+        self.k_weight = nn.Parameter(torch.ones(args.hc_mult, args.dim, dtype=torch.bfloat16), requires_grad=False)
+        self._values = None
+        self._mask = None
+
+    def forward(self, x: torch.Tensor, hash_ids=None, token_mask=None):
+        if self._values is None:
+            raise RuntimeError("Engram tables were not attached before inference")
+        count = x.shape[0] * x.shape[1]
+        values = self._values[:count].reshape(*x.shape[:2], -1)
+        kv = self.wkv(values)
+        key, value = kv.split([self.hc_mult * self.dim, self.dim], -1)
+        key = key.float().reshape(*x.shape[:2], self.hc_mult, self.dim)
+        h = x.float()
+        rstd = torch.rsqrt(h.square().mean(-1) + self.eps) * torch.rsqrt(key.square().mean(-1) + self.eps)
+        dot = (h * self.q_weight.float() * self.k_weight.float() * key).sum(-1) * rstd * self.dim**-0.5
+        gate = torch.sigmoid(torch.copysign(dot.abs().clamp_min(1e-6).sqrt(), dot))
+        mask = token_mask if token_mask is not None else self._mask[:count].reshape(*x.shape[:2])
+        gate = gate.masked_fill(~mask.unsqueeze(-1), 0)
+        return (h + gate.unsqueeze(-1) * value.float().unsqueeze(-2)).to(x.dtype)
+
+
+def _image_flags(req, start: int, length: int) -> np.ndarray:
+    flags = np.zeros(length, dtype=bool)
+    spans = [span for item in getattr(req, "mm_items", None) or () for span in item.offsets]
+    for media in getattr(req, "media", None) or ():
+        spans.append((int(media["start"]), int(media["start"]) + len(media["types"])))
+    for lo, hi in spans:
+        left = max(start, lo)
+        right = min(start + length, hi)
+        if left < right:
+            flags[left-start:right-start] = True
+    return flags
+
+
+class EngramRuntime:
+    def __init__(self, args, modules, tables, token_map, max_tokens, device, *, dummy=False):
+        self.args, self.modules, self.tables = args, modules, tables
+        self.token_map = token_map
+        self.layout = None if dummy else HashLayout.from_args(args)
+        self.dummy = dummy
+        self.device = torch.device(device)
+        self.capacity = max_tokens
+        self.host = []
+        self.host_mask = torch.ones(max_tokens, dtype=torch.bool, pin_memory=self.device.type == "cuda")
+        self.device_mask = self.host_mask.to(self.device)
+        self._copy_done = torch.cuda.Event() if self.device.type == "cuda" else None
+        self._has_copy = False
+        for module in modules:
+            width = module.hash_cols * args.engram_head_dim
+            host = torch.zeros((max_tokens, width), dtype=torch.bfloat16, pin_memory=self.device.type == "cuda")
+            self.host.append(host)
+            module._values = host.to(self.device)
+            module._mask = self.device_mask
+
+    @property
+    def pinned_bytes(self):
+        return sum(x.numel() * x.element_size() for x in self.host) + self.host_mask.numel()
+
+    @contextmanager
+    def forward_host_ctx(self, batch, use_graph):
+        if self._has_copy:
+            self._copy_done.synchronize()
+        runs, masks = [], []
+        current = batch.input_ids.detach().to("cpu", dtype=torch.int64).numpy().reshape(-1)
+        offset = 0
+        for req in batch.padded_reqs:
+            count = 1 if batch.is_decode else req.extend_len
+            start = req.device_len - 1 if batch.is_decode else req.cached_len
+            prefix_start = max(0, start - (self.args.engram_max_ngram_size - 1))
+            prefix = req.input_ids[prefix_start:start].to(dtype=torch.int64).numpy()
+            if len(prefix) != start - prefix_start:
+                raise RuntimeError("Engram history is not available for the active request")
+            ids = np.concatenate((prefix, current[offset:offset+count]))
+            image_flags = _image_flags(req, prefix_start, len(ids))
+            if not self.dummy:
+                pad_id = getattr(self.args, "engram_pad_id", 2)
+                runs.append(hash_token_run(ids, image_flags, self.token_map, self.layout, pad_id, len(prefix)))
+            masks.append(~image_flags[len(prefix):])
+            offset += count
+        if offset != len(current) or offset > self.capacity:
+            raise ValueError(f"Engram staging requires {offset} rows; capacity is {self.capacity}")
+        live = np.concatenate(masks) if masks else np.empty(0, dtype=bool)
+        self.host_mask[:offset].copy_(torch.from_numpy(live))
+        all_hashes = np.concatenate(runs) if runs else None
+        for index, (module, host) in enumerate(zip(self.modules, self.host)):
+            host[:offset].zero_()
+            if all_hashes is not None and live.any():
+                rows = self.tables[index].lookup(all_hashes[live, index])
+                host[:offset][torch.from_numpy(live)] = rows.flatten(-2)
+            module._values[:offset].copy_(host[:offset], non_blocking=True)
+        self.device_mask[:offset].copy_(self.host_mask[:offset], non_blocking=True)
+        if self._copy_done is not None:
+            self._copy_done.record()
+            self._has_copy = True
+        yield
+
+
+def prepare_engram(model, engine_config) -> int:
+    args = model._args
+    modules = [layer.engram for layer in model._transformer.layers if layer.engram is not None]
+    if not modules:
+        return 0
+    dummy = getattr(engine_config, "use_dummy_weight", False)
+    tables, mapping = [], None
+    if not dummy:
+        from freetoken.utils import download_hf_weight, load_tokenizer
+
+        folder = download_hf_weight(engine_config.model_path)
+        mapping = compressed_token_map(load_tokenizer(folder))
+        if int(mapping.max()) + 1 != args.engram_compressed_vocab_size:
+            raise ValueError("Engram tokenizer normalization does not match the checkpoint compressed vocabulary")
+        tables = [DiskEngramTable.from_checkpoint(folder, module.layer_id, args=args) for module in modules]
+        for index, table in enumerate(tables):
+            if table.num_rows != args.engram_num_embeddings[index] or table.head_dim != args.engram_head_dim:
+                raise ValueError("Engram table shape disagrees with model configuration")
+    prefill_capacity = getattr(engine_config, "max_extend_tokens", None)
+    if prefill_capacity is None:
+        prefill_capacity = min(engine_config.max_forward_len, 8192)
+    runtime = EngramRuntime(
+        args, modules, tables, mapping,
+        max(prefill_capacity, engine_config.max_running_req, engine_config.cuda_graph_max_bs or 0, 1),
+        engine_config.device if hasattr(engine_config, "device") else next(model._transformer.parameters()).device,
+        dummy=dummy,
+    )
+    model._engram_runtime = runtime
+    return runtime.pinned_bytes
+
+
+def export_engram_tables(folder: str, out_dir: str, args) -> list[str]:
+    tables_dir = os.path.join(out_dir, "engram")
+    os.makedirs(tables_dir, exist_ok=True)
+    manifest = {"version": 2, "layers": {}}
+    copied = []
+    for layer in args.engram_layer_ids:
+        table = DiskEngramTable.from_checkpoint(folder, layer, args=args)
+        entries = {"dtype": table.dtype, "block_size": table.block_size, "scale_fmt": table.scale_fmt}
+        for kind, source in (("weight", table.weights), ("scale", table.scales)):
+            relative = f"engram/layer-{layer}-{kind}.bin"
+            destination = os.path.join(out_dir, relative)
+            if os.path.realpath(source.filename) == os.path.realpath(destination):
+                raise ValueError("Engram source and destination must differ")
+            temporary = destination + ".partial"
+            remaining = source.nbytes
+            with open(source.filename, "rb") as reader, open(temporary, "wb") as writer:
+                reader.seek(source.offset)
+                while remaining:
+                    chunk = reader.read(min(16 << 20, remaining))
+                    if not chunk:
+                        raise OSError("Engram source was truncated during FTW conversion")
+                    writer.write(chunk)
+                    remaining -= len(chunk)
+            os.replace(temporary, destination)
+            entries[kind] = {"file": relative, "offset": 0, "shape": list(source.shape)}
+            copied.append(relative)
+        manifest["layers"][str(layer)] = entries
+    path = os.path.join(out_dir, "engram_tables.json")
+    with open(path + ".partial", "w", encoding="utf-8") as writer:
+        json.dump(manifest, writer, indent=2)
+    os.replace(path + ".partial", path)
+    return copied + ["engram_tables.json"]
diff --git a/python/freetoken/models/deepseek_v41/image_processor.py b/python/freetoken/models/deepseek_v41/image_processor.py
new file mode 100644
index 000000000..e914ffcd9
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/image_processor.py
@@ -0,0 +1,229 @@
+"""Image preprocessing (upstream MIT attribution: NOTICE).
+
+An image becomes a `n_vit_h x n_vit_w` patch grid for the ViT and a `n_llm_h x n_llm_w` token grid
+after the 3x3 aligner downsample, which the LLM sees as
+
+    [IMAGE_START] + ([IMAGE] * n_llm_w + [IMAGE_NEW_LINE]) * n_llm_h + [IMAGE_END]
+
+Every one of those positions carries `image_token_id` in `input_ids`; only the token type tells them
+apart. The IMAGE slots are filled with aligner rows in reading order.
+"""
+
+import base64
+import io
+import ipaddress
+import math
+import socket
+from dataclasses import dataclass
+from urllib.parse import urlsplit
+from urllib.request import HTTPRedirectHandler, build_opener
+
+import numpy as np
+import torch
+from PIL import Image, ImageOps
+
+TEXT = -1
+IMAGE_START, IMAGE, IMAGE_NEW_LINE, IMAGE_END = range(4)
+MAX_IMAGE_BYTES = 32 * 1024 * 1024
+MAX_REQUEST_IMAGES = 16
+MAX_IMAGE_PIXELS = 64 * 1024 * 1024
+
+
+@dataclass
+class ImageInput:
+    start: int
+    patches: torch.Tensor
+    n_vit_h: int
+    n_vit_w: int
+    types: torch.Tensor
+
+
+def num_image_tokens(n_llm_h: int, n_llm_w: int) -> int:
+    return n_llm_h * (n_llm_w + 1) + 2
+
+
+def llm_grid(best_height: int, best_width: int, patch_size: int, downsample_ratio: int):
+    """Token grid the aligner produces from a patch grid of this pixel size."""
+    return math.ceil((best_height // patch_size) / downsample_ratio), math.ceil(
+        (best_width // patch_size) / downsample_ratio
+    )
+
+
+def solve_resize_ratio(height, width, patch_size, downsample_ratio, max_n_token):
+    """Largest aspect-preserving pixel size whose token grid still fits in max_n_token."""
+    r = height / width
+    max_w_float = math.sqrt((max_n_token - 2) / r + 0.25) - 0.5
+    max_h_float = max_w_float * r
+    cell = patch_size * downsample_ratio
+    if max_w_float < 1.0:  # very tall: collapse to a single column
+        return (max_n_token - 2) // 2 * cell, cell
+    if max_h_float < 1.0:  # very wide: collapse to a single row
+        return cell, (max_n_token - 3) * cell
+    beta = min(math.floor(max_w_float) * cell / width, math.floor(max_h_float) * cell / height)
+    return math.floor(height * beta / patch_size) * patch_size, math.floor(width * beta / patch_size) * patch_size
+
+
+def safe_resize(height, width, best_height, best_width, patch_size, downsample_ratio, max_n_token):
+    """Shrink the pixel size until the image costs at most max_n_token LLM tokens."""
+    n_llm_h, n_llm_w = llm_grid(best_height, best_width, patch_size, downsample_ratio)
+    if num_image_tokens(n_llm_h, n_llm_w) > max_n_token:
+        best_height, best_width = solve_resize_ratio(height, width, patch_size, downsample_ratio, max_n_token)
+        n_llm_h, n_llm_w = llm_grid(best_height, best_width, patch_size, downsample_ratio)
+        assert num_image_tokens(n_llm_h, n_llm_w) <= max_n_token
+    return n_llm_h, n_llm_w, best_height, best_width
+
+
+def _validate_image_url(url: str) -> None:
+    parsed = urlsplit(url)
+    if parsed.scheme not in ("https", "http") or not parsed.hostname or parsed.username or parsed.password:
+        raise ValueError("images require an HTTP(S) URL or a base64 data URL")
+    try:
+        addresses = socket.getaddrinfo(parsed.hostname, parsed.port or (443 if parsed.scheme == "https" else 80))
+    except OSError as exc:
+        raise ValueError("could not resolve image URL") from exc
+    if not addresses or any(not ipaddress.ip_address(entry[4][0]).is_global for entry in addresses):
+        raise ValueError("image URLs must resolve to public addresses")
+
+
+class _ImageRedirectHandler(HTTPRedirectHandler):
+    def redirect_request(self, req, fp, code, msg, headers, newurl):
+        _validate_image_url(newurl)
+        return super().redirect_request(req, fp, code, msg, headers, newurl)
+
+
+def _decode_base64(data: str) -> bytes:
+    if len(data) > (MAX_IMAGE_BYTES + 2) // 3 * 4:
+        raise ValueError("image exceeds the 32 MiB upload limit")
+    try:
+        decoded = base64.b64decode(data, validate=True)
+    except (ValueError, TypeError) as exc:
+        raise ValueError("invalid base64 image") from exc
+    if len(decoded) > MAX_IMAGE_BYTES:
+        raise ValueError("image exceeds the 32 MiB upload limit")
+    return decoded
+
+
+def load_image_bytes(record) -> bytes:
+    """Load bounded API image data without allowing local filesystem paths."""
+    data = record.get("data")
+    if isinstance(data, bytes):
+        if len(data) > MAX_IMAGE_BYTES:
+            raise ValueError("image exceeds the 32 MiB upload limit")
+        return data
+    if isinstance(data, str):
+        return _decode_base64(data)
+
+    source = record.get("source")
+    if isinstance(source, dict):
+        if source.get("data") is not None:
+            return _decode_base64(source["data"])
+        if source.get("url"):
+            return load_image_bytes({"url": source["url"]})
+
+    url = record.get("url")
+    if isinstance(url, str) and url:
+        if url.startswith("data:"):
+            header, _, payload = url.partition(",")
+            if ";base64" not in header:
+                raise ValueError(f"Unsupported data URL encoding: {header}")
+            return _decode_base64(payload)
+        if url.startswith(("http://", "https://")):
+            _validate_image_url(url)
+            with build_opener(_ImageRedirectHandler()).open(url, timeout=30) as response:
+                data = response.read(MAX_IMAGE_BYTES + 1)
+            if len(data) > MAX_IMAGE_BYTES:
+                raise ValueError("image exceeds the 32 MiB download limit")
+            return data
+        raise ValueError("images require an HTTP(S) URL or a base64 data URL")
+
+    raise ValueError(f"Cannot load image from record: {list(record.keys())}")
+
+
+def plan_image_grid(width: int, height: int, args):
+    """Resize plan for an image of the given original size; a pure function of its arguments."""
+    p = args.vision_patch_size
+    if width <= 0 or height <= 0 or p <= 0 or args.vision_downsample_ratio <= 0 or args.vision_max_n_token < 4:
+        raise ValueError("invalid image dimensions or vision configuration")
+    if args.vision_max_wh_ratio is not None and width > height * args.vision_max_wh_ratio:
+        width = height * args.vision_max_wh_ratio
+    if 0 < width * height < args.vision_min_pixels:
+        ratio = (args.vision_min_pixels / (width * height)) ** 0.5
+        width = int(width * ratio)
+        height = int(height * ratio)
+    best_width = math.ceil(width / p) * p
+    best_height = math.ceil(height / p) * p
+    return safe_resize(height, width, best_height, best_width, p, args.vision_downsample_ratio, args.vision_max_n_token)
+
+
+def load_image(record, args):
+    """Load and transform one image record into ViT patches."""
+    with Image.open(io.BytesIO(load_image_bytes(record))) as source:
+        if source.width * source.height > MAX_IMAGE_PIXELS:
+            raise ValueError("image exceeds the 64 megapixel limit")
+        image = source.convert("RGB")
+    return process_image(image, args)
+
+
+def process_image(image, args):
+    """Apply the native resize and normalization to an already decoded image."""
+    p = args.vision_patch_size
+    if image.width * image.height > MAX_IMAGE_PIXELS:
+        raise ValueError("image exceeds the 64 megapixel limit")
+    image = image.convert("RGB")
+    n_llm_h, n_llm_w, best_height, best_width = plan_image_grid(image.width, image.height, args)
+    n_vit_h, n_vit_w = best_height // p, best_width // p
+    if args.vision_max_wh_ratio is not None and image.width >= args.vision_max_wh_ratio * image.height:
+        image = image.resize((best_width, best_height))
+    else:
+        image = ImageOps.pad(image, (best_width, best_height), color=(127, 127, 127))
+    x = torch.from_numpy(np.asarray(image, dtype=np.float32)).permute(2, 0, 1) / 255
+    x = (x - 0.5) / 0.5
+    patches = x.reshape(3, n_vit_h, p, n_vit_w, p).permute(1, 3, 0, 2, 4).reshape(n_vit_h * n_vit_w, 3, p, p)
+    return patches, n_vit_h, n_vit_w, n_llm_h, n_llm_w
+
+
+def image_token_types(n_llm_h: int, n_llm_w: int) -> torch.Tensor:
+    """Default layout: the aligner grid in reading order, one IMAGE_NEW_LINE per row."""
+    types = [IMAGE_START]
+    types += ([IMAGE] * n_llm_w + [IMAGE_NEW_LINE]) * n_llm_h
+    types.append(IMAGE_END)
+    return torch.tensor(types, dtype=torch.int64)
+
+
+def prepare_vl_inputs(prompt, images, tokenizer, args):
+    """Tokenize `prompt`, expanding each image placeholder token into its image span.
+
+    Returns (tokens, token_types, image_inputs). Image-span positions carry `args.image_token_id` in
+    `tokens` and are distinguished only by `token_types` (TEXT elsewhere). `image_inputs` is None when
+    the prompt has no images."""
+    from .encoding import IMAGE_PLACEHOLDER
+
+    if len(images) > MAX_REQUEST_IMAGES:
+        raise ValueError(f"at most {MAX_REQUEST_IMAGES} images are supported per request")
+
+    # The placeholder is spelled differently across tokenizer revisions, so the id comes from the
+    # config; only cross-check it when this tokenizer does know the training-time spelling.
+    image_token_id = args.image_token_id
+    placeholder_id = tokenizer.convert_tokens_to_ids(IMAGE_PLACEHOLDER)
+    if placeholder_id is not None and placeholder_id != tokenizer.unk_token_id:
+        assert placeholder_id == image_token_id, (placeholder_id, image_token_id)
+    prompt_tokens = tokenizer.encode(prompt, add_special_tokens=False)
+    num_placeholders = sum(token == image_token_id for token in prompt_tokens)
+    if num_placeholders != len(images):
+        raise ValueError(f"Found {num_placeholders} image tokens but got {len(images)} images")
+    if num_placeholders and not args.vision_enabled:
+        raise ValueError("The model config has no vision tower (vision_n_layers == 0) but the prompt contains images")
+
+    tokens, token_types, image_inputs = [], [], []
+    image_iter = iter(images)
+    for tok in prompt_tokens:
+        if tok != image_token_id:
+            tokens.append(tok)
+            token_types.append(TEXT)
+            continue
+        patches, n_vit_h, n_vit_w, n_llm_h, n_llm_w = load_image(next(image_iter), args)
+        types = image_token_types(n_llm_h, n_llm_w)
+        image_inputs.append(ImageInput(len(tokens), patches, n_vit_h, n_vit_w, types))
+        tokens += [image_token_id] * types.numel()
+        token_types += types.tolist()
+    return tokens, token_types, image_inputs or None
diff --git a/python/freetoken/models/deepseek_v41/layers.py b/python/freetoken/models/deepseek_v41/layers.py
new file mode 100644
index 000000000..f841d5098
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/layers.py
@@ -0,0 +1,70 @@
+"""Resident V4.1 projections preserve block-32 FP8 storage and activation quantization."""
+
+from __future__ import annotations
+
+import torch
+import torch.nn.functional as F
+from torch import nn
+
+
+class Linear(nn.Module):
+    def __init__(self, in_features, out_features, bias=False, kind="fp8"):
+        super().__init__()
+        self.in_features = in_features
+        self.out_features = out_features
+        self.kind = kind
+        if kind not in {"fp8", "fp32", "bf16"}:
+            raise ValueError(f"Unsupported linear storage kind: {kind}")
+        if kind == "fp8":
+            if in_features % 32 or out_features % 32:
+                raise ValueError("V4.1 FP8 dimensions must be divisible by 32")
+            dtype = torch.float8_e4m3fn
+            self.scale = nn.Parameter(torch.empty(out_features // 32, in_features // 32,
+                                                  dtype=torch.float8_e8m0fnu), requires_grad=False)
+        else:
+            dtype = torch.float32 if kind == "fp32" else torch.bfloat16
+            self.register_parameter("scale", None)
+        self.weight = nn.Parameter(torch.empty(out_features, in_features, dtype=dtype), requires_grad=False)
+        self.bias = nn.Parameter(torch.empty(out_features, dtype=dtype), requires_grad=False) if bias else None
+
+    def forward(self, x):
+        if self.kind == "fp8":
+            from freetoken.kernel.triton.dsv41.quant import block_fp8_linear
+
+            return block_fp8_linear(x, self.weight, self.scale, self.bias, block_size=32)
+        return F.linear(x, self.weight.to(x.dtype), None if self.bias is None else self.bias.to(x.dtype))
+
+
+class RMSNorm(nn.Module):
+    def __init__(self, dim: int, eps: float = 1e-20):
+        super().__init__()
+        self.eps = eps
+        self.weight = nn.Parameter(torch.ones(dim, dtype=torch.float32), requires_grad=False)
+
+    def forward(self, x):
+        if x.is_cuda:
+            from freetoken.kernel.triton.dsv4.norm import rms_norm
+
+            return rms_norm(x, self.weight, self.eps)
+        value = x.float()
+        return (value * torch.rsqrt(value.square().mean(-1, keepdim=True) + self.eps)
+                * self.weight.float()).to(x.dtype)
+
+
+class OutputHead(nn.Module):
+    def __init__(self, dim, vocab_size):
+        super().__init__()
+        self.weight = nn.Parameter(torch.empty(vocab_size, dim, dtype=torch.bfloat16), requires_grad=False)
+
+    def forward(self, x):
+        if not x.is_cuda:
+            return F.linear(x.float(), self.weight.float())
+        from freetoken.kernel.triton.dsv4.bf16_linear import bf16_linear_fp32
+
+        # The reference returns FP32 logits; keep its BF16 checkpoint head resident.
+        flat = x.reshape(-1, x.shape[-1])
+        logits = torch.stack([bf16_linear_fp32(row, self.weight) for row in flat])
+        return logits.reshape(*x.shape[:-1], self.weight.shape[0])
+
+
+__all__ = ["Linear", "RMSNorm", "OutputHead"]
diff --git a/python/freetoken/models/deepseek_v41/mm_processor.py b/python/freetoken/models/deepseek_v41/mm_processor.py
new file mode 100644
index 000000000..72fb8816d
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/mm_processor.py
@@ -0,0 +1,76 @@
+"""Native DeepSeek image spans on the shared multimodal wire and cache."""
+
+import struct
+from dataclasses import replace
+
+import torch
+
+from freetoken.message import MMItem
+from freetoken.mm import mm_pad_value
+from freetoken.mm.processor import MMProcessor, MMResult, PromptReplacement, content_hash
+
+from .args import load_args
+from .image_processor import image_token_types, process_image
+
+
+class DeepseekV41MMProcessor(MMProcessor):
+    def __init__(self, hf_config, model_path, mm):
+        super().__init__(model_path, mm)
+        self.args = load_args(hf_config)
+        if mm.image_min_tokens is not None:
+            raise ValueError("DeepSeek-V4.1 uses processor_kwargs.vision_min_pixels instead of image_min_tokens")
+        allowed = {"vision_min_pixels", "vision_max_n_token", "vision_max_wh_ratio"}
+        if set(mm.processor_kwargs) - allowed:
+            raise ValueError(f"Unsupported DeepSeek-V4.1 image processor options: {sorted(set(mm.processor_kwargs) - allowed)}")
+        overrides = dict(mm.processor_kwargs)
+        if mm.image_max_tokens is not None:
+            overrides.setdefault("vision_max_n_token", mm.image_max_tokens)
+        self.args = replace(self.args, **overrides)
+        self.placeholder = [self.args.image_token_id]
+
+    @staticmethod
+    def _item(patches, n_h, n_w, types, start=0):
+        feature = patches.to(device="cpu", dtype=torch.bfloat16).contiguous()
+        types = torch.as_tensor(types, dtype=torch.int64).tolist()
+        extra = struct.pack("<2i", n_h, n_w) + bytes(types)
+        digest = content_hash(feature, extra)
+        return MMItem(
+            modality="image", hash=digest, pad_value=mm_pad_value(digest),
+            offsets=[[start, start + len(types)]], feature=feature,
+            model_specific_data={"n_vit_h": n_h, "n_vit_w": n_w, "types": types},
+        )
+
+    def from_media(self, input_ids, media):
+        """Convert native chat-encoder output without reprocessing an image."""
+        ids = input_ids.clone()
+        items = []
+        for image in media:
+            image = image if isinstance(image, dict) else vars(image)
+            item = self._item(image["patches"], image["n_vit_h"], image["n_vit_w"],
+                              image["types"], image["start"])
+            lo, hi = item.offsets[0]
+            if lo < 0 or hi > ids.numel() or not torch.all(ids[lo:hi] == self.args.image_token_id):
+                raise ValueError("DeepSeek-V4.1 image span does not match its placeholder tokens")
+            ids[lo:hi] = item.pad_value
+            item.validate()
+            items.append(item)
+        return MMResult(ids, items, None, 0)
+
+    def process(self, images):
+        items = []
+        for image in images:
+            patches, n_h, n_w, llm_h, llm_w = process_image(image, self.args)
+            items.append(self._item(patches, n_h, n_w, image_token_types(llm_h, llm_w)))
+        return items
+
+    def prompt_replacement(self, item):
+        return PromptReplacement([self.args.image_token_id] * len(item.types))
+
+    def dummy_items(self, dtype, device):
+        ratio, patch = self.args.vision_downsample_ratio, self.args.vision_patch_size
+        return [MMItem(
+            modality="image", hash=0, pad_value=0, offsets=[[0, 4]],
+            feature=torch.zeros(ratio * ratio, 3, patch, patch, dtype=dtype, device=device),
+            model_specific_data={"n_vit_h": ratio, "n_vit_w": ratio,
+                                 "types": image_token_types(1, 1).tolist()},
+        )]
diff --git a/python/freetoken/models/deepseek_v41/model.py b/python/freetoken/models/deepseek_v41/model.py
new file mode 100644
index 000000000..3a1e4e28e
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/model.py
@@ -0,0 +1,258 @@
+"""DeepSeek-V4.1 backbone with shifted mHC mixing and image-aware MoE/Engram."""
+
+from __future__ import annotations
+
+from contextlib import contextmanager
+
+import torch
+import torch.nn.functional as F
+from torch import nn
+
+from freetoken.core import get_global_ctx
+from freetoken.models.blocks import BaseLLMModel
+
+from .attention import Attention
+from .layers import OutputHead, RMSNorm
+from .moe import MoE
+
+
+def make_identity_pre_mix(x, hc_mult):
+    pre = x.new_zeros(*x.shape[:-2], hc_mult, dtype=torch.float32)
+    pre[..., 0] = 1.0
+    return pre
+
+
+class Block(nn.Module):
+    def __init__(self, layer_id, args, *, strategy="offload", decode_target="gpu", quant_config=None):
+        super().__init__()
+        self.layer_id = layer_id
+        self.dim = args.dim
+        self.norm_eps = args.norm_eps
+        self.hc_mult = args.hc_mult
+        self.hc_sinkhorn_iters = args.hc_sinkhorn_iters
+        self.hc_eps = args.hc_eps
+        self.attn = Attention(layer_id, args)
+        self.ffn = MoE(layer_id, args, strategy=strategy, decode_target=decode_target,
+                       quant_config=quant_config)
+        self.attn_norm = RMSNorm(args.dim, args.norm_eps)
+        self.ffn_norm = RMSNorm(args.dim, args.norm_eps)
+        self.engram = None
+        if layer_id in args.engram_layer_ids:
+            from .engram import Engram
+
+            self.engram = Engram(args, layer_id)
+        mix_hc = (2 + args.hc_mult) * args.hc_mult
+        for sublayer in ("attn", "ffn"):
+            for name, shape in (("fn", (mix_hc, args.hc_mult * args.dim)),
+                                ("base", (mix_hc,)), ("scale", (3,))):
+                self.register_parameter(f"hc_{sublayer}_{name}", nn.Parameter(
+                    torch.empty(shape, dtype=torch.float32), requires_grad=False))
+
+    def hc_mixes(self, x, hc_fn, hc_scale, hc_base):
+        flat = x.flatten(-2).float()
+        mixes = F.linear(flat, hc_fn) * torch.rsqrt(flat.square().mean(-1, keepdim=True) + self.norm_eps)
+        hc = self.hc_mult
+        if x.is_cuda:
+            from freetoken.kernel.triton.dsv4.sinkhorn import hc_split_sinkhorn
+
+            pre, post, comb = hc_split_sinkhorn(mixes.reshape(-1, mixes.shape[-1]), hc_scale,
+                                               hc_base, hc, self.hc_sinkhorn_iters, self.hc_eps)
+            return (pre.reshape(*x.shape[:-2], hc), post.reshape(*x.shape[:-2], hc),
+                    comb.reshape(*x.shape[:-2], hc, hc))
+        pre = torch.sigmoid(mixes[..., :hc] * hc_scale[0] + hc_base[:hc]) + self.hc_eps
+        post = 2 * torch.sigmoid(mixes[..., hc:2 * hc] * hc_scale[1] + hc_base[hc:2 * hc])
+        comb = (mixes[..., 2 * hc:] * hc_scale[2] + hc_base[2 * hc:]).reshape(*x.shape[:-2], hc, hc)
+        comb = comb.softmax(-1) + self.hc_eps
+        comb = comb / (comb.sum(-2, keepdim=True) + self.hc_eps)
+        for _ in range(self.hc_sinkhorn_iters - 1):
+            comb = comb / (comb.sum(-1, keepdim=True) + self.hc_eps)
+            comb = comb / (comb.sum(-2, keepdim=True) + self.hc_eps)
+        return pre, post, comb
+
+    def hc_pre(self, x, pre):
+        if x.is_cuda:
+            from freetoken.kernel.triton.dsv4.hc import hc_pre_combine
+
+            return hc_pre_combine(x.reshape(-1, self.hc_mult, self.dim),
+                                  pre.reshape(-1, self.hc_mult), x.dtype).reshape(*x.shape[:-2], self.dim)
+        return (pre.unsqueeze(-1) * x.float()).sum(-2).to(x.dtype)
+
+    def hc_post(self, x, residual, post, comb):
+        if x.is_cuda:
+            from freetoken.kernel.triton.dsv4.hc import hc_post_combine
+
+            return hc_post_combine(x.reshape(-1, self.dim), residual.reshape(-1, self.hc_mult, self.dim),
+                                   post.reshape(-1, self.hc_mult),
+                                   comb.reshape(-1, self.hc_mult, self.hc_mult)).reshape(residual.shape)
+        mixed = torch.einsum("...pq,...pd->...qd", comb.float(), residual.float())
+        return (post.unsqueeze(-1) * x.float().unsqueeze(-2) + mixed).to(x.dtype)
+
+    def _forward(self, h, pre_mix, image_mask, attention):
+        residual = h
+        attn_pre, post, comb = self.hc_mixes(h, self.hc_attn_fn, self.hc_attn_scale, self.hc_attn_base)
+        x = attention(self.attn_norm(self.hc_pre(h, pre_mix)))
+        h = self.hc_post(x, residual, post, comb)
+        residual = h
+        ffn_pre, post, comb = self.hc_mixes(h, self.hc_ffn_fn, self.hc_ffn_scale, self.hc_ffn_base)
+        x = self.ffn(self.ffn_norm(self.hc_pre(h, attn_pre)), image_mask)
+        return self.hc_post(x, residual, post, comb), ffn_pre
+
+    def prefill_batched(self, h, pre_mix, image_mask, segments, positions):
+        return self._forward(h, pre_mix, image_mask,
+                             lambda x: self.attn.forward_ragged(x, segments, positions))
+
+    def decode_step(self, h, pre_mix, pos, rows, stage_cap, wctx=None):
+        return self._forward(h, pre_mix, None,
+                             lambda x: self.attn.decode_step(x, pos, rows, stage_cap, wctx))
+
+
+class Transformer(nn.Module):
+    def __init__(self, args, *, strategy="offload", decode_target="gpu", quant_config=None):
+        super().__init__()
+        if quant_config is None:
+            from .config import DeepseekV41QuantConfig
+
+            quant_config = DeepseekV41QuantConfig(args)
+        self.args = args
+        self.hc_mult = args.hc_mult
+        self.embed = nn.Embedding(args.vocab_size, args.dim, dtype=torch.bfloat16)
+        self.embed.weight.requires_grad_(False)
+        self.layers = nn.ModuleList([Block(i, args, strategy=strategy, decode_target=decode_target,
+                                            quant_config=quant_config) for i in range(args.n_layers)])
+        self.norm = RMSNorm(args.dim, args.norm_eps)
+        self.head = OutputHead(args.dim, args.vocab_size)
+        self.vision = None
+        if args.vision_enabled:
+            from .vision import Aligner, ViT
+
+            self.vision, self.aligner = ViT(args), Aligner(args)
+            for name in ("image_start", "image_end", "image_newline"):
+                self.register_parameter(name, nn.Parameter(torch.empty(args.dim, dtype=torch.bfloat16),
+                                                          requires_grad=False))
+
+    def bind(self, pool, device):
+        for layer in self.layers:
+            layer.attn.bind(pool, device)
+
+    def encode_image(self, patches, n_vit_h, n_vit_w):
+        if self.vision is None:
+            raise ValueError("This DeepSeek-V4.1 checkpoint has no vision tower")
+        return self.aligner(self.vision(patches, n_vit_h, n_vit_w), n_vit_h, n_vit_w)
+
+    def prefill_batched(self, input_ids, segments, positions, last_indices, batch=None):
+        ids = input_ids.flatten()
+        if batch is not None and getattr(batch, "mm_embeds", None) is not None:
+            ids = ids.clamp(max=self.args.vocab_size - 1)
+        flat = self.embed(ids)
+        image_mask = None
+        if batch is not None and self.vision is not None:
+            from .vision import merge_image_embeddings
+
+            flat, image_mask = merge_image_embeddings(self, batch, flat)
+        h = flat.view(1, -1, self.args.dim).unsqueeze(-2).repeat(1, 1, self.hc_mult, 1)
+        pre_mix = make_identity_pre_mix(h, self.hc_mult)
+        for layer in self.layers:
+            if layer.engram is not None:
+                h = layer.engram(h)
+            h, pre_mix = layer.prefill_batched(h, pre_mix, image_mask, segments, positions)
+        h = self.norm(self.layers[-1].hc_pre(h, pre_mix))
+        return self.head(h[0, last_indices])
+
+    def decode(self, input_ids, pos, stage_cap):
+        rows = torch.arange(input_ids.shape[0], device=input_ids.device)
+        h = self.embed(input_ids).unsqueeze(-2).repeat(1, 1, self.hc_mult, 1)
+        pre_mix = make_identity_pre_mix(h, self.hc_mult)
+        for layer in self.layers:
+            if layer.engram is not None:
+                h = layer.engram(h)
+            h, pre_mix = layer.decode_step(h, pre_mix, pos, rows, stage_cap)
+        return self.head(self.norm(self.layers[-1].hc_pre(h, pre_mix))[:, -1])
+
+
+class DeepseekV41ForCausalLM(BaseLLMModel):
+    def __init__(self, config):
+        self._config = config
+        self._args = config.dsv41_args
+        self._transformer = Transformer(self._args, strategy=config.moe_strategy,
+                                        decode_target=config.decode_target, quant_config=config.quant)
+        # The engine walks BaseOP children; resident nn.Module weights use the adapter below.
+        self._offload_layers = list(self._iter_offload_moe_layers())
+        self._bound = False
+        self._engram_runtime = None
+
+    def _ensure_bound(self):
+        if not self._bound:
+            pool = get_global_ctx().kv_cache
+            self._transformer.bind(pool, pool.device)
+            self._bound = True
+
+    def mark_for_rebind(self):
+        self._bound = False
+
+    def place_encoder_weights(self, mode):
+        if self._transformer.vision is not None:
+            self._transformer.vision.place_weights(mode)
+
+    def encode(self, item):
+        if item.modality != "image" or self._transformer.vision is None:
+            raise ValueError("DeepSeek-V4.1 supports image items only when vision is enabled")
+        from .vision import image_span_embeddings
+
+        embeddings = image_span_embeddings(self._transformer, item.feature, item.n_vit_h, item.n_vit_w, item.types)
+        if embeddings.shape[0] != item.num_tokens:
+            raise ValueError("DeepSeek-V4.1 image embedding count does not match its offsets")
+        return embeddings
+
+    def _iter_offload_moe_layers(self):
+        for layer in self._transformer.layers:
+            yield layer.ffn.experts
+
+    def load_host_tables(self, engine_config):
+        from .engram import prepare_engram
+
+        return prepare_engram(self, engine_config)
+
+    @contextmanager
+    def forward_host_ctx(self, batch, use_graph):
+        if self._engram_runtime is None:
+            yield
+        else:
+            with self._engram_runtime.forward_host_ctx(batch, use_graph):
+                yield
+
+    def state_dict(self, *, prefix="", result=None):
+        result = {} if result is None else result
+        for name, param in self._transformer.named_parameters():
+            result[f"{prefix}.{name}" if prefix else name] = param
+        return result
+
+    def load_state_dict(self, state_dict, *, prefix="", _internal=False):
+        casted = {}
+        for name, param in self._transformer.named_parameters():
+            key = f"{prefix}.{name}" if prefix else name
+            if key not in state_dict:
+                raise RuntimeError(f"Missing DeepSeek-V4.1 weight: {key}")
+            tensor = state_dict.pop(key)
+            if tensor.shape != param.shape:
+                raise ValueError(f"DeepSeek-V4.1 weight {key} has shape {tuple(tensor.shape)}; "
+                                 f"expected {tuple(param.shape)}")
+            casted[name] = tensor.to(param.dtype)
+        if state_dict and not _internal:
+            raise RuntimeError(f"Unexpected DeepSeek-V4.1 weights: {list(state_dict)[:8]}")
+        self._transformer.load_state_dict(casted, assign=True, strict=False)
+
+    def forward(self):
+        self._ensure_bound()
+        batch = get_global_ctx().batch
+        md = batch.attn_metadata
+        input_ids = batch.input_ids.long()
+        if batch.is_prefill:
+            return self._transformer.prefill_batched(input_ids, md.segments,
+                                                     batch.positions.long(), md.last_indices.long(), batch)
+        size = batch.padded_size
+        pos = batch.positions.long().view(-1)[:size]
+        stage_cap = md.stage_width - 1 if torch.cuda.is_current_stream_capturing() else int(pos.max().item())
+        return self._transformer.decode(input_ids.view(size, 1), pos, stage_cap)
+
+
+__all__ = ["Block", "Transformer", "DeepseekV41ForCausalLM", "make_identity_pre_mix"]
diff --git a/python/freetoken/models/deepseek_v41/moe.py b/python/freetoken/models/deepseek_v41/moe.py
new file mode 100644
index 000000000..9ea1af4bf
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/moe.py
@@ -0,0 +1,98 @@
+"""V4.1 bias-only routing and clamped SwiGLU, with native NVFP4 offload banks."""
+
+from __future__ import annotations
+
+import torch
+import torch.nn.functional as F
+from torch import nn
+
+from freetoken.layers import OffloadMoELayer
+
+from .layers import Linear
+
+
+class Gate(nn.Module):
+    def __init__(self, layer_id, args):
+        super().__init__()
+        self.topk = args.n_activated_experts
+        self.score_func = args.score_func
+        self.gate_temp = args.gate_temp
+        self.norm_topk_prob = args.norm_topk_prob
+        self.route_scale = args.route_scale
+        self.weight = nn.Parameter(torch.empty(args.n_routed_experts, args.dim,
+                                               dtype=torch.bfloat16), requires_grad=False)
+        self.bias = nn.Parameter(torch.empty(args.n_routed_experts, dtype=torch.float32), requires_grad=False)
+        self.bias_vl = nn.Parameter(torch.empty_like(self.bias), requires_grad=False) if args.vision_enabled else None
+
+    def forward(self, x, image_mask=None):
+        scores = F.linear(x.float(), self.weight.float()) / self.gate_temp
+        if self.score_func == "softmax":
+            scores = scores.softmax(-1)
+        elif self.score_func == "sigmoid":
+            scores = scores.sigmoid()
+        else:
+            scores = F.softplus(scores).sqrt()
+        bias = self.bias
+        if image_mask is not None and self.bias_vl is not None:
+            bias = torch.where(image_mask.reshape(-1, 1), self.bias_vl, bias)
+        indices = (scores + bias).topk(self.topk, dim=-1).indices
+        weights = scores.gather(1, indices)
+        if self.norm_topk_prob and self.topk > 1:
+            weights = weights / (weights.sum(-1, keepdim=True) + 1e-20)
+        return weights * self.route_scale, indices
+
+
+def clamped_swiglu(gate, up, limit):
+    gate, up = gate.float(), up.float()
+    if limit > 0:
+        gate, up = gate.clamp(max=limit), up.clamp(-limit, limit)
+    return F.silu(gate) * up
+
+
+class Expert(nn.Module):
+    def __init__(self, dim, inter_dim, swiglu_limit):
+        super().__init__()
+        self.w1 = Linear(dim, inter_dim)
+        self.w2 = Linear(inter_dim, dim)
+        self.w3 = Linear(dim, inter_dim)
+        self.swiglu_limit = swiglu_limit
+
+    def forward(self, x):
+        return self.w2(clamped_swiglu(self.w1(x), self.w3(x), self.swiglu_limit).to(x.dtype))
+
+
+class DSV41OffloadMoELayer(OffloadMoELayer):
+    def __init__(self, layer_id, args, *, strategy="offload", decode_target="gpu", quant_config=None):
+        if quant_config is None:
+            from .config import DeepseekV41QuantConfig
+
+            quant_config = DeepseekV41QuantConfig(args)
+        super().__init__(layer_id=layer_id, num_experts=args.n_routed_experts,
+                         top_k=args.n_activated_experts, hidden_size=args.dim,
+                         intermediate_size=args.moe_inter_dim,
+                         renormalize=args.norm_topk_prob, activation="swiglu_clamp",
+                         alpha=1.0, limit=args.swiglu_limit,
+                         strategy=strategy, decode_target=decode_target,
+                         quant_config=quant_config, prefix=f"layers.{layer_id}.ffn.experts")
+
+
+class MoE(nn.Module):
+    def __init__(self, layer_id, args, *, strategy="offload", decode_target="gpu", quant_config=None):
+        super().__init__()
+        self.dim = args.dim
+        self.gate = Gate(layer_id, args)
+        self.shared_experts = Expert(args.dim, args.moe_inter_dim, args.swiglu_limit)
+        self.experts = DSV41OffloadMoELayer(layer_id, args, strategy=strategy,
+                                          decode_target=decode_target, quant_config=quant_config)
+
+    def forward(self, x, image_mask=None):
+        shape = x.shape
+        x = x.reshape(-1, self.dim)
+        weights, indices = self.gate(x, image_mask)
+        shared = self.shared_experts(x)
+        routed = self.experts.routed_forward(x, weights.float().contiguous(),
+                                             indices.to(torch.int32).contiguous())
+        return (routed.float() + shared.float()).to(x.dtype).view(shape)
+
+
+__all__ = ["Gate", "Expert", "MoE", "DSV41OffloadMoELayer"]
diff --git a/python/freetoken/models/deepseek_v41/vision.py b/python/freetoken/models/deepseek_v41/vision.py
new file mode 100644
index 000000000..5786960db
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/vision.py
@@ -0,0 +1,233 @@
+"""DeepSeek-V4.1 vision tower; upstream MIT attribution is in this package's NOTICE."""
+
+from functools import lru_cache
+
+import torch
+import torch.nn.functional as F
+from torch import nn
+
+from freetoken.layers import BaseOP
+
+
+class _ModuleBlockAdapter(BaseOP):
+    """Expose native module parameters to the shared block streamer's tensor slots."""
+
+    def __init__(self, module):
+        self._module = module
+        for name, parameter in module.named_parameters(recurse=False):
+            setattr(self, name, parameter)
+        for name, child in module.named_children():
+            setattr(self, name, _ModuleBlockAdapter(child))
+
+    def __setattr__(self, name, value):
+        module = self.__dict__.get("_module")
+        if module is not None and name in module._parameters:
+            value = value if isinstance(value, nn.Parameter) else nn.Parameter(value, requires_grad=False)
+            setattr(module, name, value)
+        object.__setattr__(self, name, value)
+
+    def forward(self, *args):
+        return self._module(*args)
+
+
+@lru_cache(8)
+def get_vision_cos_sin(n_h: int, n_w: int, dim: int, theta: float):
+    inv_freq = 1.0 / (theta ** (torch.arange(0, dim, 2, dtype=torch.float32, device="cpu") / dim))
+    hpos = torch.arange(n_h, device="cpu").unsqueeze(1).expand(n_h, n_w)
+    wpos = torch.arange(n_w, device="cpu").unsqueeze(0).expand(n_h, n_w)
+    freqs = torch.stack([hpos, wpos], dim=-1).reshape(-1, 2, 1).float() * inv_freq
+    freqs = freqs.flatten(1)
+    return freqs.cos().unsqueeze(1), freqs.sin().unsqueeze(1)
+
+
+def apply_rotary(x: torch.Tensor, cos: torch.Tensor, sin: torch.Tensor) -> torch.Tensor:
+    dtype = x.dtype
+    x1, x2 = x.float().chunk(2, dim=-1)
+    return torch.cat([x1 * cos - x2 * sin, x2 * cos + x1 * sin], dim=-1).to(dtype)
+
+
+class RMSNorm(nn.Module):
+    def __init__(self, dim: int, eps: float = 1e-6):
+        super().__init__()
+        self.eps = eps
+        self.weight = nn.Parameter(torch.ones(dim, dtype=torch.float32))
+
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        dtype = x.dtype
+        x = x.float()
+        x = x * torch.rsqrt(x.square().mean(-1, keepdim=True) + self.eps)
+        return (self.weight * x).to(dtype)
+
+
+class PatchEmbed(nn.Module):
+    def __init__(self, args):
+        super().__init__()
+        self.proj = nn.Linear(3 * args.vision_patch_size**2, args.vision_dim, dtype=torch.bfloat16)
+
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        return self.proj(x.flatten(1))
+
+
+class Attention(nn.Module):
+    def __init__(self, args):
+        super().__init__()
+        self.n_heads = args.vision_n_heads
+        self.head_dim = args.vision_dim // args.vision_n_heads
+        self.wqkv = nn.Linear(args.vision_dim, 3 * args.vision_dim, dtype=torch.bfloat16)
+        self.wo = nn.Linear(args.vision_dim, args.vision_dim, dtype=torch.bfloat16)
+
+    def forward(self, x: torch.Tensor, cos: torch.Tensor, sin: torch.Tensor) -> torch.Tensor:
+        n = x.size(0)
+        q, k, v = (t.view(n, self.n_heads, self.head_dim) for t in self.wqkv(x).chunk(3, dim=-1))
+        q = apply_rotary(q, cos, sin)
+        k = apply_rotary(k, cos, sin)
+        o = F.scaled_dot_product_attention(
+            q.transpose(0, 1).unsqueeze(0),
+            k.transpose(0, 1).unsqueeze(0),
+            v.transpose(0, 1).unsqueeze(0),
+        )
+        return self.wo(o.squeeze(0).transpose(0, 1).reshape(n, -1))
+
+
+class MLP(nn.Module):
+    def __init__(self, args):
+        super().__init__()
+        self.w1 = nn.Linear(args.vision_dim, 2 * args.vision_inter_dim, bias=False, dtype=torch.bfloat16)
+        self.w2 = nn.Linear(args.vision_inter_dim, args.vision_dim, bias=False, dtype=torch.bfloat16)
+
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        gate, up = self.w1(x).chunk(2, dim=-1)
+        return self.w2(F.silu(gate) * up)
+
+
+class Block(nn.Module):
+    def __init__(self, args):
+        super().__init__()
+        self.norm1 = RMSNorm(args.vision_dim)
+        self.attn = Attention(args)
+        self.norm2 = RMSNorm(args.vision_dim)
+        self.mlp = MLP(args)
+
+    def forward(self, x: torch.Tensor, cos: torch.Tensor, sin: torch.Tensor) -> torch.Tensor:
+        x = x + self.attn(self.norm1(x), cos, sin)
+        return x + self.mlp(self.norm2(x))
+
+
+class ViT(nn.Module):
+    """DeepSeek ViT: full bidirectional attention over one image with 2D RoPE."""
+
+    def __init__(self, args):
+        super().__init__()
+        self.rope_dim = args.vision_dim // args.vision_n_heads // 2
+        self.rope_theta = args.vision_rope_theta
+        self.patch_embed = PatchEmbed(args)
+        self.blocks = nn.ModuleList([Block(args) for _ in range(args.vision_n_layers)])
+        self.norm = RMSNorm(args.vision_dim)
+        self._streamer = None
+        self._stream_blocks = None
+
+    def place_weights(self, mode):
+        if mode not in {"host", "gpu"}:
+            raise ValueError(f"Unsupported vision weight placement: {mode}")
+        if mode == "gpu" and self._streamer is not None:
+            self._streamer.unstream()
+            self._streamer, self._stream_blocks = None, None
+        elif mode == "host" and self._streamer is None and self.blocks:
+            from freetoken.models.weight_stream import BlockWeightStreamer
+
+            device = self.patch_embed.proj.weight.device
+            self._stream_blocks = [_ModuleBlockAdapter(block) for block in self.blocks]
+            self._streamer = BlockWeightStreamer(self._stream_blocks, device)
+
+    def forward(self, patches: torch.Tensor, n_h: int, n_w: int) -> torch.Tensor:
+        if patches.shape[0] != n_h * n_w:
+            raise ValueError("vision patch count does not match its grid")
+        x = self.patch_embed(patches)
+        cos, sin = get_vision_cos_sin(n_h, n_w, self.rope_dim, self.rope_theta)
+        cos, sin = cos.to(x.device), sin.to(x.device)
+        blocks = (self._streamer.blocks(self._stream_blocks) if self._streamer is not None
+                  else enumerate(self.blocks))
+        try:
+            for _, block in blocks:
+                x = block.forward(x, cos, sin)
+        finally:
+            if self._streamer is not None:
+                blocks.close()
+        return self.norm(x)
+
+
+class Aligner(nn.Module):
+    def __init__(self, args):
+        super().__init__()
+        self.downsample_ratio = args.vision_downsample_ratio
+        in_dim = args.vision_dim * self.downsample_ratio**2
+        self.w1 = nn.Linear(in_dim, args.dim, dtype=torch.bfloat16)
+        self.w2 = nn.Linear(args.dim, args.dim, dtype=torch.bfloat16)
+
+    def forward(self, x: torch.Tensor, n_h: int, n_w: int) -> torch.Tensor:
+        r = self.downsample_ratio
+        x = x.view(n_h, n_w, -1).permute(2, 0, 1)
+        x = F.pad(x, (0, -n_w % r, 0, -n_h % r))
+        x = F.unfold(x.unsqueeze(0), r, stride=r).squeeze(0).transpose(0, 1)
+        return self.w2(F.gelu(self.w1(x)))
+
+
+@torch.inference_mode()
+def image_span_embeddings(transformer, patches, n_h, n_w, types):
+    """Encode the full native span, including learned image delimiters and row separators."""
+    from .image_processor import IMAGE, IMAGE_END, IMAGE_NEW_LINE, IMAGE_START
+
+    weight = transformer.vision.patch_embed.proj.weight
+    types = torch.as_tensor(types, device=weight.device, dtype=torch.int64)
+    if types.ndim != 1 or torch.any((types < IMAGE_START) | (types > IMAGE_END)):
+        raise ValueError("invalid DeepSeek-V4.1 image token types")
+    features = transformer.encode_image(patches.to(device=weight.device, dtype=weight.dtype), n_h, n_w)
+    if features.shape != (int((types == IMAGE).sum()), transformer.image_start.numel()):
+        raise ValueError("vision aligner output does not match image token span")
+    span = torch.empty((types.numel(), features.shape[-1]), device=weight.device, dtype=features.dtype)
+    span[types == IMAGE] = features
+    for kind, name in ((IMAGE_START, "image_start"), (IMAGE_NEW_LINE, "image_newline"), (IMAGE_END, "image_end")):
+        span[types == kind] = getattr(transformer, name).to(features.dtype)
+    return span
+
+
+@torch.inference_mode()
+def merge_image_embeddings(transformer, batch, hidden: torch.Tensor):
+    """Scatter canonical embedding rows and compatible native media into a prefill chunk."""
+    if not batch.is_prefill:
+        return hidden, None
+    embeds = getattr(batch, "mm_embeds", None)
+    legacy = any(getattr(req, "media", None) for req in batch.reqs)
+    if embeds is None and not legacy:
+        return hidden, None
+    mask = torch.zeros(hidden.shape[0], dtype=torch.bool, device=hidden.device)
+    if embeds is not None:
+        rows = batch.mm_rows
+        if embeds.shape != (rows.numel(), hidden.shape[-1]):
+            raise ValueError("DeepSeek-V4.1 multimodal embeddings have incompatible shape")
+        hidden.index_copy_(0, rows, embeds.to(device=hidden.device, dtype=hidden.dtype))
+        mask[rows] = True
+    if not legacy:
+        return hidden, mask
+    offset = 0
+    for req in batch.reqs:
+        begin, end = req.cached_len, req.cached_len + req.extend_len
+        for item in getattr(req, "media", None) or ():
+            start = item["start"]
+            types = item["types"]
+            stop = start + types.numel()
+            left, right = max(begin, start), min(end, stop)
+            if left >= right:
+                continue
+            if "embeddings" not in item:
+                span = image_span_embeddings(transformer, item["patches"], item["n_vit_h"], item["n_vit_w"], types)
+                # The shared request object survives chunking; retain only the small CPU result.
+                item["embeddings"] = span.cpu()
+                item["patches"] = None
+            dst = slice(offset + left - begin, offset + right - begin)
+            hidden[dst] = item["embeddings"][left - start:right - start].to(hidden.device, hidden.dtype)
+            mask[dst] = True
+        offset += req.extend_len
+    if offset != hidden.shape[0]:
+        raise ValueError("prefill token count does not match request spans")
+    return hidden, mask
diff --git a/python/freetoken/models/deepseek_v41/weight.py b/python/freetoken/models/deepseek_v41/weight.py
new file mode 100644
index 000000000..9da64effd
--- /dev/null
+++ b/python/freetoken/models/deepseek_v41/weight.py
@@ -0,0 +1,200 @@
+"""Native V4.1 resident weights and NVFP4 W4A16 expert banks.
+
+The converted experts use E4M3 scales per 16 values plus a global scale.
+W4A16 intentionally does not use the checkpoint's W4A4 input_scale metadata;
+it is not bit-equivalent to the original activation-quantized reference.
+"""
+
+from __future__ import annotations
+
+import json
+import os
+import re
+import struct
+
+import safetensors
+import torch
+
+from freetoken.layers.quantization import QuantKind
+from freetoken.models.loader import drop_page_cache
+from freetoken.models.nvfp4_banks import (
+    Nvfp4ExpertSourceSpec, iter_nvfp4_expert_pieces,
+)
+from freetoken.utils import download_hf_weight
+
+from .args import load_args
+
+_EXPERT_RE = re.compile(
+    r"^layers\.(?P<layer>\d+)\.ffn\.experts\.(?P<expert>\d+)\."
+    r"(?P<proj>w1|w2|w3)\.(?P<kind>weight|weight_scale|weight_scale_2)$"
+)
+_EXPERT_PREFIX_RE = re.compile(r"^layers\.(\d+)\.ffn\.experts\.")
+_ENGRAM_TABLE_RE = re.compile(r"^layers\.\d+\.engram\.embed\.(weight|scale)$")
+_VISION_PREFIXES = ("vision.", "aligner.", "image_")
+
+
+def _weight_map(folder):
+    with open(os.path.join(folder, "model.safetensors.index.json"), encoding="utf-8") as f:
+        return json.load(f)["weight_map"]
+
+
+def read_checkpoint_headers(folder, weight_map=None):
+    """Read only safetensors JSON headers, before any expert bank is allocated."""
+    weight_map = _weight_map(folder) if weight_map is None else weight_map
+    result = {}
+    for shard in sorted(set(weight_map.values())):
+        with open(os.path.join(folder, shard), "rb") as f:
+            length_bytes = f.read(8)
+            if len(length_bytes) != 8:
+                raise ValueError(f"Truncated safetensors header: {shard}")
+            length = struct.unpack("<Q", length_bytes)[0]
+            if length <= 0 or length > 100_000_000:
+                raise ValueError(f"Invalid safetensors header length: {shard}")
+            header = json.loads(f.read(length))
+        for name, item in header.items():
+            if name == "__metadata__":
+                continue
+            if name in result:
+                raise ValueError(f"Duplicate checkpoint tensor: {name}")
+            if weight_map.get(name) != shard:
+                raise ValueError(f"Checkpoint index/shard mismatch for {name}")
+            result[name] = item
+    missing = weight_map.keys() - result.keys()
+    if missing:
+        raise ValueError(f"Checkpoint index names missing tensor {min(missing)}")
+    return result
+
+
+def validate_expert_headers(headers, config):
+    """Validate every expert component and geometry without materializing a tensor."""
+    count = 0
+    layers, experts = config.num_layers, config.num_experts
+    hidden, inter = config.hidden_size, config.moe_intermediate_size
+    if hidden % 16 or inter % 16:
+        raise ValueError("NVFP4 expert dimensions must be divisible by 16")
+    for name, item in headers.items():
+        prefix = _EXPERT_PREFIX_RE.match(name)
+        if prefix is None:
+            continue
+        match = _EXPERT_RE.fullmatch(name)
+        if match is None:
+            if name.endswith(".input_scale"):
+                continue
+            raise ValueError(f"Unrecognized V4.1 expert tensor: {name}")
+        layer, expert = int(match["layer"]), int(match["expert"])
+        if not (0 <= layer < layers and 0 <= expert < experts):
+            raise ValueError(f"V4.1 expert outside configured backbone: {name}")
+        out_dim, in_dim = (hidden, inter) if match["proj"] == "w2" else (inter, hidden)
+        kind = match["kind"]
+        shape = tuple(item["shape"])
+        expected_shape = {"weight": (out_dim, in_dim // 2),
+                          "weight_scale": (out_dim, in_dim // 16)}.get(kind)
+        expected_dtype = {"weight": "U8", "weight_scale": "F8_E4M3", "weight_scale_2": "F32"}[kind]
+        if item["dtype"] != expected_dtype or (
+                shape != expected_shape if expected_shape is not None else shape not in ((), (1,))):
+            raise ValueError(f"Malformed NVFP4 tensor {name}: {item['dtype']} {shape}; "
+                             f"expected {expected_dtype} {expected_shape or 'scalar'}")
+        count += 1
+    expected = layers * experts * 9
+    if count != expected:
+        for layer in range(layers):
+            for expert in range(experts):
+                for proj in ("w1", "w2", "w3"):
+                    for kind in ("weight", "weight_scale", "weight_scale_2"):
+                        name = f"layers.{layer}.ffn.experts.{expert}.{proj}.{kind}"
+                        if name not in headers:
+                            raise ValueError(f"Missing NVFP4 tensor: {name}")
+        raise ValueError(f"Expected {expected} NVFP4 tensors, found {count}")
+
+
+def _dequant_fp8_block(weight, scale, block=32):
+    rows, cols = weight.shape
+    if scale.shape != (rows // block, cols // block):
+        raise ValueError(f"FP8 scale shape {tuple(scale.shape)} incompatible with weight {tuple(weight.shape)}")
+    values = torch.exp2(scale.view(torch.uint8).float() - 127)
+    values = values.repeat_interleave(block, 0).repeat_interleave(block, 1)
+    return (weight.float() * values).to(torch.bfloat16)
+
+
+class _ShardReader:
+    def __init__(self, folder, weight_map, device):
+        self.folder, self.weight_map, self.device = folder, weight_map, str(device)
+        self.handles = {}
+
+    def get(self, name):
+        if name not in self.weight_map:
+            raise ValueError(f"Missing DeepSeek-V4.1 tensor: {name}")
+        shard = self.weight_map[name]
+        if shard not in self.handles:
+            self.handles[shard] = safetensors.safe_open(os.path.join(self.folder, shard),
+                                                       framework="pt", device=self.device).__enter__()
+        return self.handles[shard].get_tensor(name)
+
+    def close(self):
+        for shard, handle in self.handles.items():
+            handle.__exit__(None, None, None)
+            drop_page_cache(os.path.join(self.folder, shard))
+        self.handles.clear()
+
+
+def iter_weights(model_path, device, *, include_moe_experts=True, include_non_moe=True, include_vision=True):
+    if include_moe_experts:
+        raise ValueError("DeepSeek-V4.1 NVFP4 experts require --moe-strategy offload")
+    if not include_non_moe:
+        return
+    folder = download_hf_weight(model_path)
+    args = load_args(folder)
+    weight_map = _weight_map(folder)
+    reader = _ShardReader(folder, weight_map, device)
+    try:
+        for name in sorted(weight_map, key=lambda name: (weight_map[name], name)):
+            if name.startswith("mtp.") or _EXPERT_PREFIX_RE.match(name) or _ENGRAM_TABLE_RE.match(name):
+                continue
+            if (not include_vision or not args.vision_enabled) and name.startswith(_VISION_PREFIXES):
+                continue
+            if name.endswith(".attn.wo_a.scale"):
+                continue
+            if name.endswith(".attn.wo_a.weight"):
+                prefix = name.removesuffix(".weight")
+                yield prefix, _dequant_fp8_block(reader.get(name), reader.get(prefix + ".scale"))
+            else:
+                yield name, reader.get(name)
+    finally:
+        reader.close()
+
+
+def iter_vision_weights(model_path, device):
+    """Read the native vision stack, including its learned image delimiters."""
+    folder = download_hf_weight(model_path)
+    weight_map = _weight_map(folder)
+    reader = _ShardReader(folder, weight_map, device)
+    try:
+        for name in sorted(weight_map, key=lambda name: (weight_map[name], name)):
+            if name.startswith(_VISION_PREFIXES):
+                yield name, reader.get(name)
+    finally:
+        reader.close()
+
+
+def is_expert_tensor(name):
+    return _EXPERT_RE.fullmatch(name) is not None
+
+
+_NVFP4_SPEC = Nvfp4ExpertSourceSpec(
+    key_pattern=_EXPERT_RE, proj_to_role={"w1": "gate", "w3": "up", "w2": "down"},
+    layer_to_bank=lambda layer, config: layer if layer < config.num_layers else None,
+    desc="DeepSeek-V4.1 NVFP4 experts",
+)
+
+
+def iter_expert_pieces(model_path, config, kind, *, parallel=False, workers=8, chunk=8 << 20):
+    if kind is not QuantKind.NVFP4:
+        return None
+    folder = download_hf_weight(model_path)
+    validate_expert_headers(read_checkpoint_headers(folder), config)
+    return iter_nvfp4_expert_pieces(folder, config, _NVFP4_SPEC, parallel=parallel,
+                                    workers=workers, chunk=chunk, drop_page_cache=drop_page_cache)
+
+
+__all__ = ["iter_weights", "iter_vision_weights", "is_expert_tensor", "iter_expert_pieces",
+           "read_checkpoint_headers", "validate_expert_headers"]
diff --git a/python/freetoken/models/register.py b/python/freetoken/models/register.py
index 2a4ced10a..b5e7bc168 100644
--- a/python/freetoken/models/register.py
+++ b/python/freetoken/models/register.py
@@ -29,6 +29,7 @@ class ModelSpec:
     packed_modules_mapping: tuple[tuple[str, tuple[str, ...]], ...] = ()
     # checkpoint-name globs the family serves in bf16 although the quantization_config covers them
     unquantized_modules: tuple[str, ...] = ()
+    quant_config: str | None = None
     # "module:Class" turning the checkpoint's media into items; None: the family takes no multimodal input
     mm_processor: str | None = None
     encoders: tuple[EncoderSpec, ...] = ()
@@ -156,6 +157,13 @@ class ModelSpec:
         # the head, the KV compressors and the indexer's scorer ship bf16; the fp8 config has no modules_to_not_convert
         unquantized_modules=("head", "*.compressor.wkv", "*.compressor.wgate", "*.indexer.weights_proj"),
     ),
+    "DeepseekV41ForCausalLM": ModelSpec(
+        "freetoken.models.deepseek_v41",
+        "DeepseekV41ForCausalLM",
+        quant_config="checkpoint_quant_config",
+        mm_processor="freetoken.models.deepseek_v41.mm_processor:DeepseekV41MMProcessor",
+        encoders=(EncoderSpec("vision", "vision_config", ("image",)),),
+    ),
     "Qwen3_5MoeForConditionalGeneration": ModelSpec(
         "freetoken.models.qwen3_5_moe",
         "Qwen3_5MoeForConditionalGeneration",
@@ -320,6 +328,9 @@ def _load_attr(module_path: str, attr_name: str) -> Any:
 
 def checkpoint_quant_config(model_path: str, hf_config: Any, spec: ModelSpec):
     """The checkpoint's QuantConfig under the family's naming, or None for GGUF, whose native-quant ops the shared parser does not model yet."""
+    if spec.quant_config is not None:
+        return _load_attr(spec.module, spec.quant_config)(hf_config)
+
     from freetoken.layers.quantization import NameMap, QuantConfig
 
     if spec.parse_config == "parse_gguf_config":
diff --git a/python/freetoken/models/weight_stream.py b/python/freetoken/models/weight_stream.py
index ac529eca3..f5e5e5497 100644
--- a/python/freetoken/models/weight_stream.py
+++ b/python/freetoken/models/weight_stream.py
@@ -112,11 +112,13 @@ def blocks(self, ops: Sequence[BaseOP]) -> Iterator[tuple[int, BaseOP]]:
             slots, _ = self._layouts[i]
             for s in slots:
                 setattr(s.owner, s.attr, s.view(self.staging[buf]))
-            yield i, op
-            self.release_events[buf].record(compute)
-            self._has_release[buf] = True
-            for s in slots:
-                setattr(s.owner, s.attr, s.view(self.bank[i]))
+            try:
+                yield i, op
+            finally:
+                self.release_events[buf].record(compute)
+                self._has_release[buf] = True
+                for s in slots:
+                    setattr(s.owner, s.attr, s.view(self.bank[i]))
 
 
 __all__ = ["BlockWeightStreamer"]
diff --git a/python/freetoken/scheduler/cache.py b/python/freetoken/scheduler/cache.py
index 8f77be933..53378a90e 100644
--- a/python/freetoken/scheduler/cache.py
+++ b/python/freetoken/scheduler/cache.py
@@ -29,6 +29,11 @@ def _swa_eviction_interval() -> int:
 _SWA_RETAIN_GAP = 16
 
 
+def _has_unkeyed_media(req: PendingReq | Req) -> bool:
+    # Native media placeholders lack content hashes; MMItem pads already key the image.
+    return getattr(req, "mm_embeds", None) is not None or bool(getattr(req, "media", None))
+
+
 class CacheManager:
     def __init__(self, num_pages: int, page_size: int, page_table: torch.Tensor, type: str,
                  linear_state_pool=None, swa_pool=None, sliding_window_size=None):
@@ -93,7 +98,7 @@ def _make_prefix_cache(self, device, page_size, type):
     def match_req(self, req: PendingReq) -> MatchResult:
         input_len = req.input_len
         assert input_len > 0, "Input length must be greater than 0."
-        ids = req.input_ids[: input_len - 1]
+        ids = req.input_ids[:0] if _has_unkeyed_media(req) else req.input_ids[: input_len - 1]
         if self.is_swa:
             from freetoken.kvcache.swa_radix_cache import SWACacheHandle
             m = self.prefix_cache.match_prefix(ids)
@@ -296,6 +301,14 @@ def cache_req(self, req: Req, *, finished: bool) -> None:
         #                                           We should free it if the request has finished.
         page_indices = self.page_table[req.table_idx, : req.cached_len]
         old_handle = req.cache_handle
+        if _has_unkeyed_media(req):
+            if finished:
+                self.unlock(old_handle)
+                tail = self._padded_tail(req, old_handle.cached_len)
+                if self.swa_paged:
+                    self._free_swa(tail)
+                self._free(tail)
+            return
         insert_ids = req.input_ids[: req.cached_len]
         cached_len, new_handle = self.prefix_cache.insert_prefix(insert_ids, page_indices)
         # unlock until all operations on handle is done
@@ -336,6 +349,13 @@ def _cache_req_hybrid(self, req: Req, *, finished: bool) -> None:
         old_handle = req.cache_handle
         page_indices = self.page_table[req.table_idx, : req.cached_len]
 
+        if _has_unkeyed_media(req):
+            if finished:
+                self.unlock(old_handle)
+                self._free(page_indices[old_handle.cached_len :])
+                self._free_req_slots(req)
+            return
+
         if finished:
             # A pending freeze (the tool-call anchor, or a prefill ×64 track the request
             # finished too early to chunk-commit) is a strictly shorter prefix than the live
@@ -423,6 +443,14 @@ def _cache_req_swa(self, req: Req, *, finished: bool) -> None:
         old_handle = req.cache_handle
         page_indices = self.page_table[req.table_idx, : req.cached_len]
 
+        if _has_unkeyed_media(req):
+            if finished:
+                self.unlock(old_handle)
+                tail = self._padded_tail(req, old_handle.cached_len)
+                self._free_swa(tail)
+                self._free(tail)
+            return
+
         insert_len = align_down(req.cached_len, self.page_size)
         freed = page_indices[:0]
         if insert_len > 0:
diff --git a/python/freetoken/scheduler/mm.py b/python/freetoken/scheduler/mm.py
index 8c656c472..232332ea1 100644
--- a/python/freetoken/scheduler/mm.py
+++ b/python/freetoken/scheduler/mm.py
@@ -4,6 +4,8 @@
 
 from typing import TYPE_CHECKING, List, Tuple
 
+import torch
+
 from freetoken.utils import align_down
 
 if TYPE_CHECKING:
@@ -88,4 +90,28 @@ def plan_mm_batch(reqs, encoder_cache: EncoderCache | None) -> tuple[List[MMItem
     return jobs, plan, rows, block_ends
 
 
-__all__ = ["cut_image_spans", "mm_chunk_end", "mm_rows_after", "plan_mm_batch", "plan_mm_chunk"]
+def gather_legacy_mm_batch(reqs, image_token_id: int | None) -> tuple[List[torch.Tensor], List[int]]:
+    """Slice request-owned embeddings at this chunk's original image placeholders."""
+    parts: List[torch.Tensor] = []
+    rows: List[int] = []
+    offset = 0
+    for req in reqs:
+        embeds = getattr(req, "mm_embeds", None)
+        if embeds is not None:
+            if image_token_id is None:
+                raise ValueError("precomputed mm_embeds require an image_token_id")
+            if embeds.ndim != 2:
+                raise ValueError("precomputed mm_embeds must have shape [image_tokens, hidden]")
+            start = int((req.input_ids[:req.cached_len] == image_token_id).sum())
+            positions = (req.input_ids[req.cached_len:req.device_len] == image_token_id).nonzero().flatten()
+            end = start + positions.numel()
+            if end > embeds.shape[0]:
+                raise ValueError("image-token slots exceed precomputed mm_embeds rows")
+            if end > start:
+                parts.append(embeds[start:end])
+                rows.extend((positions + offset).tolist())
+        offset += req.extend_len
+    return parts, rows
+
+
+__all__ = ["cut_image_spans", "gather_legacy_mm_batch", "mm_chunk_end", "mm_rows_after", "plan_mm_batch", "plan_mm_chunk"]
diff --git a/python/freetoken/scheduler/prefill.py b/python/freetoken/scheduler/prefill.py
index bd8129aae..8b5c420c4 100644
--- a/python/freetoken/scheduler/prefill.py
+++ b/python/freetoken/scheduler/prefill.py
@@ -214,6 +214,8 @@ def _add_one_req(
             uid=pending_req.uid,
             cache_handle=cache_handle,
             sampling_params=pending_req.sampling_params,
+            mm_embeds=pending_req.mm_embeds,
+            media=pending_req.media,
         )
         req.mm_items = pending_req.mm_items
         req.mrope_positions_full = pending_req.mrope_positions_full
@@ -283,6 +285,8 @@ def add_one_req(self, req: UserMsg) -> None:
                 req.uid,
                 req.input_ids,
                 req.sampling_params,
+                mm_embeds=req.mm_embeds,
+                media=req.media,
                 mm_items=req.mm_items,
                 mrope_positions_full=req.mrope_positions,
                 mrope_delta=req.mrope_delta,
diff --git a/python/freetoken/scheduler/scheduler.py b/python/freetoken/scheduler/scheduler.py
index ac1bf322e..792e6298b 100644
--- a/python/freetoken/scheduler/scheduler.py
+++ b/python/freetoken/scheduler/scheduler.py
@@ -32,7 +32,7 @@
 from .config import SchedulerConfig
 from .decode import DecodeManager
 from .io import SchedulerIOMixin
-from .mm import cut_image_spans, plan_mm_batch
+from .mm import cut_image_spans, gather_legacy_mm_batch, plan_mm_batch
 from .prefill import ChunkedReq, PrefillManager
 from .status import SchedulerStatusReporter
 from .table import TableManager
@@ -739,7 +739,7 @@ def _current_cache_geometry(self) -> dict:
         config = self.config
         mc = config.model_config
         num_swa_pages = None
-        if getattr(mc, "dsv4_args", None) is not None:
+        if getattr(mc, "dsv4_args", None) is not None or getattr(mc, "dsv41_args", None) is not None:
             sizes = getattr(eng.kv_cache, "sizes", None)
             if sizes is not None:  # usable window pages = physical n_win_pages minus the dummy page
                 num_swa_pages = max(0, sizes.n_win_pages - 1)
@@ -856,8 +856,13 @@ def _gather_multimodal(self, batch: Batch) -> None:
         if plan:
             batch.mm_encoder_jobs = jobs
             batch.mm_gather_plan = plan
-            batch.mm_rows = torch.tensor(rows, dtype=torch.int64, pin_memory=True).to(self.device, non_blocking=True)
-            batch.mm_block_ends = torch.tensor(block_ends, dtype=torch.int32, pin_memory=True).to(self.device, non_blocking=True)
+            batch.mm_block_ends = torch.tensor(block_ends, dtype=torch.int32, pin_memory=torch.cuda.is_available()).to(self.device, non_blocking=True)
+        parts, legacy_rows = gather_legacy_mm_batch(batch.padded_reqs, self.config.model_config.image_token_id)
+        if parts:
+            batch.mm_embeds = torch.cat([part.to(self.device) for part in parts], dim=0)
+            rows.extend(legacy_rows)
+        if rows:
+            batch.mm_rows = torch.tensor(rows, dtype=torch.int64, pin_memory=torch.cuda.is_available()).to(self.device, non_blocking=True)
         if self._bidirectional_mm and not self._warned_cut_image and (cut := cut_image_spans(batch.padded_reqs)):
             # only a bidirectional image span loses context when cut, and only an image longer than the chunk still gets cut
             lo, hi = cut[0]
diff --git a/python/freetoken/scheduler/utils.py b/python/freetoken/scheduler/utils.py
index 83b79f2a3..fd6a3cd1e 100644
--- a/python/freetoken/scheduler/utils.py
+++ b/python/freetoken/scheduler/utils.py
@@ -17,6 +17,8 @@ class PendingReq:
     input_ids: torch.Tensor
     sampling_params: SamplingParams
     chunked_req: ChunkedReq | None = None
+    mm_embeds: torch.Tensor | None = None
+    media: list[dict] | None = None
     mm_items: list | None = None
     mrope_positions_full: torch.Tensor | None = None
     mrope_delta: int = 0
diff --git a/python/freetoken/server/anthropic_api.py b/python/freetoken/server/anthropic_api.py
index 0d60940c9..8bb74254e 100644
--- a/python/freetoken/server/anthropic_api.py
+++ b/python/freetoken/server/anthropic_api.py
@@ -117,6 +117,10 @@ async def handle_anthropic_messages(
         spec = convert_anthropic_to_genspec(
             req, model_sampling,
             reasoning_parser=getattr(state.config, "reasoning_parser", None),
+            preserve_tool_images=(
+                getattr(getattr(state.config, "model_spec", None), "model_cls", None)
+                == "DeepseekV41ForCausalLM"
+            ),
             default_max_tokens=(
                 getattr(state.config, "max_output_tokens", None) or DEFAULT_MAX_OUTPUT_TOKENS
             ),
@@ -153,7 +157,11 @@ async def handle_anthropic_count_tokens(req: AnthropicCountTokensRequest, state:
     # so it must not fall into the convert/empty-prompt ValueError branch.
     try:
         messages, template_tools, _, ctk = convert_anthropic_prompt(
-            req, reasoning_parser=getattr(state.config, "reasoning_parser", None)
+            req, reasoning_parser=getattr(state.config, "reasoning_parser", None),
+            preserve_tool_images=(
+                getattr(getattr(state.config, "model_spec", None), "model_cls", None)
+                == "DeepseekV41ForCausalLM"
+            ),
         )
     except ValueError as exc:
         return _anthropic_error_response(400, "invalid_request_error", str(exc))
@@ -182,6 +190,7 @@ async def handle_anthropic_count_tokens(req: AnthropicCountTokensRequest, state:
 def convert_anthropic_prompt(
     req: AnthropicMessagesRequest | AnthropicCountTokensRequest,
     reasoning_parser: str | None = None,
+    preserve_tool_images: bool = False,
 ) -> tuple[list[dict[str, Any]], list[dict[str, Any]] | None, list[dict[str, Any]] | None, dict[str, Any]]:
     """(messages, template_tools, parser_tools, chat_template_kwargs) — the prompt
     side of the conversion, shared by /v1/messages and /v1/messages/count_tokens so
@@ -232,17 +241,20 @@ def convert_anthropic_prompt(
                     }
                 )
             elif block.type == "tool_result":
-                text, images = _tool_result_parts(block.content)
+                content = _tool_result_content(block.content)
+                images = [p for p in content if p["type"] == "image"] if isinstance(content, list) else []
+                text = "".join(p["text"] for p in content if p["type"] == "text") if images else content
                 if msg.role == "user":
                     other.append(
                         {
                             "role": "tool",
                             "tool_call_id": block.tool_use_id or block.id or "",
-                            "content": text,
+                            "content": content if preserve_tool_images else text,
                         }
                     )
-                    # Chat templates render tool messages as text, so the images ride on the user turn that follows the tool messages (vLLM does the same).
-                    content_parts.extend(images)
+                    # Templates with text-only tool messages need images on the following user turn.
+                    if not preserve_tool_images:
+                        content_parts.extend(images)
                 else:
                     if images:
                         raise ValueError("images inside a tool_result are only accepted in a user message")
@@ -259,7 +271,7 @@ def convert_anthropic_prompt(
             else:
                 openai_msg["content"] = content_parts
         elif not tool_calls and not thinking_parts:
-            # Nothing usable in this message (e.g. image-only) — skip it.
+            # Opaque blocks such as redacted_thinking contain no model input.
             continue
         other.append(openai_msg)
 
@@ -308,9 +320,10 @@ def convert_anthropic_to_genspec(
     model_sampling: dict[str, Any],
     reasoning_parser: str | None = None,
     default_max_tokens: int = DEFAULT_MAX_OUTPUT_TOKENS,
+    preserve_tool_images: bool = False,
 ) -> GenSpec:
     messages, template_tools, parser_tools, ctk = convert_anthropic_prompt(
-        req, reasoning_parser=reasoning_parser
+        req, reasoning_parser=reasoning_parser, preserve_tool_images=preserve_tool_images
     )
     return GenSpec(
         messages=messages,
@@ -352,23 +365,24 @@ def _image_part(source: dict[str, Any] | None) -> dict[str, Any]:
     }
 
 
-def _tool_result_parts(content) -> tuple[str, list[dict[str, Any]]]:
-    """The text of a tool_result and its image blocks as image parts."""
+def _tool_result_content(content) -> str | list[dict[str, Any]]:
+    """Keep image-bearing tool results ordered for encoders that support them."""
     if content is None:
-        return "", []
+        return ""
     if isinstance(content, str):
-        return content, []
-    texts: list[str] = []
-    images: list[dict[str, Any]] = []
+        return content
+    parts: list[dict[str, Any]] = []
+    has_image = False
     for item in content:
         if isinstance(item, dict):
             if item.get("type") == "image":
-                images.append(_image_part(item.get("source")))
+                parts.append(_image_part(item.get("source")))
+                has_image = True
             else:
-                texts.append(item.get("text") or "")
+                parts.append({"type": "text", "text": item.get("text") or ""})
         else:
-            texts.append(str(item))
-    return "".join(texts), images
+            parts.append({"type": "text", "text": str(item)})
+    return parts if has_image else "".join(p["text"] for p in parts)
 
 
 # --------------------------------------------------------------------------- #
diff --git a/python/freetoken/server/api_models.py b/python/freetoken/server/api_models.py
index ffd717280..7a77e50d3 100644
--- a/python/freetoken/server/api_models.py
+++ b/python/freetoken/server/api_models.py
@@ -77,7 +77,7 @@ class ChatCompletionRequest(BaseModel):
     presence_penalty: float = 0.0
     frequency_penalty: float = 0.0
     chat_template_kwargs: dict[str, Any] = Field(default_factory=dict)
-    reasoning_effort: str | None = None
+    reasoning_effort: str | int | None = None
     # DeepSeek-wire thinking toggle ({"type": "enabled"|"disabled"}). Any so a
     # foreign shape stays ignored (extra="allow" swallowed it before this field
     # existed) instead of becoming a bare 422 at the route boundary; the handler
diff --git a/python/freetoken/server/args.py b/python/freetoken/server/args.py
index ab2fb9b74..3f2ff0d57 100644
--- a/python/freetoken/server/args.py
+++ b/python/freetoken/server/args.py
@@ -159,6 +159,15 @@ def _positive_int(value: str) -> int:
             raise argparse.ArgumentTypeError("must be >= 1")
         return n
 
+    def _window_ratio(value: str) -> float:
+        try:
+            ratio = float(value)
+        except ValueError as exc:
+            raise argparse.ArgumentTypeError("must be a number in (0, 1]") from exc
+        if not 0 < ratio <= 1:
+            raise argparse.ArgumentTypeError("must be in (0, 1]")
+        return ratio
+
     def _lazy_gpu_arg(value: str) -> tuple[str, ...]:
         from freetoken.gpu_select import gpu_arg
 
@@ -203,6 +212,8 @@ def _infer_tool_call_parser(model_path: str) -> str:
             return "qwen3_coder"
         if "qwen" in marker:
             return "qwen25"
+        if "deepseek" in marker and any(tag in marker for tag in ("v41", "v4.1", "v4_1")):
+            return "deepseekv41"
         if "deepseek" in marker and ("v4" in marker or "deepseek_v4" in marker):
             return "deepseekv32"
         if "deepseek" in marker and ("v3.2" in marker or "v32" in marker):
@@ -414,6 +425,31 @@ def _infer_reasoning_parser(model_path: str) -> str | None:
         help="Set the page size for system management.",
     )
 
+    parser.add_argument(
+        "--swa-full-tokens-ratio",
+        type=_window_ratio,
+        default=ServerArgs.swa_full_tokens_ratio,
+        help="Window-pool tokens per full-history token for paged sliding-window caches, in (0, 1].",
+    )
+
+    parser.add_argument(
+        "--kv-cache-dtype",
+        dest="kv_quant",
+        type=str,
+        default=ServerArgs.kv_quant,
+        choices=["auto", "bf16", "fp8", "nvfp4", "fp8-fp4"],
+        help=(
+            "KV-cache storage format. 'bf16' (default) stores the compute dtype; 'fp8'"
+            " stores e4m3 codes plus one fp32 scale per (token, kv head), roughly "
+            "doubling the tokens that fit in the same VRAM. Requires a compatible"
+            " Triton attention backend and a paged FULL, hybrid-SWA, QSA, or MLA/DSA pool."
+            " 'nvfp4' stores packed E2M1 with block/row scales; supports only"
+            " paged FULL, hybrid-SWA, QSA, or MLA/DSA attention with head_dim divisible by 16."
+            " 'fp8-fp4' preserves DeepSeek-V4.1's native FP8 window, FP4 compressed KV,"
+            " and MXFP4 index keys in packed storage; requires bfloat16 compute."
+        ),
+    )
+
     parser.add_argument(
         "--attention-backend",
         "--attn",
@@ -556,6 +592,7 @@ def _infer_reasoning_parser(model_path: str) -> str | None:
             "qwen3_coder",
             "mistral",
             "deepseekv32",
+            "deepseekv41",
             "gemma4",
             "glm47",
             "minimax",
diff --git a/python/freetoken/server/function_call_parser.py b/python/freetoken/server/function_call_parser.py
index 0804dcc3f..5f864e4fc 100644
--- a/python/freetoken/server/function_call_parser.py
+++ b/python/freetoken/server/function_call_parser.py
@@ -1563,13 +1563,17 @@ class DeepSeekV32Detector(BaseFormatDetector):
     Reference: https://huggingface.co/deepseek-ai/DeepSeek-V3.2
     """
 
+    _dsml_token = "｜DSML｜"
+    _block_name = "function_calls"
+    _alt_block_name = "tool_calls"
+
     def __init__(self):
         super().__init__()
-        self.dsml_token = "｜DSML｜"
-        self.bot_token = f"<{self.dsml_token}function_calls>"
-        self.eot_token = f"</{self.dsml_token}function_calls>"
-        self.alt_bot_token = f"<{self.dsml_token}tool_calls>"
-        self.alt_eot_token = f"</{self.dsml_token}tool_calls>"
+        self.dsml_token = self._dsml_token
+        self.bot_token = f"<{self.dsml_token}{self._block_name}>"
+        self.eot_token = f"</{self.dsml_token}{self._block_name}>"
+        self.alt_bot_token = f"<{self.dsml_token}{self._alt_block_name}>"
+        self.alt_eot_token = f"</{self.dsml_token}{self._alt_block_name}>"
         self.invoke_start_prefix = f"<{self.dsml_token}invoke"
         self.invoke_end_token = f"</{self.dsml_token}invoke>"
         self.param_end_token = f"</{self.dsml_token}parameter>"
@@ -1883,6 +1887,14 @@ def finish_streaming(self) -> str:
         return residual
 
 
+class DeepSeekV41Detector(DeepSeekV32Detector):
+    """V4.1 uses a space after the DSML marker and names its outer block calls."""
+
+    _dsml_token = "｜DSML｜ "
+    _block_name = "calls"
+    _alt_block_name = "calls"
+
+
 class Qwen3CoderDetector(InvokeParamStreamMixin, BaseFormatDetector):
     toolcall_opener = "<tool_call>"
     _ps_trim = "\n"
@@ -3523,6 +3535,7 @@ class FunctionCallParser:
 
     ToolCallParserEnum: Dict[str, Type[BaseFormatDetector]] = {
         "deepseekv32": DeepSeekV32Detector,
+        "deepseekv41": DeepSeekV41Detector,
         "gemma4": Gemma4Detector,
         "gpt-oss": GptOssDetector,
         "gpt_oss": GptOssDetector,
diff --git a/python/freetoken/server/generation.py b/python/freetoken/server/generation.py
index 91056e8d1..4fcaef1e4 100644
--- a/python/freetoken/server/generation.py
+++ b/python/freetoken/server/generation.py
@@ -192,9 +192,7 @@ def pick(value, key, framework):
 
 
 def render_messages(messages: list[dict[str, Any]]) -> list[dict[str, Any]]:
-    """Normalize OpenAI-shaped message dicts for the chat template: flatten text
-    content parts to a string and decode tool-call arguments from JSON. Raises
-    ValueError on a non-text content part (text-only server). Shared by all adapters."""
+    """Normalize text/tool messages while preserving ordered image content."""
     return [_render_message(m) for m in messages]
 
 
@@ -252,7 +250,15 @@ def _flatten_text_parts(parts: list[Any]) -> str | list[dict[str, Any]]:
             out.append({"type": "image", "freetoken_ref": {"kind": "url", "data": url}})
             has_image = True
         elif ptype == "image" and isinstance(part.get("freetoken_ref"), dict):
-            out.append(part)
+            out.append(dict(part))
+            has_image = True
+        elif ptype == "image":
+            source = part.get("source") or {}
+            kind = "b64" if source.get("type") == "base64" else "url"
+            data = source.get("data") if kind == "b64" else (source.get("url") or part.get("url"))
+            if not isinstance(data, str) or not data:
+                raise ValueError("image content part carries no source")
+            out.append({"type": "image", "freetoken_ref": {"kind": kind, "data": data}})
             has_image = True
         else:
             raise ValueError(f"Unsupported content part type: {ptype}")
diff --git a/python/freetoken/server/openai_api.py b/python/freetoken/server/openai_api.py
index dc2f73a97..5851d3821 100644
--- a/python/freetoken/server/openai_api.py
+++ b/python/freetoken/server/openai_api.py
@@ -167,6 +167,8 @@ async def handle_chat_completion(
     # Case/whitespace and the "off" disable synonym stay accepted here because
     # effort_toggle_kwargs normalizes and honors them downstream.
     effort = req.reasoning_effort.strip().lower() if isinstance(req.reasoning_effort, str) else None
+    if type(req.reasoning_effort) is int and not 1 <= req.reasoning_effort <= 100:
+        return create_error_response("numeric reasoning_effort must be between 1 and 100", param="reasoning_effort")
     if effort and effort not in _ACCEPTED_EFFORTS:
         return create_error_response(
             f"reasoning_effort must be one of {', '.join(_ACCEPTED_EFFORTS)}; "
diff --git a/python/freetoken/server/responses_api.py b/python/freetoken/server/responses_api.py
index 71d6eed87..28d3eb491 100644
--- a/python/freetoken/server/responses_api.py
+++ b/python/freetoken/server/responses_api.py
@@ -155,6 +155,10 @@ async def handle_responses(
         spec = convert_responses_to_genspec(
             req, model_sampling, default_max_tokens=default_max,
             reasoning_parser=getattr(state.config, "reasoning_parser", None),
+            preserve_tool_images=(
+                getattr(getattr(state.config, "model_spec", None), "model_cls", None)
+                == "DeepseekV41ForCausalLM"
+            ),
         )
         uid = await submit_generation(spec, state)
     except GenerationError as exc:
@@ -190,6 +194,7 @@ def convert_responses_to_genspec(
     model_sampling: dict[str, Any],
     default_max_tokens: int = DEFAULT_MAX_OUTPUT_TOKENS,
     reasoning_parser: str | None = None,
+    preserve_tool_images: bool = False,
 ) -> GenSpec:
     # Collect every system/developer text — the top-level `instructions` PLUS any
     # system/developer-role input items (codex sends both: a system prompt as `instructions`
@@ -206,7 +211,7 @@ def convert_responses_to_genspec(
         other.append({"role": "user", "content": req.input})
     else:
         for item in req.input:
-            for m in _convert_input_item(item):
+            for m in _convert_input_item(item, preserve_tool_images=preserve_tool_images):
                 if m.get("role") == "system":
                     system_texts.append(m.get("content") or "")
                 else:
@@ -254,7 +259,7 @@ def convert_responses_to_genspec(
     )
 
 
-def _convert_input_item(item: dict[str, Any]) -> list[dict[str, Any]]:
+def _convert_input_item(item: dict[str, Any], *, preserve_tool_images: bool = False) -> list[dict[str, Any]]:
     itype = item.get("type", "message")
     if itype == "message" or ("role" in item and "type" not in item):
         # codex sends a "developer" role (Responses instructions). Chat templates only
@@ -280,10 +285,9 @@ def _convert_input_item(item: dict[str, Any]) -> list[dict[str, Any]]:
             }
         ]
     if itype == "function_call_output":
-        output = item.get("output")
-        content = _input_content(output) if isinstance(output, list) else _stringify(output)
+        content = _tool_output(item.get("output"))
         tool_msg = {"role": "tool", "tool_call_id": item.get("call_id", ""), "content": content}
-        if isinstance(content, str):
+        if isinstance(content, str) or preserve_tool_images:
             return [tool_msg]
         # Chat templates render tool messages as text, so the images ride on a user turn after the tool message (the Anthropic path does the same).
         tool_msg["content"] = "".join(p["text"] for p in content if p["type"] == "text")
@@ -336,7 +340,7 @@ def _input_content(content: Any) -> str | list[dict[str, Any]]:
     if not isinstance(content, list):
         return _input_text(content)
     parts: list[dict[str, Any]] = []
-    has_image = False
+    has_image = any(isinstance(p, dict) and p.get("type") == "input_image" for p in content)
     for part in content:
         if isinstance(part, dict) and part.get("type") == "input_image":
             url = part.get("image_url") or part.get("url")
@@ -346,8 +350,10 @@ def _input_content(content: Any) -> str | list[dict[str, Any]]:
                 # an input_image without a url (e.g. a file_id) is not servable; fail rather than answer text-only
                 raise ValueError("input_image without image_url is not supported")
             parts.append({"type": "image", "freetoken_ref": {"kind": "url", "data": url}})
-            has_image = True
             continue
+        if has_image:
+            if not isinstance(part, dict) or part.get("type") not in ("input_text", "output_text", "text"):
+                raise ValueError("image message content must contain supported typed parts")
         parts.append({"type": "text", "text": _input_text([part])})
     if not has_image:
         return "".join(p["text"] for p in parts)
@@ -379,6 +385,14 @@ def _stringify(value: Any) -> str:
     return str(value)
 
 
+def _tool_output(value: Any) -> str | list[dict[str, Any]]:
+    if isinstance(value, list):
+        types = [part.get("type") if isinstance(part, dict) else None for part in value]
+        if "input_image" in types or (types and all(t in ("input_text", "output_text", "text") for t in types)):
+            return _input_content(value)
+    return _stringify(value)
+
+
 def _convert_tools(tools: list[dict[str, Any]] | None) -> list[dict[str, Any]] | None:
     if not tools:
         return None
diff --git a/python/freetoken/tokenizer/tokenize.py b/python/freetoken/tokenizer/tokenize.py
index 758ffd4dc..aaf11b686 100644
--- a/python/freetoken/tokenizer/tokenize.py
+++ b/python/freetoken/tokenizer/tokenize.py
@@ -4,7 +4,7 @@
 import json
 import os
 import threading
-from types import ModuleType
+from types import ModuleType, SimpleNamespace
 from typing import TYPE_CHECKING, Any, List
 
 import torch
@@ -54,6 +54,8 @@ def __init__(self, tokenizer: PreTrainedTokenizerBase, mm_processor: MMProcessor
         self.tokenizer = tokenizer
         self.mm_processor = mm_processor  # None: the model takes no images
         self._dsv4_encoder = _load_dsv4_encoder_if_needed(tokenizer)
+        self._dsv41 = bool(getattr(self._dsv4_encoder, "IS_DSV41", False))
+        self._vision_args = _dsv41_vision_args(_load_dsv41_config(tokenizer)) if self._dsv41 else None
         self._effort_profile: EffortProfile | None = None
         self._thinking_profile: ThinkingProfile | None = None
         self._effort_lock = threading.Lock()
@@ -63,6 +65,28 @@ def tokenize(self, msgs: List[TokenizeMsg]) -> List[UserMsg]:
         results: List[UserMsg] = []
         # TODO: batch tokenization
         for msg in msgs:
+            if self._dsv41 and isinstance(msg.text, list):
+                from freetoken.models.deepseek_v41.image_processor import prepare_vl_inputs
+
+                prompt, payload = _apply_dsv4_chat_encoder(
+                    self._dsv4_encoder, _dsv41_image_sources(msg.text, msg.images), msg.tools,
+                    self._sanitize_effort(msg.chat_template_kwargs or {}), return_media=True,
+                )
+                args = self.mm_processor.args if self.mm_processor is not None else self._vision_args
+                ids, _types, images = prepare_vl_inputs(prompt, payload["images"], self.tokenizer, args)
+                input_ids = torch.tensor(ids, dtype=torch.int32)
+                if self.mm_processor is not None:
+                    mm = self.mm_processor.from_media(input_ids, images or [])
+                    results.append(UserMsg(
+                        uid=msg.uid, input_ids=mm.input_ids, sampling_params=msg.sampling_params,
+                        mm_items=mm.mm_items, mrope_positions=mm.mrope_positions, mrope_delta=mm.mrope_delta,
+                    ))
+                else:
+                    msg.media = [vars(item) for item in images] if images else None
+                    results.append(UserMsg(
+                        uid=msg.uid, input_ids=input_ids, sampling_params=msg.sampling_params, media=msg.media,
+                    ))
+                continue
             prompt = self.render_prompt(msg)
             # A jinja chat template owns every special token (HF's apply_chat_template
             # tokenizes with add_special_tokens=False for the same reason): tokenizers
@@ -103,8 +127,9 @@ def render_prompt(self, msg: TokenizeMsg) -> str:
         validation, count_tokens) must quantize identically."""
         if not isinstance(msg.text, list):
             return msg.text
+        messages = _dsv41_image_sources(msg.text, msg.images) if self._dsv41 else msg.text
         return self._render(
-            msg.text, msg.tools, self._sanitize_effort(msg.chat_template_kwargs or {})
+            messages, msg.tools, self._sanitize_effort(msg.chat_template_kwargs or {})
         )
 
     def _render(
@@ -115,6 +140,8 @@ def _render(
     ) -> str:
         """Raw render, no effort sanitation — the probe needs unsupported values
         to actually reach the template so rejection is observable."""
+        if not self._dsv41 and self.mm_processor is None and _contains_images(messages):
+            raise ValueError("this model does not support image content")
         if self._dsv4_encoder is not None:
             return _apply_dsv4_chat_encoder(
                 self._dsv4_encoder, messages, tools, chat_template_kwargs
@@ -171,6 +198,10 @@ def _sanitize_effort(self, chat_template_kwargs: dict[str, Any]) -> dict[str, An
         if "reasoning_effort" not in chat_template_kwargs:
             return chat_template_kwargs
         raw = chat_template_kwargs.get("reasoning_effort")
+        if self._dsv41 and type(raw) is int:
+            if not 1 <= raw <= 100:
+                raise ValueError("DeepSeek-V4.1 reasoning_effort must be between 1 and 100")
+            return chat_template_kwargs
         mapped = quantize_effort(raw, self.effort_profile())
         if mapped == raw:
             return chat_template_kwargs
@@ -191,7 +222,89 @@ def _sanitize_effort(self, chat_template_kwargs: dict[str, Any]) -> dict[str, An
         return sanitized
 
 
+def _contains_images(value) -> bool:
+    if isinstance(value, dict):
+        return value.get("type") in ("image", "image_url", "input_image") or any(_contains_images(v) for v in value.values())
+    if isinstance(value, list):
+        return any(_contains_images(v) for v in value)
+    return False
+
+
+def _dsv41_image_sources(messages: list[dict], images: list[bytes] | None) -> list[dict]:
+    """Bind image bytes before the native encoder reorders tool-result messages."""
+    image_index = 0
+
+    def blocks(content):
+        nonlocal image_index
+        if not isinstance(content, list):
+            return content
+        result = []
+        for block in content:
+            if not isinstance(block, dict):
+                result.append(block)
+                continue
+            part = dict(block)
+            if part.get("type") in ("image", "image_url"):
+                if images is not None:
+                    if image_index >= len(images):
+                        raise ValueError("image parts and supplied image bytes do not match")
+                    part = {"type": "image", "data": images[image_index]}
+                    image_index += 1
+                elif isinstance(part.get("freetoken_ref"), dict):
+                    ref = part["freetoken_ref"]
+                    key = "data" if ref.get("kind") == "b64" else "url"
+                    part = {"type": "image", key: ref.get("data")}
+            elif part.get("type") == "tool_result":
+                part["content"] = blocks(part.get("content"))
+            result.append(part)
+        return result
+
+    rendered = []
+    for message in messages:
+        item = dict(message)
+        for key in ("content", "content_blocks"):
+            if key in item:
+                item[key] = blocks(item[key])
+        rendered.append(item)
+    if images is not None and image_index != len(images):
+        raise ValueError("image parts and supplied image bytes do not match")
+    return rendered
+
+
+def _load_dsv41_config(tokenizer) -> dict | None:
+    model_path = str(getattr(tokenizer, "name_or_path", None) or getattr(tokenizer, "_name_or_path", ""))
+    config_path = os.path.join(model_path, "config.json")
+    data = None
+    if os.path.isfile(config_path):
+        with open(config_path, encoding="utf-8") as handle:
+            data = json.load(handle)
+    elif "deepseek" in model_path.lower() and any(v in model_path.lower() for v in ("v4.1", "v41")):
+        from freetoken.utils import cached_load_hf_config
+
+        data = cached_load_hf_config(model_path).to_dict()
+    if data and (data.get("model_type") == "deepseek_v41" or "DeepseekV41ForCausalLM" in data.get("architectures", [])):
+        return data
+    return None
+
+
+def _dsv41_vision_args(data: dict):
+    vision = data.get("vision_config") or {}
+    return SimpleNamespace(
+        image_token_id=data.get("image_token_id", 129264),
+        vision_enabled=bool(vision.get("num_hidden_layers", 0)),
+        vision_patch_size=vision.get("patch_size", 14),
+        vision_downsample_ratio=vision.get("downsample_ratio", 3),
+        vision_max_wh_ratio=vision.get("max_wh_ratio"),
+        vision_min_pixels=vision.get("min_pixels", 295936),
+        vision_max_n_token=vision.get("max_image_tokens", 1024),
+    )
+
+
 def _load_dsv4_encoder_if_needed(tokenizer: PreTrainedTokenizerBase) -> ModuleType | None:
+    if _load_dsv41_config(tokenizer) is not None:
+        from freetoken.models.deepseek_v41 import encoding
+
+        return encoding
     if getattr(tokenizer, "chat_template", None):
         return None
     model_path = getattr(tokenizer, "name_or_path", None) or getattr(tokenizer, "_name_or_path", "")
@@ -215,7 +328,9 @@ def _apply_dsv4_chat_encoder(
     messages: list[dict],
     tools: list[dict] | None,
     chat_template_kwargs: dict,
-) -> str:
+    *,
+    return_media: bool = False,
+):
     rendered_messages = [dict(message) for message in messages]
     for message in rendered_messages:
         if message.get("tool_calls"):
@@ -225,10 +340,12 @@ def _apply_dsv4_chat_encoder(
 
     # No effort filtering here: the caller sanitized already, and the probe
     # needs raw values to reach the encoder's own validation.
+    extra = {"return_multi_modal_data": True} if return_media else {}
     return encoder.encode_messages(
         rendered_messages,
         thinking_mode=resolve_thinking_mode(chat_template_kwargs, tools),
         reasoning_effort=chat_template_kwargs.get("reasoning_effort"),
+        **extra,
     )
 
 
diff --git a/scripts/ftw_hotfix.py b/scripts/ftw_hotfix.py
index f9cb62d1a..74c51aac8 100644
--- a/scripts/ftw_hotfix.py
+++ b/scripts/ftw_hotfix.py
@@ -263,7 +263,8 @@ def is_checkpoint_tower_name(name: str) -> bool:
     head = name.split(".")
     if head[0] == "model" and len(head) > 1:
         head = head[1:]
-    return "vision" in head[0] or "visual" in head[0] or head[0] in _TOWER_SEGMENTS
+    return ("vision" in head[0] or "visual" in head[0] or head[0] in _TOWER_SEGMENTS
+            or head[0] in {"image_start", "image_end", "image_newline"})
 
 
 def read_tower(source: TensorSource, ftw_dir: str, checkpoint_names: list[str]) -> dict[str, torch.Tensor]:
diff --git a/tests/attention/test_dsa_kpool.py b/tests/attention/test_dsa_kpool.py
index fb856e175..bed1bbdae 100644
--- a/tests/attention/test_dsa_kpool.py
+++ b/tests/attention/test_dsa_kpool.py
@@ -49,8 +49,8 @@ def _args(num_layers=1):
     )
 
 
-@pytest.fixture()
-def harness(monkeypatch):
+@pytest.fixture(params=["none", "nvfp4"])
+def harness(monkeypatch, request):
     from freetoken.attention.dsa_indexer_kpool import Glm5NextDSABackend
     from freetoken.kvcache.dsa_pool import KpoolDSAKVCache
 
@@ -58,7 +58,7 @@ def harness(monkeypatch):
         latent_dim=LATENT, num_layers=1, num_pages=8, page_size=64,
         dtype=torch.bfloat16, device=torch.device(DEV),
         index_head_dim=DI, num_index_layers=1,
-        index_ratio=KPOOL, num_req_slots=4,
+        index_ratio=KPOOL, num_req_slots=4, kv_quant=request.param,
     )
     page_table = torch.full((4, 512), -1, dtype=torch.int32, device=DEV)
     page_table[0, :512] = torch.arange(512, dtype=torch.int32, device=DEV)
@@ -76,6 +76,23 @@ def harness(monkeypatch):
     return backend, pool, ape
 
 
+@pytest.mark.parametrize("kv_quant", ["fp8", "nvfp4"])
+def test_quantized_latent_cache_keeps_kpool_index_tiers_bf16(kv_quant):
+    from freetoken.kvcache.dsa_pool import KpoolDSAKVCache
+
+    pool = KpoolDSAKVCache(
+        latent_dim=LATENT, num_layers=1, num_pages=8, page_size=64,
+        dtype=torch.bfloat16, device=torch.device(DEV),
+        index_head_dim=DI, num_index_layers=1,
+        index_ratio=KPOOL, num_req_slots=4, kv_quant=kv_quant,
+    )
+    assert pool.latent_rows(0).dtype is torch.uint8
+    assert pool.latent_scale(0).dtype is torch.float32
+    assert pool.index_k_cache(0).dtype is torch.bfloat16
+    assert pool.tail_k(0).dtype is torch.bfloat16
+    assert pool.tail_gate(0).dtype is torch.bfloat16
+
+
 def _req(device_len, cached_len=0):
     return SimpleNamespace(
         table_idx=0, device_len=device_len, extend_len=device_len - cached_len,
@@ -113,10 +130,15 @@ def _rand_seq(total, seed=1):
 
 
 def _run(backend, batch, d, sl, ape):
+    d["kv_quant"] = backend.kvcache.kv_quant
+    backend.prepare_metadata(batch)
+    return _forward(backend, batch, d, sl, ape)
+
+
+def _forward(backend, batch, d, sl, ape):
     from freetoken.attention.dsa import DSAIndexerInputs
 
     t = batch.positions.shape[0]
-    backend.prepare_metadata(batch)
     return backend.mla_forward(
         d["q_nope"][sl], d["q_nope"].new_empty(t, H, 0),
         d["c_kv"][sl], d["c_kv"].new_empty(t, 0),
@@ -149,6 +171,14 @@ def _ref_scores(d, ape, q_idx_t, w_t, n_pools):
 def _ref_attend(d, q_t, positions):
     """Full softmax MLA over latent rows at ``positions`` for one query [H, LATENT]."""
     lat = d["c_kv"][positions].float()  # [n, LATENT]
+    if d.get("kv_quant") == "nvfp4":
+        from tests.kernels.test_kv_nvfp4 import _reference
+
+        packed, block, row = _reference(lat)
+        codes = torch.stack((packed & 15, packed >> 4), -1).flatten(-2).long()
+        grid = lat.new_tensor([0, .5, 1, 1.5, 2, 3, 4, 6,
+                               0, -.5, -1, -1.5, -2, -3, -4, -6])
+        lat = grid[codes] * block.view(torch.float8_e4m3fn).float().repeat_interleave(16, -1) * row[:, None]
     logits = q_t.float() @ lat.T * SM_SCALE  # [H, n]
     p = torch.softmax(logits, dim=-1)
     return (p @ lat).to(torch.bfloat16)
@@ -367,3 +397,52 @@ def test_padding_and_empty_batch_leave_shadow_rows_clean():
         ratio=KPOOL,
     )
     assert torch.equal(slab, before)
+
+
+def test_decode_graph_replay_tracks_slots_and_lengths(harness):
+    backend, pool, ape = harness
+    total, extra = 60, 6
+    d = _rand_seq(total + extra, seed=81)
+    _run(backend, _prefill_batch(0, total), d, slice(0, total), ape)
+    static = {k: v[total:total + 1].clone() for k, v in d.items() if isinstance(v, torch.Tensor)}
+    batch = _decode_batch(total)
+    batch.size = batch.padded_size = 1
+    backend.init_capture_graph(512, [1])
+    backend.prepare_for_capture(batch)
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(stream):
+        for _ in range(3):
+            _forward(backend, batch, static, slice(None), ape)
+    torch.cuda.current_stream().wait_stream(stream)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        out = _forward(backend, batch, static, slice(None), ape)
+    base = 0
+    for pos in range(total, total + extra):
+        if pos == total + 2:
+            from freetoken.attention.dsa import get_global_ctx
+
+            base = 256
+            pool.latent_rows(0)[base:base + 128].copy_(pool.latent_rows(0)[:128])
+            if pool.kv_quant == "nvfp4":
+                pool.latent_scale(0)[base:base + 128].copy_(pool.latent_scale(0)[:128])
+                pool.latent_block_scale(0)[base:base + 128].copy_(pool.latent_block_scale(0)[:128])
+            pool.index_k_cache(0)[64:96].copy_(pool.index_k_cache(0)[:32])
+            pool.tail_k(0)[1].copy_(pool.tail_k(0)[0])
+            pool.tail_gate(0)[1].copy_(pool.tail_gate(0)[0])
+            get_global_ctx().page_table[1, :128] = torch.arange(base, base + 128, device=DEV)
+            batch.active_table_idx.fill_(1)
+            batch.padded_reqs[0].table_idx = 1
+        for key, tensor in static.items():
+            tensor.copy_(d[key][pos:pos + 1])
+        batch.positions.fill_(pos)
+        batch.out_loc.fill_(base + pos)
+        batch.padded_reqs[0].device_len = pos + 1
+        backend.prepare_metadata(batch)
+        backend.prepare_for_replay(batch)
+        graph.replay()
+        selected = _ref_selected_positions(d, ape, d["qi"][pos], d["wi"][pos], pos)
+        ref = _ref_attend(d, d["q_nope"][pos], selected)
+        torch.testing.assert_close(out[0], ref, atol=3e-2, rtol=1e-2)
+    backend.reset_capture()
diff --git a/tests/checkpoint/test_ftw_hotfix_vision.py b/tests/checkpoint/test_ftw_hotfix_vision.py
index d4a3eac68..e9d988b7a 100644
--- a/tests/checkpoint/test_ftw_hotfix_vision.py
+++ b/tests/checkpoint/test_ftw_hotfix_vision.py
@@ -34,6 +34,11 @@ def hotfix():
     ("vision_tower.vision_model.embeddings.patch_embedding.weight", True),
     ("multi_modal_projector.linear_1.weight", True),
     ("patch_merge_mlp.linear_2.bias", True),
+    ("vision.patch_embed.proj.weight", True),
+    ("aligner.w1.weight", True),
+    ("image_start", True),
+    ("image_end", True),
+    ("image_newline", True),
     ("model.embed_audio.embedding_projection.weight", False),
     ("model.language_model.layers.0.self_attn.q_proj.weight", False),
     ("model.layers.3.mlp.experts.0.gate_proj.weight", False),
@@ -76,3 +81,57 @@ def test_every_family_with_an_encoder_has_an_encoder_only_reader():
     for spec in _MODEL_REGISTRY.values():
         if spec.encoders:
             assert callable(_load_attr(spec.module, "iter_vision_weights")), spec.module
+
+
+def test_native_v41_hotfix_reads_only_vision_tensors(hotfix, tmp_path, monkeypatch):
+    import safetensors.torch
+    import torch
+
+    from freetoken.models.deepseek_v41 import weight
+    from freetoken.models.weight import load_vision_weight
+
+    checkpoint = tmp_path / "source"
+    checkpoint.mkdir()
+    config = {"architectures": ["DeepseekV41ForCausalLM"], "model_type": "deepseek_v41", "n_layers": 1,
+              "compress_ratios": [0], "kv_source_layers": [], "index_source_layers": [],
+              "candidate_source_layer": -1, "engram_layer_ids": [], "engram_num_embeddings": [],
+              "vision_n_layers": 1, "quantization_config": {"expert_dtype": "nvfp4"}}
+    (checkpoint / "config.json").write_text(json.dumps(config))
+    tensors = {
+        "vision.patch_embed.proj.weight": torch.arange(6, dtype=torch.bfloat16).view(2, 3),
+        "vision.norm.weight": torch.arange(2, dtype=torch.float32),
+        "aligner.w1.weight": torch.ones(2, 3, dtype=torch.bfloat16),
+        "image_start": torch.full((4,), 1.0, dtype=torch.bfloat16),
+        "image_end": torch.full((4,), 2.0, dtype=torch.bfloat16),
+        "image_newline": torch.full((4,), 3.0, dtype=torch.bfloat16),
+    }
+    safetensors.torch.save_file(tensors, checkpoint / "vision.safetensors")
+    index = {name: "vision.safetensors" for name in tensors}
+    index["head.weight"] = "unavailable-text.safetensors"
+    index["layers.0.ffn.experts.0.w1.weight"] = "unavailable-experts.safetensors"
+    index["layers.0.engram.embed.weight"] = "unavailable-engram.safetensors"
+    (checkpoint / "model.safetensors.index.json").write_text(json.dumps({"weight_map": index}))
+    reads = []
+    original_get = weight._ShardReader.get
+
+    def get(reader, name):
+        reads.append(name)
+        return original_get(reader, name)
+
+    monkeypatch.setattr(weight._ShardReader, "get", get)
+    direct = dict(load_vision_weight(str(checkpoint), torch.device("cpu")))
+    assert set(direct) == set(reads) == set(tensors)
+    reads.clear()
+    ftw_like = tmp_path / "ftw"
+    ftw_like.mkdir()
+    ftw_dir = _config_dir(hotfix, checkpoint, ftw_like)
+    source = hotfix.TensorSource(None, str(checkpoint))
+    selected = [name for name in source.weight_map if hotfix.is_checkpoint_tower_name(name)]
+    got = hotfix.read_tower(source, ftw_dir, selected)
+    assert set(got) == set(tensors)
+    assert set(reads) == set(tensors)
+    for name, expected in tensors.items():
+        assert direct[name].dtype == expected.dtype
+        assert torch.equal(direct[name], expected)
+        assert got[name].dtype == expected.dtype
+        assert torch.equal(got[name], expected)
diff --git a/tests/checkpoint/test_ftw_weights.py b/tests/checkpoint/test_ftw_weights.py
index b3b3236bd..5a8dbe0c3 100644
--- a/tests/checkpoint/test_ftw_weights.py
+++ b/tests/checkpoint/test_ftw_weights.py
@@ -1,5 +1,6 @@
 """FTW replay: dropping entries by name before their bytes are read, and the vision-tower presence check."""
 
+import pytest
 import torch
 
 from freetoken.checkpoint.ftw import FTWReader, FTWWriter, ftw_tensor_names, iter_ftw_weights
@@ -54,3 +55,22 @@ def test_ftw_lacks_vision(tmp_path):
     assert not ftw_lacks_vision(str(tmp_path / "vl"))
     assert not ftw_lacks_vision(str(tmp_path))
     assert ftw_tensor_names(str(tmp_path / "vl"), "weight") == ["model.a.weight", "visual.b.weight"]
+
+
+@pytest.mark.parametrize("name", [
+    "vision_embedder.patch_embedding.weight", "vision.patch_embed.proj.weight",
+    "aligner.w1.weight", "image_start", "image_end", "image_newline",
+])
+def test_native_encoder_ftw_presence_and_text_only_read(tmp_path, monkeypatch, name):
+    _write_ftw(tmp_path, ["model.a.weight", name])
+    assert not ftw_lacks_vision(str(tmp_path))
+    reads = []
+    original = FTWReader.read_into
+
+    def spy(self, dest, entry, **kwargs):
+        reads.append(entry["name"])
+        return original(self, dest, entry, **kwargs)
+
+    monkeypatch.setattr(FTWReader, "read_into", spy)
+    got = dict(load_weight(str(tmp_path), torch.device("cpu"), include_vision=False))
+    assert list(got) == reads == ["model.a.weight"]
diff --git a/tests/engine/test_attention_backend_matrix.py b/tests/engine/test_attention_backend_matrix.py
index ef7dc7ff8..759aaab72 100644
--- a/tests/engine/test_attention_backend_matrix.py
+++ b/tests/engine/test_attention_backend_matrix.py
@@ -65,6 +65,9 @@ def _model_config(kind):
     elif kind == "dsv4":
         mc.dsv4_args = SimpleNamespace(window_size=128)
         specs = (_spec("dsv4", AttnType.DSV4, sliding_window=128),)
+    elif kind == "dsv41":
+        mc.dsv41_args = SimpleNamespace(window_size=128)
+        specs = (_spec("dsv41", AttnType.DSV41, sliding_window=128),)
     elif kind == "bsa":
         # MiniMax-M3 shape: one FULL-family group, mla=False + index dims -> BSA.
         specs = (_spec("full", AttnType.BSA, index_head_dim=128),)
@@ -114,6 +117,7 @@ def _patch_env(monkeypatch, *, major=9, flashinfer=True, sgl=True):
         ("mla", "dsa"),  # plain latent MLA
         ("dsa", "dsa"),  # MLA + DSA indexer (GLM-5.2 shape)
         ("dsv4", "dsv4_sparse"),
+        ("dsv41", "dsv41_sparse"),
         ("bsa", "m3_sparse"),  # MiniMax-M3 block-sparse GQA
         ("qsa", "qsa_sparse"),  # Qwen3.8-Flash-Next compressed-block sparse
     ],
@@ -192,12 +196,81 @@ def test_auto_dsv4_sets_window_page_size(monkeypatch):
     assert config.page_size == 128
 
 
+def test_v41_resolves_eager_window_pages_and_preserves_context(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_env(monkeypatch)
+    config = _config("dsv41", attention_backend="auto")
+    _adjust_config(config)
+    assert config.page_size == 128
+    assert config.cache_type == "swa_radix"
+    assert config.cuda_graph_max_bs == 0 and config.cuda_graph_bs == []
+    assert config.model_config.dsv41_args.max_seq_len == config.max_seq_len
+
+
+def test_v41_long_context_preserves_bounded_prefill_staging(monkeypatch):
+    from freetoken.distributed import DistributedInfo
+    from freetoken.engine.engine import _adjust_config
+    from freetoken.scheduler.config import SchedulerConfig
+
+    _patch_env(monkeypatch)
+    config = SchedulerConfig(
+        model_path="/tmp/freetoken-test-model", tp_info=DistributedInfo(rank=0, size=1),
+        dtype=torch.bfloat16, max_seq_len_override=1 << 20, max_extend_tokens=1024,
+    )
+    object.__setattr__(config, "model_config", _model_config("dsv41"))
+    _adjust_config(config)
+    assert config.max_seq_len == 1 << 20
+    assert config.max_forward_len == config.max_extend_tokens == 1024
+
+
+def test_v41_rejects_prefill_budget_that_cannot_admit_a_page(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_env(monkeypatch)
+    config = _config("dsv41", max_seq_len_override=1 << 20)
+    object.__setattr__(config, "max_extend_tokens", 127)
+    with pytest.raises(ValueError, match="max-prefill-length must fit one 128-token page"):
+        _adjust_config(config)
+
+
+@pytest.mark.parametrize("ratio", [0, -1, 1.1])
+def test_v41_rejects_invalid_window_ratio(monkeypatch, ratio):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_env(monkeypatch)
+    with pytest.raises(ValueError, match="swa_full_tokens_ratio"):
+        _adjust_config(_config("dsv41", swa_full_tokens_ratio=ratio))
+
+
+@pytest.mark.parametrize("kwargs", [{"cuda_graph_max_bs": 1}, {"cuda_graph_bs": [1]}])
+def test_v41_rejects_explicit_graph_until_supported(monkeypatch, kwargs):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_env(monkeypatch)
+    with pytest.raises(ValueError, match="eager execution"):
+        _adjust_config(_config("dsv41", **kwargs))
+
+
+def test_v41_rejects_dtype_before_allocating_quantized_pools(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_env(monkeypatch)
+    config = _config("dsv41")
+    object.__setattr__(config, "dtype", torch.float32)
+    with pytest.raises(ValueError, match="bfloat16"):
+        _adjust_config(config)
+
+
 @pytest.mark.parametrize(
     "kind, backend",
     [
         # reverse gates: type-specific backends on models without the type
         ("full", "dsa"),
         ("full", "dsv4_sparse"),
+        ("full", "dsv41_sparse"),
+        ("dsv41", "dsv4_sparse"),
+        ("dsv41", "triton"),
         ("full", "m3_sparse"),
         ("swa", "dsa"),
         ("full", "qsa_sparse"),
diff --git a/tests/engine/test_deepseek_v41_engine.py b/tests/engine/test_deepseek_v41_engine.py
new file mode 100644
index 000000000..7f832a7d6
--- /dev/null
+++ b/tests/engine/test_deepseek_v41_engine.py
@@ -0,0 +1,262 @@
+"""Protect V4.1 startup budgeting and serve an image in a fresh CUDA process."""
+
+import os
+from pathlib import Path
+import socket
+import subprocess
+import sys
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+
+@pytest.mark.parametrize("active_encoder", [False, True])
+def test_finalized_weights_and_engram_staging_precede_cache_budgets(monkeypatch, tmp_path, active_encoder):
+    import freetoken.engine.engine as engine_module
+
+    allocations = {}
+    baseline_free = 4096
+    resident_and_staging_bytes = 48 + 32 * 4 + 16 * 2 + 8 + (12 if active_encoder else 0)
+    host_table_bytes = 1024
+    cache_bytes = 64
+
+    class QuantMethod:
+        calls = 0
+
+        def finalize(self, layer):
+            self.calls += 1
+            assert layer.weight.shape == (24,)
+            layer.weight = layer.weight.repeat_interleave(2)
+            allocations["weight"] = layer.weight
+            allocations["quant_workspace"] = torch.empty(32, dtype=torch.float32, device="cpu")
+
+    class Model:
+        quant_method = QuantMethod()
+
+        def load_state_dict(self, weights):
+            self.weight = weights["weight"]
+            allocations["weight"] = self.weight
+            if active_encoder:
+                allocations["encoder_weights"] = torch.empty(256, dtype=torch.uint8, device="cpu")
+
+        def place_encoder_weights(self, mode):
+            assert mode == "host" and self.quant_method.calls == 1
+            del allocations["encoder_weights"]
+            allocations["encoder_staging"] = torch.empty(12, dtype=torch.uint8, device="cpu")
+
+        def encode(self, item):
+            raise AssertionError("encoder warmup is stubbed in the memory-budget test")
+
+        def load_host_tables(self, config):
+            assert not active_encoder or "encoder_staging" in allocations
+            allocations["engram_values"] = torch.empty(16, dtype=torch.bfloat16, device="cpu")
+            allocations["engram_mask"] = torch.empty(8, dtype=torch.bool, device="cpu")
+            return host_table_bytes
+
+    class ReachedKVPlanning(Exception):
+        pass
+
+    class Pool:
+        @staticmethod
+        def solve_num_pages(config, available_memory):
+            assert model.quant_method.calls == 1
+            assert available_memory == 3072 - resident_and_staging_bytes - cache_bytes
+            raise ReachedKVPlanning
+
+    model = Model()
+    config = SimpleNamespace(
+        model_path=str(tmp_path), tp_info=SimpleNamespace(rank=0, size=1),
+        dtype=torch.bfloat16, quant_backend="moe.nvfp4=triton", moe_strategy="offload",
+        model_config=object(), page_size=8, memory_ratio=0.75,
+        active_encoders=(SimpleNamespace(kind="vision"),) if active_encoder else (),
+        mm=SimpleNamespace(encoder_weights="host", embed_cache_device="cpu"),
+        served_modalities={"image"} if active_encoder else set(), hf_config=SimpleNamespace(),
+    )
+
+    def free_memory(self):
+        free = baseline_free - sum(t.numel() * t.element_size() for t in allocations.values())
+        return free, free
+
+    def initialize_offload_cache(self, config):
+        assert self.model.quant_method.calls == 1
+        assert self._host_tables_bytes == host_table_bytes
+        assert self._weights_bytes == resident_and_staging_bytes
+        assert self._post_weights_free == baseline_free - resident_and_staging_bytes
+        allocations["offload_cache"] = torch.empty(cache_bytes, dtype=torch.uint8, device="cpu")
+
+    # Keep the real Engine startup and finalize traversal; replace only CUDA and its consumers.
+    monkeypatch.setattr(torch.cuda, "is_initialized", lambda: False)
+    monkeypatch.setattr(torch.cuda, "Stream", lambda: object())
+    monkeypatch.setattr(torch.cuda, "set_stream", lambda stream: None)
+    monkeypatch.setattr(torch, "manual_seed", lambda seed: None)
+    monkeypatch.setattr("freetoken.gpu_select.bind_assigned_gpu", lambda rank: torch.device("cpu"))
+    monkeypatch.setattr(engine_module, "set_tp_info", lambda **kwargs: None)
+    monkeypatch.setattr(engine_module, "set_quant_backend", lambda backend: None)
+    monkeypatch.setattr(engine_module, "_adjust_ftw_quant_backend", lambda path, backend: backend)
+    monkeypatch.setattr(engine_module, "_ensure_expandable_segments", lambda: None)
+    monkeypatch.setattr(engine_module, "_adjust_config", lambda config: None)
+    monkeypatch.setattr(engine_module, "set_global_ctx", lambda ctx: None)
+    monkeypatch.setattr(engine_module, "set_rope_device", lambda device: None)
+    monkeypatch.setattr(engine_module.logger, "info_rank0", lambda *args, **kwargs: None)
+    monkeypatch.setattr(engine_module, "resolve_pool_class", lambda config: Pool)
+    monkeypatch.setattr(engine_module, "create_model", lambda config: model)
+    monkeypatch.setattr(engine_module, "state_pool_bytes", lambda config: 0)
+    monkeypatch.setattr(engine_module.Engine, "_init_communication", lambda self, config: None)
+    monkeypatch.setattr(engine_module.Engine, "_sync_get_memory", free_memory)
+    monkeypatch.setattr(engine_module.Engine, "_load_weight_state_dict", lambda self, config: {
+        "weight": torch.ones(24, dtype=torch.uint8, device="cpu"),
+    })
+    monkeypatch.setattr(engine_module.Engine, "_init_offload_moe_cache", initialize_offload_cache)
+    monkeypatch.setattr("freetoken.mm.processor.get_mm_processor", lambda *args: object())
+    monkeypatch.setattr(engine_module.Engine, "_warmup_encoders", lambda self: None)
+    with pytest.raises(ReachedKVPlanning):
+        engine_module.Engine(config)
+
+
+def _run_engine_smoke(folder, port, kv_quant, image_path):
+    import json
+    from dataclasses import asdict
+    from unittest.mock import patch
+
+    import freetoken.engine.engine as engine_module
+    from freetoken.core import Batch, Req, SamplingParams
+    from freetoken.distributed import DistributedInfo
+    from freetoken.engine.config import EngineConfig
+    from freetoken.engine.engine import Engine
+    from freetoken.layers.quantization import QuantKind, finalize_quant
+    from freetoken.moe.offload_cache import iter_offload_moe_layers
+    from freetoken.models.deepseek_v41.args import DeepseekV41Args
+    from freetoken.models.deepseek_v41.image_processor import IMAGE, IMAGE_END, IMAGE_NEW_LINE, IMAGE_START
+    from freetoken.scheduler.mm import plan_mm_batch
+
+    args = DeepseekV41Args(
+        n_layers=3, n_mtp_layers=0, compress_ratios=(0, 2, 2),
+        kv_source_layers=(1,), index_source_layers=(1,), candidate_source_layer=-1,
+        dim=64, n_heads=2, head_dim=32, rope_head_dim=16, q_lora_rank=32,
+        o_lora_rank=32, o_groups=2, window_size=8, index_n_heads=2, index_head_dim=32,
+        index_topk=3, moe_inter_dim=64, n_routed_experts=4, n_activated_experts=2,
+        vocab_size=128, hc_mult=2, engram_layer_ids=(1,), engram_num_embeddings=(17,),
+        engram_max_ngram_size=2, engram_vocab_size=17, engram_n_heads=1,
+        engram_head_dim=32, engram_compressed_vocab_size=128,
+        vision_n_layers=1, vision_dim=32, vision_n_heads=2, vision_inter_dim=32,
+        vision_patch_size=2, vision_downsample_ratio=2, image_token_id=127,
+    )
+    raw = asdict(args) | {
+        "architectures": ["DeepseekV41ForCausalLM"], "model_type": "deepseek_v41",
+        "quantization_config": {"moe_quant_algo": "NVFP4"},
+        "vision_config": {"num_hidden_layers": 1, "hidden_size": 32, "num_attention_heads": 2,
+                          "intermediate_size": 32, "patch_size": 2, "downsample_ratio": 2},
+    }
+    Path(folder, "config.json").write_text(json.dumps(raw))
+
+    class LocalEngineConfig(EngineConfig):
+        @property
+        def distributed_addr(self):
+            return f"tcp://127.0.0.1:{port}"
+
+    config = LocalEngineConfig(
+        model_path=folder, tp_info=DistributedInfo(0, 1), dtype=torch.bfloat16,
+        max_running_req=1, moe_strategy="offload", quant_backend="moe.nvfp4=triton",
+        moe_cpu_layers="", moe_cache_size=8, kv_quant=kv_quant,
+        moe_prefill_overlap=True, use_dummy_weight=True, use_pynccl=False,
+        max_seq_len_override=64, num_page_override=40, cuda_graph_max_bs=0,
+    )
+    finalized_counts = []
+
+    def checked_finalize(model):
+        count = finalize_quant(model)
+        finalized_counts.append(count)
+        return count
+
+    with patch.object(engine_module, "finalize_quant", checked_finalize):
+        engine = Engine(config)
+    assert finalized_counts == [args.n_layers]
+    assert engine.model._engram_runtime is not None
+    assert config.attention_backend == "dsv41_sparse"
+    assert config.model_config.expert_quant == "nvfp4" and config.page_size == 8
+    assert config.model_config.is_multimodal
+    assert engine.kv_cache.kv_quant == kv_quant
+    experts = list(iter_offload_moe_layers(engine.model))
+    assert len(experts) == args.n_layers
+    for expert in experts:
+        method = expert.quant_method
+        assert method.kind is QuantKind.NVFP4
+        assert method.kernel.name == "triton"
+        assert method.cfg.activation == "swiglu_clamp"
+        assert method.cfg.alpha == 1.0 and method.cfg.limit == args.swiglu_limit
+        assert method.cfg.strategy == "offload" and method.cfg.decode_target == "gpu"
+    assert engine.model._transformer.vision.patch_embed.proj.weight.dtype == torch.bfloat16
+    runtime = engine.model._engram_runtime
+    resident_bytes = sum(p.numel() * p.element_size() for p in engine.model.state_dict().values() if p.is_cuda)
+    staged_bytes = sum(module._values.numel() * module._values.element_size() for module in runtime.modules)
+    staged_bytes += runtime.device_mask.numel() * runtime.device_mask.element_size()
+    vision_streamer = engine.model._transformer.vision._streamer
+    assert vision_streamer is not None and vision_streamer.bank.is_pinned()
+    staged_bytes += vision_streamer.device_bytes
+    assert engine._weights_bytes >= resident_bytes + staged_bytes
+    sample = engine.sampler.sample
+
+    def checked_sample(logits, sampling_args):
+        assert logits.dtype == torch.float32 and torch.isfinite(logits).all()
+        return sample(logits, sampling_args)
+
+    engine.sampler.sample = checked_sample
+    engine.page_table[0, :16] = torch.arange(16, device="cuda")
+    for start in (0, 8):
+        engine.kv_cache.bind_window_pages(start, start)
+    media = [{"start": 1, "types": torch.tensor([IMAGE_START, IMAGE, IMAGE_NEW_LINE, IMAGE_END]),
+              "patches": torch.randn(4, 3, 2, 2), "n_vit_h": 2, "n_vit_w": 2}]
+    ids = torch.tensor([5, 127, 127, 127, 127, 6, 7, 8], dtype=torch.int32)
+    if image_path == "mm_items":
+        result = engine.mm_processor.from_media(ids, media)
+        req = Req(result.input_ids, 0, 0, 3, 0, SamplingParams(temperature=0), None, mm_items=result.mm_items)
+        for item in req.mm_items:
+            engine.encoder_cache.register(item.hash, req.uid, item.num_tokens)
+    else:
+        req = Req(ids, 0, 0, 3, 0, SamplingParams(temperature=0), None, media=media)
+    for phase in ("prefill", "decode"):
+        batch = Batch([req], phase)
+        batch.padded_reqs = batch.reqs
+        batch.input_ids = req.input_ids[req.cached_len:].cuda()
+        batch.positions = torch.arange(req.cached_len, req.device_len, device="cuda")
+        batch.active_table_idx = torch.tensor([0], dtype=torch.long, device="cuda")
+        batch.out_loc = engine.page_table[0, req.cached_len:req.device_len]
+        if phase == "prefill" and req.mm_items:
+            jobs, plan, rows, block_ends = plan_mm_batch([req], engine.encoder_cache)
+            batch.mm_encoder_jobs, batch.mm_gather_plan = jobs, plan
+            batch.mm_rows = torch.tensor(rows, device="cuda", dtype=torch.long)
+            batch.mm_block_ends = torch.tensor(block_ends, device="cuda", dtype=torch.int32)
+        engine.attn_backend.prepare_metadata(batch)
+        with torch.inference_mode():
+            output = engine.forward_batch(batch, engine.sampler.prepare(batch))
+        output.copy_done_event.synchronize()
+        assert output.next_tokens_cpu.shape == (1,)
+        assert 0 <= output.next_tokens_cpu.item() < 128
+        req.append_host(output.next_tokens_cpu)
+    if image_path == "mm_items":
+        assert req.mm_items[0].feature is None
+        assert not engine.encoder_cache.has(req.mm_items[0].hash)
+    else:
+        assert "embeddings" in media[0]
+    assert req.cached_len == 9 and req.device_len == 10
+    torch.cuda.synchronize()
+    torch.distributed.destroy_process_group()
+    print("V41_ENGINE_PREFILL_DECODE_IMAGE_OK")
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"], ids=["bf16", "fp8-fp4"])
+@pytest.mark.parametrize("image_path", ["legacy", "mm_items"])
+def test_engine_initialization_and_image_generation(tmp_path, kv_quant, image_path):
+    with socket.socket() as sock:
+        sock.bind(("127.0.0.1", 0))
+        port = sock.getsockname()[1]
+    result = subprocess.run([sys.executable, str(Path(__file__).resolve()), str(tmp_path), str(port), kv_quant, image_path],
+                            capture_output=True, text=True, timeout=180, env=os.environ.copy())
+    assert result.returncode == 0, result.stdout + result.stderr
+    assert "V41_ENGINE_PREFILL_DECODE_IMAGE_OK" in result.stdout
+
+
+if __name__ == "__main__":
+    _run_engine_smoke(sys.argv[1], int(sys.argv[2]), sys.argv[3], sys.argv[4])
diff --git a/tests/engine/test_kv_quant_config.py b/tests/engine/test_kv_quant_config.py
new file mode 100644
index 000000000..d31dfbc69
--- /dev/null
+++ b/tests/engine/test_kv_quant_config.py
@@ -0,0 +1,327 @@
+"""Config-time gates for ``--kv-cache-dtype`` (EngineConfig.kv_quant).
+
+fp8 KV is only half a feature: the pool has to store it AND the attention backend has
+to read the scales. Everything here must fail while the config is still a dataclass --
+after weights are resident, a wrong combination has already cost a load and (worse)
+fi/fa/trtllm would happily attend over raw e4m3 codes and produce plausible garbage.
+"""
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.attention import AttnType
+from freetoken.models.config import KVCacheGroupSpec
+
+
+def _spec(name, attn_type, *, mla=False, sliding_window=None, index_head_dim=0, index_ratio=1):
+    return KVCacheGroupSpec(
+        name=name,
+        layer_ids=(0, 1),
+        num_kv_heads=1,
+        head_dim=64,
+        sliding_window=sliding_window,
+        mla=mla,
+        index_head_dim=index_head_dim,
+        num_index_layers=2 if index_head_dim else 0,
+        index_ratio=index_ratio,
+        attn_type=attn_type,
+    )
+
+
+def _model_config(kind):
+    mc = SimpleNamespace(
+        model_type=kind,
+        single_stream_only=False,
+        is_moe=False,
+        expert_quant="none",
+        has_swa_attention=False,
+        has_linear_attention=False,
+        num_layers=4,
+        rotary_config=SimpleNamespace(max_position=1024),
+    )
+    specs = {
+        "full": (_spec("full", AttnType.FULL),),
+        "swa": (
+            _spec("full", AttnType.FULL),
+            _spec("swa", AttnType.SWA, sliding_window=128),
+        ),
+        "mla": (_spec("full", AttnType.MLA, mla=True),),
+        "dsa": (_spec("full", AttnType.DSA, mla=True, index_head_dim=128),),
+        "kpool": (_spec("full", AttnType.DSA, mla=True, index_head_dim=128, index_ratio=4),),
+        "dsv4": (_spec("dsv4", AttnType.DSV4, sliding_window=128),),
+        "dsv41": (_spec("dsv41", AttnType.DSV41, sliding_window=128),),
+        "bsa": (_spec("full", AttnType.BSA, index_head_dim=128),),
+        "qsa": (_spec("full", AttnType.QSA, index_head_dim=128, index_ratio=4),),
+    }[kind]
+    if kind == "swa":
+        mc.has_swa_attention = True
+    if kind == "dsv4":
+        mc.dsv4_args = SimpleNamespace(window_size=128)
+    if kind == "dsv41":
+        mc.dsv41_args = SimpleNamespace(window_size=128, head_dim=64, index_head_dim=128)
+    if kind in ("qsa", "kpool"):
+        mc.has_linear_attention = True
+    mc.kv_cache_group_specs = lambda: specs
+    return mc
+
+
+def _config(kind, **overrides):
+    from freetoken.distributed import DistributedInfo
+    from freetoken.engine.config import EngineConfig
+
+    config = EngineConfig(
+        model_path="/tmp/freetoken-test-model",
+        tp_info=DistributedInfo(rank=0, size=1),
+        dtype=torch.bfloat16,
+        **overrides,
+    )
+    object.__setattr__(config, "model_config", _model_config(kind))
+    return config
+
+
+def _patch_fast_machine(monkeypatch):
+    """A machine where every fast external backend is available, so an fp8 result can
+    only come from the gate and not from a missing package."""
+    from freetoken.engine import engine
+
+    monkeypatch.setattr(engine, "is_sm100_family", lambda: False)
+    monkeypatch.setattr(engine, "is_sm90_family", lambda: True)
+    monkeypatch.setattr(engine, "_flashinfer_available", lambda: True)
+    monkeypatch.setattr(engine, "_sgl_flash_attn_available", lambda: True)
+
+
+def test_kv_quant_spellings():
+    from freetoken.engine.engine import _resolve_kv_quant
+
+    assert _resolve_kv_quant("auto") == "none"
+    assert _resolve_kv_quant("bf16") == "none"
+    assert _resolve_kv_quant("FP8") == "fp8"
+    assert _resolve_kv_quant("NVFP4") == "nvfp4"
+    assert _resolve_kv_quant("FP8-FP4") == "fp8-fp4"
+    assert _resolve_kv_quant(None) == "none"
+    with pytest.raises(ValueError, match="kv-cache-dtype"):
+        _resolve_kv_quant("q8")
+
+
+def test_only_the_backends_that_read_scales_declare_fp8_support():
+    from freetoken.attention import SUPPORTED_ATTENTION_BACKENDS, attention_backend_info
+
+    fp8 = {
+        name
+        for name in SUPPORTED_ATTENTION_BACKENDS.supported_names()
+        if attention_backend_info(name).supports_fp8_kv
+    }
+    # triton serves plain paged / hybrid-SWA pools, qsa_sparse dequantizes selected
+    # rows, and dsa dequantizes MLA latent rows. External backends and sparse
+    # families without a scale path must stay out of this set.
+    assert fp8 == {"triton", "qsa_sparse", "dsa"}
+
+    nvfp4 = {
+        name
+        for name in SUPPORTED_ATTENTION_BACKENDS.supported_names()
+        if attention_backend_info(name).supports_nvfp4_kv
+    }
+    assert nvfp4 == {"triton", "qsa_sparse", "dsa"}
+
+    mixed = {
+        name
+        for name in SUPPORTED_ATTENTION_BACKENDS.supported_names()
+        if attention_backend_info(name).supports_fp8_fp4_kv
+    }
+    assert mixed == {"dsv41_sparse"}
+
+
+@pytest.mark.parametrize("backend", ["auto", "dsv41_sparse"])
+def test_native_mixed_selects_v41_backend(monkeypatch, backend):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("dsv41", attention_backend=backend, kv_quant="fp8-fp4")
+    _adjust_config(config)
+    assert config.kv_quant == "fp8-fp4"
+    assert config.attention_backend == "dsv41_sparse"
+
+
+@pytest.mark.parametrize("kind", ["full", "swa", "mla", "dsa", "kpool", "dsv4", "bsa", "qsa"])
+def test_native_mixed_rejects_other_architectures(monkeypatch, kind):
+    from freetoken.engine.engine import _adjust_config
+    from freetoken.kvcache import create_kvcache_pool
+
+    _patch_fast_machine(monkeypatch)
+    config = _config(kind, attention_backend="auto", kv_quant="fp8-fp4")
+    with pytest.raises(ValueError, match="fp8-fp4 requires DeepSeek-V4.1"):
+        _adjust_config(config)
+    with pytest.raises(ValueError, match="fp8-fp4 requires the DeepSeek-V4.1"):
+        create_kvcache_pool(config.model_config, num_pages=4, page_size=128,
+                            dtype=torch.bfloat16, device=torch.device("cpu"), kv_quant="fp8-fp4")
+
+
+@pytest.mark.parametrize("field", ["head_dim", "index_head_dim"])
+def test_native_mixed_rejects_partial_blocks_before_loading(monkeypatch, field):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("dsv41", attention_backend="auto", kv_quant="fp8-fp4")
+    setattr(config.model_config.dsv41_args, field, 72)
+    with pytest.raises(ValueError, match="divisible by 32"):
+        _adjust_config(config)
+
+
+def test_native_mixed_requires_bf16_reconstruction(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("dsv41", attention_backend="auto", kv_quant="fp8-fp4")
+    object.__setattr__(config, "dtype", torch.float16)
+    with pytest.raises(ValueError, match="requires --dtype bfloat16"):
+        _adjust_config(config)
+
+
+def test_auto_avoids_the_fast_backends_for_fp8(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    # Same machine where a 16-bit cache auto-selects the sm90 "fa,fi" tree...
+    plain = _config("full", attention_backend="auto")
+    _adjust_config(plain)
+    assert plain.kv_quant == "none"
+    assert plain.attention_backend == "fa,fi"
+
+    quantized = _config("full", attention_backend="auto", kv_quant="fp8")
+    _adjust_config(quantized)
+    assert quantized.attention_backend == "triton"
+
+
+@pytest.mark.parametrize("backend", ["fi", "fa", "trtllm", "fi,triton", "triton,fi"])
+def test_explicit_unsupported_backend_is_rejected(monkeypatch, backend):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    monkeypatch.setattr(
+        "freetoken.engine.engine.is_sm100_family", lambda: True
+    )  # let trtllm clear its own arch gate first
+    config = _config("full", attention_backend=backend, kv_quant="fp8", page_size=1)
+    with pytest.raises(ValueError, match="kv-cache-dtype fp8"):
+        _adjust_config(config)
+
+
+def test_explicit_triton_is_accepted(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("full", attention_backend="triton", kv_quant="fp8")
+    _adjust_config(config)
+    assert config.kv_quant == "fp8"
+
+
+def test_nvfp4_auto_selects_triton(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("full", attention_backend="auto", kv_quant="nvfp4")
+    _adjust_config(config)
+    assert config.attention_backend == "triton"
+
+
+@pytest.mark.parametrize("kind", ["dsv4", "dsv41", "bsa"])
+def test_nvfp4_rejects_unsupported_pools_before_allocation(monkeypatch, kind):
+    from freetoken.engine.engine import _adjust_config
+    from freetoken.kvcache import create_kvcache_pool
+
+    _patch_fast_machine(monkeypatch)
+    config = _config(kind, attention_backend="auto", kv_quant="nvfp4")
+    with pytest.raises(ValueError, match="nvfp4"):
+        _adjust_config(config)
+    with pytest.raises(ValueError, match="nvfp4"):
+        create_kvcache_pool(config.model_config, num_pages=4, page_size=1,
+                            dtype=torch.bfloat16, device=torch.device("cpu"), kv_quant="nvfp4")
+
+
+@pytest.mark.parametrize("backend", ["fi", "fa", "trtllm", "fi,triton", "triton,fi"])
+def test_nvfp4_rejects_backends_without_its_layout(monkeypatch, backend):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    monkeypatch.setattr("freetoken.engine.engine.is_sm100_family", lambda: True)
+    config = _config("full", attention_backend=backend, kv_quant="nvfp4")
+    with pytest.raises(ValueError, match="nvfp4"):
+        _adjust_config(config)
+
+
+def test_nvfp4_rejects_partial_blocks(monkeypatch):
+    from dataclasses import replace
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("full", attention_backend="auto", kv_quant="nvfp4")
+    spec = replace(config.model_config.kv_cache_group_specs()[0], head_dim=72)
+    config.model_config.kv_cache_group_specs = lambda: (spec,)
+    with pytest.raises(ValueError, match="divisible by 16"):
+        _adjust_config(config)
+
+
+def test_nvfp4_accepts_hybrid_swa(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("swa", attention_backend="auto", kv_quant="nvfp4")
+    _adjust_config(config)
+    assert config.kv_quant == "nvfp4"
+    assert config.attention_backend == "triton"
+
+
+def test_nvfp4_accepts_qsa(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("qsa", attention_backend="auto", kv_quant="nvfp4")
+    _adjust_config(config)
+    assert config.kv_quant == "nvfp4"
+    assert config.attention_backend == "qsa_sparse"
+
+
+@pytest.mark.parametrize("kind", ["mla", "dsa", "kpool"])
+@pytest.mark.parametrize("backend", ["auto", "dsa"])
+@pytest.mark.parametrize("kv_quant", ["fp8", "nvfp4"])
+def test_mla_and_dsa_select_the_scale_reading_backend(monkeypatch, kind, backend, kv_quant):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config(kind, attention_backend=backend, kv_quant=kv_quant)
+    _adjust_config(config)
+    assert config.attention_backend == "dsa"
+    assert config.page_size == (64 if kind == "kpool" else 1)
+
+
+@pytest.mark.parametrize("kind", ["dsv4", "dsv41", "bsa"])
+def test_pool_families_without_a_scale_read_path_are_rejected(monkeypatch, kind):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)  # the rejection must not depend on this box's wheels
+    config = _config(kind, attention_backend="auto", kv_quant="fp8")
+    with pytest.raises(ValueError, match="kv-cache-dtype fp8"):
+        _adjust_config(config)
+
+
+def test_qsa_keeps_fp8_available(monkeypatch):
+    """The QSA pool is the block-sparse family whose K/V rows do reach a kernel that can
+    dequantize them; the compressed index keys selection scores against are a separate,
+    always-16-bit tier, so the gate has nothing left to refuse here."""
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("qsa", attention_backend="auto", kv_quant="fp8")
+    _adjust_config(config)
+    assert config.kv_quant == "fp8"
+    assert "qsa_sparse" in config.attention_backend
+
+
+def test_swa_pool_keeps_fp8_available(monkeypatch):
+    from freetoken.engine.engine import _adjust_config
+
+    _patch_fast_machine(monkeypatch)
+    config = _config("swa", attention_backend="auto", kv_quant="fp8")
+    _adjust_config(config)
+    assert config.kv_quant == "fp8" and config.attention_backend == "triton"
diff --git a/tests/engine/test_mm_encoder.py b/tests/engine/test_mm_encoder.py
index 1c12de588..5c63f85c4 100644
--- a/tests/engine/test_mm_encoder.py
+++ b/tests/engine/test_mm_encoder.py
@@ -57,17 +57,21 @@ def test_entry_lives_until_its_consumer_gathers_the_last_row():
     assert eng.model.calls == 1
 
 
-def test_shared_image_is_encoded_once_and_sliced_per_request():
+@pytest.mark.parametrize("legacy_embeds", [False, True])
+def test_shared_image_is_encoded_once_and_sliced_per_request(legacy_embeds):
     cache = EncoderCache(storage="cpu")
     eng = _engine(cache)
     a, b = _item(h=5, n_tokens=4), _item(h=5, n_tokens=4)
     cache.register(5, 1, 4)
     cache.register(5, 2, 4)
     batch = _batch([a, b], [(1, 5, 0, 4, 4, 0), (2, 5, 0, 2, 4, 4)])
+    expected = torch.full((6, H), 5.0)
+    if legacy_embeds:
+        batch.mm_embeds = torch.full((2, H), 7.0)
+        expected = torch.cat([expected, batch.mm_embeds])
     Engine._run_mm_encoder(eng, batch)
     assert eng.model.calls == 1
-    assert batch.mm_embeds.shape == (6, H)
-    assert torch.equal(batch.mm_embeds, torch.full((6, H), 5.0))
+    assert torch.equal(batch.mm_embeds, expected)
     # request 1 consumed its image; request 2 still has rows to gather in its next chunk
     assert cache._entries[5].remaining == {2: 2}
 
diff --git a/tests/kernels/test_dsv41_aot.py b/tests/kernels/test_dsv41_aot.py
new file mode 100644
index 000000000..41f505915
--- /dev/null
+++ b/tests/kernels/test_dsv41_aot.py
@@ -0,0 +1,59 @@
+"""V4.1's native expert-copy rows must have matching prebuilt-cache spec names."""
+
+import importlib
+
+import pytest
+import torch
+
+from freetoken.kernel.aot_models import (
+    SUPPORTED_MODELS, expert_bank_row_bytes, index_variants, store_element_sizes,
+)
+
+
+ROWS = {
+    "gate_up_packed": 11_796_480,
+    "gate_up_scale": 1_474_560,
+    "gate_up_global": 9_216,
+    "down_packed": 5_898_240,
+    "down_scale": 737_280,
+    "down_global": 10_240,
+}
+
+
+def test_v41_aot_entry_matches_native_expert_layout():
+    entry = next(m for m in SUPPORTED_MODELS if m.architecture == "DeepseekV41ForCausalLM")
+    assert (entry.hidden_size, entry.moe_intermediate_size, entry.top_k) == (5120, 2304, 6)
+    assert entry.expert_formats == ("nvfp4",)
+    assert not store_element_sizes(entry) and not index_variants(entry)
+    assert expert_bank_row_bytes("nvfp4", entry.hidden_size, entry.moe_intermediate_size) == ROWS
+
+
+def test_v41_runtime_copy_names_are_in_default_aot_specs(monkeypatch):
+    from freetoken.kernel.aot import default_kernel_specs
+    from freetoken.kernel.utils import _make_name
+
+    copy = importlib.import_module("freetoken.kernel.fast_index_copy")
+    monkeypatch.setattr(copy, "load_jit", lambda name, *args, **kwargs: _make_name(name, *args))
+    claimed = {spec.name for spec in default_kernel_specs()}
+    for feature_size in ROWS.values():
+        threads, worker_size, blocks = copy.default_worker_args(feature_size)
+        runtime_name = copy._jit_fast_index_copy_module.__wrapped__(
+            feature_size=feature_size, worker_threads=threads,
+            worker_feature_size=worker_size, num_block=blocks,
+        )
+        assert runtime_name in claimed
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("feature_size", ROWS.values(), ids=ROWS.keys())
+def test_v41_native_bank_row_copy_matches_host_bytes(feature_size):
+    from freetoken.kernel.fast_index_copy import fast_index_copy_jit
+
+    src = torch.randint(0, 256, (3, feature_size), dtype=torch.uint8, pin_memory=True)
+    dst = torch.zeros((4, feature_size), dtype=torch.uint8, device="cuda")
+    src_indices = torch.tensor([2, 0], dtype=torch.int32, device="cuda")
+    dst_indices = torch.tensor([1, 3], dtype=torch.int32, device="cuda")
+    fast_index_copy_jit(dst, dst_indices, src, src_indices)
+    expected = torch.zeros_like(dst, device="cpu")
+    expected[[1, 3]] = src[[2, 0]]
+    assert torch.equal(dst.cpu(), expected)
diff --git a/tests/kernels/test_dsv41_indexer.py b/tests/kernels/test_dsv41_indexer.py
new file mode 100644
index 000000000..9277d7f3d
--- /dev/null
+++ b/tests/kernels/test_dsv41_indexer.py
@@ -0,0 +1,140 @@
+import pytest
+import torch
+
+from freetoken.kernel.triton.dsv41.indexer import index_scores, select_indices
+
+
+@pytest.mark.parametrize("device", ["cpu", pytest.param("cuda", marks=pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required"))])
+def test_paged_scores_match_reference_with_noncontiguous_pages_and_per_query_limits(device):
+    torch.manual_seed(180)
+    q = torch.randn(3, 4, 32, device=device, dtype=torch.bfloat16)
+    weights = torch.randn(3, 4, device=device, dtype=torch.bfloat16)
+    pool = torch.randn(256, 32, device=device, dtype=torch.bfloat16)
+    table = torch.stack([torch.cat((torch.arange(base, base + 64), torch.arange(0, 64)))
+                         for base in [128, 192, 256]]).to(device)
+    ids = torch.arange(64, device=device).expand(3, -1)
+    valid = torch.tensor([64, 31, 0], device=device)
+    got = index_scores(q, weights, pool, table, ids, 2, valid)
+    rows = table[:, ::2] // 2
+    keys = pool[rows]
+    dot = torch.einsum("qhd,qkd->qhk", q.float(), keys.float()).to(q.dtype)
+    ref = (dot.relu() * weights[..., None]).sum(1)
+    ref.masked_fill_(ids >= valid[:, None], -torch.inf)
+    torch.testing.assert_close(got, ref, atol=0, rtol=0)
+
+
+def reference_selection(q, weights, keys, valid, topk, candidate_count, block_size, mask=None):
+    dot = torch.einsum("qhd,kd->qhk", q.float(), keys.float()).to(q.dtype)
+    scores = (dot.relu() * weights[..., None]).sum(1)
+    scores.masked_fill_(torch.arange(keys.shape[0])[None] >= valid[:, None], -torch.inf)
+    if mask is not None:
+        scores.masked_fill_(~mask, -torch.inf)
+    picks = scores.argsort(dim=-1, descending=True, stable=True)[..., :topk].sort(-1).values
+    picks = torch.where(picks < valid[:, None], picks, -1)
+    if not candidate_count:
+        return picks, None
+    padded = torch.nn.functional.pad(scores, (0, -keys.shape[0] % block_size), value=-torch.inf)
+    block_scores = padded.unflatten(-1, (-1, block_size)).amax(-1)
+    newest = (valid - 1) // block_size
+    block_scores.masked_fill_(torch.arange(block_scores.shape[1])[None] == newest[:, None], torch.inf)
+    best = block_scores.argsort(dim=-1, descending=True, stable=True)[..., :candidate_count]
+    keep = torch.zeros_like(block_scores, dtype=torch.bool).scatter_(-1, best, block_scores.gather(-1, best) > -torch.inf)
+    return picks, keep.repeat_interleave(block_size, -1)[:, :keys.shape[0]]
+
+
+def test_streaming_topk_candidate_blocks_and_second_stage_match_dense_reference():
+    torch.manual_seed(32)
+    q = torch.randn(5, 3, 32)
+    weights = torch.rand(5, 3)
+    keys = torch.randn(77, 32)
+    table = torch.arange(77).expand(5, -1)
+    valid = torch.tensor([1, 14, 45, 76, 77])
+    got, blocks = select_indices(q, weights, keys, table, valid, 77, 1, 5,
+                                 candidate_topk=3, block_size=8, query_tile=2, key_tile=16)
+    ref, mask = reference_selection(q, weights, keys, valid, 5, 3, 8)
+    torch.testing.assert_close(got, ref)
+    expanded = torch.zeros_like(mask)
+    for row in range(5):
+        for block in blocks[row].tolist():
+            if block >= 0:
+                expanded[row, block * 8:(block + 1) * 8] = True
+    assert torch.equal(expanded, mask)
+    q2 = torch.randn_like(q)
+    got2, _ = select_indices(q2, weights, keys, table, valid, 77, 1, 5,
+                             candidates=blocks, block_size=8, query_tile=2, key_tile=16)
+    ref2, _ = reference_selection(q2, weights, keys, valid, 5, 0, 8, mask)
+    torch.testing.assert_close(got2, ref2)
+
+
+def test_empty_index_history_has_no_candidates_or_read():
+    q = torch.randn(2, 3, 32)
+    got, _ = select_indices(q, torch.ones(2, 3), torch.zeros(1, 32), torch.zeros(2, 1, dtype=torch.long),
+                             torch.zeros(2, dtype=torch.long), 0, 2, 8)
+    assert got.shape == (2, 0)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+@pytest.mark.parametrize("packed", [False, True])
+def test_million_positions_keep_score_working_memory_bounded(packed):
+    width = 1 << 20
+    q = torch.zeros(33, 4, 32, device="cuda", dtype=torch.bfloat16)
+    weights = torch.ones(33, 4, device="cuda", dtype=torch.bfloat16)
+    pool = torch.zeros(width, 17 if packed else 32, device="cuda",
+                       dtype=torch.uint8 if packed else torch.bfloat16)
+    table = torch.arange(width, device="cuda").expand(33, -1)
+    valid = torch.full((33,), width, device="cuda")
+    torch.cuda.synchronize()
+    baseline = torch.cuda.memory_allocated()
+    torch.cuda.reset_peak_memory_stats()
+    selected, candidates = select_indices(q, weights, pool, table, valid, width, 1, 512,
+                                          candidate_topk=2048, block_size=8)
+    torch.cuda.synchronize()
+    assert torch.cuda.max_memory_allocated() - baseline < 16 << 20
+    torch.testing.assert_close(selected, torch.arange(512, device="cuda").expand(33, -1))
+    assert ((candidates == (width - 1) // 8).any(-1)).all()
+
+
+@pytest.mark.parametrize("device", ["cpu", pytest.param("cuda", marks=pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required"))])
+@pytest.mark.parametrize("ratio", [1, 2])
+@pytest.mark.parametrize("heads", [4, 32])
+def test_packed_index_scores_and_candidate_selection_match_restored_keys(device, ratio, heads):
+    torch.manual_seed(917)
+    dim, total, width = 128, 257, 77
+    raw = torch.randint(0, 256, (total, dim // 2), dtype=torch.uint8)
+    scales = torch.randint(121, 128, (total, dim // 32), dtype=torch.uint8)
+    packed = torch.cat((raw, scales), -1).to(device)
+    codes = torch.stack((raw & 15, raw >> 4), -1).flatten(-2).long()
+    magnitudes = torch.tensor([0., .5, 1., 1.5, 2., 3., 4., 6.])
+    restored = (magnitudes[codes & 7] * torch.where(codes & 8 > 0, -1, 1)
+                * torch.exp2(scales.float() - 127).repeat_interleave(32, -1)).bfloat16().to(device)
+    q = torch.randn(3, heads, dim, device=device, dtype=torch.bfloat16)
+    weights = torch.rand(3, heads, device=device, dtype=torch.bfloat16)
+    rows = torch.randperm(total, device=device)[:width].expand(3, -1)
+    table = (rows * ratio)[:, :, None] + torch.arange(ratio, device=device)
+    table = table.flatten(1)
+    table[1, 8:12] = -1
+    valid = torch.tensor([77, 45, 0], device=device)
+    ids = torch.arange(width, device=device).expand(3, -1).clone()
+    ids[:, ::13] = -1
+    expected = index_scores(q, weights, restored, table, ids, ratio, valid)
+    actual = index_scores(q, weights, packed, table, ids, ratio, valid)
+    torch.testing.assert_close(actual, expected, atol=0, rtol=0)
+    options = dict(candidate_topk=3, block_size=8, query_tile=2, key_tile=16)
+    expected_ids, expected_blocks = select_indices(q, weights, restored, table, valid, width, ratio, 7, **options)
+    actual_ids, actual_blocks = select_indices(q, weights, packed, table, valid, width, ratio, 7, **options)
+    torch.testing.assert_close(actual_ids, expected_ids)
+    torch.testing.assert_close(actual_blocks, expected_blocks)
+    options.pop("candidate_topk")
+    next_q = torch.randn_like(q)
+    expected, _ = select_indices(next_q, weights, restored, table, valid, width, ratio, 7,
+                                 candidates=expected_blocks, **options)
+    actual, _ = select_indices(next_q, weights, packed, table, valid, width, ratio, 7,
+                               candidates=actual_blocks, **options)
+    torch.testing.assert_close(actual, expected)
+
+
+def test_packed_index_rejects_incomplete_scale_row():
+    q = torch.zeros(1, 4, 128)
+    with pytest.raises(ValueError, match="Packed index"):
+        index_scores(q, torch.ones(1, 4), torch.zeros(3, 65, dtype=torch.uint8),
+                     torch.arange(3)[None], torch.arange(3)[None], 1, torch.tensor([3]))
diff --git a/tests/kernels/test_dsv41_quant.py b/tests/kernels/test_dsv41_quant.py
new file mode 100644
index 000000000..2b542e1cd
--- /dev/null
+++ b/tests/kernels/test_dsv41_quant.py
@@ -0,0 +1,194 @@
+import pytest
+import torch
+
+from freetoken.kernel.triton.dsv41.quant import (
+    block_fp8_linear, fp4_roundtrip, fp8_roundtrip,
+    pack_fp4, pack_fp8, unpack_fp4, unpack_fp8,
+)
+
+
+def _fp4_reference(x, block, fmt):
+    groups = x.float().reshape(-1, block)
+    amax = groups.abs().amax(-1, keepdim=True)
+    if fmt == "e4m3":
+        scales = (amax / 6).clamp(1 / 512, 448).to(torch.float8_e4m3fn).float()
+    else:
+        scales = 2.0 ** torch.ceil(torch.log2(amax.clamp_min(6 * 2.0**-126) / 6))
+    y = (groups / scales).clamp(-6, 6)
+    thresholds = torch.tensor([.25, .75, 1.25, 1.75, 2.5, 3.5, 5.])
+    codes = torch.bucketize(y.abs().contiguous(), thresholds)
+    for index in (1, 3, 5):
+        codes[y.abs() == thresholds[index]] = index + 1
+    grid = torch.tensor([0, .5, 1, 1.5, 2, 3, 4, 6])
+    return (grid[codes] * y.sign() * scales).reshape_as(x).to(x.dtype)
+
+
+@pytest.mark.parametrize("fmt,block", [("e4m3", 16), ("e8m0", 32)])
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_fp4_roundtrip_formats(fmt, block, device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    torch.manual_seed(42)
+    x = torch.randn(35, 128, dtype=torch.bfloat16)
+    x[0].zero_()
+    x[1] *= 0.001
+    x[2] *= 10
+    x[3, :16] = torch.tensor([0, .25, .5, .75, 1, 1.25, 1.5, 1.75, 2, 2.5, 3, 3.5, 4, 5, 6, -6])
+    expected = _fp4_reference(x, block, fmt)
+    actual = fp4_roundtrip(x.to(device), block, fmt).cpu()
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+    assert torch.isfinite(actual).all()
+
+
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_fp8_roundtrip_block32(device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    torch.manual_seed(19)
+    x = torch.randn(3, 128, dtype=torch.bfloat16)
+    x[:, 32:64] *= 128
+    x[:, 96:].zero_()
+    groups = x.float().reshape(3, 4, 32)
+    scale = 2.0 ** torch.ceil(torch.log2(groups.abs().amax(-1, keepdim=True).clamp_min(1e-4) / 448))
+    expected = ((groups / scale).to(torch.float8_e4m3fn).float() * scale).reshape_as(x).to(x.dtype)
+    torch.testing.assert_close(fp8_roundtrip(x.to(device)).cpu(), expected, rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_fp4_e4m3_scale_saturates_to_representable_range(device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    x = torch.full((2, 16), 6144, dtype=torch.bfloat16, device=device)
+    x[1].neg_()
+    expected = torch.full_like(x, 2688)
+    expected[1].neg_()
+    torch.testing.assert_close(fp4_roundtrip(x), expected, rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("rows", [1, 2, 35])
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_block32_linear_independent_scales(rows, device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    torch.manual_seed(31)
+    x = torch.randn(rows, 160, dtype=torch.bfloat16)
+    w = torch.randn(96, 160).to(torch.float8_e4m3fn)
+    scale = torch.tensor([[125, 126, 127, 128, 129], [129, 125, 126, 128, 127], [127, 129, 125, 126, 128]], dtype=torch.uint8)
+    a = x.float().reshape(rows, 5, 32)
+    sa = 2.0 ** torch.ceil(torch.log2(a.abs().amax(-1, keepdim=True).clamp_min(1e-4) / 448))
+    aq = (a / sa).to(torch.float8_e4m3fn).float()
+    expected = torch.zeros(rows, 96)
+    for kb in range(5):
+        block_w = w.float()[:, kb*32:(kb+1)*32]
+        sb = (2.0 ** (scale[:, kb].float() - 127)).repeat_interleave(32)
+        expected += (aq[:, kb] @ block_w.T) * sa[:, kb] * sb
+    actual = block_fp8_linear(x.to(device), w.to(device), scale.to(device)).cpu()
+    torch.testing.assert_close(actual.float(), expected.to(torch.bfloat16).float(), rtol=0.008, atol=0.02)
+
+
+def test_invalid_scale_geometry():
+    with pytest.raises(ValueError, match="scale matrix"):
+        block_fp8_linear(torch.zeros(1, 64), torch.zeros(64, 64).to(torch.float8_e4m3fn), torch.zeros(1, 1, dtype=torch.uint8))
+
+
+@pytest.mark.parametrize("fmt,block,width", [("fp8", 32, 512), ("e4m3", 16, 512), ("e8m0", 32, 128)])
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_packed_native_rows_reconstruct_existing_roundtrip(fmt, block, width, device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    torch.manual_seed(122)
+    x = torch.randn(2, 35, width * 2, dtype=torch.bfloat16)[..., ::2]
+    x[:, 0].zero_()
+    x[:, 1] *= .001
+    x[:, 2] *= 10000
+    x[:, 3] *= 2.0**-127
+    x[:, 4, :16] = torch.tensor([0, .25, .5, .75, 1, 1.25, 1.5, 1.75, 2, 2.5, 3, 3.5, 4, 5, 6, -6])
+    source = x.to(device)
+    if fmt == "fp8":
+        packed = pack_fp8(source, block)
+        actual = unpack_fp8(packed, block)
+        expected = fp8_roundtrip(source, block)
+        groups = x.float().reshape(2, 35, -1, block)
+        scales = 2.0 ** torch.ceil(torch.log2(groups.abs().amax(-1, keepdim=True).clamp_min(1e-4) / 448))
+        reference = ((groups / scales).to(torch.float8_e4m3fn).float() * scales).reshape_as(x).bfloat16()
+        cpu_packed = pack_fp8(x, block)
+        row_bytes = width + width // block
+    else:
+        packed = pack_fp4(source, block, fmt)
+        actual = unpack_fp4(packed, block, fmt)
+        expected = fp4_roundtrip(source, block, fmt)
+        reference = _fp4_reference(x, block, fmt)
+        cpu_packed = pack_fp4(x, block, fmt)
+        row_bytes = width // 2 + width // block
+    assert packed.shape == (2, 35, row_bytes) and packed.dtype == torch.uint8
+    assert packed.is_contiguous()
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+    torch.testing.assert_close(actual.cpu(), reference, rtol=0, atol=0)
+    torch.testing.assert_close(packed.cpu(), cpu_packed, rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("fmt,block,scale_byte", [("e4m3", 16, 56), ("e8m0", 32, 127)])
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_fp4_pack_nibble_order_and_inline_scale_bytes(fmt, block, scale_byte, device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    row = [0, .5, 1, 1.5, 2, 3, 4, 6, -0., -.5, -1, -1.5, -2, -3, -4, -6]
+    x = torch.tensor(row * (block // 16), dtype=torch.bfloat16, device=device)
+    expected = torch.tensor([0x10, 0x32, 0x54, 0x76, 0x90, 0xBA, 0xDC, 0xFE] * (block // 16)
+                            + [scale_byte], dtype=torch.uint8, device=device)
+    torch.testing.assert_close(pack_fp4(x, block, fmt), expected, rtol=0, atol=0)
+    torch.testing.assert_close(unpack_fp4(expected, block, fmt), x, rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("fmt,block", [("fp8", 32), ("e4m3", 16), ("e8m0", 32)])
+@pytest.mark.parametrize("shape", [(0, 128), (2, 0, 128), (128,)])
+def test_packed_codecs_keep_empty_and_leading_shapes(fmt, block, shape):
+    x = torch.zeros(shape, dtype=torch.bfloat16)
+    if fmt == "fp8":
+        packed = pack_fp8(x, block)
+        restored = unpack_fp8(packed, block)
+    else:
+        packed = pack_fp4(x, block, fmt)
+        restored = unpack_fp4(packed, block, fmt)
+    assert restored.shape == x.shape
+    torch.testing.assert_close(restored, x, rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_packed_index_decode_handles_all_e8m0_scale_codes(device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    packed = torch.full((256, 17), 0x22, dtype=torch.uint8)
+    packed[:, -1] = torch.arange(256, dtype=torch.uint8)
+    expected = packed[:, -1:].contiguous().view(torch.float8_e8m0fnu).float().expand(-1, 32).bfloat16()
+    actual = unpack_fp4(packed.to(device), 32, "e8m0").cpu()
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0, equal_nan=True)
+
+
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+@pytest.mark.parametrize("tier", ["window", "compressed"])
+def test_packed_decode_handles_all_e4m3_codes_and_scales(device, tier):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("CUDA required")
+    codes = torch.arange(256, dtype=torch.uint8)
+    if tier == "window":
+        packed = torch.cat((codes[:, None].expand(-1, 32), torch.full((256, 1), 127, dtype=torch.uint8)), 1)
+        actual = unpack_fp8(packed.to(device)).cpu()
+        width = 32
+    else:
+        packed = torch.cat((torch.full((256, 8), 0x22, dtype=torch.uint8), codes[:, None]), 1)
+        actual = unpack_fp4(packed.to(device), 16, "e4m3").cpu()
+        width = 16
+    expected = codes.view(torch.float8_e4m3fn).float()[:, None].expand(-1, width).bfloat16()
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0, equal_nan=True)
+    assert torch.isnan(actual[[127, 255]]).all()
+
+
+@pytest.mark.parametrize("fn,args", [(pack_fp8, (torch.zeros(2, 33),)),
+                                    (pack_fp4, (torch.zeros(2, 17), 16, "e4m3")),
+                                    (unpack_fp8, (torch.zeros(2, 34, dtype=torch.uint8),)),
+                                    (unpack_fp4, (torch.zeros(2, 18), 16, "e4m3")),
+                                    (unpack_fp4, (torch.zeros(2, 18, dtype=torch.uint8), 16, "unknown"))])
+def test_packed_codecs_reject_bad_layout(fn, args):
+    with pytest.raises(ValueError):
+        fn(*args)
diff --git a/tests/kernels/test_dsv41_sparse.py b/tests/kernels/test_dsv41_sparse.py
new file mode 100644
index 000000000..a0a8062b9
--- /dev/null
+++ b/tests/kernels/test_dsv41_sparse.py
@@ -0,0 +1,156 @@
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.attention.dsv41_sparse import DSV41SparseAttnBackend
+from freetoken.kernel.triton.dsv41.sparse_attn import sparse_attn_paged
+
+
+def packed_pools(dim, device):
+    generator = torch.Generator().manual_seed(514)
+    window_codes = torch.randn(43, dim, generator=generator).to(torch.float8_e4m3fn).view(torch.uint8)
+    window_scales = torch.randint(123, 130, (43, dim // 32), generator=generator, dtype=torch.uint8)
+    window = torch.cat((window_codes, window_scales), -1).to(device)
+    restored_window = (window_codes.view(torch.float8_e4m3fn).float()
+                       * torch.exp2(window_scales.float() - 127).repeat_interleave(32, -1)).bfloat16().to(device)
+    cmp_codes = torch.randint(0, 256, (71, dim // 2), generator=generator, dtype=torch.uint8)
+    scale_grid = torch.tensor([.09375, .125, .28125, .3125, .625, 1.25, 3.5])
+    cmp_scales = scale_grid[torch.randint(0, len(scale_grid), (71, dim // 16), generator=generator)]
+    compressed = torch.cat((cmp_codes, cmp_scales.to(torch.float8_e4m3fn).view(torch.uint8)), -1).to(device)
+    code = torch.stack((cmp_codes & 15, cmp_codes >> 4), -1).flatten(-2).long()
+    values = torch.tensor([0., .5, 1., 1.5, 2., 3., 4., 6.])
+    restored_compressed = (values[code & 7] * torch.where(code & 8 > 0, -1, 1)
+                           * cmp_scales.repeat_interleave(16, -1)).bfloat16().to(device)
+    return window, compressed, restored_window, restored_compressed
+
+
+def inputs(dim, queries, device):
+    torch.manual_seed(103)
+    q = (torch.randn(2, queries, 5, dim, device=device) * .15).bfloat16()
+    sink = torch.linspace(-1, 1, 5, device=device)
+    window_ids = torch.randint(0, 43, (2, queries, 33), device=device)
+    cmp_ids = torch.randint(0, 71, (2, queries, 148), device=device)
+    ids = torch.cat((window_ids, cmp_ids), -1).int()
+    ids[..., ::7] = -1
+    ids[0, 0] = -1
+    counts = torch.full((2, queries), 117, dtype=torch.int32, device=device)
+    counts[0, 0] = 0
+    return q, sink, ids, counts
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+@pytest.mark.parametrize("dim", [32, 512])
+@pytest.mark.parametrize("queries,splits", [(1, 0), (1, 4), (3, 0), (3, 4)])
+def test_packed_sparse_matches_bf16_kernel_for_mixed_tiles_and_counts(monkeypatch, dim, queries, splits):
+    from freetoken.kernel.triton.dsv4 import sparse_attn as reference
+
+    win, cmp, restored_win, restored_cmp = packed_pools(dim, "cuda")
+    q, sink, ids, counts = inputs(dim, queries, "cuda")
+    monkeypatch.setattr(reference, "split_count", lambda *args: splits)
+    expected = reference.sparse_attn_paged(q, restored_win, restored_cmp, sink, ids, 33, dim ** -.5, counts)
+    actual = sparse_attn_paged(q, win, cmp, sink, ids, 33, dim ** -.5, counts, force_splits=splits)
+    torch.testing.assert_close(actual, expected, atol=0, rtol=0)
+    assert torch.count_nonzero(actual[0, 0]) == 0
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+@pytest.mark.parametrize("splits", [0, 4])
+def test_packed_window_only_alias_and_no_count_limit(monkeypatch, splits):
+    from freetoken.kernel.triton.dsv4 import sparse_attn as reference
+
+    win, _, restored, _ = packed_pools(512, "cuda")
+    q, sink, ids, _ = inputs(512, 1, "cuda")
+    ids = ids[..., :33]
+    monkeypatch.setattr(reference, "split_count", lambda *args: splits)
+    expected = reference.sparse_attn_paged(q, restored, restored, sink, ids, 33, 512 ** -.5)
+    actual = sparse_attn_paged(q, win, win, sink, ids, 33, 512 ** -.5, force_splits=splits)
+    torch.testing.assert_close(actual, expected, atol=0, rtol=0)
+
+
+@pytest.mark.parametrize("source", [None, 0])
+def test_cpu_backend_restores_selected_rows_and_honors_live_counts(monkeypatch, source):
+    from freetoken.kernel.triton.dsv41 import quant
+
+    win, cmp, restored_win, restored_cmp = packed_pools(32, "cpu")
+    pool = SimpleNamespace(kv_sources=[source], window_pool=[win], cmp_pool=[cmp], kv_quant="fp8-fp4")
+    backend = DSV41SparseAttnBackend.__new__(DSV41SparseAttnBackend)
+    monkeypatch.setattr(DSV41SparseAttnBackend, "pool", property(lambda self: pool))
+    q, sink, ids, counts = inputs(32, 3, "cpu")
+    if source is None:
+        ids, counts = ids[..., :33], None
+    original_unpack = quant.unpack_fp8
+    seen = []
+
+    def unpack_selected(packed, **kwargs):
+        seen.append(tuple(packed.shape))
+        assert packed.shape == (6, 33, 33)
+        return original_unpack(packed, **kwargs)
+
+    monkeypatch.setattr(quant, "unpack_fp8", unpack_selected)
+    actual = backend.attend(q, 0, ids, 33, sink, 32 ** -.5, counts, source is not None)
+    pool.kv_quant = "none"
+    pool.window_pool = [restored_win]
+    pool.cmp_pool = [restored_cmp]
+    expected = backend.attend(q, 0, ids, 33, sink, 32 ** -.5, counts, source is not None)
+    torch.testing.assert_close(actual, expected, atol=0, rtol=0)
+    assert len(seen) == 1
+
+
+def test_backend_routes_compressed_stores_without_casting(monkeypatch):
+    calls = []
+    pool = SimpleNamespace(store_compressed=lambda *args: calls.append(("attn", args)),
+                           store_indexer=lambda *args: calls.append(("idx", args)))
+    monkeypatch.setattr(DSV41SparseAttnBackend, "pool", property(lambda self: pool))
+    backend = DSV41SparseAttnBackend.__new__(DSV41SparseAttnBackend)
+    rows, packed = torch.tensor([1, 5]), torch.tensor([[255, 129], [3, 192]], dtype=torch.uint8)
+    for tier in ("attn", "idx"):
+        backend.scatter_compressed(2, tier, rows, packed)
+    assert [call[0] for call in calls] == ["attn", "idx"]
+    for _, args in calls:
+        assert args[0] is packed and args[1] == 2 and args[2] is rows
+    with pytest.raises(ValueError, match="tier"):
+        backend.scatter_compressed(2, "unknown", rows, packed)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+def test_packed_sparse_cuda_graph_reads_new_indices_counts_and_codes():
+    win, cmp, _, _ = packed_pools(512, "cuda")
+    q, sink, ids, counts = inputs(512, 1, "cuda")
+    stream = torch.cuda.Stream()
+    stream.wait_stream(torch.cuda.current_stream())
+    with torch.cuda.stream(stream):
+        for _ in range(2):
+            sparse_attn_paged(q, win, cmp, sink, ids, 33, 512 ** -.5, counts, force_splits=4)
+    torch.cuda.current_stream().wait_stream(stream)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph, stream=stream):
+        captured = sparse_attn_paged(q, win, cmp, sink, ids, 33, 512 ** -.5, counts, force_splits=4)
+    ids[1, 0, :33].fill_(3)
+    ids[1, 0, 33:].fill_(7)
+    counts[1, 0] = 71
+    win[3, :512].zero_()
+    cmp[7, :256].fill_(0x77)
+    graph.replay()
+    expected = sparse_attn_paged(q, win, cmp, sink, ids, 33, 512 ** -.5, counts, force_splits=4)
+    torch.testing.assert_close(captured, expected, atol=0, rtol=0)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+@pytest.mark.parametrize("seed", [514, 515, 516])
+@pytest.mark.parametrize("queries", [1, 32])
+def test_native_head_count_preserves_attention_with_bf16_rounding_tolerance(seed, queries):
+    from freetoken.kernel.triton.dsv4 import sparse_attn as reference
+    from freetoken.kernel.triton.dsv41.quant import pack_fp4, pack_fp8, unpack_fp4, unpack_fp8
+
+    torch.manual_seed(seed)
+    win = pack_fp8(torch.randn(2048, 512, device="cuda", dtype=torch.bfloat16) * .2)
+    cmp = pack_fp4(torch.randn(8192, 512, device="cuda", dtype=torch.bfloat16) * .2)
+    q = torch.randn(1, queries, 64, 512, device="cuda", dtype=torch.bfloat16) * .1
+    sink = torch.linspace(-1, 1, 64, device="cuda")
+    ids = torch.cat((torch.randint(0, 2048, (1, queries, 128), device="cuda"),
+                     torch.randint(0, 8192, (1, queries, 512), device="cuda")), -1).int()
+    expected = reference.sparse_attn_paged(q, unpack_fp8(win), unpack_fp4(cmp), sink, ids, 128, 512 ** -.5)
+    actual = sparse_attn_paged(q, win, cmp, sink, ids, 128, 512 ** -.5)
+    # FP32 dot layouts can differ at a final BF16 rounding boundary; one BF16 ULP is sufficient.
+    torch.testing.assert_close(actual, expected, atol=1e-6, rtol=1 / 128)
diff --git a/tests/kernels/test_e4m3_compat.py b/tests/kernels/test_e4m3_compat.py
index 8bcb8e581..1a9c57f48 100644
--- a/tests/kernels/test_e4m3_compat.py
+++ b/tests/kernels/test_e4m3_compat.py
@@ -16,6 +16,9 @@
    capability patched to sm_80/86/89/120 and every triton launch forced into
    warmup (compile-only), the full wrapper->kernel paths -- uint8 views, PDL
    gates, e4m3 branches -- must compile for the foreign arch.
+4. A cache-key guard: the @constexpr_function probes must not reference plain
+   python functions. triton's AST walk rejects that, and it disables every kernel
+   that branches on e4m3 native-ness at once, so the failure is not local.
 """
 
 from __future__ import annotations
@@ -41,6 +44,32 @@ def _native_cc() -> bool:
     return torch.cuda.get_device_capability() >= (8, 9)
 
 
+def test_constexpr_probe_never_references_a_host_function():
+    """triton hashes an @constexpr_function by walking its AST (runtime/jit.py:
+    cache_key -> record_reference), and a bare reference to a plain python function
+    raises "Unsupported function referenced: <function e4m3_native ...>". Making
+    e4m3_native_cx defer to the host probe did exactly that, and it took out EVERY
+    kernel that branches on e4m3 native-ness at once (here: the PLE gather, inside
+    CUDA graph capture). Modules (``target_info.cuda_capability_geq``) survive the
+    walk, functions do not -- so unifying the two probes has to go the other way: the
+    code that owns a buffer follows the buffer (tests/kernels/test_kv_fp8.py)."""
+    import inspect
+
+    from freetoken.kernel.triton import e4m3_compat
+
+    lines = inspect.getsource(e4m3_compat).splitlines()
+    start = next(i for i, ln in enumerate(lines) if ln.startswith("def e4m3_native_cx("))
+    body = []
+    for line in lines[start + 1:]:
+        if line and not line.startswith((" ", "\t")):
+            break
+        body.append(line)
+    assert "e4m3_native(" not in "\n".join(body), (
+        "a constexpr_function may not call a host function: triton's cache-key walk "
+        "rejects it and every e4m3 kernel stops compiling"
+    )
+
+
 # ======================================================================================
 # 1. Primitives vs the native fp8 unit (needs sm_89+ hardware for the reference).
 # ======================================================================================
@@ -113,6 +142,49 @@ def k(x_ptr, y_ptr, N, BLOCK: tl.constexpr):
         assert int(((y != ref) & ~(y.isnan() & ref.isnan())).sum()) == 0
 
 
+def test_kv_tile_scaled16_agrees_with_the_f32_loader():
+    """kv_load_e4m3_tile_scaled16 is kv_load_e4m3_tile_f32 with its last two steps --
+    the widen to fp32 and the ``* 256.0`` -- left for the caller to fold into the
+    per-(token, kv_head) dequant scale it has to apply anyway. That fold is only legal
+    if the two loaders agree on every code, so pin it: all 256, NaN patterns included.
+
+    Second assertion pins the property that lets the 16-bit tile be narrowed to a bf16
+    compute dtype for free: every value carries at most the code's own 3 mantissa bits
+    and lies in +-1.75, so bf16 holds it exactly. That is what makes scaling AFTER the
+    dot strictly more accurate than upstream's scale-then-narrow.
+
+    Third pins masked lanes at zero in both, so a masked tile contributes nothing once
+    the scale is applied to the dot output instead of to the tile."""
+    import triton
+    import triton.language as tl
+
+    from freetoken.kernel.triton.e4m3_compat import (
+        KV_TILE_SCALE,
+        kv_load_e4m3_tile_f32,
+        kv_load_e4m3_tile_scaled16,
+    )
+
+    @triton.jit
+    def k(v_ptr, wide_ptr, narrow_ptr, N, BLOCK: tl.constexpr):
+        offs = tl.arange(0, BLOCK)
+        mask = offs < N
+        tl.store(wide_ptr + offs, kv_load_e4m3_tile_f32(v_ptr + offs, mask))
+        tl.store(narrow_ptr + offs, kv_load_e4m3_tile_scaled16(v_ptr + offs, mask))
+
+    n = 256
+    v = torch.zeros(2 * n, dtype=torch.uint8, device="cuda")
+    v[:n] = torch.arange(n, dtype=torch.uint8, device="cuda")
+    wide = torch.empty(2 * n, dtype=torch.float32, device="cuda")
+    narrow = torch.empty(2 * n, dtype=torch.float16, device="cuda")
+    k[(1,)](v, wide, narrow, n, BLOCK=2 * n)
+
+    assert torch.equal(narrow.to(torch.float32) * KV_TILE_SCALE.value, wide)
+    assert torch.equal(
+        narrow.to(torch.bfloat16).to(torch.float32), narrow.to(torch.float32)
+    )
+    assert not wide[n:].any() and not narrow[n:].any()
+
+
 # ======================================================================================
 # Shared emit path: every affected wrapper, deterministic inputs.
 # ======================================================================================
@@ -228,6 +300,60 @@ def _emit_all(path: str) -> None:
     out["nv_moe_prefill"] = f32(fused_experts_nvfp4(
         hid4, gup, gus4, gug, dnp, dns4, dng, tw, tids, S))
 
+    # fp8 KV cache: the fused quantize+scatter writer the pools call in store_kv, and the
+    # QSA sparse reader that turns codes + row scales back into operands. Codes are plain
+    # bytes on every arch (kv_quant.kv_codes_dtype), so comparing them as BYTES is what
+    # pins a native run and a forced-EMU run to the very same encoding, and the two
+    # QSA runs -- one over codes, one over the same real numbers pre-rounded into bf16 --
+    # must produce identical bits: the reader casts back to the query dtype before the
+    # dot, so any decode difference lands on the comparison instead of hiding in a
+    # tolerance.
+    from freetoken.kernel.triton.kv_quant import (
+        alloc_codes, codes_to_f32, quantize_kv_to_cache,
+    )
+    from freetoken.kernel.triton.qsa import qsa_sparse_paged_attention
+
+    slots, kvh, hd, page = 128, 2, 64, 64
+    kv_in = (torch.randn(slots, kvh * hd, device=dev, dtype=torch.float32) * 4.0).to(
+        torch.bfloat16
+    )
+    k_codes = alloc_codes((slots, kvh, hd), dev)
+    v_codes = alloc_codes((slots, kvh, hd), dev)
+    k_sc = torch.zeros((slots, kvh), dtype=torch.float32, device=dev)
+    v_sc = torch.zeros_like(k_sc)
+    quantize_kv_to_cache(
+        k=kv_in,
+        v=kv_in.flip(-1).contiguous(),
+        out_loc=torch.arange(slots, dtype=torch.int32, device=dev),
+        k_cache=k_codes,
+        v_cache=v_codes,
+        k_scale=k_sc,
+        v_scale=v_sc,
+    )
+    out["kvfp8_codes"] = k_codes.view(torch.uint8).to(torch.int16).cpu()
+    out["kvfp8_scale"] = k_sc.cpu()
+
+    pages = slots // page
+    kc4, vc4 = k_codes.view(pages, page, kvh, hd), v_codes.view(pages, page, kvh, hd)
+    q = torch.randn(2, 2 * kvh, hd, device=dev, dtype=torch.bfloat16)
+    sel = (
+        torch.arange(2 * page, dtype=torch.int32, device=dev)[None, :]
+        .repeat(2, 1)
+        .contiguous()
+    )
+    table = torch.arange(pages, dtype=torch.int32, device=dev)[None, :].contiguous()
+    t2r = torch.zeros(2, dtype=torch.int32, device=dev)
+    out["kvfp8_qsa"] = f32(qsa_sparse_paged_attention(
+        q, kc4, vc4, sel, table, t2r, k_scale=k_sc, v_scale=v_sc))
+    # The very same numbers, pre-rounded into a bf16 cache. The reader casts its
+    # dequantized operands to the query dtype before tl.dot, so the two runs must end up
+    # bit-identical (asserted where launches really execute: tests/kernels/test_qsa_fp8.py
+    # -- the compile gate below runs warmup-only, where outputs are never written).
+    out["kvfp8_qsa_bf16"] = f32(qsa_sparse_paged_attention(
+        q,
+        (codes_to_f32(kc4) * k_sc.view(pages, page, kvh, 1)).to(torch.bfloat16),
+        (codes_to_f32(vc4) * v_sc.view(pages, page, kvh, 1)).to(torch.bfloat16),
+        sel, table, t2r))
     torch.save(out, path)
 
 
diff --git a/tests/kernels/test_kv_fp8.py b/tests/kernels/test_kv_fp8.py
new file mode 100644
index 000000000..817646886
--- /dev/null
+++ b/tests/kernels/test_kv_fp8.py
@@ -0,0 +1,263 @@
+"""FP8 (e4m3) KV quantization: the store kernel against an independent oracle.
+
+Expectations come from a brute-force nearest-code search over an e4m3 table decoded
+from first principles (sign / exponent / mantissa), NOT from the kernels' own
+rounding helpers -- so a drift in ``round_e4m3`` or in the new ``e4m3_f32_to_u8``
+encoder fails here instead of being blessed by itself.
+"""
+
+from __future__ import annotations
+
+import pytest
+import torch
+
+if not torch.cuda.is_available():  # pragma: no cover
+    pytest.skip("CUDA required", allow_module_level=True)
+
+from freetoken.distributed import set_tp_info, try_get_tp_info
+from freetoken.kernel.triton.kv_quant import (
+    KV_SCALE_DTYPE,
+    alloc_codes,
+    codes_to_f32,
+    kv_codes_dtype,
+    quantize_kv_to_cache,
+)
+
+DEV = torch.device("cuda")
+FP8_MAX = 448.0
+CODE_448 = 0x7E
+
+
+def _init_tp() -> None:
+    if try_get_tp_info() is None:
+        set_tp_info(rank=0, size=1)
+
+
+@pytest.fixture(autouse=True)
+def _tp():
+    _init_tp()
+
+
+def _e4m3_table() -> dict[int, float]:
+    """Every finite e4m3fn value, decoded by hand from its bit fields.
+
+    code = S EEEE MMM, exponent bias 7: normal (E>0) -> (-1)^S * 2^(E-7) * (1+M/8);
+    subnormal (E==0) -> (-1)^S * 2^-6 * (M/8). E==15 & M==7 (0x7F/0xFF) is the NaN
+    pattern and is dropped, which caps the format at +-448.
+    """
+    values: dict[int, float] = {}
+    for code in range(256):
+        sign = -1.0 if (code >> 7) & 1 else 1.0
+        exp, mant = (code >> 3) & 0x0F, code & 0x07
+        if exp == 15 and mant == 7:
+            continue
+        values[code] = sign * (
+            (2.0**-6) * (mant / 8.0) if exp == 0 else (2.0 ** (exp - 7)) * (1.0 + mant / 8.0)
+        )
+    return values
+
+
+E4M3_VALUES = _e4m3_table()
+CODES = torch.tensor(sorted(E4M3_VALUES), dtype=torch.int32)
+GRID = torch.tensor([E4M3_VALUES[c] for c in CODES.tolist()], dtype=torch.float64)
+NEGATIVE_ZERO = 0x80
+
+
+def _as_bytes(t: torch.Tensor) -> torch.Tensor:
+    """A code buffer as raw bytes (the fp8 view on sm_89+, uint8 below it)."""
+    return t if t.dtype == torch.uint8 else t.view(torch.uint8)
+
+
+def _canonical_zero(codes: torch.Tensor) -> torch.Tensor:
+    """Fold -0.0 (0x80) onto +0.0 (0x00).
+
+    The two store paths legitimately disagree on zero's sign bit: sm_89+ converts with
+    ``x.to(fp8e4nv)`` (keeps it), the emulated path rounds first and
+    ``round_e4m3(-0.0)`` is documented to return +0.0. They decode to the same number,
+    so a test that compares bytes must not care which one it got.
+    """
+    return torch.where(codes == NEGATIVE_ZERO, torch.zeros_like(codes), codes)
+
+
+def _ref_codes(x: torch.Tensor) -> torch.Tensor:
+    """Independent encoder: nearest grid value by brute force, ties resolved to the
+    EVEN code (RNE). ``x`` is any shape; returns int32 codes."""
+    grid = GRID.to(x.device)
+    codes = CODES.to(x.device)
+    dist = (x.to(torch.float64).reshape(-1, 1) - grid.unsqueeze(0)).abs()
+    near = dist == dist.min(dim=-1, keepdim=True).values
+    big = 1 << 30
+    codes = codes.unsqueeze(0).expand_as(dist)
+    any_code = torch.where(near, codes, torch.full_like(codes, big))
+    even = near & (codes % 2 == 0)
+    even_code = torch.where(even, codes, torch.full_like(codes, big))
+    has_even = (even_code < big).any(dim=-1)
+    return torch.where(has_even, even_code.min(dim=-1).values, any_code.min(dim=-1).values)
+
+
+def _store(rows_k: torch.Tensor, rows_v: torch.Tensor):
+    """Quantize ``[T, heads, dim]`` rows into a fresh code buffer, returning
+    ``(k_codes, v_codes, k_scales, v_scales)``."""
+    _init_tp()
+    tokens, heads, dim = rows_k.shape
+    k_cache = alloc_codes((tokens, heads, dim), DEV)
+    v_cache = alloc_codes((tokens, heads, dim), DEV)
+    k_scale = torch.zeros((tokens, heads), dtype=KV_SCALE_DTYPE, device=DEV)
+    v_scale = torch.zeros((tokens, heads), dtype=KV_SCALE_DTYPE, device=DEV)
+    quantize_kv_to_cache(
+        k=rows_k.reshape(tokens, -1),
+        v=rows_v.reshape(tokens, -1),
+        out_loc=torch.arange(tokens, dtype=torch.int32, device=DEV),
+        k_cache=k_cache,
+        v_cache=v_cache,
+        k_scale=k_scale,
+        v_scale=v_scale,
+    )
+    torch.cuda.synchronize()
+    return k_cache, v_cache, k_scale, v_scale
+
+
+def test_scale_is_amax_over_e4m3_max():
+    torch.manual_seed(0)
+    k = (torch.randn(4, 2, 64, device=DEV, dtype=torch.bfloat16) * 3.0).to(torch.bfloat16)
+    v = torch.randn_like(k)
+    _, _, k_scale, v_scale = _store(k, v)
+    # The kernel widens to fp32 before the amax/divide, so the reference must too:
+    # a bf16 intermediate would round the expected scale and hide a precision bug.
+    torch.testing.assert_close(
+        k_scale, k.abs().to(torch.float32).amax(dim=-1) / FP8_MAX, rtol=1e-6, atol=0
+    )
+    torch.testing.assert_close(
+        v_scale, v.abs().to(torch.float32).amax(dim=-1) / FP8_MAX, rtol=1e-6, atol=0
+    )
+
+
+def test_codes_match_the_reference_quantizer_and_reconstruction_is_close():
+    torch.manual_seed(1)
+    tokens, heads, dim = 8, 3, 128
+    # Feed the rows as a qkv slice, the way the attention backends really do: the
+    # row pitch is then wider than the row, which the store kernel must honour.
+    qkv = torch.randn(tokens, heads * dim * 3, device=DEV, dtype=torch.bfloat16)
+    qkv[:, heads * dim : 2 * heads * dim] *= 5.0  # K: large magnitude
+    qkv[:, 2 * heads * dim :] *= 0.01  # V: subnormal end of the e4m3 grid
+    _, k_rows, v_rows = qkv.split(heads * dim, dim=-1)
+    k = k_rows.view(tokens, heads, dim)
+    v = v_rows.view(tokens, heads, dim).clamp(-FP8_MAX, FP8_MAX)
+    k_cache, v_cache, k_scale, v_scale = _store(k, v)
+
+    for rows, cache, scale in ((k, k_cache, k_scale), (v, v_cache, v_scale)):
+        f32 = rows.to(torch.float32)
+        ref_scale = f32.abs().amax(dim=-1, keepdim=True) / FP8_MAX
+        expected = _ref_codes((f32 / ref_scale).clamp(-FP8_MAX, FP8_MAX))
+        got = _canonical_zero(_as_bytes(cache).reshape(-1).to(torch.int32))
+        expected = _canonical_zero(expected)
+        assert torch.equal(got, expected), (
+            f"{int((got != expected).sum())} code mismatches of {got.numel()}"
+        )
+        deq = codes_to_f32(cache) * scale.unsqueeze(-1)
+        err = (deq - f32).abs().max(dim=-1).values
+        assert torch.all(err <= 0.08 * f32.abs().amax(dim=-1)), float(err.max())
+
+
+def test_encoder_inverts_the_grid_through_the_scale_one_path():
+    """Pack every e4m3 grid value into a row that also holds 448.0: the row scale is
+    then exactly 1.0, so each stored byte IS the encoder's answer for that value."""
+    dim, per_row = 256, 255
+    pairs = sorted(E4M3_VALUES.items())
+    tokens = -(-len(pairs) // per_row)
+    rows = torch.zeros(tokens, 1, dim, dtype=torch.float32)
+    expected = torch.zeros(tokens, dim, dtype=torch.uint8)
+    for t in range(tokens):
+        rows[t, 0, 0] = FP8_MAX  # the amax anchor
+        expected[t, 0] = CODE_448
+        for j in range(per_row):
+            i = t * per_row + j
+            if i >= len(pairs):
+                break
+            code, value = pairs[i]
+            rows[t, 0, j + 1] = value
+            expected[t, j + 1] = code
+
+    k_cache, _, k_scale, _ = _store(
+        rows.to(DEV, dtype=torch.bfloat16),
+        torch.zeros(tokens, 1, dim, dtype=torch.bfloat16, device=DEV),
+    )
+    assert torch.equal(k_scale[:, 0], torch.ones_like(k_scale[:, 0]))
+    got = _canonical_zero(_as_bytes(k_cache)[:, 0, :])
+    want = _canonical_zero(expected.to(DEV))
+    bad = got != want
+    assert not bad.any().item(), (
+        f"{int(bad.sum())} of {want.numel()} grid values round-tripped wrong; "
+        f"first at {bad.nonzero()[0].tolist()}: expected "
+        f"{want[bad][0].item():#x} got {got[bad][0].item():#x}"
+    )
+
+
+def test_zero_row_stays_finite_and_exact():
+    k = torch.zeros(2, 2, 32, device=DEV, dtype=torch.bfloat16)
+    k_cache, _, k_scale, _ = _store(k, k.clone())
+    assert torch.isfinite(k_scale).all()
+    assert (k_scale > 0).all(), "an all-zero row must still store a usable scale"
+    assert (codes_to_f32(k_cache) == 0).all()
+
+
+def test_codes_are_plain_bytes_and_the_kernel_decode_matches_torch():
+    """The KV codec never puts an fp8 type in front of Triton.
+
+    Codes live in a uint8 buffer on EVERY architecture and the kernel widens them with
+    the software decoder, while the expectation below is torch's OWN e4m3 cast of those
+    very bytes. That pins the one claim the design rests on: byte for byte, the
+    software decode reads what a native fp8 unit would -- which is what lets the
+    quantized cache behave identically on GPUs where the fp8 type is illegal.
+    """
+    import triton
+    import triton.language as tl
+
+    from freetoken.kernel.triton.e4m3_compat import kv_load_e4m3_tile_f32
+
+    rows = torch.randn(5, 2, 64, device=DEV, dtype=torch.bfloat16) * 3.0
+    k_cache, _, _, _ = _store(rows, rows.clone())
+    assert kv_codes_dtype() is torch.uint8, "keep the fp8 type out of kernel signatures"
+    assert k_cache.dtype is torch.uint8 and k_cache.element_size() == 1
+    want = codes_to_f32(k_cache)  # torch reinterprets these bytes as e4m3 and casts
+
+    @triton.jit
+    def read_out(codes_ptr, out_ptr, n, BLOCK: tl.constexpr):
+        offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK)
+        got = kv_load_e4m3_tile_f32(codes_ptr + offs, offs < n)
+        tl.store(out_ptr + offs, got, mask=offs < n)
+
+    n = k_cache.numel()
+    out = torch.zeros(n, dtype=torch.float32, device=DEV)
+    read_out[(triton.cdiv(n, 256),)](k_cache, out, n, BLOCK=256)
+    flat = want.reshape(-1)
+    assert torch.equal(out, flat), (
+        f"{int((out != flat).sum())} of {n} codes decode differently from torch's cast"
+    )
+
+
+def test_kv_codec_has_no_arch_or_dtype_branch():
+    """The two rejected designs had one thing in common: they chose an arm.
+
+    The compile-time fp8-native probe answers a question the allocator already
+    answered -- and on one box answered wrongly -- while a test against the pointer's
+    element type is NOT pruned by triton, so the dead arm still gets type-checked.
+    That is how an int mask fill ended up in front of an fp8 pointer, twice. Both
+    codecs are straight-line now; pin that, plus the identifiers of the two rejected
+    designs, so neither creeps back in as a "fast path".
+    """
+    import inspect
+
+    from freetoken.kernel.triton import kv_quant
+    from freetoken.kernel.triton.e4m3_compat import kv_load_e4m3_tile_f32
+
+    for obj in (kv_load_e4m3_tile_f32, kv_quant._kv_quant_scatter_kernel):
+        src = inspect.getsource(getattr(obj, "fn", obj))
+        for banned in ("e4m3_native", "dtype.element_ty"):
+            assert banned not in src, f"{banned} is back in {obj.__name__}"
+        body = src.split('"""')[-1].splitlines()
+        arms = [
+            line.strip() for line in body
+            if line.strip().startswith(("if ", "elif ", "else"))
+        ]
+        assert not arms, f"{obj.__name__} must not branch: {arms}"
diff --git a/tests/kernels/test_kv_nvfp4.py b/tests/kernels/test_kv_nvfp4.py
new file mode 100644
index 000000000..662b44bf2
--- /dev/null
+++ b/tests/kernels/test_kv_nvfp4.py
@@ -0,0 +1,349 @@
+"""NVFP4 KV against independent torch rounding and dense attention references."""
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.distributed import set_tp_info, try_get_tp_info
+from freetoken.kvcache.mha_pool import MHAKVCache
+
+pytestmark = pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+
+
+def _pool(dim=128, heads=2, slots=96, layer_ids=None):
+    if try_get_tp_info() is None:
+        set_tp_info(rank=0, size=1)
+    return MHAKVCache(heads, 4, dim, slots // 4, 4, torch.bfloat16,
+                      torch.device("cuda"), layer_ids=layer_ids, kv_quant="nvfp4")
+
+
+def _reference(x):
+    shape = x.shape
+    x = x.float().reshape(*shape[:-1], -1, 16)
+    row = x.abs().flatten(-2).amax(-1).clamp_min(1e-10) / 2688.0
+    block = (x.abs().amax(-1) / (6 * row[..., None])).clamp_max(448).to(torch.float8_e4m3fn)
+    denom = block.float() * row[..., None]
+    normalized = torch.where(denom[..., None] > 0, x / denom[..., None], 0)
+    grid = torch.tensor([0, .5, 1, 1.5, 2, 3, 4, 6], device=x.device)
+    distance = (normalized.abs()[..., None] - grid).abs()
+    nearest = distance == distance.amin(-1, keepdim=True)
+    codes = torch.arange(8, device=x.device).expand_as(distance)
+    # Prefer even codes on ties, independently of the kernel's threshold encoding.
+    rank = torch.where(nearest, codes % 2 * 8 + codes, 32)
+    code = rank.argmin(-1) | ((normalized < 0).long() * 8)
+    code = code.reshape(shape)
+    packed = (code[..., ::2] | (code[..., 1::2] << 4)).to(torch.uint8)
+    return packed, block.view(torch.uint8), row
+
+
+def _decode(pool, which, layer=1):
+    codes = getattr(pool, f"{which}_cache")(layer).flatten(0, 1)
+    block = getattr(pool, f"{which}_block_scale")(layer).view(torch.float8_e4m3fn).float()
+    row = getattr(pool, f"{which}_scale")(layer)
+    code = torch.stack((codes & 15, codes >> 4), -1).flatten(-2).long()
+    grid = torch.tensor([0, .5, 1, 1.5, 2, 3, 4, 6,
+                         0, -.5, -1, -1.5, -2, -3, -4, -6], device=codes.device)
+    return grid[code] * block.repeat_interleave(16, -1) * row[..., None]
+
+
+def _decode_latent(pool, layer=0):
+    codes = pool.latent_rows(layer)
+    code = torch.stack((codes & 15, codes >> 4), -1).flatten(-2).long()
+    grid = torch.tensor([0, .5, 1, 1.5, 2, 3, 4, 6,
+                         0, -.5, -1, -1.5, -2, -3, -4, -6], device=codes.device)
+    block = pool.latent_block_scale(layer).view(torch.float8_e4m3fn).float()
+    return grid[code] * block.repeat_interleave(16, -1) * pool.latent_scale(layer)[:, None]
+
+
+@pytest.mark.parametrize("rope", [0, 64])
+@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16])
+def test_latent_scatter_matches_reference_and_preserves_prefix(rope, dtype):
+    from freetoken.kvcache.dsa_pool import MLAKVCache
+
+    torch.manual_seed(73)
+    dim = 512 + rope
+    pool = MLAKVCache(dim, 4, 2, 64, dtype, torch.device("cuda"),
+                      layer_ids=(1, 3), kv_quant="nvfp4")
+    x = torch.randn(7, dim + 32, device="cuda", dtype=dtype)[:, :dim]
+    x[0].zero_()
+    x[1, :16] *= 100
+    x[2] *= 1e-5
+    loc = torch.tensor([64, 3, 127, 14, 6, 90, 31], device="cuda")
+    pool.store_kv(x[:4, :512], x[:4, 512:], loc[:4], 3)
+    before = [v.clone() for v in (pool.latent_rows(3), pool.latent_scale(3),
+                                  pool.latent_block_scale(3))]
+    pool.store_kv(x[4:, :512], x[4:, 512:], loc[4:], 3)
+    pool.store_kv(x[:0, :512], x[:0, 512:], loc[:0], 3)
+    for got, saved in zip((pool.latent_rows(3), pool.latent_scale(3),
+                           pool.latent_block_scale(3)), before):
+        torch.testing.assert_close(got[loc[:4]], saved[loc[:4]], rtol=0, atol=0)
+    packed, block, row = _reference(x)
+    torch.testing.assert_close(pool.latent_rows(3)[loc], packed)
+    torch.testing.assert_close(pool.latent_block_scale(3)[loc], block)
+    torch.testing.assert_close(pool.latent_scale(3)[loc], row)
+    assert torch.count_nonzero(pool.latent_rows(1)) == 0
+    assert torch.isfinite(_decode_latent(pool, 3)).all()
+    assert pool.k_cache(3).data_ptr() == pool.v_cache(3).data_ptr()
+    # Reused physical slots replace codes and both scales together.
+    x[4:].mul_(0.125)
+    pool.store_kv(x[4:, :512], x[4:, 512:], loc[4:], 3)
+    packed, block, row = _reference(x[4:])
+    torch.testing.assert_close(pool.latent_rows(3)[loc[4:]], packed)
+    torch.testing.assert_close(pool.latent_block_scale(3)[loc[4:]], block)
+    torch.testing.assert_close(pool.latent_scale(3)[loc[4:]], row)
+
+
+@pytest.mark.parametrize("rope", [0, 64])
+@pytest.mark.parametrize("splits", [0, 4])
+@pytest.mark.parametrize("queries,broadcast", [(1, False), (5, False), (5, True)])
+@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16])
+def test_sparse_mla_nvfp4_matches_restored_reference(rope, splits, queries, broadcast, dtype):
+    from freetoken.kernel.triton.glm_dsa_sparse import glm_dsa_sparse_attn
+    from freetoken.kvcache.dsa_pool import MLAKVCache
+
+    torch.manual_seed(74)
+    dim, n = 512 + rope, 67
+    pool = MLAKVCache(dim, 1, 2, 64, torch.bfloat16, torch.device("cuda"), kv_quant="nvfp4")
+    x = torch.randn(n, dim, device="cuda", dtype=torch.bfloat16)
+    x *= torch.linspace(.2, 2, n, device="cuda")[:, None]
+    if rope:
+        x[:, 512:] *= 3
+    loc = torch.randperm(128, device="cuda")[:n]
+    pool.store_kv(x[:, :512], x[:, 512:], loc, 0)
+    q = torch.randn(2, queries, 19, dim, device="cuda", dtype=dtype)
+    sel = loc.repeat(2, 1 if broadcast else queries, 1).to(torch.int32)
+    sel[1] = sel[1].flip(-1)
+    sel[..., 5] = -1
+    cnt = torch.full((2, queries), n, device="cuda", dtype=torch.int32)
+    cnt[0, 0] = 0
+    cnt[1, 0] = 13
+    out = glm_dsa_sparse_attn(
+        q, pool.latent_rows(0), sel, .04, counts=cnt, d_v=512,
+        pool_scale=pool.latent_scale(0), pool_block_scale=pool.latent_block_scale(0),
+        kv_quant="nvfp4", force_splits=splits,
+    )
+    decoded = _decode_latent(pool)
+    ref = torch.zeros_like(out, dtype=torch.float32)
+    for b in range(2):
+        for m in range(queries):
+            rows = sel[b, 0 if broadcast else m, :int(cnt[b, m])]
+            rows = rows[rows >= 0].long()
+            if rows.numel():
+                kv = decoded[rows]
+                ref[b, m] = (q[b, m].float() @ kv.T * .04).softmax(-1) @ kv[:, :512]
+    torch.testing.assert_close(out.float(), ref, atol=2e-2, rtol=1e-2)
+
+
+@pytest.mark.parametrize("dim", [16, 48, 64, 128, 256, 512])
+@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16])
+def test_scatter_matches_independent_reference(dim, dtype):
+    torch.manual_seed(41)
+    pool = _pool(dim)
+    # K is a projection slice; V is independently contiguous.
+    k = torch.randn(5, 4 * dim, device="cuda", dtype=dtype)[:, :2 * dim]
+    v = torch.randn(5, 2 * dim, device="cuda", dtype=dtype)
+    k[0].zero_()
+    k[1, :16] *= 100
+    v[2] *= 1e-5
+    loc = torch.tensor([7, 1, 95, 31, 12], device="cuda", dtype=torch.int64)
+    pool.store_kv(k, v, loc, 1)
+    for which, source in (("k", k), ("v", v)):
+        packed, block, row = _reference(source.reshape(5, 2, dim))
+        torch.testing.assert_close(getattr(pool, f"{which}_cache")(1).flatten(0, 1)[loc], packed)
+        torch.testing.assert_close(getattr(pool, f"{which}_block_scale")(1)[loc], block)
+        torch.testing.assert_close(getattr(pool, f"{which}_scale")(1)[loc], row)
+        assert torch.isfinite(_decode(pool, which)).all()
+        assert torch.count_nonzero(_decode(pool, which)[0]) == 0
+
+
+def test_e2m1_grid_and_round_to_even_boundaries():
+    positive = torch.tensor([0, .25, .5, .75, 1, 1.25, 1.5, 1.75,
+                             2, 2.5, 3, 3.5, 4, 5, 6], device="cuda")
+    values = torch.cat((positive, -positive))
+    values = torch.cat((values, values.nextafter(torch.full_like(values, float("inf"))),
+                        values.nextafter(torch.full_like(values, -float("inf")))))
+    pool = _pool(dim=32, slots=192)
+    rows = torch.zeros(values.numel(), 2, 32, device="cuda")
+    rows[:, :, 0] = values[:, None]
+    rows[:, :, 15] = 6
+    rows[:, :, 31] = 2688  # Forces row_scale=1 and the first block_scale=1.
+    loc = torch.arange(values.numel(), device="cuda", dtype=torch.int32)
+    pool.store_kv(rows.flatten(1), rows.flatten(1), loc, 1)
+    packed, block, row = _reference(rows)
+    torch.testing.assert_close(pool.k_cache(1).flatten(0, 1)[loc], packed)
+    torch.testing.assert_close(pool.k_block_scale(1)[loc], block)
+    torch.testing.assert_close(pool.k_scale(1)[loc], row)
+
+
+@pytest.mark.parametrize("dim", [64, 128, 256, 512])
+@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16])
+@pytest.mark.parametrize("mode", ["paged", "decode", "extend", "split"])
+def test_attention_reads_packed_cache(dim, dtype, mode):
+    from freetoken.kernel.triton.attention import (
+        decode_paged_attention, extend_paged_attention, paged_attention,
+    )
+
+    torch.manual_seed(42)
+    pool = _pool(dim)
+    n, prefix = 67, 35
+    k, v = [torch.randn(n, 2 * dim, device="cuda", dtype=dtype) for _ in range(2)]
+    # Distinct token/head and block scales expose accidental FP8 rescaling of NVFP4.
+    row_gain = torch.linspace(0.25, 2.0, n * 2, device="cuda").view(n, 2, 1)
+    block_gain = torch.linspace(0.5, 1.5, dim // 16, device="cuda").repeat_interleave(16)
+    k.view(n, 2, dim).mul_(row_gain * block_gain)
+    v.view(n, 2, dim).mul_(row_gain.flip(0) * block_gain.flip(0))
+    loc = torch.randperm(95, device="cuda")[:n] + 1
+    pool.store_kv(k, v, loc, 1)
+    tokens = 1 if mode in ("paged", "decode") else n - prefix
+    q = torch.randn(tokens, 6, dim, device="cuda", dtype=dtype)
+    indptr = torch.tensor([0, n], device="cuda", dtype=torch.int32)
+    pos = torch.arange(n - tokens, n, device="cuda")
+    args = dict(q=q, k_cache=pool.k_cache(1).flatten(0, 1),
+                v_cache=pool.v_cache(1).flatten(0, 1), k_scale=pool.k_scale(1),
+                v_scale=pool.v_scale(1), k_block_scale=pool.k_block_scale(1),
+                v_block_scale=pool.v_block_scale(1), kv_quant="nvfp4", sm_scale=dim ** -.5)
+    kd, vd = _decode(pool, "k")[loc], _decode(pool, "v")[loc]
+    if mode == "split":
+        kd[prefix:] = k[prefix:].view(-1, 2, dim).float()
+        vd[prefix:] = v[prefix:].view(-1, 2, dim).float()
+    # Tensor-core paths round restored K/V to the query dtype before dot products.
+    if mode != "paged":
+        kd, vd = kd.to(q.dtype).float(), vd.to(q.dtype).float()
+    kd, vd = [x.repeat_interleave(3, 1).transpose(0, 1) for x in (kd, vd)]
+    score = torch.einsum("thd,hnd->thn", q.float(), kd) * dim ** -.5
+    score.masked_fill_(torch.arange(n, device="cuda")[None, None, :] > pos[:, None, None], -float("inf"))
+    ref = torch.einsum("thn,hnd->thd", score.softmax(-1), vd).to(q.dtype)
+    if mode == "paged":
+        actual = paged_attention(**args, indptr=indptr, indices=loc,
+                                 q_to_req=torch.zeros(tokens, device="cuda", dtype=torch.int32), q_positions=pos)
+    elif mode == "decode":
+        actual = decode_paged_attention(**args, indptr=indptr, indices=loc, q_positions=pos,
+            attn_logits=torch.empty(1, 6, 8, dim, device="cuda"),
+            attn_lse=torch.empty(1, 6, 8, device="cuda"),
+            num_kv_splits=torch.tensor([8], device="cuda", dtype=torch.int32), max_kv_splits=8)
+    else:
+        extra = {} if mode == "extend" else dict(k_extend=k[prefix:].view(-1, 2, dim), v_extend=v[prefix:].view(-1, 2, dim))
+        actual = extend_paged_attention(**args, **extra,
+            qo_indptr=torch.tensor([0, tokens], device="cuda", dtype=torch.int32),
+            kv_indptr=indptr, kv_indices=loc,
+            prefix_lens=torch.tensor([prefix], device="cuda", dtype=torch.int32), max_q_len=tokens)
+    torch.testing.assert_close(actual, ref, atol=0.008, rtol=0.025)
+
+
+def test_pool_budget_rebuild_and_layer_mapping():
+    from freetoken.kvcache.base import spec_kv_bytes_per_token
+    from freetoken.models.config import KVCacheGroupSpec
+
+    pool = _pool(layer_ids=(1, 3))
+    spec = KVCacheGroupSpec(name="full", layer_ids=(1, 3), num_kv_heads=2, head_dim=128, sliding_window=None)
+    cfg = SimpleNamespace(kv_quant="nvfp4", dtype=torch.bfloat16, tp_info=SimpleNamespace(size=1))
+    assert spec_kv_bytes_per_token(spec, cfg) == pool.unit_bytes()[0] == 2 * 2 * 2 * 76
+    with pytest.raises(KeyError):
+        pool.k_block_scale(0)
+    for pages in (32, 8):
+        pool.rebuild(pages)
+        assert pool.k_cache(3).shape == (pages, 4, 2, 64)
+        assert pool.k_block_scale(3).shape == (pages * 4, 2, 8)
+        assert pool.unit_bytes()[0] == 608
+        assert torch.count_nonzero(_decode(pool, "k", 3)) == 0
+
+
+def test_store_cuda_graph_replay_changes_slots():
+    pool = _pool()
+    k = torch.randn(2, 256, device="cuda", dtype=torch.bfloat16)
+    loc = torch.tensor([1, 2], device="cuda", dtype=torch.int32)
+    pool.store_kv(k, k, loc, 1)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        pool.store_kv(k, k, loc, 1)
+    k.mul_(2)
+    loc.copy_(torch.tensor([4, 9], device="cuda", dtype=torch.int32))
+    graph.replay()
+    expected, _, _ = _reference(k.view(2, 2, 128))
+    torch.testing.assert_close(pool.k_cache(1).flatten(0, 1)[loc], expected)
+
+
+def test_hybrid_swa_pool_packs_both_groups_and_rebuilds(monkeypatch):
+    from freetoken.distributed.info import DistributedInfo
+    from freetoken.kvcache.hybrid_swa_pool import HybridSWAKVCache
+    from freetoken.models.config import KVCacheGroupSpec
+
+    monkeypatch.setattr(
+        "freetoken.kvcache.hybrid_swa_pool.get_tp_info",
+        lambda: DistributedInfo(rank=0, size=1),
+    )
+    groups = (
+        KVCacheGroupSpec("full", (1,), 2, 64, None),
+        KVCacheGroupSpec("swa", (0,), 2, 128, 32),
+    )
+    pool = HybridSWAKVCache(
+        groups=groups,
+        num_layers=2,
+        num_full_pages=12,
+        page_size=1,
+        num_swa_tokens=18,
+        dtype=torch.bfloat16,
+        device=torch.device("cuda"),
+        kv_quant="nvfp4",
+    )
+    assert pool.k_cache(1).shape == (12, 1, 2, 32)
+    assert pool.k_cache(0).shape == (18, 1, 2, 64)
+    assert pool.k_block_scale(1).shape == (12, 2, 4)
+    assert pool.k_block_scale(0).shape == (18, 2, 8)
+    assert pool.unit_bytes() == (160, 304)
+
+    full_loc = torch.tensor([4, 9], device="cuda", dtype=torch.int32)
+    swa_loc = torch.tensor([2, 7], device="cuda", dtype=torch.int32)
+    pool.alloc_swa(swa_loc)
+    for layer, dim, loc, cache_loc in ((1, 64, full_loc, full_loc), (0, 128, swa_loc, None)):
+        rows = torch.randn(2, 2 * dim, device="cuda", dtype=torch.bfloat16)
+        pool.store_kv(rows, rows, loc, layer)
+        if cache_loc is None:
+            cache_loc = pool.translate_loc_from_full_to_swa(loc)
+        expected, block, row = _reference(rows.view(2, 2, dim))
+        codes = pool.k_cache(layer).flatten(0, 1)[cache_loc]
+        torch.testing.assert_close(codes, expected)
+        torch.testing.assert_close(pool.k_block_scale(layer)[cache_loc], block)
+        torch.testing.assert_close(pool.k_scale(layer)[cache_loc], row)
+
+    pool.rebuild(num_full_pages=6, num_swa_tokens=10)
+    assert pool.k_cache(1).shape == (6, 1, 2, 32)
+    assert pool.k_cache(0).shape == (10, 1, 2, 64)
+    assert pool.unit_bytes() == (160, 304)
+    assert pool.swa_available_size() == 9
+
+
+def test_backend_decode_graph_replays_new_kv_and_page_tables(monkeypatch):
+    from freetoken.attention.triton import TritonAttentionBackend, TritonMetadata
+    from freetoken.kernel.triton.attention import paged_attention
+
+    pool = _pool()
+    monkeypatch.setattr("freetoken.attention.triton.get_global_ctx",
+                        lambda: SimpleNamespace(kv_cache=pool))
+    backend = TritonAttentionBackend(SimpleNamespace(num_qo_heads=6, head_dim=128))
+    q = torch.randn(2, 6, 128, device="cuda", dtype=torch.bfloat16)
+    k, v = [torch.randn(2, 256, device="cuda", dtype=q.dtype) for _ in range(2)]
+    loc = torch.tensor([7, 11], device="cuda", dtype=torch.int32)
+    indptr = torch.tensor([0, 1, 2], device="cuda", dtype=torch.int32)
+    positions = torch.zeros(2, device="cuda", dtype=torch.int32)
+    indices = loc.clone()
+    req = torch.arange(2, device="cuda", dtype=torch.int32)
+    metadata = TritonMetadata(cu_seqlens_q_gpu=indptr, indptr=indptr, indices=indices,
+        q_to_req=req, q_positions=positions, is_decode=True, prefix_lens=positions,
+        max_q_len=1)
+    batch = SimpleNamespace(out_loc=loc, attn_metadata=metadata)
+    backend.forward(q, k, v, 1, batch)
+    graph = torch.cuda.CUDAGraph()
+    with torch.cuda.graph(graph):
+        actual = backend.forward(q, k, v, 1, batch)
+    for new_slots in ([31, 44], [7, 11]):
+        k.normal_()
+        v.normal_()
+        q.normal_()
+        loc.copy_(torch.tensor(new_slots, device="cuda", dtype=torch.int32))
+        indices.copy_(loc)
+        graph.replay()
+        ref = paged_attention(q, _decode(pool, "k").to(q.dtype),
+            _decode(pool, "v").to(q.dtype), indptr, indices, req, positions, 128 ** -.5)
+        torch.testing.assert_close(actual, ref, atol=0.008, rtol=0.025)
diff --git a/tests/kernels/test_mrope.py b/tests/kernels/test_mrope.py
index f8402f875..9868cbf18 100644
--- a/tests/kernels/test_mrope.py
+++ b/tests/kernels/test_mrope.py
@@ -18,6 +18,52 @@
     "interleaved_glm": [0, 1, 2] * 10 + [0, 1],
 }
 
+YARN = (("rope_type", "yarn"), ("factor", 4.0), ("original_max_position_embeddings", 262144))
+
+
+def _hf_yarn_parameters(device):
+    from transformers.modeling_rope_utils import ROPE_INIT_FUNCTIONS
+
+    class Config:
+        head_dim = hidden_size = HEAD
+        num_attention_heads = 1
+        max_position_embeddings = 1048576
+        rope_parameters = {"rope_theta": 1e7, "partial_rotary_factor": .25, **dict(YARN)}
+
+        def standardize_rope_params(self):
+            pass
+
+    return ROPE_INIT_FUNCTIONS["yarn"](Config(), device=torch.device(device))
+
+
+def test_yarn_mrope_keeps_scaled_main_and_index_frequencies():
+    from freetoken.layers.rotary import MRotaryEmbedding, get_rope
+
+    get_rope.cache_clear()
+    kwargs = dict(rotary_dim=ROT, max_position=257, base=1e7, rope_scaling=YARN)
+    plain = get_rope(head_dim=HEAD, **kwargs)
+    main = get_rope(head_dim=HEAD, mrope_section=SECTION, **kwargs)
+    index = get_rope(head_dim=128, mrope_section=SECTION, **kwargs)
+    assert isinstance(main, MRotaryEmbedding) and isinstance(index, MRotaryEmbedding)
+    inv_freq, amplitude = _hf_yarn_parameters("cpu")
+    angles = torch.outer(torch.arange(257, dtype=torch.float32), inv_freq)
+    expected = torch.cat((angles.cos(), angles.sin()), dim=-1) * amplitude
+    for rope in (plain, main, index):
+        torch.testing.assert_close(rope._cos_sin_cache, expected, rtol=0, atol=1e-6)
+    torch.testing.assert_close(main._cos_sin_cache, plain._cos_sin_cache, rtol=0, atol=0)
+    torch.testing.assert_close(index._cos_sin_cache, plain._cos_sin_cache, rtol=0, atol=0)
+    assert main._section_table.tolist() == REFERENCE_TABLES["interleaved"]
+    assert amplitude > 1.0
+    get_rope.cache_clear()
+
+
+def test_mrope_rejects_partial_proportional_cache_layout():
+    from freetoken.layers.rotary import get_rope
+
+    with pytest.raises(ValueError, match="partial proportional"):
+        get_rope(head_dim=HEAD, rotary_dim=ROT, max_position=8, base=1e7,
+                 rope_scaling=(("rope_type", "proportional"),), mrope_section=SECTION)
+
 
 def test_section_tables():
     from freetoken.layers.rotary import build_section_table
@@ -112,3 +158,34 @@ def test_kernel_matches_torch_fallback():
     apply_mrope_torch_fallback(pos3, q2, k2, HEAD, cache, sec)
     assert torch.allclose(q.float(), q2.float(), atol=1e-2, rtol=1e-2)
     assert torch.allclose(k.float(), k2.float(), atol=1e-2, rtol=1e-2)
+
+
+@cuda
+@pytest.mark.parametrize("head_dim", [128, HEAD], ids=["index", "attention"])
+def test_yarn_mrope_cuda_matches_hf_frequencies_and_preserves_partial_tail(head_dim):
+    from freetoken.layers.rotary import get_rope
+
+    get_rope.cache_clear()
+    with torch.device("cuda"):
+        rope = get_rope(head_dim=head_dim, rotary_dim=ROT, max_position=4096, base=1e7,
+                        rope_scaling=YARN, mrope_section=SECTION)
+    gen = torch.Generator(device="cuda").manual_seed(29)
+    positions = torch.randint(0, 4096, (3, 37), device="cuda", dtype=torch.int32, generator=gen)
+    q = torch.randn(37, 3 * head_dim, device="cuda", dtype=torch.bfloat16, generator=gen)
+    k = torch.randn(37, head_dim, device="cuda", dtype=torch.bfloat16, generator=gen)
+    expected = [q.clone(), k.clone()]
+    inv_freq, amplitude = _hf_yarn_parameters("cuda")
+    axes = torch.tensor(REFERENCE_TABLES["interleaved"], device="cuda")
+    angles = positions[axes].T.float() * inv_freq
+    cos, sin = (angles.cos() * amplitude).unsqueeze(1), (angles.sin() * amplitude).unsqueeze(1)
+    for tensor in expected:
+        value = tensor.view(37, -1, head_dim)
+        lo, hi = value[..., :ROT // 2].float(), value[..., ROT // 2:ROT].float()
+        value[..., :ROT // 2] = (lo * cos - hi * sin).to(value.dtype)
+        value[..., ROT // 2:ROT] = (hi * cos + lo * sin).to(value.dtype)
+    rope.forward(positions, q, k)
+    for actual, reference in zip((q, k), expected):
+        torch.testing.assert_close(actual, reference, rtol=.02, atol=.02)
+        torch.testing.assert_close(actual.view(37, -1, head_dim)[..., ROT:],
+                                    reference.view(37, -1, head_dim)[..., ROT:], rtol=0, atol=0)
+    get_rope.cache_clear()
diff --git a/tests/kernels/test_qsa_fp8.py b/tests/kernels/test_qsa_fp8.py
new file mode 100644
index 000000000..1fa6669f8
--- /dev/null
+++ b/tests/kernels/test_qsa_fp8.py
@@ -0,0 +1,278 @@
+"""QSA sparse attention reading an fp8 KV pool (kernel/triton/qsa/attend.py).
+
+The oracle is the SAME data in a bf16 cache, not a tolerance. The kernel widens the e4m3
+codes to fp32, multiplies by the row scale, and casts back to the query dtype before
+``tl.dot`` -- so a cache holding ``c * s`` and a cache holding ``c`` with scale ``s`` feed
+the matmul bit-identical operands, and the two runs must agree bit for bit. Anything wrong
+in scale indexing (the page/offset -> slot arithmetic), in the K-vs-V broadcast direction,
+or in the masked-fill values shows up as a mismatch instead of slop inside an epsilon.
+
+Two quantizer variants are covered:
+  * hand-made codes with power-of-two scales, which also keeps the real values on the e4m3
+    grid, so the fp8 buffer, the bf16 buffer and the test agree on the numbers;
+  * the fused store kernel the pools actually call (``kv_quant.quantize_kv_to_cache``,
+    amax/448 scales) over ordinary gaussian rows -- whose quantization error against the
+    ORIGINAL rows is bounded separately, because that part is quality, not exactness.
+
+Each parametrization also covers a different launcher profile: a small
+``rows * kv_heads`` pushes it into split-K (partials + merge kernel), a large one takes
+the direct-write path.
+"""
+
+from __future__ import annotations
+
+import pytest
+import torch
+
+from freetoken.kernel.triton.kv_quant import (
+    FP8,
+    alloc_codes,
+    codes_to_f32,
+    kv_codes_dtype,
+    quantize_kv_to_cache,
+)
+from freetoken.kernel.triton.qsa import qsa_sparse_paged_attention
+
+pytestmark = pytest.mark.skipif(
+    not torch.cuda.is_available(), reason="Triton attention needs CUDA"
+)
+
+PAGE = 64  # the page size the qsa_sparse backend registers
+HEAD_DIM = 64
+
+
+def _on_grid(shape, device, generator) -> torch.Tensor:
+    """Values of the form +-m * 2^e with m in [8, 16): four significant bits, i.e. every
+    one of them is representable in e4m3 AND in bf16, so encoding is lossless and both
+    caches hold the very same real numbers."""
+    mantissa = torch.randint(
+        8, 16, shape, device=device, generator=generator, dtype=torch.int32
+    )
+    exponent = torch.randint(
+        -6, 5, shape, device=device, generator=generator, dtype=torch.int32
+    )
+    sign = torch.where(
+        torch.randint(0, 2, shape, device=device, generator=generator).bool(), 1.0, -1.0
+    )
+    return sign * mantissa.to(torch.float32) * torch.pow(2.0, exponent.to(torch.float32))
+
+
+def _code_buffer(values_f32: torch.Tensor) -> torch.Tensor:
+    """Encode e4m3-exact fp32 values into a buffer of the pool's code dtype."""
+    codes = values_f32.to(FP8)
+    if kv_codes_dtype() is torch.uint8:
+        codes = codes.view(torch.uint8)
+    assert codes.dtype == kv_codes_dtype(), codes.dtype
+    return codes.contiguous()
+
+
+def _layout(rows: int, topk: int, num_req: int):
+    """(indices, block_table, token_to_req) over three 64-token pages.
+
+    Tokens 32..31+topk straddle pages 0 and 1, and the two requests map their logical
+    pages to physical ones in OPPOSITE order -- a page-table or page/offset slip cannot
+    hide behind a symmetric fixture.
+    """
+    device = "cuda"
+    block_table = torch.tensor(
+        [[0, 1, 2], [2, 1, 0]], dtype=torch.int32, device=device
+    )[:num_req].contiguous()
+    indices = (
+        torch.arange(topk, dtype=torch.int32, device=device)[None, :] + 32
+    ).repeat(rows, 1).contiguous()
+    token_to_req = (
+        torch.arange(rows, dtype=torch.int32, device=device) % num_req
+    ).contiguous()
+    return indices, block_table, token_to_req
+
+
+def _run(codes, scales, q, indices, block_table, token_to_req):
+    return qsa_sparse_paged_attention(
+        q,
+        codes[0],
+        codes[1],
+        indices,
+        block_table,
+        token_to_req,
+        k_scale=scales[0],
+        v_scale=scales[1],
+    )
+
+
+
+# rows 1 x 1 kv head -> base_programs 1  -> BLOCK_N 16, 4 tiles -> NUM_SPLITS 4 (split-K).
+# rows 16 x 2        -> base_programs 32 -> BLOCK_N 64, 1 tile  -> NUM_SPLITS 1 (direct).
+@pytest.mark.parametrize(("rows", "kv_heads", "topk"), [(1, 1, 64), (16, 2, 64)])
+def test_fp8_codes_match_the_bf16_cache_bit_for_bit(rows, kv_heads, topk):
+    torch.manual_seed(3)
+    device = torch.device("cuda")
+    num_pages, num_query_heads = 3, 2 * kv_heads
+    slots = num_pages * PAGE
+
+    k_values = _on_grid((num_pages, PAGE, kv_heads, HEAD_DIM), device, None)
+    v_values = _on_grid((num_pages, PAGE, kv_heads, HEAD_DIM), device, None)
+    # Alternate the scale exponent by slot -- a scale broadcast along the wrong axis then
+    # changes the answer instead of cancelling out -- and give K and V opposite parities.
+    parity = torch.arange(slots, device=device, dtype=torch.float32) % 2
+    k_scale = (
+        torch.pow(2.0, parity * 5 - 3).unsqueeze(-1).expand(slots, kv_heads).contiguous()
+    )
+    v_scale = (
+        torch.pow(2.0, (1 - parity) * 4 - 2).unsqueeze(-1).expand(slots, kv_heads).contiguous()
+    )
+    k_ref = (k_values * k_scale.view(num_pages, PAGE, kv_heads, 1)).to(torch.bfloat16)
+    v_ref = (v_values * v_scale.view(num_pages, PAGE, kv_heads, 1)).to(torch.bfloat16)
+
+    q = torch.randn(rows, num_query_heads, HEAD_DIM, device=device, dtype=torch.bfloat16)
+    indices, block_table, token_to_req = _layout(rows, topk, num_req=2)
+
+    got = _run(
+        (_code_buffer(k_values), _code_buffer(v_values)),
+        (k_scale, v_scale),
+        q,
+        indices,
+        block_table,
+        token_to_req,
+    )
+    want = qsa_sparse_paged_attention(
+        q, k_ref, v_ref, indices, block_table, token_to_req
+    )
+    assert torch.equal(got, want), (
+        "fp8 QSA attend diverged from the same data in a bf16 cache (max diff "
+        f"{(got.float() - want.float()).abs().max().item():.3e})"
+    )
+
+
+@pytest.mark.parametrize(("rows", "kv_heads", "topk"), [(1, 1, 64), (16, 2, 64)])
+def test_pool_writer_codes_match_the_bf16_cache_bit_for_bit(rows, kv_heads, topk):
+    """Ordinary rows through the fused quantize+scatter writer store_kv calls: the
+    codes/scales it produces, decoded on the host with torch's own e4m3 cast (never the
+    decoder under test), must feed the dot identical operands."""
+    torch.manual_seed(5)
+    device = torch.device("cuda")
+    num_pages, num_query_heads = 3, 2 * kv_heads
+    slots = num_pages * PAGE
+
+    k_rows = torch.randn(slots, kv_heads * HEAD_DIM, device=device, dtype=torch.float32)
+    v_rows = torch.randn(slots, kv_heads * HEAD_DIM, device=device, dtype=torch.float32)
+    # Wide amplitude spread: every 8th row carries a 64x outlier. That is what a per-row
+    # amax scale exists to absorb, and what a scale read from the wrong slot blows up on.
+    outlier = (torch.arange(slots, device=device) % 8 == 0).unsqueeze(-1)
+    k_rows = k_rows * torch.where(outlier, 64.0, 1.0)
+    v_rows = v_rows * torch.where(outlier.flip(0), 32.0, 0.5)
+
+    k_flat = alloc_codes((slots, kv_heads, HEAD_DIM), device)
+    v_flat = alloc_codes((slots, kv_heads, HEAD_DIM), device)
+    k_scale = torch.zeros((slots, kv_heads), dtype=torch.float32, device=device)
+    v_scale = torch.zeros_like(k_scale)
+    quantize_kv_to_cache(
+        k=k_rows,
+        v=v_rows,
+        out_loc=torch.arange(slots, dtype=torch.int32, device=device),
+        k_cache=k_flat,
+        v_cache=v_flat,
+        k_scale=k_scale,
+        v_scale=v_scale,
+    )
+    torch.cuda.synchronize()
+
+    k_ref = (codes_to_f32(k_flat) * k_scale.unsqueeze(-1)).to(torch.bfloat16)
+    v_ref = (codes_to_f32(v_flat) * v_scale.unsqueeze(-1)).to(torch.bfloat16)
+    q = torch.randn(rows, num_query_heads, HEAD_DIM, device=device, dtype=torch.bfloat16)
+    indices, block_table, token_to_req = _layout(rows, topk, num_req=2)
+
+    shape = (num_pages, PAGE, kv_heads, HEAD_DIM)
+    got = _run(
+        (k_flat.view(shape), v_flat.view(shape)),
+        (k_scale, v_scale),
+        q,
+        indices,
+        block_table,
+        token_to_req,
+    )
+    want = qsa_sparse_paged_attention(
+        q, k_ref.view(shape), v_ref.view(shape), indices, block_table, token_to_req
+    )
+    assert torch.equal(got, want), (
+        "fused-writer codes diverged (max diff "
+        f"{(got.float() - want.float()).abs().max().item():.3e})"
+    )
+
+    # Quantization QUALITY, bounded per row: e4m3's half-ulp is 2^-4 of the binade a
+    # value lands in, and the writer scales each row so its amax sits at 448.
+    for source, codes, scales in ((k_rows, k_flat, k_scale), (v_rows, v_flat, v_scale)):
+        assert torch.isfinite(scales).all() and (scales > 0).all()
+        decoded = codes_to_f32(codes) * scales.unsqueeze(-1)
+        original = source.view(slots, kv_heads, HEAD_DIM)
+        amax = original.abs().amax(dim=-1, keepdim=True)
+        rel = ((decoded - original).abs() / amax.clamp_min(1e-9)).max()
+        assert rel.item() <= 0.07, f"e4m3 grid error {rel.item():.4f} above its half-ulp"
+
+
+@pytest.mark.parametrize("which", ["k_only", "v_only", "bf16_cache", "wrong_shape"])
+def test_scale_arguments_are_validated(which):
+    """A half-supplied or mismatched scale pair must fail loudly: attending over raw
+    codes while believing they are bf16 is the exact failure mode this feature cannot be
+    allowed to have, and it produces plausible garbage rather than an error."""
+    device = "cuda"
+    kv_heads, rows = 1, 1
+    k = torch.zeros(2, PAGE, kv_heads, HEAD_DIM, device=device, dtype=kv_codes_dtype())
+    v = torch.zeros_like(k)
+    scale = torch.ones(2 * PAGE, kv_heads, dtype=torch.float32, device=device)
+    q = torch.zeros(rows, 2 * kv_heads, HEAD_DIM, device=device, dtype=torch.bfloat16)
+    indices = torch.zeros(rows, 8, dtype=torch.int32, device=device)
+    block_table = torch.zeros(1, 2, dtype=torch.int32, device=device)
+    token_to_req = torch.zeros(rows, dtype=torch.int32, device=device)
+
+    kwargs = {"k_scale": scale, "v_scale": scale}
+    if which == "k_only":
+        kwargs.pop("v_scale")
+    elif which == "v_only":
+        kwargs.pop("k_scale")
+    elif which == "bf16_cache":
+        k, v = k.to(torch.bfloat16), v.to(torch.bfloat16)
+    else:
+        kwargs["k_scale"] = scale[:-1]
+
+    with pytest.raises(ValueError, match="QSA"):
+        qsa_sparse_paged_attention(
+            q, k, v, indices, block_table, token_to_req, **kwargs
+        )
+
+
+def test_qsa_scoring_refuses_fp8_operands():
+    """The indexer's scoring dot is 16-bit only, and the wrapper has to say so.
+
+    In the field an e4m3 ``q_index`` -- produced by a backend that took its scratch dtype
+    from a pool reporting its STORE dtype -- surfaced as ``CompilationError: Unsupported
+    rhs dtype fp8e4nv`` inside CUDA-graph capture: a dead scheduler and a stopped API
+    server, forty seconds after the pool allocated fine. Same mistake, now stopped at
+    the call with a name on it.
+    """
+    from freetoken.kernel.triton.qsa import qsa_mqa_paged
+
+    device = "cuda"
+    rows, heads, dim, pages, cmp_page = 2, 2, 64, 2, 16
+    good = {
+        "q": torch.zeros(rows, heads, dim, device=device, dtype=torch.bfloat16),
+        "k_cache": torch.zeros(
+            pages, cmp_page, 1, dim, device=device, dtype=torch.bfloat16
+        ),
+        "page_table": torch.zeros(1, pages, dtype=torch.int32, device=device),
+        "token_to_req": torch.zeros(rows, dtype=torch.int32, device=device),
+        "query_positions": torch.arange(rows, dtype=torch.int32, device=device),
+        "sequence_lengths": torch.zeros(1, dtype=torch.int32, device=device),
+        "compress_ratio": cmp_page,
+        # Zero columns -> the wrapper returns before launching. This control proves the
+        # new check still lets the dtype it exists to allow through.
+        "logits": torch.zeros(rows, 0, dtype=torch.float32, device=device),
+        "visible_blocks": torch.zeros(rows, dtype=torch.int32, device=device),
+    }
+    qsa_mqa_paged(**good)
+
+    for name in ("q", "k_cache"):
+        bad = dict(good)
+        bad[name] = alloc_codes(tuple(good[name].shape), device)
+        with pytest.raises(ValueError, match="16-bit only"):
+            qsa_mqa_paged(**bad)
+
diff --git a/tests/kernels/test_qsa_nvfp4.py b/tests/kernels/test_qsa_nvfp4.py
new file mode 100644
index 000000000..591c6f1d4
--- /dev/null
+++ b/tests/kernels/test_qsa_nvfp4.py
@@ -0,0 +1,141 @@
+"""QSA sparse attention over the packed NVFP4 KV tier."""
+
+from __future__ import annotations
+
+import pytest
+import torch
+
+from freetoken.kernel.triton.kv_nvfp4 import quantize_nvfp4_to_cache
+from freetoken.kernel.triton.qsa import qsa_sparse_paged_attention
+
+pytestmark = pytest.mark.skipif(
+    not torch.cuda.is_available(), reason="Triton attention needs CUDA"
+)
+
+PAGE = 64
+HEAD_DIM = 64
+
+
+def _layout(rows: int, topk: int):
+    block_table = torch.tensor(
+        [[0, 1, 2], [2, 1, 0]], dtype=torch.int32, device="cuda"
+    ).contiguous()
+    indices = (
+        torch.arange(topk, dtype=torch.int32, device="cuda")[None, :] + 32
+    ).repeat(rows, 1).contiguous()
+    token_to_req = torch.arange(rows, dtype=torch.int32, device="cuda") % 2
+    return indices, block_table, token_to_req.contiguous()
+
+
+def _decode(codes: torch.Tensor, row_scale: torch.Tensor, block_scale: torch.Tensor):
+    """Decode with torch, independently of the Triton read path under test."""
+    lo = (codes & 15).to(torch.long)
+    hi = (codes >> 4).to(torch.long)
+    code = torch.stack((lo, hi), dim=-1).reshape(*codes.shape[:-1], HEAD_DIM)
+    magnitude = torch.tensor(
+        [0.0, 0.5, 1.0, 1.5, 2.0, 3.0, 4.0, 6.0], device=codes.device
+    )[code & 7]
+    value = torch.where(code < 8, magnitude, -magnitude)
+    block = block_scale.view(torch.float8_e4m3fn).to(torch.float32)
+    return (value * block.repeat_interleave(16, dim=-1) * row_scale.unsqueeze(-1)).to(
+        torch.bfloat16
+    )
+
+
+@pytest.mark.parametrize(("rows", "kv_heads"), [(1, 1), (16, 2)])
+def test_qsa_nvfp4_matches_its_bf16_decode(rows, kv_heads):
+    """Page-table indirection must address the same packed slot and both scale tiers."""
+    torch.manual_seed(19)
+    pages, query_heads, slots = 3, 2 * kv_heads, 3 * PAGE
+    k = torch.randn(slots, kv_heads * HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    v = torch.randn_like(k) * 0.25
+    out_loc = torch.arange(slots, dtype=torch.int32, device="cuda")
+    code_shape = (slots, kv_heads, HEAD_DIM // 2)
+    block_shape = (slots, kv_heads, HEAD_DIM // 16)
+    k_codes = torch.zeros(code_shape, dtype=torch.uint8, device="cuda")
+    v_codes = torch.zeros_like(k_codes)
+    k_scale = torch.zeros((slots, kv_heads), dtype=torch.float32, device="cuda")
+    v_scale = torch.zeros_like(k_scale)
+    k_block = torch.zeros(block_shape, dtype=torch.uint8, device="cuda")
+    v_block = torch.zeros_like(k_block)
+    quantize_nvfp4_to_cache(
+        k, v, out_loc, k_codes, v_codes, k_scale, v_scale, k_block, v_block
+    )
+    torch.cuda.synchronize()
+
+    q = torch.randn(rows, query_heads, HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    indices, block_table, token_to_req = _layout(rows, topk=64)
+    packed_shape = (pages, PAGE, kv_heads, HEAD_DIM // 2)
+    dense_shape = (pages, PAGE, kv_heads, HEAD_DIM)
+    got = qsa_sparse_paged_attention(
+        q,
+        k_codes.view(packed_shape),
+        v_codes.view(packed_shape),
+        indices,
+        block_table,
+        token_to_req,
+        k_scale=k_scale,
+        v_scale=v_scale,
+        kv_quant="nvfp4",
+        k_block_scale=k_block,
+        v_block_scale=v_block,
+    )
+    want = qsa_sparse_paged_attention(
+        q,
+        _decode(k_codes, k_scale, k_block).view(dense_shape),
+        _decode(v_codes, v_scale, v_block).view(dense_shape),
+        indices,
+        block_table,
+        token_to_req,
+    )
+    assert torch.equal(got, want), (
+        "NVFP4 QSA attend diverged from its bf16 decode (max diff "
+        f"{(got.float() - want.float()).abs().max().item():.3e})"
+    )
+
+
+def test_qsa_nvfp4_requires_both_block_scale_tensors():
+    q = torch.zeros(1, 2, HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    codes = torch.zeros(2, PAGE, 1, HEAD_DIM // 2, device="cuda", dtype=torch.uint8)
+    scales = torch.ones(2 * PAGE, 1, device="cuda", dtype=torch.float32)
+    indices = torch.zeros(1, 8, device="cuda", dtype=torch.int32)
+    block_table = torch.zeros(1, 2, device="cuda", dtype=torch.int32)
+    token_to_req = torch.zeros(1, device="cuda", dtype=torch.int32)
+
+    with pytest.raises(ValueError, match="NVFP4"):
+        qsa_sparse_paged_attention(
+            q,
+            codes,
+            codes,
+            indices,
+            block_table,
+            token_to_req,
+            k_scale=scales,
+            v_scale=scales,
+            kv_quant="nvfp4",
+        )
+
+
+def test_qsa_nvfp4_splitk_reads_value_rows():
+    """A zero query makes the split-K result the selected V-row average."""
+    torch.manual_seed(23)
+    slots = 3 * PAGE
+    v = torch.randn(slots, HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    codes = torch.zeros((slots, 1, HEAD_DIM // 2), dtype=torch.uint8, device="cuda")
+    row = torch.zeros((slots, 1), dtype=torch.float32, device="cuda")
+    block = torch.zeros((slots, 1, HEAD_DIM // 16), dtype=torch.uint8, device="cuda")
+    quantize_nvfp4_to_cache(
+        v, v, torch.arange(slots, dtype=torch.int32, device="cuda"), codes, codes,
+        row, row, block, block,
+    )
+    q = torch.zeros(1, 2, HEAD_DIM, device="cuda", dtype=torch.bfloat16)
+    indices, block_table, token_to_req = _layout(1, topk=64)
+    args = (indices, block_table, token_to_req)
+    got = qsa_sparse_paged_attention(
+        q, codes.view(3, PAGE, 1, HEAD_DIM // 2), codes.view(3, PAGE, 1, HEAD_DIM // 2),
+        *args, k_scale=row, v_scale=row, kv_quant="nvfp4", k_block_scale=block,
+        v_block_scale=block,
+    )
+    ref = _decode(codes, row, block).view(3, PAGE, 1, HEAD_DIM)
+    want = qsa_sparse_paged_attention(q, ref, ref, *args)
+    assert torch.equal(got, want)
diff --git a/tests/kernels/test_triton_attention.py b/tests/kernels/test_triton_attention.py
index a29f621dd..ea86fec5b 100644
--- a/tests/kernels/test_triton_attention.py
+++ b/tests/kernels/test_triton_attention.py
@@ -49,6 +49,34 @@ def _reference_paged_attention(
     return torch.stack(outs, dim=0)
 
 
+def _mm_test_pool(monkeypatch, heads, head_dim, num_slots, kv_quant):
+    from freetoken.kvcache.mha_pool import MHAKVCache
+
+    monkeypatch.setattr(
+        "freetoken.kvcache.mha_pool.get_tp_info", lambda: SimpleNamespace(size=1)
+    )
+    return MHAKVCache(
+        heads, 1, head_dim, num_slots, 1, torch.bfloat16,
+        torch.device("cuda"), kv_quant=kv_quant,
+    )
+
+
+def _restored_mha_cache(pool, which):
+    codes = getattr(pool, f"{which}_cache")(0).flatten(0, 1)
+    if pool.kv_quant == "none":
+        return codes.float()
+    row_scale = getattr(pool, f"{which}_scale")(0)
+    if pool.kv_quant == "fp8":
+        return codes.view(torch.float8_e4m3fn).float() * row_scale[..., None]
+    code = torch.stack((codes & 15, codes >> 4), -1).flatten(-2).long()
+    values = torch.tensor(
+        [0, .5, 1, 1.5, 2, 3, 4, 6, 0, -.5, -1, -1.5, -2, -3, -4, -6],
+        device=codes.device,
+    )
+    block_scale = getattr(pool, f"{which}_block_scale")(0).view(torch.float8_e4m3fn).float()
+    return values[code] * block_scale.repeat_interleave(16, -1) * row_scale[..., None]
+
+
 def test_triton_backend_passes_attention_sinks_to_paged_kernel(monkeypatch):
     from freetoken.attention import AttentionSpec
     from freetoken.attention.triton import TritonAttentionBackend, TritonMetadata
@@ -70,6 +98,14 @@ def k_cache(self, layer_id):
         def v_cache(self, layer_id):
             return self.v
 
+        # A 16-bit pool answers None here. The backend reads it on every path since the
+        # fp8 store landed, so the double answers it too.
+        def k_scale(self, layer_id):
+            return None
+
+        def v_scale(self, layer_id):
+            return None
+
     kv_cache = FakeKVCache()
     monkeypatch.setattr(
         "freetoken.attention.triton.get_global_ctx",
@@ -517,7 +553,10 @@ def test_extend_triton_attention_matches_reference(
 @pytest.mark.skipif(not torch.cuda.is_available(), reason="Triton attention needs CUDA")
 @pytest.mark.parametrize("use_split_inputs", [False, True])
 @pytest.mark.parametrize("sliding_window", [None, 48])
-def test_extend_triton_attention_with_bidirectional_blocks_matches_reference(use_split_inputs: bool, sliding_window: int | None):
+@pytest.mark.parametrize("kv_quant", ["none", "fp8", "nvfp4"])
+def test_extend_triton_attention_with_bidirectional_blocks_matches_reference(
+    monkeypatch, use_split_inputs: bool, sliding_window: int | None, kv_quant: str,
+):
     """Rows of an image span see the span's later keys across q tiles, text rows stay causal, the window still bounds the past."""
     from freetoken.kernel.triton.attention import extend_paged_attention
 
@@ -531,9 +570,13 @@ def test_extend_triton_attention_with_bidirectional_blocks_matches_reference(use
     q = torch.randn(total_q, num_q_heads, head_dim, device=device, dtype=torch.bfloat16)
     k_cache = torch.randn(total_kv, num_kv_heads, head_dim, device=device, dtype=torch.bfloat16)
     v_cache = torch.randn(total_kv, num_kv_heads, head_dim, device=device, dtype=torch.bfloat16)
+    row_gain = torch.linspace(.25, 2, total_kv * num_kv_heads, device=device).view(total_kv, num_kv_heads, 1)
+    block_gain = torch.linspace(.5, 1.5, head_dim // 16, device=device).repeat_interleave(16)
+    k_cache.mul_(row_gain * block_gain)
+    v_cache.mul_(row_gain.flip(0) * block_gain.flip(0))
     qo_indptr = torch.tensor([0] + extend_lens, dtype=torch.int32, device=device).cumsum_(0)
     kv_indptr = torch.tensor([0] + seq_lens, dtype=torch.int32, device=device).cumsum_(0)
-    indices = torch.arange(total_kv, dtype=torch.int32, device=device)
+    indices = torch.randperm(total_kv, device=device).to(torch.int32)
     prefix_lens = torch.tensor(cached_lens, dtype=torch.int32, device=device)
     q_to_req = torch.empty(total_q, dtype=torch.int32, device=device)
     q_positions = torch.empty(total_q, dtype=torch.int64, device=device)
@@ -548,16 +591,29 @@ def test_extend_triton_attention_with_bidirectional_blocks_matches_reference(use
     k_extend = torch.cat([k_cache[kv_indptr[i] + cached_lens[i] : kv_indptr[i + 1]] for i in range(2)])
     v_extend = torch.cat([v_cache[kv_indptr[i] + cached_lens[i] : kv_indptr[i + 1]] for i in range(2)])
     sm_scale = head_dim**-0.5
+    pool = _mm_test_pool(monkeypatch, num_kv_heads, head_dim, total_kv, kv_quant)
+    pool.store_kv(k_cache.flatten(1), v_cache.flatten(1), indices, 0)
+    k_ref, v_ref = [_restored_mha_cache(pool, which) for which in ("k", "v")]
+    if use_split_inputs:
+        for i, cached_len in enumerate(cached_lens):
+            slots = indices[kv_indptr[i] + cached_len : kv_indptr[i + 1]].long()
+            k_ref[slots] = k_extend[qo_indptr[i] : qo_indptr[i + 1]].float()
+            v_ref[slots] = v_extend[qo_indptr[i] : qo_indptr[i + 1]].float()
+    if kv_quant == "nvfp4":
+        k_ref, v_ref = [value.to(q.dtype).float() for value in (k_ref, v_ref)]
 
     actual = extend_paged_attention(
-        q, k_cache, v_cache, qo_indptr, kv_indptr, indices, prefix_lens, max(extend_lens), sm_scale,
+        q, pool.k_cache(0).flatten(0, 1), pool.v_cache(0).flatten(0, 1),
+        qo_indptr, kv_indptr, indices, prefix_lens, max(extend_lens), sm_scale,
         sliding_window=sliding_window, block_ends=block_ends,
         k_extend=k_extend if use_split_inputs else None, v_extend=v_extend if use_split_inputs else None,
+        k_scale=pool.k_scale(0), v_scale=pool.v_scale(0), kv_quant=kv_quant,
+        k_block_scale=pool.k_block_scale(0), v_block_scale=pool.v_block_scale(0),
     )
     expected = _reference_paged_attention(
-        q, k_cache, v_cache, kv_indptr, indices, q_to_req, q_positions, sm_scale, sliding_window, block_ends=block_ends
+        q, k_ref, v_ref, kv_indptr, indices, q_to_req, q_positions, sm_scale, sliding_window, block_ends=block_ends
     )
-    causal_only = _reference_paged_attention(q, k_cache, v_cache, kv_indptr, indices, q_to_req, q_positions, sm_scale, sliding_window)
+    causal_only = _reference_paged_attention(q, k_ref, v_ref, kv_indptr, indices, q_to_req, q_positions, sm_scale, sliding_window)
 
     torch.testing.assert_close(actual.float(), expected.float(), atol=2e-2, rtol=2e-2)
     assert not torch.allclose(expected.float(), causal_only.float(), atol=2e-2, rtol=2e-2)
@@ -666,6 +722,36 @@ def test_select_extend_tile_is_shared_memory_aware(head_dim, smem_optin, expecte
     assert _select_extend_tile(head_dim, block_d, smem_optin) == expected
 
 
+@pytest.mark.parametrize(
+    ("head_dim", "smem_optin", "expected_16bit", "expected_fp8"),
+    [
+        # consumer opt-in smem (sm_86/sm_89 ~99KB): a 1-byte cache buys BLOCK_N 32 -> 64
+        (256, 101376, (64, 32), (64, 64)),
+        # where the fast tile already fits, the cache dtype changes nothing
+        (256, 232448, (128, 64), (128, 64)),
+        # unknown budget stays conservative for both
+        (256, 0, (64, 32), (64, 32)),
+    ],
+)
+def test_select_extend_tile_uses_kv_cache_element_size(
+    head_dim, smem_optin, expected_16bit, expected_fp8
+):
+    """The q tile is always 2 bytes/element but K and V follow the cache, so charging
+    K/V at 2 bytes regardless makes an fp8 cache run a smaller tile than it has shared
+    memory for. On an RTX 3070 (sm_86, 99KB opt-in, head_dim 256) that was BLOCK_N 32
+    where 64 fits, worth ~9% of prefill time at 99k context.
+
+    The 16-bit column is the no-regression half: ``kv_bytes=2`` reproduces the previous
+    budget identically, since (M + 2N) * D * 2 == (M*2 + 2N*2) * D."""
+    import triton
+
+    from freetoken.kernel.triton.attention import _select_extend_tile
+
+    block_d = triton.next_power_of_2(head_dim)
+    assert _select_extend_tile(head_dim, block_d, smem_optin, 2) == expected_16bit
+    assert _select_extend_tile(head_dim, block_d, smem_optin, 1) == expected_fp8
+
+
 @pytest.mark.skipif(not torch.cuda.is_available(), reason="Triton attention needs CUDA")
 def test_triton_backend_stores_kv_and_matches_reference(monkeypatch):
     from freetoken.attention import AttentionSpec
@@ -688,6 +774,14 @@ def k_cache(self, layer_id):
         def v_cache(self, layer_id):
             return self.v
 
+        # A 16-bit pool answers None here. The backend reads it on every path since the
+        # fp8 store landed, so the double answers it too.
+        def k_scale(self, layer_id):
+            return None
+
+        def v_scale(self, layer_id):
+            return None
+
     device = torch.device("cuda")
     head_dim = 256
     page_table = torch.tensor([[0, 1], [2, 3]], dtype=torch.int32, device=device)
@@ -737,38 +831,29 @@ def v_cache(self, layer_id):
 
 
 @pytest.mark.skipif(not torch.cuda.is_available(), reason="Triton attention needs CUDA")
-def test_triton_backend_applies_the_batch_block_ends_only_when_the_spec_asks(monkeypatch):
+@pytest.mark.parametrize("kv_quant", ["none", "fp8", "nvfp4"])
+def test_triton_backend_applies_the_batch_block_ends_only_when_the_spec_asks(monkeypatch, kv_quant):
     from freetoken.attention import AttentionSpec
     from freetoken.attention.triton import TritonAttentionBackend
 
-    class FakeKVCache:
-        def __init__(self, device: torch.device, head_dim: int):
-            self.device = device
-            self.dtype = torch.bfloat16
-            self.k = torch.randn(3, 1, head_dim, device=device, dtype=torch.bfloat16)
-            self.v = torch.randn(3, 1, head_dim, device=device, dtype=torch.bfloat16)
-
-        def store_kv(self, k, v, out_loc, layer_id):
-            self.k[out_loc.to(torch.long)] = k.view(k.shape[0], 1, -1)
-            self.v[out_loc.to(torch.long)] = v.view(v.shape[0], 1, -1)
-
-        def k_cache(self, layer_id):
-            return self.k
-
-        def v_cache(self, layer_id):
-            return self.v
-
+    torch.manual_seed(19)
     device = torch.device("cuda")
     head_dim = 256
-    kv_cache = FakeKVCache(device, head_dim)
-    ctx = SimpleNamespace(kv_cache=kv_cache, page_table=torch.tensor([[0, 1, 2]], dtype=torch.int32, device=device))
+    kv_cache = _mm_test_pool(monkeypatch, 1, head_dim, 4, kv_quant)
+    page_table = torch.tensor([[2, 0, 3, 1]], dtype=torch.int32, device=device)
+    kv_cache.store_kv(
+        torch.randn(4, head_dim, device=device, dtype=torch.bfloat16),
+        torch.randn(4, head_dim, device=device, dtype=torch.bfloat16),
+        page_table.flatten(), 0,
+    )
+    ctx = SimpleNamespace(kv_cache=kv_cache, page_table=page_table)
     monkeypatch.setattr("freetoken.attention.triton.get_global_ctx", lambda: ctx)
     backend = TritonAttentionBackend(SimpleNamespace())
     # one request: cached token 0, chunk tokens 1 and 2 forming one image span [1, 3)
     batch = SimpleNamespace(
         padded_reqs=[SimpleNamespace(extend_len=2, device_len=3, cached_len=1, table_idx=0)],
         positions=torch.tensor([1, 2], dtype=torch.int64, device=device),
-        out_loc=torch.tensor([1, 2], dtype=torch.int32, device=device),
+        out_loc=page_table[0, 1:3],
         mm_block_ends=torch.tensor([3, 3], dtype=torch.int32, device=device),
     )
     # bf16 takes the extend kernel path; fp32 would fall back to the paged kernel, which has no block mask
@@ -777,15 +862,38 @@ def v_cache(self, layer_id):
     v = torch.randn(2, head_dim, device=device, dtype=torch.bfloat16)
     backend.prepare_metadata(batch)
     md = batch.attn_metadata
-    ref = lambda ends: _reference_paged_attention(
-        q, kv_cache.k, kv_cache.v, md.indptr, md.indices, md.q_to_req, md.q_positions, head_dim**-0.5, None, block_ends=ends
-    )
+    def ref(ends, split_inputs=True):
+        k_ref, v_ref = [_restored_mha_cache(kv_cache, which) for which in ("k", "v")]
+        if split_inputs:
+            k_ref[batch.out_loc.long()] = k.view(-1, 1, head_dim).float()
+            v_ref[batch.out_loc.long()] = v.view(-1, 1, head_dim).float()
+        if kv_quant == "nvfp4":
+            k_ref, v_ref = [value.to(q.dtype).float() for value in (k_ref, v_ref)]
+        return _reference_paged_attention(
+            q, k_ref, v_ref, md.indptr, md.indices, md.q_to_req, md.q_positions,
+            head_dim**-0.5, None, block_ends=ends,
+        )
+
     causal = backend.forward(q, k, v, layer_id=0, batch=batch, attn_spec=AttentionSpec(sm_scale=head_dim**-0.5))
     torch.testing.assert_close(causal.float(), ref(None).float(), atol=2e-2, rtol=2e-2)
     blocks = backend.forward(q, k, v, layer_id=0, batch=batch, attn_spec=AttentionSpec(sm_scale=head_dim**-0.5, bidirectional_mm_blocks=True))
     torch.testing.assert_close(blocks.float(), ref(batch.mm_block_ends).float(), atol=2e-2, rtol=2e-2)
     assert not torch.allclose(blocks[0].float(), causal[0].float(), atol=2e-2, rtol=2e-2)  # token 1 now sees token 2
 
+    batch = SimpleNamespace(
+        padded_reqs=[SimpleNamespace(extend_len=1, device_len=4, cached_len=3, table_idx=0)],
+        positions=torch.tensor([3], dtype=torch.int64, device=device),
+        out_loc=page_table[0, 3:4],
+        mm_block_ends=None,
+    )
+    q = torch.randn(1, 2, head_dim, device=device, dtype=torch.bfloat16)
+    k = torch.randn(1, head_dim, device=device, dtype=torch.bfloat16)
+    v = torch.randn(1, head_dim, device=device, dtype=torch.bfloat16)
+    backend.prepare_metadata(batch)
+    md = batch.attn_metadata
+    decoded = backend.forward(q, k, v, layer_id=0, batch=batch, attn_spec=AttentionSpec(bidirectional_mm_blocks=True))
+    torch.testing.assert_close(decoded.float(), ref(None, split_inputs=False).float(), atol=2e-2, rtol=2e-2)
+
 
 @pytest.mark.skipif(not torch.cuda.is_available(), reason="Triton attention needs CUDA")
 def test_triton_backend_replay_metadata_uses_capture_buffers(monkeypatch):
@@ -884,3 +992,271 @@ def test_triton_metadata_keeps_full_indices_and_optional_swa_indices(monkeypatch
     assert metadata.indices.tolist() == [10, 11, 20, 21, 22]
     assert metadata.swa_indices is not None
     assert metadata.swa_indices.tolist() == [110, 111, 120, 121, 122]
+
+
+def _fp8_cache(k_rows: torch.Tensor, v_rows: torch.Tensor):
+    """Quantize ``[slots, heads, dim]`` KV into codes + scales, in the layout the
+    attention kernels expect."""
+    from freetoken.kernel.triton.kv_quant import alloc_codes, quantize_kv_to_cache
+
+    slots, heads, dim = k_rows.shape
+    k_cache = alloc_codes((slots, heads, dim), k_rows.device)
+    v_cache = alloc_codes((slots, heads, dim), k_rows.device)
+    k_scale = torch.zeros((slots, heads), dtype=torch.float32, device=k_rows.device)
+    v_scale = torch.zeros((slots, heads), dtype=torch.float32, device=k_rows.device)
+    quantize_kv_to_cache(
+        k=k_rows.reshape(slots, heads * dim),
+        v=v_rows.reshape(slots, heads * dim),
+        out_loc=torch.arange(slots, dtype=torch.int32, device=k_rows.device),
+        k_cache=k_cache,
+        v_cache=v_cache,
+        k_scale=k_scale,
+        v_scale=v_scale,
+    )
+    torch.cuda.synchronize()
+    return k_cache, v_cache, k_scale, v_scale
+
+
+def _dequantized(codes: torch.Tensor, scale: torch.Tensor) -> torch.Tensor:
+    """Host-side decode (torch's e4m3 cast), so the reference never reuses the
+    software decoder the kernel is being tested against."""
+    from freetoken.kernel.triton.kv_quant import codes_to_f32
+
+    return codes_to_f32(codes) * scale.unsqueeze(-1)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="Triton attention needs CUDA")
+@pytest.mark.parametrize(("head_dim", "num_kv_heads"), [(64, 4), (128, 2)])
+def test_decode_paged_attention_decodes_fp8_scales(head_dim: int, num_kv_heads: int):
+    """fp8 decode must equal the SAME data dequantized by hand.
+
+    Not compared against the bf16 cache: the fp8-vs-bf16 gap is quantization error,
+    already bounded in tests/kernels/test_kv_fp8.py. What this pins is that the
+    kernel applies the right scale to the right row -- a dropped or mis-indexed
+    scale is off by a factor of amax/448, which no tolerance hides.
+    """
+    from freetoken.kernel.triton.attention import decode_paged_attention
+
+    torch.manual_seed(11)
+    device = torch.device("cuda")
+    batch, num_q_heads, max_kv_splits = 2, 8, 8
+    seq_lens = [6, 9]
+    total_kv = sum(seq_lens)
+    q = torch.randn(batch, num_q_heads, head_dim, device=device, dtype=torch.bfloat16)
+    # Row magnitudes far from 1, so a missing scale cannot pass by luck.
+    k_rows = (torch.randn(total_kv, num_kv_heads, head_dim, device=device) * 8.0).to(
+        torch.bfloat16
+    )
+    v_rows = (torch.randn(total_kv, num_kv_heads, head_dim, device=device) * 0.05).to(
+        torch.bfloat16
+    )
+    indptr = torch.tensor([0, seq_lens[0], total_kv], dtype=torch.int32, device=device)
+    indices = torch.arange(total_kv, dtype=torch.int32, device=device)
+    q_positions = torch.tensor(
+        [seq_lens[0] - 1, seq_lens[1] - 1], dtype=torch.int64, device=device
+    )
+    q_to_req = torch.arange(batch, dtype=torch.int32, device=device)
+    attn_logits = torch.empty(
+        batch, num_q_heads, max_kv_splits, head_dim, dtype=torch.float32, device=device
+    )
+    attn_lse = torch.empty(batch, num_q_heads, max_kv_splits, dtype=torch.float32, device=device)
+    num_kv_splits = torch.full((batch,), max_kv_splits, dtype=torch.int32, device=device)
+    sm_scale = head_dim**-0.5
+
+    k_codes, v_codes, k_scale, v_scale = _fp8_cache(k_rows, v_rows)
+    actual = decode_paged_attention(
+        q,
+        k_codes,
+        v_codes,
+        indptr,
+        indices,
+        q_positions,
+        attn_logits,
+        attn_lse,
+        num_kv_splits,
+        max_kv_splits,
+        sm_scale,
+        k_scale=k_scale,
+        v_scale=v_scale,
+    )
+    expected = _reference_paged_attention(
+        q,
+        _dequantized(k_codes, k_scale),
+        _dequantized(v_codes, v_scale),
+        indptr,
+        indices,
+        q_to_req,
+        q_positions,
+        sm_scale,
+        None,
+    )
+    # The kernel rounds the decoded operands to the compute dtype before tl.dot; the
+    # reference stays in fp32, hence the same tolerance the bf16 tests already use.
+    torch.testing.assert_close(actual.float(), expected.float(), atol=2e-2, rtol=2e-2)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="Triton attention needs CUDA")
+def test_paged_attention_decodes_fp8_scales():
+    """The non-tl.dot fallback kernel takes the same scales (it serves head_dim > 256
+    with a short prefill, and every decode batch when the grouped path is not used)."""
+    from freetoken.kernel.triton.attention import paged_attention
+
+    torch.manual_seed(13)
+    device = torch.device("cuda")
+    head_dim, num_kv_heads, num_q_heads = 64, 2, 4
+    seq_lens = [5, 3]
+    total_kv = sum(seq_lens)
+    q = torch.randn(total_kv, num_q_heads, head_dim, device=device, dtype=torch.bfloat16)
+    k_rows = (torch.randn(total_kv, num_kv_heads, head_dim, device=device) * 16.0).to(
+        torch.bfloat16
+    )
+    v_rows = (torch.randn(total_kv, num_kv_heads, head_dim, device=device) * 0.02).to(
+        torch.bfloat16
+    )
+    indptr = torch.tensor([0, seq_lens[0], total_kv], dtype=torch.int32, device=device)
+    indices = torch.arange(total_kv, dtype=torch.int32, device=device)
+    q_to_req = torch.tensor(
+        [0] * seq_lens[0] + [1] * seq_lens[1], dtype=torch.int32, device=device
+    )
+    q_positions = torch.cat(
+        [
+            torch.arange(seq_lens[0], dtype=torch.int64, device=device),
+            torch.arange(seq_lens[1], dtype=torch.int64, device=device),
+        ]
+    )
+    k_codes, v_codes, k_scale, v_scale = _fp8_cache(k_rows, v_rows)
+    actual = paged_attention(
+        q,
+        k_codes,
+        v_codes,
+        indptr,
+        indices,
+        q_to_req,
+        q_positions,
+        head_dim**-0.5,
+        k_scale=k_scale,
+        v_scale=v_scale,
+    )
+    expected = _reference_paged_attention(
+        q,
+        _dequantized(k_codes, k_scale),
+        _dequantized(v_codes, v_scale),
+        indptr,
+        indices,
+        q_to_req,
+        q_positions,
+        head_dim**-0.5,
+        None,
+    )
+    torch.testing.assert_close(actual.float(), expected.float(), atol=2e-2, rtol=2e-2)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="Triton attention needs CUDA")
+@pytest.mark.parametrize("use_split_inputs", [False, True])
+def test_extend_paged_attention_decodes_fp8_scales(use_split_inputs: bool):
+    """Prefill over an fp8 cache.
+
+    With ``use_split_inputs`` the kernel reads a request's own new tokens from
+    ``k_extend`` (compute dtype, never quantized) and only the cached prefix from the
+    fp8 codes; without it every row is served from the codes. The reference mirrors
+    that split, so a scale leaking onto the extend path -- or failing to apply to the
+    cache path -- cannot pass.
+    """
+    from freetoken.kernel.triton.attention import extend_paged_attention
+
+    torch.manual_seed(14)
+    device = torch.device("cuda")
+    head_dim, num_kv_heads, num_q_heads = 64, 2, 8
+    cached_lens, extend_lens = [4, 2], [3, 2]
+    seq_lens = [c + e for c, e in zip(cached_lens, extend_lens)]
+    total_q, total_kv = sum(extend_lens), sum(seq_lens)
+    q = torch.randn(total_q, num_q_heads, head_dim, device=device, dtype=torch.bfloat16)
+    k_extend = (torch.randn(total_q, num_kv_heads, head_dim, device=device) * 3.0).to(
+        torch.bfloat16
+    )
+    v_extend = (torch.randn(total_q, num_kv_heads, head_dim, device=device) * 3.0).to(
+        torch.bfloat16
+    )
+    # Magnitudes far from 1: a dropped scale is then off by orders of magnitude.
+    k_cache = (torch.randn(total_kv, num_kv_heads, head_dim, device=device) * 0.3).to(
+        torch.bfloat16
+    )
+    v_cache = (torch.randn(total_kv, num_kv_heads, head_dim, device=device) * 3.0).to(
+        torch.bfloat16
+    )
+    qo_indptr = torch.tensor([0] + extend_lens, dtype=torch.int32, device=device).cumsum_(0)
+    kv_indptr = torch.tensor([0] + seq_lens, dtype=torch.int32, device=device).cumsum_(0)
+    # KV positions are logical, but FP8 codes and their scales are addressed by the
+    # physical slots from the page table. A contiguous table masks a regression that
+    # looks scales up with the logical position instead of the slot.
+    indices = torch.tensor(
+        [10, 3, 8, 1, 9, 0, 7, 2, 6, 4, 5], dtype=torch.int32, device=device
+    )
+    prefix_lens = torch.tensor(cached_lens, dtype=torch.int32, device=device)
+    q_to_req = torch.empty(total_q, dtype=torch.int32, device=device)
+    q_positions = torch.empty(total_q, dtype=torch.int64, device=device)
+    q_off = kv_off = 0
+    for req_idx, (cached_len, extend_len) in enumerate(zip(cached_lens, extend_lens)):
+        q_to_req[q_off : q_off + extend_len].fill_(req_idx)
+        q_positions[q_off : q_off + extend_len] = torch.arange(
+            cached_len, cached_len + extend_len, dtype=torch.int64, device=device
+        )
+        # The step's own tokens are what the engine stores at the tail of the span.
+        k_cache[kv_off + cached_len : kv_off + cached_len + extend_len] = k_extend[
+            q_off : q_off + extend_len
+        ]
+        v_cache[kv_off + cached_len : kv_off + cached_len + extend_len] = v_extend[
+            q_off : q_off + extend_len
+        ]
+        q_off += extend_len
+        kv_off += cached_len + extend_len
+    sm_scale = head_dim**-0.5
+
+    num_slots = total_kv + 1
+    k_slots = torch.zeros(
+        num_slots, num_kv_heads, head_dim, dtype=k_cache.dtype, device=device
+    )
+    v_slots = torch.zeros_like(k_slots)
+    k_slots[indices.to(torch.long)] = k_cache
+    v_slots[indices.to(torch.long)] = v_cache
+    k_codes, v_codes, k_scale, v_scale = _fp8_cache(k_slots, v_slots)
+    k_ref = _dequantized(k_codes, k_scale).clone()
+    v_ref = _dequantized(v_codes, v_scale).clone()
+    if use_split_inputs:
+        q_off = kv_off = 0
+        for cached_len, extend_len in zip(cached_lens, extend_lens):
+            current_slots = indices[
+                kv_off + cached_len : kv_off + cached_len + extend_len
+            ].to(torch.long)
+            k_ref[current_slots] = k_extend[q_off : q_off + extend_len].float()
+            v_ref[current_slots] = v_extend[q_off : q_off + extend_len].float()
+            q_off += extend_len
+            kv_off += cached_len + extend_len
+
+    actual = extend_paged_attention(
+        q=q,
+        k_cache=k_codes,
+        v_cache=v_codes,
+        qo_indptr=qo_indptr,
+        kv_indptr=kv_indptr,
+        kv_indices=indices,
+        prefix_lens=prefix_lens,
+        max_q_len=max(extend_lens),
+        sm_scale=sm_scale,
+        k_extend=k_extend if use_split_inputs else None,
+        v_extend=v_extend if use_split_inputs else None,
+        k_scale=k_scale,
+        v_scale=v_scale,
+    )
+    expected = _reference_paged_attention(
+        q,
+        k_ref,
+        v_ref,
+        kv_indptr,
+        indices,
+        q_to_req,
+        q_positions,
+        sm_scale,
+        None,
+    )
+    torch.testing.assert_close(actual.float(), expected.float(), atol=3e-2, rtol=3e-2)
diff --git a/tests/kvcache/test_dsa_pool.py b/tests/kvcache/test_dsa_pool.py
index 7f4b1c709..8bcdb8d5e 100644
--- a/tests/kvcache/test_dsa_pool.py
+++ b/tests/kvcache/test_dsa_pool.py
@@ -16,13 +16,14 @@
 LATENT, IDX_DIM = 80, 32  # latent 64+16: kernel spans need pow-2 dims (512+64 in the real model)
 
 
-def _pool(num_pages: int):
+def _pool(num_pages: int, kv_quant: str = "none"):
     from freetoken.kvcache.dsa_pool import DSAKVCache
 
     return DSAKVCache(
         latent_dim=LATENT, num_layers=2, num_pages=num_pages, page_size=1,
         dtype=torch.bfloat16, device=torch.device("cuda"),
         index_head_dim=IDX_DIM, num_index_layers=1,
+        kv_quant=kv_quant,
     )
 
 
@@ -91,6 +92,39 @@ def test_sparse_decode_reads_grown_pool_through_backend_kernels():
     assert (o[0, 0].float() - ref_o).abs().max().item() < 2e-2
 
 
+def test_fp8_latent_rows_decode_with_their_scales():
+    """The DSA kernel must read codes and scales from matching physical rows."""
+    from freetoken.kernel.triton.glm_dsa_sparse import glm_dsa_sparse_attn
+    from freetoken.kernel.triton.kv_quant import codes_to_f32
+
+    torch.manual_seed(7)
+    pool = _pool(128, kv_quant="fp8")
+    rows = torch.randperm(128, device="cuda")[:96].to(torch.int32)
+    c_kv = torch.randn(96, LATENT - 16, device="cuda", dtype=torch.bfloat16)
+    k_rope = torch.randn(96, 16, device="cuda", dtype=torch.bfloat16)
+    pool.store_kv(c_kv, k_rope, rows, layer_id=0)
+    assert pool.latent_rows(0).dtype is torch.uint8
+    assert pool.latent_scale(0).dtype is torch.float32
+
+    q = torch.randn(1, 1, 4, LATENT, device="cuda", dtype=torch.bfloat16)
+    sel = rows[:64].view(1, 1, -1)
+    cnt = torch.tensor([[64]], device="cuda", dtype=torch.int32)
+    out = glm_dsa_sparse_attn(
+        q, pool.latent_rows(0), sel, 0.1, counts=cnt, d_v=LATENT - 16,
+        pool_scale=pool.latent_scale(0),
+    )
+    out_split = glm_dsa_sparse_attn(
+        q, pool.latent_rows(0), sel, 0.1, counts=cnt, d_v=LATENT - 16,
+        pool_scale=pool.latent_scale(0), force_splits=4,
+    )
+    decoded = codes_to_f32(pool.latent_rows(0)) * pool.latent_scale(0).unsqueeze(-1)
+    picked = decoded[sel.view(-1).long()]
+    logits = (q[0, 0].float() @ picked.T) * 0.1
+    ref = logits.softmax(-1) @ picked[:, : LATENT - 16]
+    assert (out[0, 0].float() - ref).abs().max().item() < 2e-2
+    assert (out_split.float() - ref).abs().max().item() < 2e-2
+
+
 def test_mla_pool_selected_by_group_spec():
     """The factory keys MLA/DSA pools off the attention-group spec, never the model
     payload -- and zeroed index dims (the dense ablation) fall back to MLAKVCache."""
@@ -120,6 +154,10 @@ def cfg(index_dim, n_idx):
     dsa = create_kvcache_pool(model_config=cfg(IDX_DIM, 1), num_pages=8, page_size=1,
                               device=torch.device("cuda"), dtype=torch.bfloat16)
     assert isinstance(dsa, DSAKVCache)
+    fp8_dsa = create_kvcache_pool(model_config=cfg(IDX_DIM, 1), num_pages=8, page_size=1,
+                                  device=torch.device("cuda"), dtype=torch.bfloat16,
+                                  kv_quant="fp8")
+    assert isinstance(fp8_dsa, DSAKVCache) and fp8_dsa.store_dtype is torch.uint8
     mla = create_kvcache_pool(model_config=cfg(0, 0), num_pages=8, page_size=1,
                               device=torch.device("cuda"), dtype=torch.bfloat16)
     assert isinstance(mla, MLAKVCache) and not isinstance(mla, DSAKVCache)
@@ -136,3 +174,57 @@ def test_rebuild_shrink_and_engine_wiring():
     pool.rebuild_from_config(config=None, num_pages=63)
     assert pool.latent_rows(0).shape[0] == 64  # 63 + 1 dummy page
     assert pool.index_k_cache(0).shape[0] == 64
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8", "nvfp4"])
+@pytest.mark.parametrize("ratio", [1, 4])
+def test_latent_budget_matches_allocations_and_rebuild(kv_quant, ratio):
+    from types import SimpleNamespace
+    from freetoken.attention import AttnType
+    from freetoken.kvcache.dsa_pool import DSAKVCache, KpoolDSAKVCache
+    from freetoken.models.config import KVCacheGroupSpec
+
+    cls = KpoolDSAKVCache if ratio > 1 else DSAKVCache
+    spec = KVCacheGroupSpec(
+        name="full", layer_ids=(3, 7), num_kv_heads=1, head_dim=512, sliding_window=None,
+        mla=True, index_head_dim=128, num_index_layers=2, index_ratio=ratio,
+        attn_type=AttnType.DSA,
+    )
+    cfg = SimpleNamespace(
+        kv_quant=kv_quant, dtype=torch.bfloat16, tp_info=SimpleNamespace(size=1),
+        max_running_req=3, page_size=64,
+        model_config=SimpleNamespace(kv_cache_group_specs=lambda: (spec,)),
+    )
+    extra = dict(index_ratio=ratio, num_req_slots=4) if ratio > 1 else {}
+    pool = cls(512, 8, 2, 64, torch.bfloat16, torch.device("cuda"),
+               index_head_dim=128, num_index_layers=2, layer_ids=(3, 7),
+               kv_quant=kv_quant, **extra)
+    page_bytes, fixed, page_size, _ = cls.kv_cost(cfg)
+    for pages in (2, 5, 1):
+        pool.rebuild_from_config(cfg, pages)
+        allocated_pages = pages + 1
+        buffers = (pool._kv_buffer, pool._scale_buffer, pool._block_scale_buffer,
+                   pool._index_k_buffer)
+        if ratio > 1:
+            buffers += (pool._tail_k, pool._tail_gate)
+            assert pool.cmp_scratch_base == allocated_pages * 64 // ratio
+            assert pool.tail_k(0).dtype == torch.bfloat16
+        actual = sum(b.numel() * b.element_size() for b in buffers if b is not None)
+        assert actual == allocated_pages * page_bytes + fixed
+        assert pool.unit_bytes() == (page_bytes // page_size, 0)
+        assert pool.latent_rows(7).shape == (allocated_pages * 64, 256 if kv_quant == "nvfp4" else 512)
+        loc = torch.tensor([allocated_pages * 64 - 1], device="cuda")
+        x = torch.ones(1, 512, device="cuda", dtype=torch.bfloat16)
+        pool.store_kv(x, x[:, :0], loc, 7)
+        if kv_quant == "nvfp4":
+            from tests.kernels.test_kv_nvfp4 import _decode_latent
+
+            torch.testing.assert_close(_decode_latent(pool, 7)[loc], x.float())
+        assert pool.k_cache(7).data_ptr() == pool.v_cache(7).data_ptr()
+
+
+def test_nvfp4_latent_rejects_partial_blocks():
+    from freetoken.kvcache.dsa_pool import MLAKVCache
+
+    with pytest.raises(ValueError, match="divisible by 16"):
+        MLAKVCache(72, 1, 1, 1, torch.bfloat16, torch.device("cpu"), kv_quant="nvfp4")
diff --git a/tests/kvcache/test_dsv41_pool.py b/tests/kvcache/test_dsv41_pool.py
new file mode 100644
index 000000000..616fa5ec5
--- /dev/null
+++ b/tests/kvcache/test_dsv41_pool.py
@@ -0,0 +1,334 @@
+import json
+from pathlib import Path
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.kvcache.dsv41_cost_model import (
+    _dsv41_pool_sizes, dsv41_auto_cost_model, dsv41_pool_sizes, dsv41_pool_bytes,
+    dsv41_solve_num_pages, dsv41_unit_bytes, source_layers,
+)
+from freetoken.kvcache.dsv41_layout import dsv41_row_bytes
+from freetoken.kvcache.dsv41_paged_pool import DSV41PagedKVCache
+
+
+def args():
+    return SimpleNamespace(n_layers=5, compress_ratios=(0, 2, 2, 1, 1),
+                           kv_source_layers=(1, 3), index_source_layers=(1, 3, 4),
+                           head_dim=32, index_head_dim=32, window_size=128)
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_only_sources_allocate_compressed_keys_and_pool_bytes_are_exact(kv_quant):
+    a = args()
+    sizes = dsv41_pool_sizes(13, a, .4)
+    pool = DSV41PagedKVCache(sizes, a, torch.device("cpu"), n_scratch=3, kv_quant=kv_quant)
+    assert pool.kv_sources == (None, 1, 1, 3, 3)
+    assert pool.index_sources == (None, 1, 1, 3, 4)
+    assert [i for i, p in enumerate(pool.cmp_pool) if p is not None] == [1, 3]
+    assert [i for i, p in enumerate(pool.idx_pool) if p is not None] == [1, 3]
+    assert [i for i, p in enumerate(pool.state_ring) if p is not None] == [1]
+    assert pool.total_bytes() == dsv41_pool_bytes(sizes, a, 3, kv_quant)
+    assert pool.dtype == torch.bfloat16
+    assert pool.state_ring[1].buffer.dtype == torch.float32
+    expected_widths = (32, 32, 32) if kv_quant == "none" else (33, 18, 17)
+    stored = (pool.window_pool[0], pool.cmp_pool[1], pool.idx_pool[1])
+    assert tuple(t.shape[1] for t in stored) == expected_widths
+    storage_dtype = torch.bfloat16 if kv_quant == "none" else torch.uint8
+    assert all(t.dtype == storage_dtype for t in stored)
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_two_request_pending_pairs_have_disjoint_state_after_page_recycling(kv_quant):
+    a = args()
+    pool = DSV41PagedKVCache(dsv41_pool_sizes(8, a, 1), a, torch.device("cpu"), kv_quant=kv_quant)
+    pool._init_paged_state(2, True)
+    full = torch.arange(256)
+    pool.alloc_swa(full)
+    window = pool.translate_full_to_window(torch.tensor([0, 128]))
+    state = pool.state_loc(window, 2, 128)
+    expected = torch.stack((torch.full((64,), 3.), torch.full((64,), -7.)))
+    pool.set_state(1, state, expected)
+    assert state[0] != state[1]
+    torch.testing.assert_close(pool.get_state(1, state), expected)
+    pool.free_swa(full[:128])
+    pool.alloc_swa(torch.arange(256, 384))
+    assert pool.translate_full_to_window(torch.tensor([0])).item() == -1
+    torch.testing.assert_close(pool.get_state(1, state[1:]), expected[1:])
+
+
+@pytest.mark.parametrize("field,value", [("compress_ratios", (0, 2, 1, 1, 1)),
+                                         ("kv_source_layers", (1, 4)),
+                                         ("index_source_layers", (3, 4))])
+def test_reject_malformed_source_geometry(field, value):
+    a = args()
+    setattr(a, field, value)
+    with pytest.raises(ValueError):
+        source_layers(a)
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_solver_respects_exact_source_owned_cost_and_maximal_fit(kv_quant):
+    a = args()
+    config = SimpleNamespace(model_config=SimpleNamespace(dsv41_args=a), max_seq_len=8192,
+                             max_running_req=2, cache_type="swa_radix", swa_full_tokens_ratio=.2,
+                             swa_num_pages_override=None, kv_quant=kv_quant)
+    budget = dsv41_pool_bytes(_dsv41_pool_sizes(config, 64), a, 3, kv_quant) + 100
+    usable = dsv41_solve_num_pages(config, budget)
+    assert dsv41_pool_bytes(_dsv41_pool_sizes(config, usable + 1), a, 3, kv_quant) <= budget
+    assert dsv41_pool_bytes(_dsv41_pool_sizes(config, usable + 2), a, 3, kv_quant) > budget
+
+
+@pytest.mark.parametrize("ratio", [.001, .2, .333, 1.])
+@pytest.mark.parametrize("window_pages", [None, 28, 57])
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_auto_budget_includes_dummy_floor_and_window_rounding(ratio, window_pages, kv_quant):
+    a = args()
+    config = SimpleNamespace(model_config=SimpleNamespace(dsv41_args=a), max_seq_len=1 << 20,
+                             max_running_req=2, cache_type="swa_radix", swa_full_tokens_ratio=ratio,
+                             swa_num_pages_override=window_pages, kv_quant=kv_quant)
+    per_page, fixed, P, reserve = dsv41_auto_cost_model(config)
+    for usable in (reserve // P, 41, 100, 8192):
+        exact = dsv41_pool_bytes(_dsv41_pool_sizes(config, usable + 1), a, 3, kv_quant)
+        assert usable * per_page + fixed >= exact
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_rebuild_budget_keeps_auxiliary_fixed_allocation(kv_quant):
+    from freetoken.kvcache.base import CacheRebuildRejected
+
+    a = args()
+    sizes = dsv41_pool_sizes(32, a, .5)
+    pool = DSV41PagedKVCache(sizes, a, torch.device("cpu"), n_scratch=3, kv_quant=kv_quant)
+    config = SimpleNamespace(model_config=SimpleNamespace(dsv41_args=a), max_seq_len=4096,
+                             max_running_req=2, cache_type="swa_radix", swa_full_tokens_ratio=.5,
+                             swa_num_pages_override=None, memory_ratio=1., kv_quant=kv_quant)
+    kwargs = dict(num_pages=None, target_moe=0, per_expert_bytes=0,
+                  baseline_free=pool.total_bytes() + 4096, weights_bytes=0, current_num_pages=31)
+    pool.validate_rebuild(config, extra_fixed_bytes=4096, **kwargs)
+    with pytest.raises(CacheRebuildRejected, match="exceeding budget"):
+        pool.validate_rebuild(config, extra_fixed_bytes=4097, **kwargs)
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_smallest_solved_pool_honors_usable_page_floor(kv_quant):
+    a = args()
+    config = SimpleNamespace(model_config=SimpleNamespace(dsv41_args=a), max_seq_len=8192,
+                             max_running_req=2, cache_type="swa_radix", swa_full_tokens_ratio=.2,
+                             swa_num_pages_override=None, kv_quant=kv_quant)
+    floor = DSV41PagedKVCache.min_kv_tokens(config) // a.window_size
+    budget = dsv41_pool_bytes(_dsv41_pool_sizes(config, floor + 1), a, 3, kv_quant)
+    assert dsv41_solve_num_pages(config, budget) == floor
+    with pytest.raises(ValueError, match="window working set"):
+        dsv41_solve_num_pages(config, budget - 1)
+
+
+def _target_config(tokens, ratio, kv_quant="none"):
+    from freetoken.models.deepseek_v41.config import parse_config
+
+    fixture = Path(__file__).parents[1] / "models/fixtures/deepseek_v41_nvfp4_config.json"
+    return SimpleNamespace(model_config=parse_config(json.loads(fixture.read_text())),
+                           max_seq_len=tokens, max_running_req=2, cache_type="swa_radix",
+                           swa_full_tokens_ratio=ratio, swa_num_pages_override=None, kv_quant=kv_quant)
+
+
+@pytest.mark.parametrize("tokens,ratio,window_pages,bf16_bytes,packed_bytes", [
+    (32768, .2, 52, 379_465_736, 171_409_848),
+    (131072, .2, 205, 1_500_745_736, 677_061_048),
+    (1048576, .2, 1639, 11_997_630_472, 5_412_839_864),
+    (1048576, .02, 164, 4_228_132_872, 1_389_134_264),
+    (1048576, .01, 82, 3_796_201_480, 1_165_443_512),
+])
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_target_kv_budget_dimensions(tokens, ratio, window_pages, bf16_bytes, packed_bytes, kv_quant):
+    config = _target_config(tokens, ratio, kv_quant)
+    sizes = _dsv41_pool_sizes(config, tokens // 128 + 1)
+    assert sizes.n_win_pages == window_pages
+    expected_bytes = bf16_bytes if kv_quant == "none" else packed_bytes
+    assert dsv41_pool_bytes(sizes, config.model_config.dsv41_args, 3, kv_quant) == expected_bytes
+    expected_units = (3208, 41152) if kv_quant == "none" else (898, 21312)
+    assert dsv41_unit_bytes(config.model_config.dsv41_args, 128, kv_quant) == expected_units
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_target_long_context_auto_plan_fits_conservative_30_gib_baseline(kv_quant):
+    from freetoken.engine.cache_budget import resolve_moe_cache_auto
+    from freetoken.moe.expert_banks import bank_bytes_estimate
+
+    config = _target_config(1048576, .02, kv_quant)
+    bank_bytes = bank_bytes_estimate(config.model_config)
+    assert bank_bytes == 306_063_605_760
+    a = config.model_config.dsv41_args
+    assert sum(a.engram_num_embeddings) * (a.engram_head_dim + a.engram_head_dim // 32) == 202_758_032_400
+    expert_bytes = bank_bytes // (40 * 384)
+    per_page, fixed, P, floor = dsv41_auto_cost_model(config)
+    weights = 12_190_720_448 + 25_166_848
+    slots, pages, overlap = resolve_moe_cache_auto(
+        baseline_free=30 << 30, weights_bytes=weights, memory_ratio=.9,
+        cache_per_page=per_page, fixed_cache_size=fixed, per_expert_bytes=expert_bytes,
+        num_experts=384, total_experts=40 * 384, prefill_overlap=False,
+        kv_reserve_tokens=max(1048576, floor), page_size=P, max_slots=None,
+    )
+    assert slots >= 384 and pages * P >= 1048576 and not overlap
+    exact_kv = dsv41_pool_bytes(_dsv41_pool_sizes(config, pages + 1), a, 3, kv_quant)
+    assert weights + slots * expert_bytes + exact_kv <= int(.9 * (30 << 30))
+
+
+def _packed_config():
+    return SimpleNamespace(model_config=SimpleNamespace(dsv41_args=args()), max_seq_len=8192,
+                           max_running_req=2, cache_type="swa_radix", swa_full_tokens_ratio=.5,
+                           swa_num_pages_override=None, kv_quant="fp8-fp4", page_size=128,
+                           num_page_override=None, memory_ratio=1.)
+
+
+def test_factory_builds_inline_packed_rows_and_prices_the_dummy_page():
+    from freetoken.kvcache import create_kv_pool
+
+    config = _packed_config()
+    pool = create_kv_pool(config, 31, torch.device("cpu"), torch.bfloat16)
+    assert pool.kv_quant == "fp8-fp4"
+    assert pool.sizes.full_token == 4096
+    assert pool.sizes.n_win_pages == 21
+    assert pool.n_scratch == 3
+    assert pool.total_bytes() == 702_554
+    assert pool.total_bytes() == dsv41_pool_bytes(pool.sizes, pool.args, 3, "fp8-fp4")
+    dummy_full = torch.arange(pool.sizes.full_token - 128, pool.sizes.full_token)
+    dummy_window = torch.arange(pool.sizes.n_win_slots - 128, pool.sizes.n_win_slots)
+    torch.testing.assert_close(pool.translate_full_to_window(dummy_full), dummy_window)
+    assert pool.full_to_window[-1] == -1
+    assert torch.count_nonzero(pool.window_pool[0][dummy_window]) == 0
+
+
+def _byte_rows(width, starts):
+    return ((torch.arange(width)[None, :] + torch.tensor(starts)[:, None]) % 256).to(torch.uint8)
+
+
+def test_inline_codes_and_scales_follow_window_reuse_without_touching_other_request():
+    from freetoken.kvcache import create_kv_pool
+
+    pool = create_kv_pool(_packed_config(), 31, torch.device("cpu"), torch.bfloat16)
+    pool.alloc_swa(torch.arange(256))
+    slots = pool.translate_full_to_window(torch.tensor([0, 128]))
+    first = _byte_rows(33, [200, 127])
+    pool.store_window(first, 0, slots)
+    torch.testing.assert_close(pool.window_pool[0][slots], first)
+    pool.free_swa(torch.arange(128))
+    pool.alloc_swa(torch.arange(256, 384))
+    recycled = pool.translate_full_to_window(torch.tensor([256]))
+    assert recycled[0] == slots[0]
+    replacement = _byte_rows(33, [251])
+    pool.store_window(replacement, 0, recycled)
+    torch.testing.assert_close(pool.window_pool[0][recycled], replacement)
+    torch.testing.assert_close(pool.window_pool[0][slots[1:]], first[1:])
+    assert pool.full_to_window[0] == -1
+
+
+@pytest.mark.parametrize("layer", [1, 3])
+@pytest.mark.parametrize("tier,width", [("compressed", 18), ("indexer", 17)])
+def test_packed_source_rows_and_each_decode_scratch_are_independent(layer, tier, width):
+    from freetoken.kvcache import create_kv_pool
+
+    pool = create_kv_pool(_packed_config(), 31, torch.device("cpu"), torch.bfloat16)
+    backing = pool.cmp_pool[layer] if tier == "compressed" else pool.idx_pool[layer]
+    base = pool.cmp_scratch_base[layer] if tier == "compressed" else pool.idx_scratch_base[layer]
+    writer = getattr(pool, "store_" + tier)
+    main_rows = torch.tensor([0, base - 1])
+    main_values = _byte_rows(width, [17, 201])
+    writer(main_values, layer, main_rows)
+    scratch_rows = base + torch.arange(3)
+    scratch_values = _byte_rows(width, [77, 154, 231])
+    writer(scratch_values, layer, scratch_rows)
+    torch.testing.assert_close(backing[main_rows], main_values)
+    torch.testing.assert_close(backing[scratch_rows], scratch_values)
+    assert torch.count_nonzero(backing[1]) == 0
+    other = 3 if layer == 1 else 1
+    other_backing = pool.cmp_pool[other] if tier == "compressed" else pool.idx_pool[other]
+    assert torch.count_nonzero(other_backing) == 0
+
+
+@pytest.mark.parametrize("method,layer,width", [("store_window", 0, 33),
+                                               ("store_compressed", 1, 18),
+                                               ("store_indexer", 1, 17)])
+@pytest.mark.parametrize("bad", ["dtype", "width", "row_count", "rank"])
+def test_packed_writes_reject_malformed_rows_before_modifying_storage(method, layer, width, bad):
+    a = args()
+    pool = DSV41PagedKVCache(dsv41_pool_sizes(8, a, 1), a, torch.device("cpu"), kv_quant="fp8-fp4")
+    values = _byte_rows(width, [200])
+    if bad == "dtype":
+        values = values.to(torch.bfloat16)
+    elif bad == "width":
+        values = values[:, :-1]
+    elif bad == "row_count":
+        values = values.expand(2, -1)
+    else:
+        values = values.unsqueeze(0)
+    with pytest.raises(ValueError, match="row|uint8"):
+        getattr(pool, method)(values, layer, torch.tensor([0]))
+    assert torch.count_nonzero(pool.window_pool[0]) == 0
+    assert torch.count_nonzero(pool.cmp_pool[1]) == 0
+    assert torch.count_nonzero(pool.idx_pool[1]) == 0
+
+
+@pytest.mark.parametrize("method,width", [("store_compressed", 18), ("store_indexer", 17)])
+def test_reusing_layer_cannot_write_a_source_pool(method, width):
+    a = args()
+    pool = DSV41PagedKVCache(dsv41_pool_sizes(8, a, 1), a, torch.device("cpu"), kv_quant="fp8-fp4")
+    with pytest.raises(ValueError, match="source layer"):
+        getattr(pool, method)(_byte_rows(width, [1]), 2, torch.tensor([0]))
+
+
+def test_packed_rebuild_reallocates_every_inline_row_and_preserves_storage_format():
+    from freetoken.kvcache import create_kv_pool
+    from freetoken.kvcache.base import CacheRebuildRejected
+
+    config = _packed_config()
+    pool = create_kv_pool(config, 31, torch.device("cpu"), torch.bfloat16)
+    table = torch.zeros(3, 8192, dtype=torch.int32)
+    pool.attach_page_table(table)
+    for pages in (63, 31):
+        previous_window, previous_cmp, previous_index = pool.window_pool[0], pool.cmp_pool[1], pool.idx_pool[1]
+        previous_window.fill_(231)
+        previous_cmp.fill_(154)
+        previous_index.fill_(77)
+        pool.rebuild_from_config(config, pages)
+        assert pool.window_pool[0] is not previous_window
+        assert pool.cmp_pool[1] is not previous_cmp
+        assert pool.idx_pool[1] is not previous_index
+        assert pool.full_loc_map is table
+        assert pool.kv_quant == "fp8-fp4" and pool.dtype == torch.bfloat16
+        assert pool.total_bytes() == dsv41_pool_bytes(pool.sizes, pool.args, 3, "fp8-fp4")
+        assert torch.count_nonzero(pool.window_pool[0]) == 0
+        assert torch.count_nonzero(pool.cmp_pool[1]) == 0
+        assert torch.count_nonzero(pool.idx_pool[1]) == 0
+        assert pool.cmp_scratch_base[1] == pool.sizes.full_token // 2
+        assert pool.idx_scratch_base[3] == pool.sizes.full_token
+        dummy = torch.tensor([pool.sizes.full_token - 128])
+        assert pool.translate_full_to_window(dummy).item() == pool.sizes.n_win_slots - 128
+    previous_window = pool.window_pool[0]
+    config.kv_quant = "none"
+    with pytest.raises(CacheRebuildRejected, match="restart"):
+        pool.rebuild_from_config(config, 63)
+    with pytest.raises(CacheRebuildRejected, match="restart"):
+        pool.validate_rebuild(config, num_pages=63, target_moe=0, per_expert_bytes=0,
+                              baseline_free=1 << 30, weights_bytes=0, current_num_pages=31)
+    assert pool.window_pool[0] is previous_window
+    with pytest.raises(AttributeError):
+        pool.kv_quant = "none"
+
+
+@pytest.mark.parametrize("field", ["head_dim", "index_head_dim"])
+def test_packed_layout_rejects_incompatible_channel_blocks(field):
+    a = args()
+    setattr(a, field, 72)
+    with pytest.raises(ValueError, match="divisible by 32"):
+        dsv41_row_bytes(a, "fp8-fp4")
+    with pytest.raises(ValueError, match="divisible by 32"):
+        DSV41PagedKVCache(dsv41_pool_sizes(8, a, 1), a, torch.device("cpu"), kv_quant="fp8-fp4")
+
+
+@pytest.mark.parametrize("kv_quant", ["fp8", "nvfp4", "invalid"])
+def test_pool_rejects_other_storage_formats(kv_quant):
+    a = args()
+    with pytest.raises(ValueError, match="storage format"):
+        DSV41PagedKVCache(dsv41_pool_sizes(8, a, 1), a, torch.device("cpu"), kv_quant=kv_quant)
diff --git a/tests/kvcache/test_kv_cache_rebuild.py b/tests/kvcache/test_kv_cache_rebuild.py
index a5ec94a08..ec5259277 100644
--- a/tests/kvcache/test_kv_cache_rebuild.py
+++ b/tests/kvcache/test_kv_cache_rebuild.py
@@ -64,6 +64,13 @@ def test_mla_and_dsa_rebuild_from_config_and_unit_bytes():
     # the index slab's per-token bytes ride on top of the latent slab's, each floored on its own
     assert dsa.unit_bytes() == (layers * latent * 2 + n_idx * idx_dim * 2, 0)
 
+    fp8 = MLAKVCache(latent_dim=latent, num_layers=layers, num_pages=8, page_size=1,
+                     dtype=torch.bfloat16, device=torch.device("cpu"), kv_quant="fp8")
+    fp8.rebuild_from_config(config=None, num_pages=20)
+    assert fp8.latent_rows(0).dtype is torch.uint8
+    # One byte per latent element plus one FP32 scale for each latent layer.
+    assert fp8.unit_bytes() == (layers * latent + layers * 4, 0)
+
 
 def _hybrid_groups():
     from freetoken.models.config import KVCacheGroupSpec
diff --git a/tests/kvcache/test_mha_pool_fp8.py b/tests/kvcache/test_mha_pool_fp8.py
new file mode 100644
index 000000000..8d74b0edd
--- /dev/null
+++ b/tests/kvcache/test_mha_pool_fp8.py
@@ -0,0 +1,270 @@
+"""The fp8 KV pool: code buffer + per-(token, head) scales, sized and rebuilt together.
+
+The pool-side half of --kv-cache-dtype fp8. The interesting failures are the silent
+ones: a scale buffer that misses the layer_ids remap, a rebuild that resizes the codes
+but not the scales, or a ``unit_bytes`` that drifts from ``kv_cost`` (which is what the
+VRAM budget, the cache sliders and the runtime rebuild all divide by).
+"""
+
+from __future__ import annotations
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.distributed import set_tp_info, try_get_tp_info
+from freetoken.models.config import KVCacheGroupSpec
+
+if not torch.cuda.is_available():  # pragma: no cover
+    pytest.skip("CUDA required", allow_module_level=True)
+
+from freetoken.kernel.triton.kv_quant import kv_codes_dtype
+
+DEV = torch.device("cuda")
+HEADS, DIM, LAYERS, PAGES, PAGE_SIZE = 4, 64, 3, 6, 8
+
+
+def _init_tp() -> None:
+    if try_get_tp_info() is None:
+        set_tp_info(rank=0, size=1)
+
+
+@pytest.fixture(autouse=True)
+def _tp():
+    _init_tp()
+
+
+def _pool(kv_quant="fp8", num_pages=PAGES, layer_ids=None, num_layers=LAYERS):
+    from freetoken.kvcache.mha_pool import MHAKVCache
+
+    _init_tp()
+    return MHAKVCache(
+        num_kv_heads=HEADS,
+        num_layers=num_layers,
+        head_dim=DIM,
+        num_pages=num_pages,
+        page_size=PAGE_SIZE,
+        dtype=torch.bfloat16,
+        device=DEV,
+        layer_ids=layer_ids,
+        kv_quant=kv_quant,
+    )
+
+
+def test_kv_store_dtype_selection():
+    from freetoken.kvcache.mha_pool import _kv_store_dtype
+
+    assert _kv_store_dtype(torch.bfloat16, "none") is torch.bfloat16
+    assert _kv_store_dtype(torch.bfloat16, "fp8") is kv_codes_dtype()
+    with pytest.raises(ValueError, match="kv_quant"):
+        _kv_store_dtype(torch.bfloat16, "q6")
+
+
+def test_fp8_pool_keeps_geometry_and_adds_scale_views():
+    pool = _pool()
+    assert pool.kv_quant == "fp8"
+    # Same shape as the 16-bit pool -- only the element type changed.
+    assert pool._kv_buffer.shape == (2, LAYERS, PAGES, PAGE_SIZE, HEADS, DIM)
+    assert pool._kv_buffer.dtype == kv_codes_dtype()
+    assert pool.store_dtype == kv_codes_dtype()
+    # ``dtype`` stays the COMPUTE dtype (kvcache/base.py). Backends size their scratch
+    # with it -- reporting codes here hands e4m3 to a 16-bit tl.dot, which does not fail
+    # until CUDA-graph capture (exactly how QSA's indexer died in the field).
+    assert pool.dtype == torch.bfloat16
+    slots = PAGES * PAGE_SIZE
+    for layer in range(LAYERS):
+        assert pool.k_scale(layer).shape == (slots, HEADS)
+        assert pool.v_scale(layer).shape == (slots, HEADS)
+        assert pool.k_scale(layer).dtype == torch.float32
+    # A 16-bit pool exposes no scales at all.
+    assert _pool(kv_quant="none").k_scale(0) is None
+
+
+def test_layer_ids_remap_applies_to_scales_too():
+    # Hybrid GDN models back only their full-attention layers; a scale view that
+    # forgot the remap would hand layer 7's rows to layer 2's attention.
+    # (1, 3) names a subset of a model, so layer id 3 has to be inside it: LAYERS is 3, which
+    # makes 3 one past the end, so this pool is told the depth the ids imply.
+    layer_ids = (1, 3)
+    pool = _pool(layer_ids=layer_ids, num_layers=4)
+    assert pool._kv_buffer.shape[1] == 2
+    with pytest.raises(KeyError):
+        pool.k_scale(0)
+def test_store_kv_scatters_codes_and_scales():
+    from freetoken.kernel.triton.kv_quant import codes_to_f32
+
+    tokens = 5
+    torch.manual_seed(7)
+    rows = torch.randn(tokens, HEADS * DIM, device=DEV, dtype=torch.bfloat16) * 2.0
+    out_loc = torch.tensor(
+        [3, 0, PAGES * PAGE_SIZE - 1, 40, 17], device=DEV, dtype=torch.int32
+    )
+    pool = _pool()
+    pool.store_kv(rows, rows.clone(), out_loc, layer_id=0)
+    torch.cuda.synchronize()
+
+    codes = codes_to_f32(pool.k_cache(0).view(-1, HEADS, DIM))
+    scale = pool.k_scale(0)
+    f32 = rows.view(tokens, HEADS, DIM).to(torch.float32)
+    deq = codes[out_loc.long()] * scale[out_loc.long()].unsqueeze(-1)
+    amax = f32.abs().amax(dim=-1, keepdim=True)
+    # The scales are per (token, head), so the bound is relative to each row's max.
+    err = ((deq - f32).abs() / amax.clamp_min(1e-6)).max()
+    assert float(err) < 0.08, float(err)
+    # Rows nobody wrote must stay exactly zero -- the scatter touched only out_loc.
+    untouched = torch.ones(PAGES * PAGE_SIZE, dtype=torch.bool, device=DEV)
+    untouched[out_loc.long()] = False
+    assert (codes[untouched] == 0).all()
+    assert (scale[untouched] == 0).all()
+
+
+def test_rebuild_resizes_codes_and_scales_together():
+    pool = _pool()
+    before = id(pool)
+    pool.rebuild(11)
+    assert id(pool) == before  # identity preserved (backends cache the object)
+    assert pool._kv_buffer.shape == (2, LAYERS, 11, PAGE_SIZE, HEADS, DIM)
+    slots = 11 * PAGE_SIZE
+    assert pool.k_scale(0).shape == (slots, HEADS)
+    assert pool.k_cache(0).shape[0] == 11
+    # Per-token cost is page-count invariant, codes and scales alike.
+    assert pool.unit_bytes() == _pool().unit_bytes()
+
+
+def _sizing_config(kv_quant):
+    from freetoken.attention import AttnType
+
+    spec = KVCacheGroupSpec(
+        name="full",
+        layer_ids=tuple(range(LAYERS)),
+        num_kv_heads=HEADS,
+        head_dim=DIM,
+        sliding_window=None,
+        attn_type=AttnType.FULL,
+    )
+    mc = SimpleNamespace(
+        has_swa_attention=False,
+        has_linear_attention=False,
+        num_layers=LAYERS,
+        num_kv_heads=HEADS,
+        head_dim=DIM,
+        kv_cache_group_specs=lambda: (spec,),
+    )
+    return SimpleNamespace(
+        model_config=mc,
+        page_size=PAGE_SIZE,
+        dtype=torch.bfloat16,
+        tp_info=SimpleNamespace(size=1),
+        kv_quant=kv_quant,
+    )
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8"])
+def test_unit_bytes_matches_the_cost_model_that_sized_the_pool(kv_quant):
+    """The budget solve and the live pool must agree byte for byte."""
+    from freetoken.kvcache.base import spec_kv_bytes_per_token
+    from freetoken.kvcache.mha_pool import MHAKVCache
+
+    config = _sizing_config(kv_quant)
+    (spec,) = config.model_config.kv_cache_group_specs()
+    per_token = spec_kv_bytes_per_token(spec, config)
+    assert MHAKVCache.kv_cost(config)[0] == per_token * PAGE_SIZE
+
+    pool = _pool(kv_quant=kv_quant)
+    assert pool.unit_bytes() == (per_token, 0)
+
+
+def test_fp8_lands_just_above_half_the_bytes():
+    """Half the code bytes, plus a scale sidecar of 4 B per (token, slab, layer, head)."""
+    plain, quantized = _pool("none").unit_bytes()[0], _pool("fp8").unit_bytes()[0]
+    scales = 2 * LAYERS * HEADS * 4
+    assert quantized == plain // 2 + scales, (plain, quantized, scales)
+
+
+def test_latent_kv_pool_accepts_fp8():
+    """MLA carries one latent row per token, so it uses one FP32 scale per row."""
+    from freetoken.attention import AttnType
+    from freetoken.kvcache import create_kvcache_pool
+
+    spec = KVCacheGroupSpec(
+        name="full",
+        layer_ids=(0, 1),
+        num_kv_heads=HEADS,
+        head_dim=DIM,
+        sliding_window=None,
+        mla=True,
+        attn_type=AttnType.MLA,
+    )
+    mc = SimpleNamespace(
+        has_swa_attention=False,
+        has_linear_attention=False,
+        num_layers=2,
+        num_kv_heads=HEADS,
+        head_dim=DIM,
+        kv_cache_group_specs=lambda: (spec,),
+    )
+    fp8 = create_kvcache_pool(
+        model_config=mc,
+        num_pages=4,
+        page_size=1,
+        dtype=torch.bfloat16,
+        device=DEV,
+        kv_quant="fp8",
+    )
+    assert fp8.kv_quant == "fp8" and fp8.store_dtype is torch.uint8
+    pool = create_kvcache_pool(
+        model_config=mc,
+        num_pages=4,
+        page_size=1,
+        dtype=torch.bfloat16,
+        device=DEV,
+        kv_quant="none",
+    )
+    assert pool.kv_quant == "none"
+
+
+def test_hybrid_swa_pool_also_separates_compute_and_store_dtype(monkeypatch):
+    """The hybrid-SWA pool is the other family that stores codes, so it must report the
+    same pair. Any backend sizing scratch off ``dtype`` would take e4m3 home with it
+    here too -- what that looked like in practice is QSA's indexer (kvcache/base.py)."""
+    from freetoken.distributed.info import DistributedInfo
+    from freetoken.kvcache.hybrid_swa_pool import HybridSWAKVCache
+    from freetoken.models.config import KVCacheGroupSpec
+
+    monkeypatch.setattr(
+        "freetoken.kvcache.hybrid_swa_pool.get_tp_info",
+        lambda: DistributedInfo(rank=0, size=1),
+    )
+    groups = (
+        KVCacheGroupSpec(
+            name="full", layer_ids=(2, 5), num_kv_heads=2, head_dim=DIM,
+            sliding_window=None,
+        ),
+        KVCacheGroupSpec(
+            name="swa", layer_ids=(0, 1, 3, 4), num_kv_heads=2, head_dim=DIM,
+            sliding_window=32,
+        ),
+    )
+
+    def build(kv_quant: str):
+        return HybridSWAKVCache(
+            groups=groups,
+            num_layers=6,
+            num_full_pages=PAGES,
+            page_size=PAGE_SIZE,
+            num_swa_tokens=PAGES * PAGE_SIZE,
+            dtype=torch.bfloat16,
+            device=DEV,
+            kv_quant=kv_quant,
+        )
+
+    quantized, plain = build("fp8"), build("none")
+    assert quantized.dtype is torch.bfloat16
+    assert plain.dtype is torch.bfloat16
+    assert quantized.store_dtype == kv_codes_dtype()
+    assert plain.store_dtype is torch.bfloat16
+    # ...while the buffers really did shrink: the two properties must not be aliases.
+    assert quantized.k_cache(2).element_size() == 1
+    assert plain.k_cache(2).element_size() == 2
+    assert quantized.k_scale(2) is not None and plain.k_scale(2) is None
diff --git a/tests/kvcache/test_qsa_pool.py b/tests/kvcache/test_qsa_pool.py
index dea722f63..327200cfb 100644
--- a/tests/kvcache/test_qsa_pool.py
+++ b/tests/kvcache/test_qsa_pool.py
@@ -37,7 +37,8 @@ def _tp(monkeypatch):
     )
 
 
-def _pool(num_pages=4, page_size=64, index_ratio=4, num_req_slots=4, ring_capacity=None):
+def _pool(num_pages=4, page_size=64, index_ratio=4, num_req_slots=4, ring_capacity=None,
+          kv_quant="none", mrope=False):
     return QSAKVCache(
         num_kv_heads=2,
         num_layers=8,
@@ -52,6 +53,8 @@ def _pool(num_pages=4, page_size=64, index_ratio=4, num_req_slots=4, ring_capaci
         num_req_slots=num_req_slots,
         ring_capacity=ring_capacity,
         layer_ids=(1, 3, 5, 7),
+        kv_quant=kv_quant,
+        mrope=mrope,
     )
 
 
@@ -185,23 +188,42 @@ def test_kv_cost_prices_ring_and_scratch_as_fixed():
     assert fixed == 4 * row * (QSAKVCache.ring_capacity_for(4) + 1)
 
 
-def test_unit_bytes_matches_the_cost_model():
+@pytest.mark.parametrize("kv_quant", ["none", "fp8", "nvfp4"])
+@pytest.mark.parametrize("mrope", [False, True])
+def test_unit_bytes_matches_the_cost_model(kv_quant, mrope):
     spec = _spec()
     config = _config(spec)
-    pool = _pool()
+    config.kv_quant = kv_quant
+    config.model_config.model_is_mrope = mrope
+    pool = _pool(kv_quant=kv_quant, mrope=mrope)
     kv_bytes, swa_bytes = pool.unit_bytes()
     assert swa_bytes == 0
     # the scratch rows and the ring must NOT inflate the per-token slider
-    assert kv_bytes == spec_kv_bytes_per_token(spec, config)
+    assert kv_bytes == spec_kv_bytes_per_token(spec, config) + (12 if mrope else 0)
     assert kv_bytes * 64 == QSAKVCache.kv_cost(config)[0]
-
-
-def test_resolve_pool_class_and_factory():
+    for num_pages in (4, 9):
+        if num_pages != 4:
+            pool.rebuild(num_pages)
+        per_page, fixed, _, _ = QSAKVCache.kv_cost(config)
+        buffers = [pool._kv_buffer, pool._scale_buffer, pool._block_scale_buffer,
+                   pool._cmp_k_buffer, pool._pending_ring, pool._rope_positions]
+        actual_bytes = sum(t.numel() * t.element_size() for t in buffers if t is not None)
+        assert actual_bytes == per_page * num_pages + fixed
+        assert pool.unit_bytes() == (kv_bytes, 0)
+        assert pool.cmp_k_cache(0).dtype == pool.pending_ring(0).dtype == torch.bfloat16
+        if mrope:
+            assert pool.rope_positions.shape == (num_pages * 64, 3)
+            assert pool.rope_positions.dtype == torch.int32
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "nvfp4"])
+@pytest.mark.parametrize("mrope", [False, True])
+def test_resolve_pool_class_and_factory(kv_quant, mrope):
     from freetoken.kvcache import create_kvcache_pool, resolve_pool_class
 
     spec = _spec()
     mc = SimpleNamespace(
-        model_is_mrope=False,
+        model_is_mrope=mrope,
         num_layers=8, has_swa_attention=False, has_linear_attention=True,
         num_kv_heads=2, head_dim=64, dsv4_args=None,
     )
@@ -209,11 +231,18 @@ def test_resolve_pool_class_and_factory():
     assert resolve_pool_class(mc) is QSAKVCache
 
     pool = create_kvcache_pool(
-        mc, num_pages=4, page_size=64, dtype=torch.bfloat16, device=DEV, num_req_slots=4
+        mc, num_pages=4, page_size=64, dtype=torch.bfloat16, device=DEV, num_req_slots=4,
+        kv_quant=kv_quant,
     )
     assert isinstance(pool, QSAKVCache)
     assert pool._kv_buffer.shape[1] == 4  # not the model's 8 layers
     assert pool.cmp_k_cache(0).shape == (4 * 64 // 4 + 4, 32)
+    assert pool.kv_quant == kv_quant
+    assert (pool._rope_positions is not None) == mrope
+    if kv_quant == "nvfp4":
+        assert pool.k_cache(1).dtype == torch.uint8
+        assert pool.k_cache(1).shape[-1] == 32
+        assert pool.k_block_scale(1).shape[-1] == 4
 
     with pytest.raises(ValueError, match="num_req_slots"):
         create_kvcache_pool(mc, num_pages=4, page_size=64, dtype=torch.bfloat16, device=DEV)
diff --git a/tests/kvcache/test_qsa_pool_fp8.py b/tests/kvcache/test_qsa_pool_fp8.py
new file mode 100644
index 000000000..c0a91f862
--- /dev/null
+++ b/tests/kvcache/test_qsa_pool_fp8.py
@@ -0,0 +1,221 @@
+"""The QSA pool under ``--kv-cache-dtype fp8``: quantized K/V, 16-bit index tiers.
+
+Only the paged K/V slabs change. The compressed index slab, the per-request ring and
+their scratch rows must stay 16-bit whatever the KV store does -- block selection scores
+against them and the score kernel asserts their dtype -- while the byte account the
+startup budget, the cache sliders and the runtime rebuild all divide by has to keep
+telling the two halves apart. A drift here does not crash: it silently buys the wrong
+number of pages.
+"""
+
+from __future__ import annotations
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.attention import AttnType
+from freetoken.kvcache.base import spec_kv_bytes_per_token
+from freetoken.models.config import KVCacheGroupSpec
+
+if not torch.cuda.is_available():  # pragma: no cover
+    pytest.skip("CUDA required", allow_module_level=True)
+
+from freetoken.kernel.triton.kv_quant import codes_to_f32, kv_codes_dtype
+from freetoken.kvcache.qsa_pool import QSAKVCache
+
+DEV = torch.device("cuda")
+PAGE_SIZE = 64  # the page size the qsa_sparse backend registers
+LAYER_IDS = (1, 3, 5, 7)
+HEADS, DIM, INDEX_DIM, INDEX_LAYERS, RATIO = 2, 64, 32, 4, 4
+
+
+@pytest.fixture(autouse=True)
+def _tp(monkeypatch):
+    from freetoken.distributed.info import DistributedInfo
+
+    monkeypatch.setattr(
+        "freetoken.kvcache.mha_pool.get_tp_info",
+        lambda: DistributedInfo(rank=0, size=1),
+    )
+
+
+def _pool(kv_quant="fp8", num_pages=4, num_req_slots=4):
+    return QSAKVCache(
+        num_kv_heads=HEADS,
+        num_layers=8,
+        head_dim=DIM,
+        num_pages=num_pages,
+        page_size=PAGE_SIZE,
+        dtype=torch.bfloat16,
+        device=DEV,
+        index_head_dim=INDEX_DIM,
+        num_index_layers=INDEX_LAYERS,
+        index_ratio=RATIO,
+        num_req_slots=num_req_slots,
+        layer_ids=LAYER_IDS,
+        kv_quant=kv_quant,
+    )
+
+
+def _spec():
+    return KVCacheGroupSpec(
+        name="full",
+        layer_ids=LAYER_IDS,
+        num_kv_heads=HEADS,
+        head_dim=DIM,
+        sliding_window=None,
+        index_head_dim=INDEX_DIM,
+        num_index_layers=INDEX_LAYERS,
+        index_ratio=RATIO,
+        attn_type=AttnType.QSA,
+    )
+
+
+def _config(kv_quant, *, page_size=PAGE_SIZE, max_running_req=3):
+    mc = SimpleNamespace(num_layers=8, has_swa_attention=False, has_linear_attention=True)
+    mc.kv_cache_group_specs = lambda: (_spec(),)
+    return SimpleNamespace(
+        model_config=mc,
+        page_size=page_size,
+        dtype=torch.bfloat16,
+        tp_info=SimpleNamespace(size=1),
+        max_running_req=max_running_req,
+        kv_quant=kv_quant,
+    )
+
+
+def test_fp8_replaces_the_kv_slab_and_adds_scale_views():
+    pool = _pool()
+    bf16 = _pool(kv_quant="none")
+    assert pool.kv_quant == "fp8"
+    # Same geometry as the 16-bit pool -- only the element type changed, because the
+    # attend kernels index codes exactly the way they index values.
+    assert pool._kv_buffer.shape == bf16._kv_buffer.shape
+    assert pool._kv_buffer.dtype == kv_codes_dtype()
+    # dtype = compute dtype (what the backend sizes its INDEXER scratch with),
+    # store_dtype = what the buffer holds. Swapping them is the bug that compiled an
+    # e4m3 operand into QSA's scoring dot and died at graph capture.
+    assert pool.dtype is torch.bfloat16
+    assert pool.store_dtype == kv_codes_dtype()
+    assert bf16.dtype is torch.bfloat16 and bf16.store_dtype is torch.bfloat16
+    assert pool.k_cache(3).shape == (4, PAGE_SIZE, HEADS, DIM)
+    assert pool.k_cache(3).element_size() == 1
+    slots = 4 * PAGE_SIZE
+    assert pool.k_scale(3).shape == (slots, HEADS)
+    assert pool.k_scale(3).dtype is torch.float32
+    assert pool.v_scale(3).shape == (slots, HEADS)
+    # Zero-filled: e4m3 has NaN bit patterns, and the dummy page / unwritten tail rows
+    # must never read back as one.
+    assert pool.k_scale(3).abs().sum().item() == 0.0
+    assert codes_to_f32(pool.k_cache(3)).abs().sum().item() == 0.0
+    # A 16-bit pool keeps answering None, so the backends' k_scale(...) pass-through is
+    # the only branch that ever differs between the two.
+    assert bf16.k_scale(3) is None and bf16.v_scale(3) is None
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8", "nvfp4"])
+def test_index_tiers_stay_16_bit_whatever_the_kv_store_does(kv_quant):
+    """Block selection is quantization-agnostic by construction: it reads the compressed
+    index keys, not the KV rows, so fp8 must not touch these three buffers."""
+    pool = _pool(kv_quant=kv_quant)
+    assert pool.cmp_k_cache(0).dtype is torch.bfloat16
+    assert pool.pending_ring(0).dtype is torch.bfloat16
+    assert pool.cmp_k_cache(0).shape == (4 * PAGE_SIZE // RATIO + 4, INDEX_DIM)
+    # ...and the byte account says so too: only the 1-byte KV codes got cheaper.
+    spec = _spec()
+    cost = spec_kv_bytes_per_token(spec, _config(kv_quant))
+    plain = spec_kv_bytes_per_token(spec, _config("none"))
+    kv_layers = len(LAYER_IDS)
+    kv_16bit = 2 * DIM * HEADS * 2 * kv_layers          # two slabs, 16-bit codes
+    scale_term = 2 * kv_layers * HEADS * 4              # one fp32 scale per (slot, head)
+    index_term = INDEX_DIM * INDEX_LAYERS * 2 // RATIO  # the untouched 16-bit slab
+    assert plain == kv_16bit + index_term
+    if kv_quant == "fp8":
+        assert cost == kv_16bit // 2 + scale_term + index_term
+    elif kv_quant == "nvfp4":
+        block_term = 2 * kv_layers * HEADS * (DIM // 16)
+        assert cost == kv_16bit // 4 + scale_term + block_term + index_term
+    else:
+        assert cost == plain
+
+
+def test_unit_bytes_and_kv_cost_still_agree_when_quantized():
+    """The pool's own allocation and the budget model must divide the same way -- the
+    scale sidecar is priced in base.spec_kv_bytes_per_token, not here."""
+    spec, config = _spec(), _config("fp8")
+    pool = _pool()
+    kv_bytes, swa_bytes = pool.unit_bytes()
+    assert swa_bytes == 0
+    assert kv_bytes == spec_kv_bytes_per_token(spec, config)
+    assert kv_bytes * PAGE_SIZE == QSAKVCache.kv_cost(config)[0]
+    # The 16-bit pool's per-token figure is the reference: codes halve the KV term, the
+    # scale sidecar and the untouched index slab keep the total above a clean half.
+    plain_kv = spec_kv_bytes_per_token(spec, _config("none"))
+    assert plain_kv // 2 < kv_bytes < plain_kv
+
+
+def test_rebuild_resizes_codes_and_scales_together():
+    pool = _pool(num_pages=4)
+    ident = id(pool)
+    pool.rebuild(16)
+    assert id(pool) == ident
+    assert pool.k_cache(1).shape == (16, PAGE_SIZE, HEADS, DIM)
+    assert pool.k_scale(1).shape == (16 * PAGE_SIZE, HEADS)
+    assert pool.k_scale(1).abs().sum().item() == 0.0
+    assert pool.cmp_k_cache(0).shape == (16 * PAGE_SIZE // RATIO + 4, INDEX_DIM)
+
+
+def test_store_kv_writes_the_slot_the_attend_kernel_will_read():
+    """out_loc numbering (page * page_size + offset) is the contract between the fused
+    writer and the attend kernel's scale slot arithmetic -- this is that round trip."""
+    torch.manual_seed(0)
+    pool = _pool(num_pages=5)
+    slots = 5 * PAGE_SIZE
+    rows = (0, 1, 63, 64, 255, 256)  # page boundaries included: 63/64 and 255/256
+    k = torch.randn(len(rows), HEADS * DIM, device=DEV, dtype=torch.bfloat16) * 3.0
+    v = torch.randn(len(rows), HEADS * DIM, device=DEV, dtype=torch.bfloat16) * 0.25
+    out_loc = torch.tensor(rows, dtype=torch.int32, device=DEV)
+
+    pool.store_kv(k, v, out_loc, layer_id=3)
+    torch.cuda.synchronize()
+
+    codes = codes_to_f32(pool.k_cache(3).view(slots, HEADS, DIM))
+    decoded_k = codes[out_loc] * pool.k_scale(3)[out_loc].unsqueeze(-1)
+    decoded_v = codes_to_f32(pool.v_cache(3).view(slots, HEADS, DIM))[out_loc] * pool.v_scale(3)[
+        out_loc
+    ].unsqueeze(-1)
+    for want, got in ((k, decoded_k), (v, decoded_v)):
+        w = want.view(len(rows), HEADS, DIM).to(torch.float32)
+        amax = w.abs().amax(dim=-1, keepdim=True)
+        assert (((got - w).abs() / amax).max().item()) < 0.07
+    # Rows nobody wrote stay zero rather than NaN -- the dummy page depends on it.
+    untouched = torch.tensor([r for r in range(64) if r not in rows], device=DEV)
+    assert pool.k_scale(3)[untouched].abs().sum().item() == 0.0
+
+
+def test_factory_threads_kv_quant_into_the_qsa_pool():
+    from freetoken.kvcache import create_kvcache_pool
+
+    mc = SimpleNamespace(
+        num_layers=8, has_swa_attention=False, has_linear_attention=True,
+        num_kv_heads=HEADS, head_dim=DIM, dsv4_args=None,
+    )
+    mc.kv_cache_group_specs = lambda: (_spec(),)
+    pool = create_kvcache_pool(
+        mc, num_pages=4, page_size=PAGE_SIZE, dtype=torch.bfloat16, device=DEV,
+        num_req_slots=4, kv_quant="fp8",
+    )
+    assert isinstance(pool, QSAKVCache) and pool.kv_quant == "fp8"
+    assert pool.k_cache(1).element_size() == 1 and pool.k_scale(1) is not None
+
+
+def test_nvfp4_replaces_only_qsa_kv_tiers():
+    pool = _pool(kv_quant="nvfp4")
+    slots = 4 * PAGE_SIZE
+    assert pool.k_cache(3).shape == (4, PAGE_SIZE, HEADS, DIM // 2)
+    assert pool.k_block_scale(3).shape == (slots, HEADS, DIM // 16)
+    assert pool.v_block_scale(3).dtype is torch.uint8
+    assert pool.cmp_k_cache(0).dtype is torch.bfloat16
+
diff --git a/tests/models/fixtures/deepseek_v41_libertai_nvfp4_config.json b/tests/models/fixtures/deepseek_v41_libertai_nvfp4_config.json
new file mode 100644
index 000000000..468e498a0
--- /dev/null
+++ b/tests/models/fixtures/deepseek_v41_libertai_nvfp4_config.json
@@ -0,0 +1,177 @@
+{
+  "architectures": [
+    "DeepseekV41ForCausalLM"
+  ],
+  "model_type": "deepseek_v41",
+  "dtype": "bfloat16",
+  "transformers_version": "5.6.0",
+  "bos_token_id": 0,
+  "eos_token_id": 1,
+  "pad_token_id": 2,
+  "image_token_id": 129264,
+  "quantization_config": {
+    "quant_method": "fp8",
+    "activation_scheme": "dynamic",
+    "weight_block_size": [
+      32,
+      32
+    ],
+    "scale_fmt": "ue8m0",
+    "expert_dtype": "nvfp4",
+    "expert_block_size": 16,
+    "expert_scale_fmt": "e4m3",
+    "expert_global_scale": true,
+    "engram_dtype": "fp4",
+    "engram_block_size": 32,
+    "engram_scale_fmt": "ue8m0",
+    "repacked_by": "LibertAI/dsv41_fp4_stream.py"
+  },
+  "text_config": {
+    "model_type": "deepseek_v41_text",
+    "vocab_size": 129280,
+    "hidden_size": 5120,
+    "moe_intermediate_size": 2304,
+    "num_hidden_layers": 40,
+    "num_attention_heads": 64,
+    "num_key_value_heads": 1,
+    "head_dim": 512,
+    "qk_rope_head_dim": 64,
+    "q_lora_rank": 1280,
+    "o_lora_rank": 1024,
+    "o_groups": 8,
+    "hidden_act": "silu",
+    "swiglu_limit": 10.0,
+    "rms_norm_eps": 1e-20,
+    "attention_bias": false,
+    "attention_dropout": 0.0,
+    "initializer_range": 0.02,
+    "use_cache": true,
+    "tie_word_embeddings": false,
+    "max_position_embeddings": 1048576,
+    "rope_theta": 10000,
+    "rope_scaling": {
+      "rope_type": "yarn",
+      "factor": 16,
+      "beta_fast": 32,
+      "beta_slow": 1,
+      "original_max_position_embeddings": 65536
+    },
+    "n_routed_experts": 384,
+    "n_shared_experts": 1,
+    "num_experts_per_tok": 6,
+    "scoring_func": "sqrtsoftplus",
+    "topk_method": "noaux_tc",
+    "norm_topk_prob": true,
+    "routed_scaling_factor": 1.5,
+    "sliding_window": 128,
+    "compress_ratios": [
+      0,
+      0,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      0,
+      0,
+      0
+    ],
+    "compress_rope_theta": 160000,
+    "kv_source_layer_ids": [
+      2,
+      8,
+      14,
+      20
+    ],
+    "index_source_layer_ids": [
+      2,
+      8,
+      14,
+      20,
+      24,
+      28,
+      32,
+      36
+    ],
+    "index_n_heads": 32,
+    "index_head_dim": 128,
+    "index_topk": 512,
+    "candidate_source_layer_id": 20,
+    "candidate_topk_blocks": 2048,
+    "candidate_block_size": 8,
+    "hc_mult": 4,
+    "hc_sinkhorn_iters": 20,
+    "hc_eps": 1e-06,
+    "engram_layer_ids": [
+      1,
+      14
+    ],
+    "engram_num_embeddings": [
+      384006168,
+      384016682
+    ],
+    "engram_max_ngram_size": 4,
+    "engram_vocab_size": 16000000,
+    "engram_n_heads": 8,
+    "engram_head_dim": 256,
+    "engram_pad_token_id": 2,
+    "engram_compressed_vocab_size": 99092,
+    "num_nextn_predict_layers": 3,
+    "dspark_block_size": 5,
+    "dspark_noise_token_id": 128799,
+    "dspark_target_layer_ids": [
+      37,
+      38,
+      39
+    ],
+    "dspark_markov_rank": 256,
+    "dspark_n_routed_experts": 128,
+    "dspark_num_experts_per_tok": 3
+  },
+  "vision_config": {
+    "model_type": "deepseek_v41_vision",
+    "num_hidden_layers": 32,
+    "hidden_size": 1024,
+    "num_attention_heads": 16,
+    "intermediate_size": 2816,
+    "patch_size": 14,
+    "rope_theta": 10000,
+    "downsample_ratio": 3,
+    "max_image_tokens": 1024,
+    "min_pixels": 295936,
+    "max_wh_ratio": null
+  }
+}
\ No newline at end of file
diff --git a/tests/models/fixtures/deepseek_v41_nvfp4_config.json b/tests/models/fixtures/deepseek_v41_nvfp4_config.json
new file mode 100644
index 000000000..b42385da3
--- /dev/null
+++ b/tests/models/fixtures/deepseek_v41_nvfp4_config.json
@@ -0,0 +1,365 @@
+{
+  "architectures": [
+    "DeepseekV41ForCausalLM"
+  ],
+  "bos_token_id": 0,
+  "dtype": "bfloat16",
+  "eos_token_id": 1,
+  "image_token_id": 129264,
+  "model_type": "deepseek_v41",
+  "pad_token_id": 2,
+  "quantization_config": {
+    "activation_scheme": "dynamic",
+    "config_groups": {
+      "group_0": {
+        "input_activations": {
+          "dynamic": false,
+          "group_size": 16,
+          "num_bits": 4,
+          "type": "float"
+        },
+        "targets": [
+          "Linear"
+        ],
+        "weights": {
+          "dynamic": false,
+          "group_size": 16,
+          "num_bits": 4,
+          "type": "float"
+        }
+      }
+    },
+    "expert_dtype": "fp4",
+    "group_size": 16,
+    "ignore": [
+      "*.attn.*",
+      "*.ffn.shared_experts.*",
+      "head",
+      "mtp.*"
+    ],
+    "kv_cache_quant_algo": null,
+    "moe_quant_algo": "NVFP4",
+    "producer": {
+      "name": "modelopt",
+      "version": "dsv4-nvfp4-experts"
+    },
+    "quant_algo": "MIXED_PRECISION",
+    "quant_method": "fp8",
+    "quantized_layers": {
+      "layers.0.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.1.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.10.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.11.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.12.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.13.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.14.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.15.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.16.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.17.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.18.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.19.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.2.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.20.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.21.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.22.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.23.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.24.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.25.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.26.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.27.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.28.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.29.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.3.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.30.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.31.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.32.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.33.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.34.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.35.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.36.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.37.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.38.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.39.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.4.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.5.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.6.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.7.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.8.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      },
+      "layers.9.ffn.experts": {
+        "group_size": 16,
+        "quant_algo": "NVFP4"
+      }
+    },
+    "scale_fmt": "ue8m0",
+    "weight_block_size": [
+      32,
+      32
+    ]
+  },
+  "text_config": {
+    "attention_bias": false,
+    "attention_dropout": 0.0,
+    "candidate_block_size": 8,
+    "candidate_source_layer_id": 20,
+    "candidate_topk_blocks": 2048,
+    "compress_ratios": [
+      0,
+      0,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      2,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      1,
+      0,
+      0,
+      0
+    ],
+    "compress_rope_theta": 160000,
+    "dspark_block_size": 5,
+    "dspark_markov_rank": 256,
+    "dspark_n_routed_experts": 128,
+    "dspark_noise_token_id": 128799,
+    "dspark_num_experts_per_tok": 3,
+    "dspark_target_layer_ids": [
+      37,
+      38,
+      39
+    ],
+    "engram_compressed_vocab_size": 99092,
+    "engram_head_dim": 256,
+    "engram_layer_ids": [
+      1,
+      14
+    ],
+    "engram_max_ngram_size": 4,
+    "engram_n_heads": 8,
+    "engram_num_embeddings": [
+      384006168,
+      384016682
+    ],
+    "engram_pad_token_id": 2,
+    "engram_vocab_size": 16000000,
+    "hc_eps": 1e-06,
+    "hc_mult": 4,
+    "hc_sinkhorn_iters": 20,
+    "head_dim": 512,
+    "hidden_act": "silu",
+    "hidden_size": 5120,
+    "index_head_dim": 128,
+    "index_n_heads": 32,
+    "index_source_layer_ids": [
+      2,
+      8,
+      14,
+      20,
+      24,
+      28,
+      32,
+      36
+    ],
+    "index_topk": 512,
+    "initializer_range": 0.02,
+    "kv_source_layer_ids": [
+      2,
+      8,
+      14,
+      20
+    ],
+    "max_position_embeddings": 1048576,
+    "model_type": "deepseek_v41_text",
+    "moe_intermediate_size": 2304,
+    "n_routed_experts": 384,
+    "n_shared_experts": 1,
+    "norm_topk_prob": true,
+    "num_attention_heads": 64,
+    "num_experts_per_tok": 6,
+    "num_hidden_layers": 40,
+    "num_key_value_heads": 1,
+    "num_nextn_predict_layers": 3,
+    "o_groups": 8,
+    "o_lora_rank": 1024,
+    "q_lora_rank": 1280,
+    "qk_rope_head_dim": 64,
+    "rms_norm_eps": 1e-20,
+    "rope_scaling": {
+      "beta_fast": 32,
+      "beta_slow": 1,
+      "factor": 16,
+      "original_max_position_embeddings": 65536,
+      "rope_type": "yarn"
+    },
+    "rope_theta": 10000,
+    "routed_scaling_factor": 1.5,
+    "scoring_func": "sqrtsoftplus",
+    "sliding_window": 128,
+    "swiglu_limit": 10.0,
+    "tie_word_embeddings": false,
+    "topk_method": "noaux_tc",
+    "use_cache": true,
+    "vocab_size": 129280
+  },
+  "transformers_version": "5.6.0",
+  "vision_config": {
+    "downsample_ratio": 3,
+    "hidden_size": 1024,
+    "intermediate_size": 2816,
+    "max_image_tokens": 1024,
+    "max_wh_ratio": null,
+    "min_pixels": 295936,
+    "model_type": "deepseek_v41_vision",
+    "num_attention_heads": 16,
+    "num_hidden_layers": 32,
+    "patch_size": 14,
+    "rope_theta": 10000
+  }
+}
diff --git a/tests/models/qwen4_exp/common.py b/tests/models/qwen4_exp/common.py
index 0a52fb0e0..f3bddd7a9 100644
--- a/tests/models/qwen4_exp/common.py
+++ b/tests/models/qwen4_exp/common.py
@@ -154,6 +154,7 @@ def __init__(
         device: str = "cuda",
         dtype: torch.dtype = torch.bfloat16,
         page_size: int = 64,
+        kv_quant: str = "none",
     ) -> None:
         from freetoken.attention.qsa_sparse import QSASparseAttnBackend
         from freetoken.kvcache import create_kvcache_pool
@@ -170,6 +171,7 @@ def __init__(
             dtype=dtype,
             device=self.device,
             num_req_slots=self.num_req_slots,
+            kv_quant=kv_quant,
         )
         self.page_table = torch.zeros(
             (self.num_req_slots, num_pages * page_size), dtype=torch.int32, device=self.device
diff --git a/tests/models/qwen4_exp/test_config.py b/tests/models/qwen4_exp/test_config.py
index be532eeca..0358239f8 100644
--- a/tests/models/qwen4_exp/test_config.py
+++ b/tests/models/qwen4_exp/test_config.py
@@ -182,6 +182,20 @@ def test_vision_turns_on_mrope_and_the_tower():
 def test_text_only_keeps_the_1d_rope():
     config = parse_config(_hf_config())
     assert not config.is_multimodal and not config.model_is_mrope and config.rotary_config.mrope_section is None
+
+
+def test_released_yarn_config_preserves_scaling_when_vision_adds_mrope():
+    hf = _hf_config()
+    hf.text_config.max_position_embeddings = 1048576
+    hf.text_config.rope_parameters.update(rope_type="yarn", factor=4.0,
+                                          original_max_position_embeddings=262144)
+    plain = parse_config(hf).rotary_config
+    hf.vision_config = _vision_config()
+    multimodal = parse_config(hf).rotary_config
+    assert plain.scaling == multimodal.scaling
+    assert multimodal.scaling["rope_type"] == "yarn" and multimodal.scaling["factor"] == 4.0
+    assert multimodal.rotary_dim == 64 and multimodal.max_position == 1048576
+    assert multimodal.mrope_section == [11, 11, 10] and multimodal.mrope_layout == "interleaved"
 # the merged-projection prefixes the model asks the QuantConfig about (attention.py / gdn.py)
 DENSE_PREFIXES = (
     "model.layers.3.self_attn.qkv_proj", "model.layers.3.self_attn.o_proj",
diff --git a/tests/models/qwen4_exp/test_qsa_backend.py b/tests/models/qwen4_exp/test_qsa_backend.py
index 236610791..cb819508a 100644
--- a/tests/models/qwen4_exp/test_qsa_backend.py
+++ b/tests/models/qwen4_exp/test_qsa_backend.py
@@ -5,17 +5,21 @@
     exactly the causal prefix and the layer output must match ``TorchDenseQSAReference`` (fp32)
     and a flashinfer dense run over the same pool;
 (b) chunked prefill at unaligned cut points equals one-shot prefill (the dual-source compress);
-(c) a captured decode replay equals the eager decode step.
+(c) a captured decode replay equals the eager decode step;
+(d) an fp8 KV pool (``--kv-cache-dtype fp8``) keeps block selection bit-identical to the
+    16-bit run -- only the selected K/V rows are read back as e4m3 codes -- and the layer
+    output stays within quantization error of it.
 """
 
 from __future__ import annotations
 
+import math
 from types import SimpleNamespace
 
 import pytest
 import torch
 
-from .common import Fixture, requires_cuda, parsed_config, selection_spy
+from .common import Fixture, hf_config, requires_cuda, parsed_config, selection_spy
 
 QSA_LAYER = 3
 
@@ -249,3 +253,182 @@ def test_two_qsa_layers_keep_separate_slab_slots(monkeypatch):
 
     slab = fixture.pool.cmp_k_cache
     assert not torch.equal(slab(0), slab(1))
+
+
+def _prefill_under_kv(monkeypatch, config, kv_quant: str, lengths):
+    """One prefill of the QSA layer under a given KV store.
+
+    Each call builds its own Fixture on purpose: a Fixture owns the global ctx (pool,
+    page table, backend), so two KV stores cannot share one scenario. The weight seed
+    (``Fixture.layer``) and the input seed (``_inputs``) are fixed, so the two runs differ
+    ONLY in how the K/V rows are stored.
+    """
+    fixture = Fixture(config, num_pages=128, kv_quant=kv_quant)
+    attn = fixture.layer(QSA_LAYER)
+    seen = selection_spy(monkeypatch, fixture.backend)
+    inputs = _inputs(fixture, lengths)
+    x = torch.cat([row[:n] for row, n in zip(inputs, lengths)])
+    reqs = [fixture.req(i, 0, n) for i, n in enumerate(lengths)]
+    batch = fixture.batch(reqs, "prefill")
+    out = attn.forward(x, batch)
+    # the selection lives in a scratch buffer the next forward overwrites
+    return fixture, out.clone(), seen["indices"].clone(), batch.positions.clone()
+
+
+@requires_cuda
+def test_fp8_kv_pool_keeps_selection_and_output(monkeypatch):
+    """--kv-cache-dtype fp8 through the real layer: e4m3 codes + per-row scales in, same
+    answer out to within quantization error -- and, because block selection scores 16-bit
+    compressed index keys that fp8 never touches, the SAME selection bit for bit."""
+    config = parsed_config()
+    lengths = [2051, 1000, 137]  # every complete block is selected here
+
+    plain, plain_out, plain_idx, _ = _prefill_under_kv(monkeypatch, config, "none", lengths)
+    quant, quant_out, quant_idx, positions = _prefill_under_kv(
+        monkeypatch, config, "fp8", lengths
+    )
+
+    # The tripwire for the field failure: the backend sizes its indexer scratch with
+    # pool.dtype, which must stay the COMPUTE dtype even when store_dtype is e4m3. An
+    # fp8 q_index compiles into qsa_mqa_paged's dot and dies at graph capture.
+    assert quant.backend.dtype is torch.bfloat16
+    assert plain.backend.dtype is torch.bfloat16
+    assert quant.pool.store_dtype != torch.bfloat16
+    assert quant.pool.kv_quant == "fp8" and plain.pool.kv_quant == "none"
+    assert quant.pool.k_cache(QSA_LAYER).element_size() == 1
+    assert quant.pool.v_cache(QSA_LAYER).element_size() == 1
+    assert plain.pool.k_scale(QSA_LAYER) is None and plain.pool.v_scale(QSA_LAYER) is None
+    pages, page_size, kv_heads = quant.pool.k_cache(QSA_LAYER).shape[:3]
+    assert quant.pool.k_scale(QSA_LAYER).shape == (pages * page_size, kv_heads)
+    assert quant.pool.k_scale(QSA_LAYER).dtype is torch.float32
+
+    for pool in (plain.pool, quant.pool):
+        assert pool.cmp_k_cache(0).dtype is torch.bfloat16
+    assert torch.equal(quant_idx, plain_idx), (
+        "quantizing the KV rows changed which blocks the indexer selected -- the index "
+        "tier is supposed to be 16-bit in both runs"
+    )
+    _assert_selection_is_causal_prefix(quant_idx, positions)
+
+    # Looser than the 2e-2 the 16-bit run needs against the same reference: e4m3 carries
+    # four significant bits, so ~1e-2 relative per stored element is the floor here.
+    torch.testing.assert_close(quant_out.float(), plain_out.float(), rtol=4e-2, atol=4e-2)
+
+
+def _mrope_config(rope_type):
+    from freetoken.distributed import set_tp_info, try_get_tp_info
+    from freetoken.models.qwen4_exp.config import parse_config
+
+    if try_get_tp_info() is None:
+        set_tp_info(rank=0, size=1)
+    hf = hf_config(budget=32, max_position=512)
+    hf.text_config.rope_parameters.update(mrope_section=[11, 11, 10], mrope_interleaved=True)
+    if rope_type == "yarn":
+        hf.text_config.rope_parameters.update(rope_type="yarn", factor=4.0,
+                                             original_max_position_embeddings=262144)
+    hf.vision_config = SimpleNamespace(
+        hidden_size=32, depth=1, num_heads=4, intermediate_size=64, patch_size=2,
+        temporal_patch_size=2, spatial_merge_size=2, num_position_embeddings=64,
+        out_hidden_size=hf.text_config.hidden_size, in_channels=3,
+    )
+    config = parse_config(hf)
+    assert config.model_is_mrope and config.rotary_config.mrope_layout == "interleaved"
+    return config
+
+
+def _image_positions(length, start, height, width, device):
+    positions = torch.arange(length, dtype=torch.int32, device=device).repeat(3, 1)
+    pixels = torch.arange(height * width, dtype=torch.int32, device=device)
+    end = start + pixels.numel()
+    positions[0, start:end] = start
+    positions[1, start:end] = start + pixels // width
+    positions[2, start:end] = start + pixels % width
+    positions[:, end:] += max(height, width) - pixels.numel()
+    return positions
+
+
+def _index_mrope_reference(positions, rope_type):
+    # Independently spell out Qwen's [11, 11, 10] interleaving and partial-rope frequencies.
+    axes = torch.tensor([1 if i % 3 == 1 else 2 if i % 3 == 2 and i < 30 else 0
+                         for i in range(32)], device=positions.device)
+    inv = 1.0 / (1e7 ** (torch.arange(0, 64, 2, device=positions.device).float() / 64))
+    amplitude = 1.0
+    if rope_type == "yarn":
+        low = math.floor(32 * math.log(262144 / (64 * math.pi)) / math.log(1e7))
+        high = math.ceil(32 * math.log(262144 / (2 * math.pi)) / math.log(1e7))
+        blend = ((torch.arange(32, device=positions.device).float() - low) / (high - low)).clamp(0, 1)
+        inv = inv * (1 - 0.75 * blend)
+        amplitude = 1 + 0.1 * math.log(4)
+    phases = positions[axes].t().float() * inv
+    return torch.cat((phases.cos(), phases.sin()), dim=-1) * amplitude
+
+
+def _mrope_under_kv(monkeypatch, kv_quant, chunked, rope_type):
+    fixture = Fixture(_mrope_config(rope_type), num_pages=16, max_running_req=2, kv_quant=kv_quant)
+    attn = fixture.layer(QSA_LAYER)
+    lengths, steps = [83, 59], 3
+    inputs = _inputs(fixture, lengths, extra=steps)
+    full_positions = [_image_positions(n + steps, start, h, w, fixture.device)
+                      for n, start, h, w in ((83, 29, 3, 7), (59, 13, 4, 5))]
+    cuts = [37, 19] if chunked else lengths
+    reqs = [fixture.req(i, 0, cut) for i, cut in enumerate(cuts)]
+    snapshots = []
+    with monkeypatch.context() as patch:
+        seen = selection_spy(patch, fixture.backend)
+
+        def forward(phase):
+            batch = fixture.batch(reqs, phase)
+            batch.mrope_positions = torch.cat([p[:, r.cached_len:r.device_len]
+                                                for p, r in zip(full_positions, reqs)], dim=1)
+            batch.get_attn_positions = lambda: batch.mrope_positions
+            x = torch.cat([row[r.cached_len:r.device_len] for row, r in zip(inputs, reqs)])
+            out = attn.forward(x, batch)
+            md = batch.attn_metadata
+            group_positions = torch.cat([p[:, torch.arange(r.cached_len, r.device_len, device=fixture.device) // 4 * 4]
+                                          for p, r in zip(full_positions, reqs)], dim=1)
+            torch.testing.assert_close(md.q_rope_cache, _index_mrope_reference(batch.mrope_positions, rope_type),
+                                       rtol=1e-6, atol=1e-6)
+            torch.testing.assert_close(md.k_rope_cache, _index_mrope_reference(group_positions, rope_type),
+                                       rtol=1e-6, atol=1e-6)
+            for r, positions in zip(reqs, full_positions):
+                slots = fixture.page_table[r.table_idx, :r.device_len].long()
+                assert torch.equal(fixture.pool.rope_positions[slots], positions[:, :r.device_len].t())
+            snapshots.append(dict(
+                out=out.float().cpu(), indices=seen["indices"].cpu(),
+                # Non-closing groups collide in unread scratch rows; only the persistent slab is defined.
+                compressed=fixture.pool.cmp_k_cache(0)[:fixture.pool.cmp_scratch_base].cpu().clone(),
+                ring=fixture.pool.pending_ring(0).cpu().clone(),
+                rope=fixture.pool.rope_positions.cpu().clone(),
+            ))
+
+        forward("prefill")
+        if chunked:
+            reqs = [fixture.req(i, cut, n) for i, (cut, n) in enumerate(zip(cuts, lengths))]
+            forward("prefill")
+        for _ in range(steps):
+            for req in reqs:
+                fixture.step(req)
+            forward("decode")
+    assert fixture.pool.kv_quant == kv_quant
+    assert fixture.pool.k_cache(QSA_LAYER).dtype == (torch.uint8 if kv_quant == "nvfp4" else fixture.dtype)
+    return snapshots
+
+
+@requires_cuda
+@pytest.mark.parametrize("rope_type", ["default", "yarn"])
+@pytest.mark.parametrize("chunked", [False, True], ids=["prefill-decode", "image-cut-decode"])
+def test_mrope_nvfp4_kv_keeps_index_and_rope_state(monkeypatch, chunked, rope_type):
+    """Image cuts split ratio-4 groups; NVFP4 changes only K/V, including later decode writes."""
+    plain = _mrope_under_kv(monkeypatch, "none", chunked, rope_type)
+    quant = _mrope_under_kv(monkeypatch, "nvfp4", chunked, rope_type)
+    assert len(plain) == len(quant)
+    for step, (expected, actual) in enumerate(zip(plain, quant)):
+        for key in ("indices", "compressed", "ring", "rope"):
+            assert torch.equal(actual[key], expected[key]), f"{key} changed at step {step}"
+        assert torch.isfinite(actual["out"]).all()
+        relative_rmse = ((actual["out"] - expected["out"]).square().mean()
+                         / expected["out"].square().mean()).sqrt()
+        cosine = torch.nn.functional.cosine_similarity(actual["out"].flatten(), expected["out"].flatten(), dim=0)
+        assert relative_rmse < 0.18, f"NVFP4 relative RMSE {relative_rmse.item():.4f} at step {step}"
+        assert cosine > 0.98, f"NVFP4 cosine {cosine.item():.4f} at step {step}"
+
diff --git a/tests/models/qwen4_exp/test_skeleton.py b/tests/models/qwen4_exp/test_skeleton.py
index ac598e65d..d5fe01c3b 100644
--- a/tests/models/qwen4_exp/test_skeleton.py
+++ b/tests/models/qwen4_exp/test_skeleton.py
@@ -502,3 +502,144 @@ def test_decoder_stack_prefill_and_decode(monkeypatch):
         decode_logits = model.forward()
     assert decode_logits.shape == (len(prompts), config.vocab_size)
     assert torch.isfinite(decode_logits.float()).all()
+
+
+@requires_cuda
+@pytest.mark.parametrize("checkpoint", ["radixark", "nvidia"])
+@torch.inference_mode()
+def test_nvfp4_experts_and_kv_preserve_full_model_continuation(checkpoint, monkeypatch):
+    """Real GDN/QSA/PLE and NVFP4 MoE across ragged chunks, ratio-4 groups and page 64."""
+    from copy import deepcopy
+    from dataclasses import replace
+
+    from freetoken.attention.qsa_sparse import QSASparseAttnBackend
+    from freetoken.kvcache.linear_state_pool import LinearStatePool
+    from freetoken.kvcache.qsa_pool import QSAKVCache
+    from freetoken.layers.quantization import NameMap, QuantBackend, QuantConfig, QuantKind
+    from freetoken.layers.quantization import finalize_quant
+    from freetoken.models.qwen4_exp.model import Qwen4ExpForConditionalGeneration
+    from freetoken.models.qwen4_exp.ple import PLE_CONV_STATE, PLE_NGRAM_STATE
+    from freetoken.models.register import get_model_spec
+    from freetoken.moe.expert_banks import build_expert_banks
+    from freetoken.moe.offload_cache import OffloadMoeCache, attach_offload_moe_cache, iter_offload_moe_layers
+    from freetoken.utils.torch_utils import torch_dtype
+
+    from .common import Fixture, NVIDIA_NVFP4, RADIXARK_NVFP4
+
+    hf = toy_hf_config()
+    hf.quantization_config = deepcopy(RADIXARK_NVFP4 if checkpoint == "radixark" else NVIDIA_NVFP4)
+    spec = get_model_spec(hf.architectures[0])
+    quant = QuantConfig.from_hf(hf, name_map=NameMap(
+        roots=spec.checkpoint_roots, segments=spec.checkpoint_segments,
+        packed=spec.packed_modules_mapping), unquantized=spec.unquantized_modules)
+    config = replace(parse_config(hf), quant=quant, moe_strategy="offload", decode_target="gpu")
+    monkeypatch.setattr("freetoken.layers.quantization.quant_backend._QUANT_BACKEND",
+                        QuantBackend.parse("moe.nvfp4=triton"))
+    device, dtype = torch.device("cuda"), torch.bfloat16
+    with torch.device(device), torch_dtype(dtype):
+        model = Qwen4ExpForConditionalGeneration(config)
+    gen = torch.Generator(device=device).manual_seed(63)
+    _fill(model, gen, scale=.03)
+    assert finalize_quant(model) > config.num_layers
+    experts = list(iter_offload_moe_layers(model))
+    assert len(experts) == config.num_layers
+    assert all(layer.quant_method.kind is QuantKind.NVFP4 for layer in experts)
+    assert all(layer.quant_method.kernel.name == "triton" for layer in experts)
+    torch.manual_seed(64)
+    banks = build_expert_banks(experts[0].quant_method, config.num_layers, None,
+                               device=device, dummy=True)
+    cache = OffloadMoeCache(config.num_layers, config.num_experts, 2 * config.num_experts,
+                            device, quant_format=banks.quant_format, layout=banks.layout)
+    cache.set_bank_sources(banks.sources)
+    attach_offload_moe_cache(model, cache)
+    for ple in model.model.ple_layers:
+        multipliers, sizes, offsets = hash_constants(config.qwen4_args)
+        emb = ple.ple_embedding
+        emb.layer_multipliers.copy_(multipliers)
+        emb.ngram_heads_vocab_sizes.copy_(sizes)
+        emb.ngram_heads_offsets.copy_(offsets)
+        table = torch.randn(int(offsets[-1] + sizes[-1]), config.qwen4_args.ngram_head_dim,
+                             generator=gen, device=device, dtype=dtype) * .03
+        emb.attach_table(GpuResidentTable(table, dtype=dtype))
+
+    tokens = [((torch.arange(length) * 3 + offset) % 500).long()
+              for length, offset in ((67, 13), (71, 101))]
+    tokens[0][29], tokens[1][33] = EOS, EOS
+
+    selected = []
+    select = QSASparseAttnBackend._select
+
+    def capture_selection(backend, index, metadata, slot):
+        indices = select(backend, index, metadata, slot)
+        selected[:] = [indices.clone()]
+        return indices
+
+    monkeypatch.setattr(QSASparseAttnBackend, "_select", capture_selection)
+
+    def run(chunks, kv_quant):
+        fixture = Fixture(config, num_pages=8, max_running_req=3, kv_quant=kv_quant)
+        pool = fixture.pool
+        assert isinstance(pool, QSAKVCache) and pool.kv_quant == kv_quant
+        states = LinearStatePool(config.linear_attention_group(), 4, dtype, device,
+                                  slot_states=config.slot_states)
+        fixture.ctx.linear_state_pool = states
+        cache.reset()
+        reqs = [SimpleNamespace(table_idx=i + 1, linear_slot_idx=i + 1, cached_len=0,
+                                device_len=0, extend_len=0, mamba_ping_pong=None) for i in range(2)]
+        for phase, lengths in chunks:
+            for req, length in zip(reqs, lengths):
+                fixture.allocate(req.table_idx, req.device_len, length)
+                req.cached_len, req.device_len = req.device_len, length
+                req.extend_len = length - req.cached_len
+            batch = fixture.batch(reqs, phase)
+            batch.input_ids = torch.cat([ids[req.cached_len:req.device_len]
+                                         for ids, req in zip(tokens, reqs)]).to(device)
+            batch.linear_table_idx = torch.tensor([1, 2], dtype=torch.int32, device=device)
+            batch.fla_metadata = None
+            with fixture.ctx.forward_batch(batch):
+                logits = model.forward()
+            assert logits.shape == (2, config.vocab_size) and torch.isfinite(logits).all()
+            assert not torch.equal(logits[0], logits[1])
+            context = states.slot_state(PLE_NGRAM_STATE)[1:3]
+            expected = torch.stack([ids[length - 2:length] for ids, length in zip(tokens, lengths)]).to(device)
+            torch.testing.assert_close(context.long(), expected, rtol=0, atol=0)
+        locations = torch.cat([fixture.page_table[req.table_idx, :req.device_len] for req in reqs]).long()
+        for kind in ("k", "v"):
+            codes = getattr(pool, f"{kind}_cache")(3)
+            if kv_quant == "none":
+                assert codes.dtype == dtype and codes.shape[-1] == config.head_dim
+                continue
+            scales = getattr(pool, f"{kind}_scale")(3)[locations]
+            blocks = getattr(pool, f"{kind}_block_scale")(3)[locations]
+            assert codes.dtype == torch.uint8 and codes.shape[-1] == config.head_dim // 2
+            assert codes.view(-1, config.num_kv_heads, config.head_dim // 2)[locations].any()
+            assert scales.dtype == torch.float32 and torch.isfinite(scales).all() and (scales > 0).all()
+            assert blocks.dtype == torch.uint8 and blocks.shape[-1] == config.head_dim // 16
+            assert torch.isfinite(blocks.view(torch.float8_e4m3fn).float()).all() and blocks.any()
+        assert pool.cmp_k_cache(0).dtype == pool.pending_ring(0).dtype == dtype
+        grouped_locations = torch.cat([fixture.page_table[req.table_idx, :req.device_len // 4 * 4:4] // 4
+                                       for req in reqs]).long()
+        compressed = pool.cmp_k_cache(0)[grouped_locations].clone()
+        assert compressed.abs().max() > 0 and torch.isfinite(compressed).all()
+        rec = states.recurrent_states[:, 1:3].clone()
+        conv = states.conv_states[:, 1:3].clone()
+        ple_conv = states.slot_state(PLE_CONV_STATE, config.qwen4_args.ple_layer_ids[0])[1:3].clone()
+        assert rec.abs().max() > 0 and ple_conv.abs().max() > 0
+        return logits.clone(), rec, conv, ple_conv, compressed, selected[0]
+
+    full = [("prefill", (67, 71))]
+    chunks = [("prefill", (31, 35)), ("prefill", (63, 67))]
+    chunks += [("decode", (64 + step, 68 + step)) for step in range(4)]
+    bf16 = [run(schedule, "none") for schedule in (full, chunks)]
+    nvfp4 = [run(schedule, "nvfp4") for schedule in (full, chunks)]
+
+    for dense, packed in zip(bf16, nvfp4):
+        # QSA is last, so changing its KV storage cannot change the earlier states or indexer.
+        for actual, expected in zip(packed[1:], dense[1:]):
+            torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+    for reference, continued in (bf16, nvfp4):
+        for actual, expected in zip(continued[1:4], reference[1:4]):
+            torch.testing.assert_close(actual.float(), expected.float(), rtol=.02, atol=.002)
+    torch.testing.assert_close(bf16[1][0].float(), bf16[0][0].float(), rtol=.02, atol=.002)
+    # GDN's prefill/decode rounding can cross FP4 bins; logits need a separate error budget.
+    torch.testing.assert_close(nvfp4[1][0].float(), nvfp4[0][0].float(), rtol=.02, atol=.006)
diff --git a/tests/models/test_deepseek_v41_config.py b/tests/models/test_deepseek_v41_config.py
new file mode 100644
index 000000000..b0bdde08d
--- /dev/null
+++ b/tests/models/test_deepseek_v41_config.py
@@ -0,0 +1,139 @@
+"""Captured NVFP4 configs: s-zaizen 179b7cda and LibertAIDAI dfce15b9."""
+
+import json
+from pathlib import Path
+
+import pytest
+
+from freetoken.models.deepseek_v41.args import load_args
+from freetoken.models.deepseek_v41.config import parse_config
+
+
+@pytest.fixture
+def checkpoint_config():
+    return json.loads((Path(__file__).parent / "fixtures/deepseek_v41_nvfp4_config.json").read_text())
+
+
+@pytest.fixture
+def libertai_config():
+    return json.loads((Path(__file__).parent / "fixtures/deepseek_v41_libertai_nvfp4_config.json").read_text())
+
+
+def test_target_mixed_quant_and_multimodal_config(checkpoint_config):
+    config = parse_config(checkpoint_config)
+    args = config.dsv41_args
+    assert (config.num_layers, config.hidden_size, config.num_experts, config.num_experts_per_tok) == (40, 5120, 384, 6)
+    assert config.expert_quant == "nvfp4"
+    assert config.weight_block_size == (32, 32)
+    assert (config.hidden_act, config.hidden_act_alpha, config.swiglu_limit) == ("swiglu_clamp", 1.0, 10.0)
+    assert config.is_multimodal and args.vision_enabled
+    assert args.vision_n_layers == 32 and config.image_token_id == 129264
+    assert args.engram_layer_ids == (1, 14)
+    assert (args.engram_dtype, args.engram_block_size, args.engram_scale_fmt) == ("fp8", 32, "ue8m0")
+    assert args.kv_source_layers == (2, 8, 14, 20)
+    assert args.index_source_layers == (2, 8, 14, 20, 24, 28, 32, 36)
+    assert len(args.compress_ratios) == 43 and args.n_mtp_layers == 3
+    assert config.attention_groups[0].layer_ids == tuple(range(40))
+    assert config.rotary_config.max_position == 1048576
+    assert config.rms_norm_eps == 1e-20
+
+
+def test_parse_needs_no_local_inference_file(checkpoint_config):
+    from freetoken.utils.hf import RawConfigShim
+
+    hf_config = RawConfigShim(checkpoint_config, _name_or_path="s-zaizen/DeepSeek-V4.1-Flash-NVFP4")
+    assert parse_config(hf_config) == parse_config(checkpoint_config)
+
+
+def test_explicit_vision_disable_preserves_native_quantization(checkpoint_config):
+    from freetoken.models.register import checkpoint_quant_config, get_model_spec
+    from freetoken.layers.quantization import QuantKind
+
+    checkpoint_config["vision_config"] = None
+    config = parse_config(checkpoint_config)
+    assert not config.is_multimodal and not config.dsv41_args.vision_enabled
+    spec = get_model_spec("DeepseekV41ForCausalLM")
+    assert spec.encoders[0].modalities == ("image",) and spec.mm_processor is not None
+    quant = checkpoint_quant_config("unused", checkpoint_config, spec)
+    assert quant.scheme_for_name("layers.2.attn.wq_a").kind is QuantKind.FP8_BLOCK
+    assert quant.scheme_for_name("layers.2.ffn.experts").kind is QuantKind.NVFP4
+
+
+def test_libertai_native_nvfp4_with_packed_engram(libertai_config):
+    from freetoken.utils.hf import RawConfigShim
+
+    config = parse_config(libertai_config)
+    args = config.dsv41_args
+    assert config.expert_quant == "nvfp4" and config.weight_block_size == (32, 32)
+    assert (config.num_layers, config.hidden_size, config.num_experts) == (40, 5120, 384)
+    assert config.is_multimodal and args.vision_n_layers == 32
+    assert (args.engram_dtype, args.engram_block_size, args.engram_scale_fmt) == ("fp4", 32, "ue8m0")
+    assert args.engram_num_embeddings == (384006168, 384016682)
+    assert "moe_quant_algo" not in libertai_config["quantization_config"]
+    assert "quantized_layers" not in libertai_config["quantization_config"]
+    assert parse_config(RawConfigShim(libertai_config)) == config
+
+
+@pytest.mark.parametrize("key,value,match", [
+    ("expert_dtype", "fp4", "NVFP4"),
+    ("expert_block_size", 32, "group_size"),
+    ("expert_scale_fmt", "ue8m0", "E4M3"),
+    ("expert_global_scale", False, "global scale"),
+    ("engram_dtype", "bf16", "FP8 or FP4"),
+    ("engram_block_size", 16, "block_size=32"),
+    ("engram_scale_fmt", "e4m3", "UE8M0"),
+])
+def test_reject_incompatible_native_storage(libertai_config, key, value, match):
+    libertai_config["quantization_config"][key] = value
+    with pytest.raises(ValueError, match=match):
+        parse_config(libertai_config)
+
+
+def test_engram_quant_metadata_overrides_text_storage(libertai_config, tmp_path):
+    libertai_config["text_config"]["engram_dtype"] = "fp8"
+    (tmp_path / "config.json").write_text(json.dumps(libertai_config))
+    assert load_args(tmp_path).engram_dtype == "fp4"
+    assert load_args(tmp_path, engram_dtype="fp8").engram_dtype == "fp8"
+
+
+def test_legacy_fp4_label_alone_does_not_imply_nvfp4(checkpoint_config):
+    del checkpoint_config["quantization_config"]["moe_quant_algo"]
+    del checkpoint_config["quantization_config"]["quantized_layers"]
+    with pytest.raises(ValueError, match="NVFP4"):
+        parse_config(checkpoint_config)
+
+
+@pytest.mark.parametrize("key,value,match", [
+    ("moe_quant_algo", "MXFP4", "NVFP4"),
+    ("group_size", 32, "group_size"),
+    ("weight_block_size", [128, 128], "32x32"),
+    ("scale_fmt", "float", "UE8M0"),
+])
+def test_reject_incompatible_quantization(checkpoint_config, key, value, match):
+    checkpoint_config["quantization_config"][key] = value
+    with pytest.raises(ValueError, match=match):
+        parse_config(checkpoint_config)
+
+
+def test_reject_missing_layer_quantization(checkpoint_config):
+    del checkpoint_config["quantization_config"]["quantized_layers"]["layers.39.ffn.experts"]
+    with pytest.raises(ValueError, match="backbone layer 39"):
+        parse_config(checkpoint_config)
+
+
+def test_infer_nvfp4_from_per_layer_quantization(checkpoint_config):
+    del checkpoint_config["quantization_config"]["moe_quant_algo"]
+    assert parse_config(checkpoint_config).expert_quant == "nvfp4"
+
+
+@pytest.mark.parametrize("field,value", [
+    ("kv_source_layer_ids", [2, 8, 14]),
+    ("index_source_layer_ids", [8, 14, 20]),
+    ("compress_ratios", [0] * 39),
+    ("engram_num_embeddings", [4]),
+    ("gate_temp", 0),
+])
+def test_invalid_attention_or_engram_layout(checkpoint_config, field, value):
+    checkpoint_config["text_config"][field] = value
+    with pytest.raises(ValueError):
+        load_args(checkpoint_config)
diff --git a/tests/models/test_deepseek_v41_engram.py b/tests/models/test_deepseek_v41_engram.py
new file mode 100644
index 000000000..c5bf6ec0b
--- /dev/null
+++ b/tests/models/test_deepseek_v41_engram.py
@@ -0,0 +1,404 @@
+import json
+import math
+from types import SimpleNamespace
+
+import numpy as np
+import pytest
+import torch
+from safetensors.torch import save_file
+
+from freetoken.models.deepseek_v41.engram import (
+    DiskEngramTable, Engram, EngramRuntime, HashLayout, compressed_token_map,
+    export_engram_tables, hash_token_run, prepare_engram,
+)
+
+
+def _args():
+    return SimpleNamespace(
+        engram_layer_ids=(1, 14), engram_max_ngram_size=3, engram_n_heads=2,
+        engram_vocab_size=10, engram_num_embeddings=(60, 120),
+        engram_compressed_vocab_size=7, engram_head_dim=32, engram_pad_id=2,
+        engram_dtype="fp8", engram_block_size=32, engram_scale_fmt="ue8m0",
+    )
+
+
+def _scalar_hash(ids, images, mapping, layout, pad):
+    output = []
+    for position in range(len(ids)):
+        row = []
+        for layer in range(len(layout.layer_ids)):
+            values, blocked = [], False
+            for shift in range(layout.max_ngram_size):
+                source = position - shift
+                blocked |= source < 0 or (source >= 0 and images[source])
+                values.append(int(mapping[pad if blocked else ids[source]]))
+            running = values[0] * int(layout.multipliers[layer, 0])
+            per_layer = []
+            for shift in range(1, layout.max_ngram_size):
+                running ^= values[shift] * int(layout.multipliers[layer, shift])
+                for head in range(layout.heads):
+                    column = (shift - 1) * layout.heads + head
+                    per_layer.append(running % int(layout.primes[layer, column]) + int(layout.offsets[layer, column]))
+            row.append(per_layer)
+        output.append(row)
+    return np.asarray(output)
+
+
+def test_hash_matches_scalar_and_chunk_boundary():
+    args = _args()
+    layout = HashLayout.from_args(args)
+    mapping = np.arange(7, dtype=np.int64)
+    ids = np.array([1, 3, 5, 6, 4, 1, 3, 6])
+    images = np.array([False, False, True, True, False, False, False, False])
+    expected = _scalar_hash(ids, images, mapping, layout, args.engram_pad_id)
+    np.testing.assert_array_equal(hash_token_run(ids, images, mapping, layout, args.engram_pad_id), expected)
+    for boundary in range(1, len(ids)):
+        start = max(0, boundary - 2)
+        actual = hash_token_run(ids[start:], images[start:], mapping, layout, args.engram_pad_id, boundary-start)
+        np.testing.assert_array_equal(actual, expected[boundary:])
+    alternate = ids.copy()
+    alternate[:2] = [6, 4]
+    changed = hash_token_run(alternate, images, mapping, layout, args.engram_pad_id)
+    np.testing.assert_array_equal(changed[4:], expected[4:])
+
+
+def test_hash_rejects_wrong_bucket_geometry():
+    args = _args()
+    args.engram_num_embeddings = (61, 120)
+    with pytest.raises(ValueError, match="table sizes"):
+        HashLayout.from_args(args)
+
+
+def test_content_pad_ids_use_image_boundaries_before_vocabulary_lookup():
+    from freetoken.models.deepseek_v41.engram import _image_flags
+
+    args, mapping = _args(), np.arange(7, dtype=np.int64)
+    layout = HashLayout.from_args(args)
+    ids = np.array([1, 3, 5, 6, 4, 1, 3, 6])
+    req = SimpleNamespace(mm_items=[SimpleNamespace(offsets=[[2, 4]])])
+    flags = _image_flags(req, 0, len(ids))
+    expected = hash_token_run(ids, flags, mapping, layout, args.engram_pad_id)
+    ids[2:4] = [1_000_100, 1_000_100]
+    np.testing.assert_array_equal(hash_token_run(ids, flags, mapping, layout, args.engram_pad_id), expected)
+    np.testing.assert_array_equal(_image_flags(req, 3, 4), [True, False, False, False])
+    with pytest.raises(ValueError, match="outside the tokenizer"):
+        hash_token_run(ids, np.zeros_like(flags), mapping, layout, args.engram_pad_id)
+
+
+def test_compressed_tokens_match_training_normalization():
+    texts = [" The", "the", "THE", "\u00e9", "e", "\uff25", " ", "\t", "",
+             "a\r\n\tb", "a b", "\ufffd", "\ufffd"]
+
+    class Tokenizer:
+        def __len__(self):
+            return len(texts)
+
+        @property
+        def backend_tokenizer(self):
+            return self
+
+        def decode(self, ids, *, skip_special_tokens):
+            assert skip_special_tokens is False
+            return texts[ids[0]]
+
+        def id_to_token(self, token_id):
+            return f"<byte-{token_id}>"
+
+    np.testing.assert_array_equal(compressed_token_map(Tokenizer()),
+                                  [0, 0, 0, 1, 1, 1, 2, 2, 3, 4, 4, 5, 6])
+
+
+def test_disk_table_decodes_all_e8m0_codes():
+    weights = torch.ones(256, 32).to(torch.float8_e4m3fn)
+    codes = torch.arange(256, dtype=torch.uint8).view(256, 1)
+    table = DiskEngramTable(weights.view(torch.uint8).numpy(), codes.numpy())
+    expected = (weights.float() * codes.view(torch.float8_e8m0fnu).float()).to(torch.bfloat16)
+    torch.testing.assert_close(table.lookup(np.arange(256)), expected, rtol=0, atol=0, equal_nan=True)
+
+
+def test_disk_table_rejects_partial_scale_group():
+    with pytest.raises(ValueError, match="per 32"):
+        DiskEngramTable(np.zeros((4, 33), dtype=np.uint8), np.zeros((4, 1), dtype=np.uint8))
+
+
+def _checkpoint(folder, dtype="fp8"):
+    args = _args()
+    args.engram_dtype = dtype
+    tensors = {}
+    for layer, rows in zip(args.engram_layer_ids, args.engram_num_embeddings):
+        w = (torch.arange(rows * 32).reshape(rows, 32) % 9 - 4).float().to(torch.float8_e4m3fn)
+        if dtype == "fp4":
+            w = (torch.arange(rows * 16).reshape(rows, 16) % 256).to(torch.uint8)
+        s = torch.full((rows, 1), 128, dtype=torch.uint8).view(torch.float8_e8m0fnu)
+        tensors[f"layers.{layer}.engram.embed.weight"] = w
+        tensors[f"layers.{layer}.engram.embed.scale"] = s
+    save_file(tensors, str(folder / "tables.safetensors"))
+    (folder / "model.safetensors.index.json").write_text(json.dumps({"weight_map": {k: "tables.safetensors" for k in tensors}}))
+    return args, tensors
+
+
+def _reference_rows(tensors, layer, ids):
+    weight = tensors[f"layers.{layer}.engram.embed.weight"]
+    scale = tensors[f"layers.{layer}.engram.embed.scale"]
+    if weight.dtype != torch.uint8:
+        return (weight.float().reshape(len(weight), -1, 32) * scale.float()[..., None]).flatten(-2)[ids]
+    result = []
+    for row in ids.reshape(-1).tolist():
+        values = []
+        for column in range(weight.shape[1] * 2):
+            code = int(weight[row, column // 2]) >> (4 * (column % 2)) & 15
+            exponent, fraction = (code & 7) >> 1, code & 1
+            magnitude = fraction * .5 if exponent == 0 else math.ldexp(1 + fraction * .5, exponent - 1)
+            exponent_scale = int(scale.view(torch.uint8)[row, column // 32]) - 127
+            value = math.ldexp(magnitude, exponent_scale) * (-1 if code & 8 else 1)
+            values.append(value)
+        result.append(values)
+    return torch.tensor(result).reshape(*ids.shape, weight.shape[1] * 2)
+
+
+@pytest.mark.parametrize("dtype", ["fp8", "fp4"])
+def test_disk_table_and_standalone_export(tmp_path, dtype):
+    from freetoken.checkpoint.convert import _copy_metadata
+
+    source, dest = tmp_path / "source", tmp_path / "ftw"
+    source.mkdir()
+    dest.mkdir()
+    args, tensors = _checkpoint(source, dtype)
+    metadata = {"config.json": '{"model_type": "deepseek_v41"}',
+                "inference/config.json": '{"dim": 5120}',
+                "encoding/encoding.py": "IS_DSV41 = True\n",
+                "tokenizer.json": "{}"}
+    for name, text in metadata.items():
+        path = source / name
+        path.parent.mkdir(parents=True, exist_ok=True)
+        path.write_text(text)
+    ids = np.array([[3, 1, 3, 7], [5, 2, 0, 1]])
+    for layer in args.engram_layer_ids:
+        table = DiskEngramTable.from_checkpoint(str(source), layer, args=args)
+        assert isinstance(table.weights, np.memmap) and not table.weights.flags.writeable
+        assert isinstance(table.scales, np.memmap) and not table.scales.flags.writeable
+        expected = _reference_rows(tensors, layer, torch.from_numpy(ids))
+        torch.testing.assert_close(table.lookup(ids).float(), expected, rtol=0, atol=0)
+    copied = _copy_metadata(str(source), str(dest))
+    assert set(copied) == set(metadata)
+    export_engram_tables(str(source), str(dest), args)
+    manifest = json.loads((dest / "engram_tables.json").read_text())
+    assert manifest["version"] == 2
+    for layer in args.engram_layer_ids:
+        info = manifest["layers"][str(layer)]
+        assert info["dtype"] == dtype and info["block_size"] == 32 and info["scale_fmt"] == "ue8m0"
+        for kind in ("weight", "scale"):
+            original = tensors[f"layers.{layer}.engram.embed.{kind}"].view(torch.uint8).numpy().tobytes()
+            assert (dest / info[kind]["file"]).read_bytes() == original
+    source.rename(tmp_path / "unavailable")
+    for name, text in metadata.items():
+        assert (dest / name).read_text() == text
+    for layer in args.engram_layer_ids:
+        restored = DiskEngramTable.from_checkpoint(str(dest), layer, args=args)
+        expected = _reference_rows(tensors, layer, torch.from_numpy(ids))
+        torch.testing.assert_close(restored.lookup(ids).float(), expected, rtol=0, atol=0)
+
+
+def test_legacy_fp8_manifest_remains_loadable(tmp_path):
+    source, dest = tmp_path / "source", tmp_path / "ftw"
+    source.mkdir()
+    args, tensors = _checkpoint(source)
+    export_engram_tables(str(source), str(dest), args)
+    manifest_path = dest / "engram_tables.json"
+    manifest = json.loads(manifest_path.read_text())
+    manifest["version"] = 1
+    for info in manifest["layers"].values():
+        for field in ("dtype", "block_size", "scale_fmt"):
+            info.pop(field)
+    manifest_path.write_text(json.dumps(manifest))
+    ids = torch.tensor([1, 7, 3])
+    actual = DiskEngramTable.from_checkpoint(str(dest), 1, args=args).lookup(ids.numpy())
+    torch.testing.assert_close(actual.float(), _reference_rows(tensors, 1, ids), rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("dtype", ["fp8", "fp4"])
+def test_table_rejects_config_format_mismatch(tmp_path, dtype):
+    source, dest = tmp_path / "source", tmp_path / "ftw"
+    source.mkdir()
+    args, _ = _checkpoint(source, dtype)
+    export_engram_tables(str(source), str(dest), args)
+    args.engram_dtype = "fp4" if dtype == "fp8" else "fp8"
+    for folder in (source, dest):
+        with pytest.raises(ValueError, match="format disagrees"):
+            DiskEngramTable.from_checkpoint(str(folder), 1, args=args)
+
+
+def test_fp4_table_uses_per_block_scale_and_low_nibble_first():
+    codes = torch.arange(256, dtype=torch.uint8).repeat(2, 1)
+    scale_codes = torch.arange(111, 127, dtype=torch.uint8).repeat(2, 1)
+    tensors = {"layers.1.engram.embed.weight": codes,
+               "layers.1.engram.embed.scale": scale_codes.view(torch.float8_e8m0fnu)}
+    ids = torch.tensor([[1, 0, 1]])
+    table = DiskEngramTable(codes.numpy(), scale_codes.numpy(), dtype="fp4")
+    expected = _reference_rows(tensors, 1, ids).to(torch.bfloat16)
+    torch.testing.assert_close(table.lookup(ids.numpy()), expected, rtol=0, atol=0)
+    with torch.device("meta"):
+        actual = table.lookup(ids.numpy())
+    assert actual.device.type == "cpu"
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+    assert table.lookup(np.empty((0, 2), dtype=np.int64)).shape == (0, 2, 512)
+
+
+def test_fp4_lookup_reads_only_unique_requested_rows():
+    class RequestedRowsOnly:
+        dtype = np.dtype("uint8")
+        ndim = 2
+
+        def __init__(self, width, fill):
+            self.shape = (384006168, width)
+            self.fill = fill
+            self.reads = []
+
+        def __getitem__(self, ids):
+            np.testing.assert_array_equal(ids, [1, 7, 1000])
+            self.reads.append(ids.copy())
+            return np.full((len(ids), self.shape[1]), self.fill, dtype=np.uint8)
+
+        def __array__(self, *args, **kwargs):
+            raise AssertionError("full Engram table materialization")
+
+    weight, scale = RequestedRowsOnly(128, 0x32), RequestedRowsOnly(8, 128)
+    table = DiskEngramTable(weight, scale, dtype="fp4")
+    actual = table.lookup(np.array([[1000, 1, 7, 1]]))
+    expected = torch.tensor([2., 3.], dtype=torch.bfloat16).repeat(128).expand(1, 4, 256)
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+    assert len(weight.reads) == len(scale.reads) == 1
+
+
+@pytest.mark.parametrize("dtype,block,fmt", [("nvfp4", 32, "ue8m0"),
+                                           ("fp4", 16, "ue8m0"), ("fp4", 32, "e4m3")])
+def test_disk_table_rejects_unsupported_format(dtype, block, fmt):
+    with pytest.raises(ValueError, match="block-32 UE8M0"):
+        DiskEngramTable(np.zeros((4, 128), dtype=np.uint8), np.zeros((4, 8), dtype=np.uint8),
+                        dtype=dtype, block_size=block, scale_fmt=fmt)
+
+
+@pytest.mark.parametrize("shape", [(4, 4), (4, 16), (3, 8)])
+def test_fp4_table_rejects_scale_shape_for_wrong_logical_width(shape):
+    with pytest.raises(ValueError, match="per 32 weight channels"):
+        DiskEngramTable(np.zeros((4, 128), dtype=np.uint8), np.zeros(shape, dtype=np.uint8), dtype="fp4")
+
+
+def test_prepare_engram_loads_fp4_tables_without_expanding_storage(tmp_path, monkeypatch):
+    import freetoken.models.deepseek_v41.engram as engram_module
+    import freetoken.utils
+
+    args, tensors = _checkpoint(tmp_path, "fp4")
+    layers = [SimpleNamespace(engram=SimpleNamespace(layer_id=i, hash_cols=4)) for i in args.engram_layer_ids]
+    model = SimpleNamespace(_args=args, _transformer=SimpleNamespace(layers=layers))
+    config = SimpleNamespace(model_path=str(tmp_path), device="cpu", use_dummy_weight=False,
+                             max_extend_tokens=3, max_running_req=1, cuda_graph_max_bs=0)
+    monkeypatch.setattr(freetoken.utils, "download_hf_weight", lambda path: path)
+    monkeypatch.setattr(freetoken.utils, "load_tokenizer", lambda path: object())
+    monkeypatch.setattr(engram_module, "compressed_token_map", lambda tokenizer: np.arange(7))
+    assert prepare_engram(model, config) == 3 * (2 * 4 * 32 * 2 + 1)
+    for layer_id, table in zip(args.engram_layer_ids, model._engram_runtime.tables):
+        assert table.dtype == "fp4" and table.head_dim == 32
+        assert isinstance(table.weights, np.memmap) and isinstance(table.scales, np.memmap)
+        assert table.weights.nbytes + table.scales.nbytes == table.num_rows * 17
+        ids = torch.tensor([1, 3, 5])
+        torch.testing.assert_close(table.lookup(ids.numpy()).float(), _reference_rows(tensors, layer_id, ids),
+                                   rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("dtype", ["fp8", "fp4"])
+def test_runtime_isolates_requests_and_images(tmp_path, dtype):
+    args, _ = _checkpoint(tmp_path, dtype)
+    tables = [DiskEngramTable.from_checkpoint(str(tmp_path), layer) for layer in args.engram_layer_ids]
+    modules = [SimpleNamespace(hash_cols=4) for _ in tables]
+    runtime = EngramRuntime(args, modules, tables, np.arange(7), 12, "cpu")
+    first = SimpleNamespace(input_ids=torch.tensor([1, 3, 5, 4]), cached_len=2, extend_len=2,
+                            device_len=4, media=[{"start": 2, "types": [0]}])
+    second = SimpleNamespace(input_ids=torch.tensor([6, 1]), cached_len=0, extend_len=2, device_len=2, media=None)
+    batch = SimpleNamespace(input_ids=torch.tensor([5, 4, 6, 1]), padded_reqs=[first, second], is_decode=False)
+    with runtime.forward_host_ctx(batch, False):
+        assert not modules[0]._mask[0]
+        assert modules[0]._mask[1:4].all()
+        assert not modules[0]._values[0].any()
+        expected_ids = hash_token_run(np.array([6, 1]), np.zeros(2, bool), np.arange(7), runtime.layout, 2)
+        torch.testing.assert_close(modules[0]._values[2:4], tables[0].lookup(expected_ids[:, 0]).flatten(-2))
+
+
+def test_runtime_decode_uses_current_gpu_token_and_keeps_image_boundary(tmp_path):
+    args, _ = _checkpoint(tmp_path)
+    tables = [DiskEngramTable.from_checkpoint(str(tmp_path), layer) for layer in args.engram_layer_ids]
+    modules = [SimpleNamespace(hash_cols=4) for _ in tables]
+    runtime = EngramRuntime(args, modules, tables, np.arange(7), 3, "cpu")
+    # Under overlap, the current token is in batch.input_ids before append_host runs.
+    req = SimpleNamespace(input_ids=torch.tensor([1, 3, 5, 4]), device_len=5,
+                          media=[{"start": 1, "types": [0, 1, 3]}])
+    dummy = SimpleNamespace(input_ids=torch.tensor([0]), device_len=1, media=None)
+    batch = SimpleNamespace(input_ids=torch.tensor([6, 0]), padded_reqs=[req, dummy], is_decode=True)
+    with runtime.forward_host_ctx(batch, True):
+        expected = hash_token_run(np.array([1, 3, 5, 4, 6]), np.array([0, 1, 1, 1, 0]),
+                                  np.arange(7), runtime.layout, 2)[-1:]
+        for index, table in enumerate(tables):
+            torch.testing.assert_close(modules[index]._values[:1], table.lookup(expected[:, index]).flatten(-2))
+        assert modules[0]._mask[:2].all()
+
+
+@pytest.mark.parametrize("prefill,decode,graph,expected", [(7, 3, 16, 16), (9, 16, 3, 16), (None, 4, 4, 8192)])
+def test_engram_staging_capacity_follows_batch_budget(prefill, decode, graph, expected):
+    args = _args()
+    layers = [SimpleNamespace(engram=SimpleNamespace(hash_cols=4)) for _ in args.engram_layer_ids]
+    model = SimpleNamespace(_args=args, _transformer=SimpleNamespace(layers=layers))
+    config = SimpleNamespace(use_dummy_weight=True, device="cpu", max_forward_len=1 << 20,
+                             max_running_req=decode, cuda_graph_max_bs=graph)
+    if prefill is not None:
+        config.max_extend_tokens = prefill
+    pinned_bytes = prepare_engram(model, config)
+    assert model._engram_runtime.capacity == expected
+    assert pinned_bytes == expected * (len(layers) * 4 * 32 * 2 + 1)
+
+
+def test_engram_staging_rejects_larger_batch_before_writing():
+    args = _args()
+    module = SimpleNamespace(hash_cols=4)
+    runtime = EngramRuntime(args, [module], [], None, 2, "cpu", dummy=True)
+    req = SimpleNamespace(input_ids=torch.tensor([1, 3, 5]), cached_len=0, extend_len=3, media=None)
+    batch = SimpleNamespace(input_ids=req.input_ids, padded_reqs=[req], is_decode=False)
+    with pytest.raises(ValueError, match="requires 3 rows; capacity is 2"):
+        with runtime.forward_host_ctx(batch, False):
+            pass
+    assert not module._values.any()
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+@pytest.mark.parametrize("dtype", ["fp8", "fp4"])
+def test_engram_staging_updates_captured_graph(tmp_path, dtype):
+    args, _ = _checkpoint(tmp_path, dtype)
+    args.dim, args.hc_mult, args.norm_eps = 32, 2, 1e-6
+    module = Engram(args, 1).cuda()
+    torch.manual_seed(7)
+    with torch.no_grad():
+        module.wkv.weight.copy_(torch.randn(module.wkv.weight.shape, device="cuda") * .1)
+        module.wkv.scale.view(torch.uint8).fill_(127)
+    table = DiskEngramTable.from_checkpoint(str(tmp_path), 1)
+    runtime = EngramRuntime(args, [module], [table], np.arange(7), 2, "cuda")
+    hidden = torch.randn(1, 1, 2, 32, device="cuda", dtype=torch.bfloat16)
+    req = SimpleNamespace(input_ids=torch.tensor([1, 3]), device_len=3, media=None)
+    batch = SimpleNamespace(input_ids=torch.tensor([6], device="cuda"), padded_reqs=[req], is_decode=True)
+    with torch.inference_mode():
+        with runtime.forward_host_ctx(batch, False):
+            expected = module(hidden).clone()
+        stream = torch.cuda.Stream()
+        stream.wait_stream(torch.cuda.current_stream())
+        with torch.cuda.stream(stream):
+            for _ in range(3):
+                module(hidden)
+        torch.cuda.current_stream().wait_stream(stream)
+        graph = torch.cuda.CUDAGraph()
+        with torch.cuda.graph(graph, stream=stream):
+            output = module(hidden)
+        with runtime.forward_host_ctx(batch, True):
+            graph.replay()
+        torch.testing.assert_close(output, expected)
+        req.media = [{"start": 2, "types": [0]}]
+        with runtime.forward_host_ctx(batch, True):
+            graph.replay()
+        torch.testing.assert_close(output, hidden, rtol=0, atol=0)
diff --git a/tests/models/test_deepseek_v41_model.py b/tests/models/test_deepseek_v41_model.py
new file mode 100644
index 000000000..e3d8d83f5
--- /dev/null
+++ b/tests/models/test_deepseek_v41_model.py
@@ -0,0 +1,285 @@
+"""V4.1 mHC passes each pre-mix to the following sublayer, including the head."""
+
+import json
+from pathlib import Path
+
+import pytest
+import torch
+from torch import nn
+
+from freetoken.models.deepseek_v41.model import Block, Transformer, make_identity_pre_mix
+
+
+def _hc_block():
+    block = Block.__new__(Block)
+    nn.Module.__init__(block)
+    block.dim, block.hc_mult = 2, 2
+    block.norm_eps, block.hc_eps, block.hc_sinkhorn_iters = 1e-20, 1e-6, 3
+    return block
+
+
+def test_shifted_pre_mix_and_final_collapse():
+    block = _hc_block()
+    block.attn_norm, block.ffn_norm, block.ffn = nn.Identity(), nn.Identity(), nn.Identity()
+    block.ffn = type("ImageAwareIdentity", (nn.Module,), {"forward": lambda self, x, mask: x})()
+    for name in ("hc_attn_fn", "hc_attn_scale", "hc_attn_base", "hc_ffn_fn", "hc_ffn_scale", "hc_ffn_base"):
+        setattr(block, name, None)
+    mix_iter = iter([
+        (torch.tensor([[[0., 1.]]]), torch.tensor([[[1., 2.]]]), torch.eye(2).view(1, 1, 2, 2)),
+        (torch.tensor([[[.25, .75]]]), torch.tensor([[[2., 3.]]]), torch.eye(2).view(1, 1, 2, 2)),
+    ])
+    block.hc_mixes = lambda *args: next(mix_iter)
+    h = torch.tensor([[[[1., 2.], [3., 4.]]]])
+    seen = []
+    result, next_pre = block._forward(h, make_identity_pre_mix(h, 2), None,
+                                      lambda x: seen.append(x.clone()) or x)
+    torch.testing.assert_close(seen[0], torch.tensor([[[1., 2.]]]))
+    torch.testing.assert_close(result, torch.tensor([[[[12., 20.], [20., 32.]]]]))
+    torch.testing.assert_close(block.hc_pre(result, next_pre), torch.tensor([[[18., 29.]]]))
+
+
+@pytest.mark.parametrize("device", ["cpu", "cuda"])
+def test_hc_sinkhorn_and_combination_match_float_reference(device):
+    if device == "cuda" and not torch.cuda.is_available():
+        pytest.skip("requires CUDA")
+    torch.manual_seed(7)
+    block = _hc_block()
+    x = torch.randn(2, 3, 2, 2, device=device, dtype=torch.bfloat16)
+    fn = torch.randn(8, 4, device=device)
+    scale = torch.tensor([.3, -.2, .5], device=device)
+    base = torch.randn(8, device=device)
+    pre, post, comb = block.hc_mixes(x, fn, scale, base)
+    normalized = x.flatten(-2).double()
+    projected = (normalized @ fn.double().T) / (normalized.square().mean(-1, keepdim=True) + block.norm_eps).sqrt()
+    pre_ref = (projected[..., :2] * scale[0].double() + base[:2].double()).sigmoid() + block.hc_eps
+    post_ref = 2 * (projected[..., 2:4] * scale[1].double() + base[2:4].double()).sigmoid()
+    comb_ref = (projected[..., 4:] * scale[2].double() + base[4:].double()).view(2, 3, 2, 2).softmax(-1) + block.hc_eps
+    comb_ref /= comb_ref.sum(-2, keepdim=True) + block.hc_eps
+    for _ in range(block.hc_sinkhorn_iters - 1):
+        comb_ref /= comb_ref.sum(-1, keepdim=True) + block.hc_eps
+        comb_ref /= comb_ref.sum(-2, keepdim=True) + block.hc_eps
+    for actual, expected in ((pre, pre_ref), (post, post_ref), (comb, comb_ref)):
+        torch.testing.assert_close(actual.double(), expected, atol=1e-6, rtol=1e-5)
+    value = torch.randn(2, 3, 2, device=device, dtype=torch.bfloat16)
+    reference = post_ref.unsqueeze(-1) * value.double().unsqueeze(-2)
+    reference += torch.einsum("...pq,...pd->...qd", comb_ref, x.double())
+    torch.testing.assert_close(block.hc_post(value, x, post, comb), reference.bfloat16(), atol=.02, rtol=.008)
+    torch.testing.assert_close(block.hc_pre(x, pre), (pre_ref.unsqueeze(-1) * x.double()).sum(-2).bfloat16())
+
+
+def test_target_builds_metadata_without_allocating_experts_or_engram_tables():
+    from freetoken.distributed import set_tp_info, try_get_tp_info
+    from freetoken.layers.quantization import finalize_quant, QuantKind
+    from freetoken.models.deepseek_v41.config import parse_config
+    from freetoken.models.deepseek_v41.model import DeepseekV41ForCausalLM
+    from freetoken.moe.offload_cache import iter_offload_moe_layers
+
+    if try_get_tp_info() is None:
+        set_tp_info(0, 1)
+    raw = json.loads((Path(__file__).parent / "fixtures/deepseek_v41_nvfp4_config.json").read_text())
+    with torch.device("meta"):
+        adapter = DeepseekV41ForCausalLM(parse_config(raw))
+    model = adapter._transformer
+    experts = list(iter_offload_moe_layers(adapter))
+    assert len(experts) == 40
+    assert all(expert is layer.ffn.experts for expert, layer in zip(experts, model.layers))
+    assert all(expert.quant_method.kind is QuantKind.NVFP4 for expert in experts)
+    assert finalize_quant(adapter) == 40
+    params = dict(model.named_parameters())
+    assert sum(p.numel() * p.element_size() for p in params.values()) == 12_190_720_448
+    assert "head.weight" in params
+    assert not any("hc_head" in name or ".ffn.experts." in name or ".engram.embed." in name for name in params)
+    assert len(model.layers) == 40 and model.vision is not None
+    assert params["layers.0.ffn.gate.bias_vl"].shape == (384,)
+    assert params["layers.1.engram.wkv.weight"].shape == (25600, 6144)
+    assert params["layers.20.attn.compressor.wkv.weight"].shape == (512, 5120)
+    assert "layers.21.attn.compressor.wkv.weight" not in params
+
+
+def _cuda_stack():
+    import numpy as np
+    from dataclasses import asdict
+
+    from freetoken.attention.dsv41_sparse import DSV41SparseAttnBackend
+    from freetoken.core import Context, get_global_ctx, set_global_ctx
+    from freetoken.distributed import set_tp_info, try_get_tp_info
+    from freetoken.kvcache.dsv41_cost_model import dsv41_pool_sizes
+    from freetoken.kvcache.dsv41_paged_pool import DSV41PagedKVCache
+    from freetoken.models.deepseek_v41.args import DeepseekV41Args
+    from freetoken.models.deepseek_v41.config import parse_config
+    from freetoken.models.deepseek_v41.engram import DiskEngramTable, EngramRuntime
+    from freetoken.models.deepseek_v41.model import DeepseekV41ForCausalLM
+    from freetoken.moe.offload_cache import OffloadMoeCache
+
+    if try_get_tp_info() is None:
+        set_tp_info(0, 1)
+    args = DeepseekV41Args(
+        n_layers=5, n_mtp_layers=0, compress_ratios=(0, 2, 2, 1, 1),
+        kv_source_layers=(1, 3), index_source_layers=(1, 3, 4),
+        candidate_source_layer=3, candidate_topk_blocks=2, candidate_block_size=2,
+        dim=64, n_heads=2, head_dim=32, rope_head_dim=16, q_lora_rank=32,
+        o_lora_rank=32, o_groups=2, window_size=8, index_n_heads=2, index_head_dim=32,
+        index_topk=3, moe_inter_dim=64, n_routed_experts=4, n_activated_experts=2,
+        vocab_size=128, hc_mult=2, engram_layer_ids=(1,), engram_num_embeddings=(17,),
+        engram_max_ngram_size=2, engram_vocab_size=17, engram_n_heads=1,
+        engram_head_dim=32, engram_compressed_vocab_size=128,
+        vision_n_layers=1, vision_dim=32, vision_n_heads=2, vision_inter_dim=32,
+        vision_patch_size=2, vision_downsample_ratio=2, image_token_id=127,
+    )
+    config = parse_config(asdict(args) | {"quantization_config": {"moe_quant_algo": "NVFP4"}})
+    with torch.device("cuda"):
+        model = DeepseekV41ForCausalLM(config)
+    torch.manual_seed(12)
+    state = {}
+    for name, param in model.state_dict().items():
+        if param.dtype == torch.float8_e8m0fnu:
+            value = torch.full(param.shape, 1 / 32, device="cuda").to(param.dtype)
+        elif "norm" in name and name.endswith("weight") or name.endswith(("q_weight", "k_weight")):
+            value = torch.ones(param.shape, dtype=param.dtype, device="cuda")
+        else:
+            value = (torch.randn(param.shape, device="cuda") * .12).to(param.dtype)
+        state[name] = value
+    model.load_state_dict(state)
+    pool = DSV41PagedKVCache(dsv41_pool_sizes(40, args, 1., P=8), args,
+                            torch.device("cuda"), P=8, n_scratch=3)
+    page_table = torch.empty(2, 64, dtype=torch.long, device="cuda")
+    for row in range(2):
+        page_table[row] = torch.arange(row * 64, (row + 1) * 64, device="cuda").view(-1, 8).flip(0).flatten()
+    pool.attach_page_table(page_table)
+    for base in range(0, 128, 8):
+        pool.bind_window_pages(base, base)
+    try:
+        ctx = get_global_ctx()
+    except AssertionError:
+        ctx = Context(page_size=8)
+        set_global_ctx(ctx)
+    ctx.kv_cache = pool
+    ctx.attn_backend = backend = DSV41SparseAttnBackend(config)
+    cache = OffloadMoeCache(args.n_layers, args.n_routed_experts, args.n_routed_experts,
+                           torch.device("cuda"), quant_format="nvfp4")
+    sources = {}
+    for name, shape, dtype in (
+        ("gate_up_packed", (4, 128, 32), torch.uint8),
+        ("gate_up_scale", (4, 128, 4), torch.float8_e4m3fn),
+        ("gate_up_global", (4, 128), torch.float16),
+        ("down_packed", (4, 64, 32), torch.uint8),
+        ("down_scale", (4, 64, 4), torch.float8_e4m3fn),
+        ("down_global", (4, 64), torch.float16),
+    ):
+        sources[name] = []
+        for _ in range(args.n_layers):
+            values = torch.randint(0, 256, shape, dtype=dtype) if dtype == torch.uint8 else torch.full(shape, .0625).to(dtype)
+            sources[name].append(values.pin_memory())
+    cache.set_bank_sources(sources)
+    cache.reset()
+    for layer in model._iter_offload_moe_layers():
+        layer.offload_cache = cache
+    modules = [layer.engram for layer in model._transformer.layers if layer.engram is not None]
+    weights = torch.randn(17, 32).to(torch.float8_e4m3fn).view(torch.uint8).numpy()
+    table = DiskEngramTable(weights, np.full((17, 1), 126, dtype=np.uint8))
+    model._engram_runtime = EngramRuntime(args, modules, [table], np.arange(128), 64, "cuda")
+    return model, ctx, backend, pool, cache
+
+
+def _make_batch(ids, *, cached=0, row=0, decode=False, media=None):
+    from freetoken.core import Batch, Req, SamplingParams
+
+    req = Req(torch.tensor(ids, dtype=torch.int32), row, cached, 4, row,
+              SamplingParams(), None, media=media)
+    batch = Batch([req], "decode" if decode else "prefill")
+    batch.padded_reqs = batch.reqs
+    batch.input_ids = req.input_ids[-1:].cuda() if decode else req.input_ids[cached:].cuda()
+    batch.positions = torch.tensor([len(ids) - 1], device="cuda") if decode else torch.arange(cached, len(ids), device="cuda")
+    batch.active_table_idx = torch.tensor([row], dtype=torch.long, device="cuda")
+    return batch
+
+
+def _forward_batch(model, ctx, backend, ids, **kwargs):
+    batch = _make_batch(ids, **kwargs)
+    backend.prepare_metadata(batch)
+    with torch.inference_mode(), ctx.forward_batch(batch), model.forward_host_ctx(batch, False):
+        result = model.forward()
+    torch.cuda.synchronize()
+    return result
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_cuda_whole_model_nvfp4_csa2_engram_prefill_decode_and_images():
+    from freetoken.models.deepseek_v41.image_processor import IMAGE, IMAGE_END, IMAGE_NEW_LINE, IMAGE_START
+
+    model, ctx, backend, pool, cache = _cuda_stack()
+    ids = list(range(5, 18))
+    full = _forward_batch(model, ctx, backend, ids, row=0)
+    _forward_batch(model, ctx, backend, ids[:-1], row=1)
+    decode = _forward_batch(model, ctx, backend, ids, cached=len(ids) - 1, row=1, decode=True)
+    assert full.dtype == decode.dtype == torch.float32
+    assert torch.isfinite(full).all() and torch.isfinite(decode).all()
+    torch.testing.assert_close(decode, full, atol=.025, rtol=.025)
+    image_ids = ids.copy()
+    image_ids[1:5] = [127] * 4
+    image = {"start": 1, "types": torch.tensor([IMAGE_START, IMAGE, IMAGE_NEW_LINE, IMAGE_END]),
+             "patches": torch.randn(4, 3, 2, 2), "n_vit_h": 2, "n_vit_w": 2}
+    image_result = _forward_batch(model, ctx, backend, image_ids, row=0, media=[image])
+    assert torch.isfinite(image_result).all()
+    assert "embeddings" in image and image["patches"] is None
+    _forward_batch(model, ctx, backend, image_ids[:3], row=1, media=[image])
+    chunked = _forward_batch(model, ctx, backend, image_ids, cached=3, row=1, media=[image])
+    torch.testing.assert_close(chunked, image_result, atol=.025, rtol=.025)
+    assert not torch.allclose(image_result, full, atol=.005, rtol=.005)
+    different_image = {"start": 1, "types": image["types"],
+                       "patches": torch.randn(4, 3, 2, 2), "n_vit_h": 2, "n_vit_w": 2}
+    changed = _forward_batch(model, ctx, backend, image_ids, row=0, media=[different_image])
+    assert not torch.equal(changed, image_result)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@pytest.mark.parametrize("target", ["cpu", "hybrid"])
+def test_cuda_whole_model_cpu_and_hybrid_nvfp4_decode(target):
+    from freetoken.moe.cpu_executor import CpuMoeExecutor, compiled_extension_supports
+
+    if not compiled_extension_supports("swiglu_clamp"):
+        pytest.skip("CPU extension needs the clamped SwiGLU epilogue")
+    model, ctx, backend, _, cache = _cuda_stack()
+    ids = list(range(5, 18))
+    reference = _forward_batch(model, ctx, backend, ids, row=0)
+    _forward_batch(model, ctx, backend, ids[:-1], row=1)
+    cache.decode_target = target
+    cache.cpu_layer_ids = frozenset(range(5)) if target == "cpu" else frozenset()
+    cache.hybrid_max_fetch = 1
+    executor = CpuMoeExecutor(cache, top_k=2, activation="swiglu_clamp",
+                              apply_router_weight_on_input=False, num_threads=2,
+                              max_tokens=2, device=torch.device("cuda"),
+                              swiglu_alpha=1.0, swiglu_limit=10.0)
+    cache.set_cpu_executor(executor)
+    result = _forward_batch(model, ctx, backend, ids, cached=len(ids)-1, row=1, decode=True)
+    assert torch.isfinite(result).all()
+    torch.testing.assert_close(result, reference, atol=.025, rtol=.025)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_cuda_whole_model_batched_requests_keep_engram_and_kv_separate():
+    from freetoken.core import Batch
+
+    model, ctx, backend, _, _ = _cuda_stack()
+    sequences = [list(range(5, 18)), list(range(30, 43))]
+    references = [_forward_batch(model, ctx, backend, ids, row=row) for row, ids in enumerate(sequences)]
+    requests = [_make_batch(ids[:-1], row=row).reqs[0] for row, ids in enumerate(sequences)]
+    batch = Batch(requests, "prefill")
+    batch.padded_reqs = requests
+    batch.input_ids = torch.cat([req.input_ids for req in requests]).cuda()
+    batch.positions = torch.arange(12, device="cuda").repeat(2)
+    batch.active_table_idx = torch.arange(2, dtype=torch.long, device="cuda")
+    backend.prepare_metadata(batch)
+    with torch.inference_mode(), ctx.forward_batch(batch), model.forward_host_ctx(batch, False):
+        model.forward()
+    requests = [_make_batch(ids, cached=12, row=row, decode=True).reqs[0] for row, ids in enumerate(sequences)]
+    batch = Batch(requests, "decode")
+    batch.padded_reqs = requests
+    batch.input_ids = torch.tensor([ids[-1] for ids in sequences], dtype=torch.int32, device="cuda")
+    batch.positions = torch.full((2,), 12, device="cuda")
+    batch.active_table_idx = torch.arange(2, dtype=torch.long, device="cuda")
+    backend.prepare_metadata(batch)
+    with torch.inference_mode(), ctx.forward_batch(batch), model.forward_host_ctx(batch, False):
+        result = model.forward()
+    torch.cuda.synchronize()
+    torch.testing.assert_close(result, torch.cat(references), atol=.025, rtol=.025)
diff --git a/tests/models/test_deepseek_v41_moe.py b/tests/models/test_deepseek_v41_moe.py
new file mode 100644
index 000000000..fabb9f86b
--- /dev/null
+++ b/tests/models/test_deepseek_v41_moe.py
@@ -0,0 +1,66 @@
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.models.deepseek_v41.moe import DSV41OffloadMoELayer, Gate, clamped_swiglu
+
+
+@pytest.mark.parametrize("score_func", ["softmax", "sigmoid", "sqrtsoftplus"])
+@pytest.mark.parametrize("topk,norm", [(1, True), (2, True), (2, False)])
+def test_router_uses_temperature_and_bias_only_for_selection(score_func, topk, norm):
+    args = SimpleNamespace(n_activated_experts=topk, score_func=score_func, gate_temp=2.0,
+                           norm_topk_prob=norm, route_scale=1.5, n_routed_experts=3,
+                           dim=2, vision_enabled=True)
+    gate = Gate(0, args)
+    gate.weight.data.copy_(torch.tensor([[1., 0.], [0., 1.], [-1., 1.]]))
+    gate.bias.data.copy_(torch.tensor([0., 10., 0.]))
+    gate.bias_vl.data.copy_(torch.tensor([0., 0., 10.]))
+    x = torch.tensor([[2., -1.], [2., -1.]], dtype=torch.bfloat16)
+    mask = torch.tensor([False, True])
+    scores = torch.tensor([[1., -.5, -1.5], [1., -.5, -1.5]])
+    if score_func == "softmax":
+        scores = scores.exp() / scores.exp().sum(-1, keepdim=True)
+    elif score_func == "sigmoid":
+        scores = 1 / (1 + (-scores).exp())
+    else:
+        scores = torch.logaddexp(torch.zeros_like(scores), scores).sqrt()
+    expected_ids = torch.tensor([[1, 0], [2, 0]])[:, :topk]
+    expected = scores.gather(1, expected_ids)
+    if norm and topk > 1:
+        expected = expected / (expected.sum(-1, keepdim=True) + 1e-20)
+    actual, ids = gate(x, mask)
+    assert torch.equal(ids, expected_ids)
+    torch.testing.assert_close(actual, expected * 1.5)
+
+
+def test_swiglu_clamps_gate_only_above_and_up_both_sides():
+    gate = torch.tensor([-12., 12., 12., 2.])
+    up = torch.tensor([12., -12., 2., 3.])
+    clipped_gate = torch.tensor([-12., 10., 10., 2.])
+    clipped_up = torch.tensor([10., -10., 2., 3.])
+    expected = clipped_gate / (1 + (-clipped_gate).exp()) * clipped_up
+    torch.testing.assert_close(clamped_swiglu(gate, up, 10.), expected)
+
+
+@pytest.mark.parametrize("strategy,decode_target", [("offload", "gpu"), ("cpu", "cpu"), ("hybrid", "hybrid")])
+def test_native_expert_method_receives_clamp_and_execution_policy(strategy, decode_target):
+    from freetoken.distributed import set_tp_info, try_get_tp_info
+    from freetoken.layers.quantization import QuantBackend, QuantKind, set_quant_backend
+
+    if try_get_tp_info() is None:
+        set_tp_info(0, 1)
+    set_quant_backend(QuantBackend.parse("moe.nvfp4=triton"))
+    args = SimpleNamespace(n_layers=1, index_source_layers=(), engram_layer_ids=(),
+                           n_routed_experts=2, n_activated_experts=1, dim=32, moe_inter_dim=32,
+                           norm_topk_prob=True, swiglu_limit=10.0)
+    layer = DSV41OffloadMoELayer(0, args, strategy=strategy, decode_target=decode_target)
+    method = layer.quant_method
+    assert method.kind is QuantKind.NVFP4 and method.kernel.name == "triton"
+    assert (method.cfg.activation, method.cfg.alpha, method.cfg.limit) == ("swiglu_clamp", 1.0, 10.0)
+    assert (layer.alpha, layer.limit) == (1.0, 10.0)
+    assert (method.cfg.strategy, method.cfg.decode_target) == (strategy, decode_target)
+    assert not method.scheme.has("input_scale")
+    assert set(method.layout()) == {"gate_up", "gate_up_scale", "gate_up_global",
+                                    "down", "down_scale", "down_global"}
+    assert not layer.state_dict()
diff --git a/tests/models/test_deepseek_v41_vision.py b/tests/models/test_deepseek_v41_vision.py
new file mode 100644
index 000000000..6cc55b93d
--- /dev/null
+++ b/tests/models/test_deepseek_v41_vision.py
@@ -0,0 +1,238 @@
+from types import SimpleNamespace
+import base64
+import io
+
+import pytest
+import torch
+from PIL import Image
+
+from freetoken.models.deepseek_v41.image_processor import (
+    IMAGE, IMAGE_END, IMAGE_NEW_LINE, IMAGE_START, image_token_types,
+    load_image, load_image_bytes, num_image_tokens, plan_image_grid,
+)
+from freetoken.models.deepseek_v41.vision import Aligner, ViT, merge_image_embeddings
+
+
+def _args(**overrides):
+    values = dict(vision_patch_size=2, vision_dim=8, vision_n_heads=2,
+                  vision_inter_dim=12, vision_n_layers=2, vision_rope_theta=10000.0,
+                  vision_downsample_ratio=2, dim=6, vision_min_pixels=16,
+                  vision_max_n_token=24, vision_max_wh_ratio=None)
+    return SimpleNamespace(**(values | overrides))
+
+
+@pytest.mark.parametrize("width,height", [(1, 10000), (10000, 1), (1, 1), (37, 53), (2000, 3000)])
+def test_image_resize_respects_token_budget(width, height):
+    args = _args()
+    h, w, pixels_h, pixels_w = plan_image_grid(width, height, args)
+    assert h > 0 and w > 0
+    assert pixels_h % args.vision_patch_size == pixels_w % args.vision_patch_size == 0
+    assert num_image_tokens(h, w) <= args.vision_max_n_token
+
+
+def test_image_patches_normalization_and_layout():
+    buffer = io.BytesIO()
+    Image.new("RGB", (8, 4), (255, 0, 0)).save(buffer, format="PNG")
+    encoded = base64.b64encode(buffer.getvalue()).decode()
+    patches, nh, nw, lh, lw = load_image({"url": "data:image/png;base64," + encoded}, _args())
+    assert (nh, nw, lh, lw) == (2, 4, 1, 2)
+    assert patches.shape == (8, 3, 2, 2)
+    assert patches.dtype == torch.float32
+    torch.testing.assert_close(patches[:, 0], torch.ones_like(patches[:, 0]))
+    torch.testing.assert_close(patches[:, 1:], -torch.ones_like(patches[:, 1:]))
+    assert image_token_types(lh, lw).tolist() == [IMAGE_START, IMAGE, IMAGE, IMAGE_NEW_LINE, IMAGE_END]
+
+
+@pytest.mark.parametrize("url", ["/etc/passwd", "file:///etc/passwd", "http://127.0.0.1/test", "http://[::1]/test"])
+def test_api_image_loader_rejects_local_sources(url):
+    with pytest.raises(ValueError):
+        load_image_bytes({"url": url})
+
+
+def test_aligner_channel_and_pixel_order():
+    torch.manual_seed(12)
+    args = _args()
+    aligner = Aligner(args).float()
+    values = torch.randn(3, 5, args.vision_dim)
+    rows = []
+    for h in range(0, 3, 2):
+        for w in range(0, 5, 2):
+            block = torch.zeros(args.vision_dim, 2, 2)
+            crop = values[h:h + 2, w:w + 2].permute(2, 0, 1)
+            block[:, :crop.shape[1], :crop.shape[2]] = crop
+            rows.append(block.flatten())
+    reference = aligner.w2(torch.nn.functional.gelu(aligner.w1(torch.stack(rows))))
+    torch.testing.assert_close(aligner(values.reshape(-1, args.vision_dim), 3, 5), reference)
+
+
+def test_vit_attention_matches_explicit_bidirectional_reference():
+    torch.manual_seed(20)
+    args = _args()
+    model = ViT(args).float()
+    patches = torch.randn(6, 3, 2, 2)
+    from freetoken.models.deepseek_v41.vision import apply_rotary, get_vision_cos_sin
+
+    x = model.patch_embed(patches)
+    cos, sin = get_vision_cos_sin(2, 3, model.rope_dim, model.rope_theta)
+    for block in model.blocks:
+        q, k, v = [t.reshape(6, args.vision_n_heads, -1).transpose(0, 1)
+                   for t in block.attn.wqkv(block.norm1(x)).chunk(3, -1)]
+        q = apply_rotary(q.transpose(0, 1), cos, sin).transpose(0, 1)
+        k = apply_rotary(k.transpose(0, 1), cos, sin).transpose(0, 1)
+        scores = q @ k.transpose(-1, -2) / block.attn.head_dim ** 0.5
+        attention = (scores.softmax(-1) @ v).transpose(0, 1).reshape(6, -1)
+        x = x + block.attn.wo(attention)
+        x = x + block.mlp(block.norm2(x))
+    torch.testing.assert_close(model(patches, 2, 3), model.norm(x), atol=1e-6, rtol=1e-5)
+
+
+def test_image_embedding_scatter_across_every_chunk_boundary():
+    class Tower:
+        def __init__(self):
+            self.calls = 0
+            self.vision = SimpleNamespace(patch_embed=SimpleNamespace(proj=SimpleNamespace(weight=torch.zeros(1))))
+            self.image_start = torch.full((3,), 10.0)
+            self.image_newline = torch.full((3,), 20.0)
+            self.image_end = torch.full((3,), 30.0)
+
+        def encode_image(self, patches, nh, nw):
+            self.calls += 1
+            return torch.arange(12, dtype=torch.float32).reshape(4, 3)
+
+    types = image_token_types(2, 2)
+    span = torch.stack([torch.full((3,), 10.0), *torch.arange(6.).reshape(2, 3),
+                        torch.full((3,), 20.0), *torch.arange(6., 12.).reshape(2, 3),
+                        torch.full((3,), 20.0), torch.full((3,), 30.0)])
+    expected = torch.cat((torch.zeros(2, 3), span, torch.zeros(2, 3)))
+    for boundary in range(1, expected.shape[0]):
+        model = Tower()
+        media = [dict(start=2, patches=torch.zeros(4, 3, 2, 2), n_vit_h=2, n_vit_w=2, types=types)]
+        outputs, masks = [], []
+        for start, stop in ((0, boundary), (boundary, expected.shape[0])):
+            req = SimpleNamespace(cached_len=start, extend_len=stop-start, media=media)
+            batch = SimpleNamespace(is_prefill=True, reqs=[req])
+            h, mask = merge_image_embeddings(model, batch, torch.zeros(stop-start, 3))
+            outputs.append(h)
+            masks.append(mask)
+        torch.testing.assert_close(torch.cat(outputs), expected)
+        assert torch.cat(masks).tolist() == [False] * 2 + [True] * types.numel() + [False] * 2
+        assert model.calls == 1
+        assert media[0]["patches"] is None
+        assert media[0]["types"] is types
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+def test_vit_bf16_cuda_matches_cpu():
+    torch.manual_seed(31)
+    cpu = ViT(_args()).float().eval()
+    gpu = ViT(_args()).to(device="cuda", dtype=torch.bfloat16).eval()
+    gpu.load_state_dict(cpu.state_dict())
+    patches = torch.randn(6, 3, 2, 2)
+    with torch.no_grad():
+        expected = cpu(patches, 2, 3)
+        actual = gpu(patches.to(device="cuda", dtype=torch.bfloat16), 2, 3)
+    torch.testing.assert_close(actual.float().cpu(), expected, rtol=0.04, atol=0.025)
+
+
+def test_vision_parameters_preserve_checkpoint_dtypes():
+    vision = ViT(_args())
+    aligner = Aligner(_args())
+    for name, parameter in vision.named_parameters():
+        assert parameter.dtype == (torch.float32 if "norm" in name else torch.bfloat16)
+    assert all(p.dtype == torch.bfloat16 for p in aligner.parameters())
+
+
+def _processor():
+    from freetoken.mm.config import MultimodalConfig
+    from freetoken.models.deepseek_v41.mm_processor import DeepseekV41MMProcessor
+
+    config = {"image_token_id": 99, "vision_config": {
+        "num_hidden_layers": 2, "hidden_size": 8, "num_attention_heads": 2,
+        "intermediate_size": 12, "patch_size": 2, "downsample_ratio": 2,
+        "min_pixels": 16, "max_image_tokens": 24,
+    }}
+    return DeepseekV41MMProcessor(config, "unused", MultimodalConfig())
+
+
+def test_native_media_and_shared_processor_produce_identical_content_keys():
+    from freetoken.models.deepseek_v41.image_processor import ImageInput
+
+    processor = _processor()
+    buffer = io.BytesIO()
+    Image.new("RGB", (8, 4), (255, 0, 0)).save(buffer, format="PNG")
+    raw = buffer.getvalue()
+    result = processor.apply(torch.tensor([11, 99, 12], dtype=torch.int32), [raw])
+    patches, nh, nw, lh, lw = load_image({"data": raw}, processor.args)
+    types = image_token_types(lh, lw)
+    legacy_ids = torch.tensor([11] + [99] * len(types) + [12], dtype=torch.int32)
+    converted = processor.from_media(legacy_ids, [ImageInput(1, patches, nh, nw, types)])
+    assert torch.equal(converted.input_ids, result.input_ids)
+    assert converted.mm_items[0].hash == result.mm_items[0].hash
+    assert torch.equal(converted.mm_items[0].feature, result.mm_items[0].feature)
+    assert result.mm_items[0].offsets == [[1, 6]]
+    assert result.mm_items[0].types == types.tolist()
+    assert result.mrope_positions is None and result.mrope_delta == 0
+    changed_grid = processor._item(patches, nw, nh, types, 1)
+    assert changed_grid.hash != result.mm_items[0].hash
+    assert legacy_ids.tolist() == [11, 99, 99, 99, 99, 99, 12]
+
+
+def test_precomputed_mm_rows_scatter_with_image_router_mask():
+    hidden = torch.zeros(5, 3)
+    embeddings = torch.tensor([[1., 2., 3.], [4., 5., 6.]])
+    batch = SimpleNamespace(is_prefill=True, reqs=[SimpleNamespace()],
+                            mm_embeds=embeddings, mm_rows=torch.tensor([1, 3]))
+    actual, mask = merge_image_embeddings(None, batch, hidden)
+    torch.testing.assert_close(actual[[1, 3]], embeddings)
+    assert actual[[0, 2, 4]].count_nonzero() == 0
+    assert mask.tolist() == [False, True, False, True, False]
+
+
+def test_streamer_adapter_rebinds_native_parameters_without_changing_keys():
+    from freetoken.models.deepseek_v41.vision import _ModuleBlockAdapter
+    from freetoken.models.weight_stream import _slots
+
+    block = ViT(_args()).blocks[0]
+    names = set(block.state_dict())
+    adapter = _ModuleBlockAdapter(block)
+    slots, size = _slots(adapter)
+    row = torch.zeros(size, dtype=torch.uint8)
+    for slot in slots:
+        value = slot.view(row)
+        value.fill_(.25)
+        setattr(slot.owner, slot.attr, value)
+    assert set(block.state_dict()) == names
+    assert all(torch.all(p == .25) for p in block.parameters())
+    assert all(p.untyped_storage().data_ptr() == row.untyped_storage().data_ptr() for p in block.parameters())
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA")
+@torch.inference_mode()
+def test_native_vision_host_streaming_matches_resident_across_repeated_images(monkeypatch):
+    torch.manual_seed(71)
+    vision = ViT(_args(vision_n_layers=3)).cuda().eval()
+    patches = [torch.randn(6, 3, 2, 2, device="cuda", dtype=torch.bfloat16) for _ in range(2)]
+    expected = [vision(patch, 2, 3).clone() for patch in patches]
+    weights = {name: p.cpu().clone() for name, p in vision.named_parameters()}
+    vision.place_weights("host")
+    streamer = vision._streamer
+    assert streamer.bank.is_pinned() and streamer.staging.shape[0] == 2
+    assert all(p.device.type == "cpu" for block in vision.blocks for p in block.parameters())
+    assert vision.patch_embed.proj.weight.is_cuda
+    with monkeypatch.context() as mp:
+        def fail_block(*args):
+            raise RuntimeError("interrupted vision block")
+
+        mp.setattr(vision.blocks[1], "forward", fail_block)
+        with pytest.raises(RuntimeError, match="interrupted vision block"):
+            vision(patches[0], 2, 3)
+    assert all(p.device.type == "cpu" for block in vision.blocks for p in block.parameters())
+    for _ in range(2):
+        for patch, reference in zip(patches, expected):
+            torch.testing.assert_close(vision(patch, 2, 3), reference, rtol=0, atol=0)
+    assert all(p.device.type == "cpu" for block in vision.blocks for p in block.parameters())
+    for name, p in vision.named_parameters():
+        torch.testing.assert_close(p.cpu(), weights[name], rtol=0, atol=0)
+    vision.place_weights("gpu")
+    assert vision._streamer is None and all(p.is_cuda for p in vision.parameters())
+    torch.testing.assert_close(vision(patches[0], 2, 3), expected[0], rtol=0, atol=0)
diff --git a/tests/models/test_deepseek_v41_weight.py b/tests/models/test_deepseek_v41_weight.py
new file mode 100644
index 000000000..a9bc267bf
--- /dev/null
+++ b/tests/models/test_deepseek_v41_weight.py
@@ -0,0 +1,135 @@
+"""Small native NVFP4 shards exercise loading without the released weight payloads."""
+
+import json
+from types import SimpleNamespace
+
+import pytest
+import safetensors.torch
+import torch
+
+from freetoken.layers.quantization import MoEConfig, Nvfp4MoEMethod, QuantKind
+from freetoken.layers.quantization.scheme import nvfp4_scheme
+from freetoken.models.deepseek_v41 import weight
+from freetoken.moe.expert_banks import build_expert_banks
+from freetoken.moe.expert_pieces import iter_expert_pieces
+
+
+@pytest.fixture(params=[True, False], ids=["modelopt_input_scale", "native_no_input_scale"])
+def tiny_shards(tmp_path, request):
+    config = SimpleNamespace(num_layers=1, num_experts=2, hidden_size=32, moe_intermediate_size=32,
+                             architectures=["DeepseekV41ForCausalLM"])
+    tensors, globals_ = {}, {}
+    for expert in range(2):
+        for proj_idx, proj in enumerate(("w1", "w2", "w3"), 1):
+            name = f"layers.0.ffn.experts.{expert}.{proj}"
+            tensors[name + ".weight"] = torch.full((32, 16), expert * 16 + proj_idx, dtype=torch.uint8)
+            tensors[name + ".weight_scale"] = torch.full((32, 2), proj_idx, dtype=torch.float32).to(torch.float8_e4m3fn)
+            globals_[name + ".weight_scale_2"] = torch.tensor((expert + 1) * proj_idx / 8)
+            if request.param:
+                globals_[name + ".input_scale"] = torch.tensor(99.0)
+    globals_["mtp.0.ffn.experts.0.w1.weight"] = torch.ones(1, 1, dtype=torch.uint8)
+    globals_["layers.0.engram.embed.weight"] = (torch.ones(1, 32).to(torch.float8_e4m3fn) if request.param
+                                               else torch.ones(1, 16, dtype=torch.uint8))
+    globals_["layers.0.engram.embed.scale"] = torch.ones(1, 1).to(torch.float8_e8m0fnu)
+    safetensors.torch.save_file(tensors, tmp_path / "bulk.safetensors")
+    safetensors.torch.save_file(globals_, tmp_path / "global.safetensors")
+    index = {key: "bulk.safetensors" for key in tensors} | {key: "global.safetensors" for key in globals_}
+    (tmp_path / "model.safetensors.index.json").write_text(json.dumps({"weight_map": index}))
+    return tmp_path, config
+
+
+def _load_banks(folder, config, *, parallel=False, layer_sink=None):
+    method = Nvfp4MoEMethod(MoEConfig(num_experts=config.num_experts, hidden=config.hidden_size,
+                                     intermediate=config.moe_intermediate_size, top_k=1,
+                                     scheme=nvfp4_scheme(input_scale=False), strategy="offload",
+                                     activation="swiglu_clamp", alpha=1.0, limit=10.0), "triton")
+    pieces = iter_expert_pieces(str(folder), config, QuantKind.NVFP4, parallel=parallel,
+                                workers=2, chunk=4096)
+    return build_expert_banks(method, config.num_layers, pieces, device=torch.device("cpu"),
+                              layer_sink=layer_sink).sources
+
+
+def test_validate_and_load_split_global_scales(tiny_shards):
+    folder, config = tiny_shards
+    completed = []
+    banks = _load_banks(folder, config, layer_sink=lambda layer, banks: completed.append(layer))
+    assert completed == [0]
+    for expert in range(2):
+        assert torch.all(banks["gate_up"][0][expert, :32] == expert * 16 + 1)
+        assert torch.all(banks["gate_up"][0][expert, 32:] == expert * 16 + 3)
+        assert torch.all(banks["down"][0][expert] == expert * 16 + 2)
+        assert torch.all(banks["gate_up_scale"][0][expert, :32].float() == 1)
+        assert torch.all(banks["gate_up_scale"][0][expert, 32:].float() == 3)
+        assert torch.all(banks["down_scale"][0][expert].float() == 2)
+        assert torch.all(banks["gate_up_global"][0][expert, :32] == (expert + 1) / 8)
+        assert torch.all(banks["gate_up_global"][0][expert, 32:] == (expert + 1) * 3 / 8)
+        assert torch.all(banks["down_global"][0][expert] == (expert + 1) / 4)
+
+
+def test_serial_parallel_bank_bytes_match(tiny_shards):
+    folder, config = tiny_shards
+    sink = lambda layer, banks: None
+    serial = _load_banks(folder, config, layer_sink=sink)
+    parallel = _load_banks(folder, config, parallel=True, layer_sink=sink)
+    for name in serial:
+        assert torch.equal(serial[name][0].view(torch.uint8), parallel[name][0].view(torch.uint8)), name
+
+
+def test_wrong_expert_kind_does_not_decode_nvfp4_as_another_format(tiny_shards):
+    folder, config = tiny_shards
+    assert weight.iter_expert_pieces(str(folder), config, QuantKind.MXFP4) is None
+
+
+def test_expert_stream_validates_headers_before_creating_iterator(tiny_shards, monkeypatch):
+    folder, config = tiny_shards
+    headers = weight.read_checkpoint_headers(folder)
+    del headers["layers.0.ffn.experts.0.w1.weight"]
+    monkeypatch.setattr(weight, "read_checkpoint_headers", lambda folder: headers)
+    with pytest.raises(ValueError, match="Missing NVFP4 tensor"):
+        weight.iter_expert_pieces(str(folder), config, QuantKind.NVFP4)
+
+
+@pytest.mark.parametrize("mutation,match", [
+    ("missing", "Missing NVFP4 tensor"), ("dtype", "Malformed NVFP4 tensor"),
+    ("shape", "Malformed NVFP4 tensor"), ("expert", "outside configured backbone"),
+])
+def test_bad_expert_headers_fail_before_allocation(tiny_shards, mutation, match):
+    folder, config = tiny_shards
+    headers = weight.read_checkpoint_headers(folder)
+    key = "layers.0.ffn.experts.0.w1.weight_scale"
+    if mutation == "missing":
+        del headers[key]
+    elif mutation == "dtype":
+        headers[key]["dtype"] = "F8_E8M0"
+    elif mutation == "shape":
+        headers[key]["shape"] = [32, 1]
+    else:
+        headers[key.replace("experts.0", "experts.2")] = headers.pop(key)
+    with pytest.raises(ValueError, match=match):
+        weight.validate_expert_headers(headers, config)
+
+
+@pytest.mark.parametrize("include_vision", [False, True])
+def test_resident_loader_preserves_native_keys_and_excludes_external_tables(tiny_shards, include_vision):
+    folder, _ = tiny_shards
+    config = {"n_layers": 1, "compress_ratios": [0], "kv_source_layers": [],
+              "index_source_layers": [], "candidate_source_layer": -1,
+              "engram_layer_ids": [], "engram_num_embeddings": [], "vision_n_layers": 1}
+    (folder / "config.json").write_text(json.dumps(config))
+    tensors = {"head.weight": torch.arange(1024).view(32, 32).to(torch.bfloat16),
+               "layers.0.attn.wo_a.weight": torch.ones(32, 32).to(torch.float8_e4m3fn),
+               "layers.0.attn.wo_a.scale": torch.tensor([[128]], dtype=torch.uint8).view(torch.float8_e8m0fnu),
+               "vision.patch_embed.proj.weight": torch.ones(2, 3, dtype=torch.bfloat16),
+               "aligner.w1.weight": torch.ones(2, 3, dtype=torch.bfloat16),
+               "image_start": torch.ones(32, dtype=torch.bfloat16)}
+    safetensors.torch.save_file(tensors, folder / "resident.safetensors")
+    index = weight._weight_map(folder)
+    index.update({name: "resident.safetensors" for name in tensors})
+    (folder / "model.safetensors.index.json").write_text(json.dumps({"weight_map": index}))
+    resident = dict(weight.iter_weights(str(folder), "cpu", include_moe_experts=False, include_vision=include_vision))
+    expected = {"head.weight", "layers.0.attn.wo_a"}
+    if include_vision:
+        expected |= {"vision.patch_embed.proj.weight", "aligner.w1.weight", "image_start"}
+    assert set(resident) == expected
+    assert torch.all(resident["layers.0.attn.wo_a"] == 2)
+    assert torch.equal(resident["head.weight"], tensors["head.weight"])
diff --git a/tests/models/test_dsv41_attention.py b/tests/models/test_dsv41_attention.py
new file mode 100644
index 000000000..cba40f83d
--- /dev/null
+++ b/tests/models/test_dsv41_attention.py
@@ -0,0 +1,271 @@
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.attention.dsv4_sparse import DSV4AttnMetadata
+from freetoken.attention.dsv41_sparse import DSV41SparseAttnBackend
+from freetoken.core import Context, get_global_ctx, set_global_ctx
+from freetoken.kernel.triton.dsv41.quant import fp4_roundtrip, fp8_roundtrip, unpack_fp4, unpack_fp8
+from freetoken.kvcache.dsv41_cost_model import dsv41_pool_sizes
+from freetoken.kvcache.dsv41_paged_pool import DSV41PagedKVCache
+from freetoken.models.deepseek_v41.attention import Attention, apply_rope
+
+
+def _args():
+    return SimpleNamespace(n_layers=5, compress_ratios=(0, 2, 2, 1, 1),
+        kv_source_layers=(1, 3), index_source_layers=(1, 3, 4),
+        candidate_source_layer=3, candidate_topk_blocks=2, candidate_block_size=2,
+        dim=64, n_heads=2, head_dim=32, rope_head_dim=16, q_lora_rank=32,
+        o_lora_rank=32, o_groups=2, window_size=8, index_n_heads=2, index_head_dim=32,
+        index_topk=3, norm_eps=1e-20, rope_theta=10000., compress_rope_theta=160000.,
+        original_seq_len=65536, rope_factor=16., beta_fast=32, beta_slow=1)
+
+
+def _stack(device="cpu", weights=None, kv_quant="none"):
+    args = _args()
+    pool = DSV41PagedKVCache(dsv41_pool_sizes(40, args, 1, P=8), args, torch.device(device),
+                           P=8, n_scratch=3, kv_quant=kv_quant)
+    table = torch.empty(2, 64, dtype=torch.long, device=device)
+    for row in range(2):
+        table[row] = torch.arange(row * 64, (row + 1) * 64, device=device).view(-1, 8).flip(0).flatten()
+    pool.attach_page_table(table)
+    for base in range(0, 128, 8):
+        pool.bind_window_pages(base, base)
+    try:
+        ctx = get_global_ctx()
+    except AssertionError:
+        ctx = Context(page_size=8)
+        set_global_ctx(ctx)
+    ctx.kv_cache = pool
+    ctx.attn_backend = backend = DSV41SparseAttnBackend(SimpleNamespace(dsv41_args=args))
+    layers = [Attention(i, args).to(device) for i in range(args.n_layers)]
+    if weights is None:
+        torch.manual_seed(8)
+        for layer in layers:
+            for name, p in layer.named_parameters():
+                if p.dtype == torch.float8_e8m0fnu:
+                    p.data.copy_(torch.full(p.shape, 1 / 16, device=device).to(p.dtype))
+                elif "norm.weight" in name:
+                    p.data.fill_(1)
+                else:
+                    p.data.copy_((torch.randn(p.shape, device=device) * .15).to(p.dtype))
+    else:
+        for layer, state in zip(layers, weights):
+            layer.load_state_dict(state)
+    for layer in layers:
+        layer.bind(pool, torch.device(device))
+    return ctx, backend, pool, layers
+
+
+def _reference(layers, inputs):
+    out = []
+    length = inputs.shape[1]
+    positions = torch.arange(length, device=inputs.device)
+    shared_kv = shared_keys = shared_picks = candidates = None
+    for layer, x in zip(layers, inputs):
+        qr = layer.q_norm(layer.wq_a(x))
+        q = apply_rope(layer.wq_b(qr).unflatten(-1, (layer.n_heads, layer.head_dim)), positions, layer.inv_freq)
+        kv = fp8_roundtrip(apply_rope(layer.kv_norm(layer.wkv(x)), positions, layer.inv_freq), block_size=32)
+        ratio = layer.ratio
+        if layer.is_kv_source:
+            comp = layer.compressor
+            if ratio == 2:
+                count = length // 2 * 2
+                raw = comp.wkv(x[:count].float()).view(-1, 2, layer.head_dim)
+                scores = comp.wgate(x[:count].float()).view_as(raw)
+                latent = comp.norm((raw * scores.softmax(1)).sum(1).to(x.dtype))
+            else:
+                latent = comp.norm(comp.wkv(x))
+            compressed_pos = torch.arange(length // ratio, device=x.device) * ratio
+            idx = layer.indexer
+            shared_keys = fp4_roundtrip(apply_rope(idx.k_norm(idx.wk(latent)), compressed_pos, layer.inv_freq),
+                                        block_size=32, scale_format="e8m0")
+            shared_kv = fp4_roundtrip(apply_rope(latent, compressed_pos, layer.inv_freq),
+                                      block_size=16, scale_format="e4m3")
+        if layer.is_index_source:
+            idx = layer.indexer
+            iq = fp4_roundtrip(apply_rope(idx.wq_b(qr).unflatten(-1, (idx.n_heads, idx.head_dim)),
+                                          positions, layer.inv_freq), block_size=32, scale_format="e8m0")
+            weights = idx.weights_proj(x) * idx.scale
+            dots = torch.einsum("qhd,kd->qhk", iq.float(), shared_keys.float()).to(iq.dtype)
+            score = (dots.relu() * weights[..., None]).sum(1)
+            visible = (positions + 1) // ratio
+            score.masked_fill_(torch.arange(shared_keys.shape[0], device=x.device)[None] >= visible[:, None], -torch.inf)
+            if layer.layer_id == layer.args.candidate_source_layer:
+                block_size = layer.args.candidate_block_size
+                blocks = torch.nn.functional.pad(score, (0, -score.shape[-1] % block_size), value=-torch.inf)
+                blocks = blocks.unflatten(-1, (-1, block_size)).amax(-1)
+                newest = (visible - 1) // block_size
+                blocks.masked_fill_(torch.arange(blocks.shape[-1], device=x.device)[None] == newest[:, None], torch.inf)
+                selected = blocks.argsort(dim=-1, descending=True, stable=True)[..., :layer.args.candidate_topk_blocks]
+                keep = torch.zeros_like(blocks, dtype=torch.bool).scatter_(-1, selected, blocks.gather(-1, selected) > -torch.inf)
+                candidates = keep.repeat_interleave(block_size, -1)[:, :score.shape[-1]]
+            elif layer.layer_id > layer.args.candidate_source_layer:
+                score.masked_fill_(~candidates, -torch.inf)
+            shared_picks = score.argsort(dim=-1, descending=True, stable=True)[..., :layer.args.index_topk].sort(-1).values
+            shared_picks = torch.where(shared_picks < visible[:, None], shared_picks, -1)
+        token_outputs = []
+        for position in range(length):
+            selected = kv[max(0, position - layer.window_size + 1):position + 1]
+            if ratio:
+                indices = shared_picks[position]
+                selected = torch.cat((selected, shared_kv[indices[indices >= 0]]), 0)
+            logits = q[position].float() @ selected.float().T * layer.softmax_scale
+            probs = torch.cat((logits, layer.attn_sink[:, None]), -1).softmax(-1)[:, :-1]
+            token_outputs.append((probs @ selected.float()).to(x.dtype))
+        value = apply_rope(torch.stack(token_outputs), positions, layer.inv_freq, inverse=True)
+        grouped = value.view(length, layer.n_groups, -1)
+        projected = torch.einsum("tgd,grd->tgr", grouped, layer.wo_a.view(layer.n_groups, layer.o_lora_rank, -1))
+        out.append(layer.wo_b(projected.flatten(1)))
+    return torch.stack(out)
+
+
+@pytest.mark.parametrize("device", ["cpu", pytest.param("cuda", marks=pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required"))])
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_source_sharing_and_two_stage_attention_match_direct_reference(device, kv_quant):
+    _, _, _, layers = _stack(device, kv_quant=kv_quant)
+    torch.manual_seed(2)
+    inputs = torch.randn(5, 17, 64, device=device).bfloat16()
+    expected = _reference(layers, inputs)
+    actual = torch.stack([layer.forward_ragged(x[None], [(0, 17, 0, 0)], torch.arange(17, device=device))[0]
+                          for layer, x in zip(layers, inputs)])
+    torch.testing.assert_close(actual, expected, rtol=.02 if device == "cuda" else 0, atol=.002 if device == "cuda" else 0)
+
+
+@pytest.mark.parametrize("cut", [1, 2, 7, 8, 9, 16])
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_prefill_chunks_preserve_partial_pairs_and_shared_sources(cut, kv_quant):
+    _, _, _, layers = _stack(kv_quant=kv_quant)
+    weights = [layer.state_dict() for layer in layers]
+    torch.manual_seed(31)
+    inputs = torch.randn(5, 19, 64).bfloat16()
+    expected = torch.stack([layer.forward_ragged(x[None], [(0, 19, 0, 0)], torch.arange(19))[0]
+                            for layer, x in zip(layers, inputs)])
+    _, _, _, layers = _stack(weights=weights, kv_quant=kv_quant)
+    first = [layer.forward_ragged(x[None, :cut], [(0, cut, 0, 0)], torch.arange(cut))[0]
+             for layer, x in zip(layers, inputs)]
+    second = [layer.forward_ragged(x[None, cut:], [(0, 19 - cut, 0, cut)], torch.arange(cut, 19))[0]
+              for layer, x in zip(layers, inputs)]
+    actual = torch.stack([torch.cat((a, b)) for a, b in zip(first, second)])
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("kv_quant", ["none", "fp8-fp4"])
+def test_batched_decode_is_request_isolated_and_matches_full_prefill(kv_quant):
+    _, _, _, layers = _stack(kv_quant=kv_quant)
+    weights = [layer.state_dict() for layer in layers]
+    torch.manual_seed(34)
+    inputs = torch.randn(2, 5, 17, 64).bfloat16()
+    expected = [_reference(layers, sample) for sample in inputs]
+    ctx, backend, pool, layers = _stack(weights=weights, kv_quant=kv_quant)
+    for row in range(2):
+        for layer, x in zip(layers, inputs[row]):
+            layer.forward_ragged(x[None, :7], [(0, 7, row, 0)], torch.arange(7))
+    actual = [[] for _ in range(2)]
+    for position in range(7, 17):
+        md = DSV4AttnMetadata(last_indices=torch.arange(2), full_snap=pool.full_loc_map.clone(), window_ar=torch.arange(8))
+        batch = SimpleNamespace(attn_metadata=md)
+        with ctx.forward_batch(batch):
+            rows, positions = torch.arange(2), torch.full((2,), position)
+            step = [layer.decode_step(inputs[:, i, position:position + 1], positions, rows, position)[:, 0]
+                    for i, layer in enumerate(layers)]
+        stacked = torch.stack(step)
+        for row in range(2):
+            actual[row].append(stacked[:, row])
+    for row in range(2):
+        torch.testing.assert_close(torch.stack(actual[row], 1), expected[row][:, 7:], rtol=0, atol=0)
+
+
+@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required")
+def test_cuda_packed_prefill_and_batched_decode_match_bf16_across_boundaries():
+    torch.manual_seed(47)
+    inputs = torch.randn(2, 5, 18, 64, device="cuda").bfloat16()
+    starts, steps = (7, 15), 3
+
+    def run(kv_quant, weights=None):
+        ctx, backend, pool, layers = _stack("cuda", weights=weights, kv_quant=kv_quant)
+        segments = [(0, starts[0], 0, 0), (starts[0], starts[1], 1, 0)]
+        positions = torch.cat([torch.arange(length, device="cuda") for length in starts])
+        selected, prefill = [], []
+
+        def record_indices(layer):
+            if layer.ratio:
+                selected.append({key: value.clone() for key, value in backend.shared_indices.items()})
+
+        for i, layer in enumerate(layers):
+            x = torch.cat([inputs[row, i, :length] for row, length in enumerate(starts)])
+            prefill.append(layer.forward_ragged(x[None], segments, positions)[0])
+            record_indices(layer)
+        prefill = torch.stack(prefill)
+        outputs = [[prefill[:, offset:offset + length]]
+                   for offset, length in ((0, starts[0]), (starts[0], starts[1]))]
+        rows = torch.arange(2, device="cuda")
+        for step in range(steps):
+            positions = torch.tensor([start + step for start in starts], device="cuda")
+            md = DSV4AttnMetadata(last_indices=rows, full_snap=pool.full_loc_map.clone(),
+                                 window_ar=torch.arange(8, device="cuda"))
+            decoded = []
+            with ctx.forward_batch(SimpleNamespace(attn_metadata=md)):
+                for i, layer in enumerate(layers):
+                    x = torch.stack([inputs[row, i, start + step] for row, start in enumerate(starts)])
+                    decoded.append(layer.decode_step(x[:, None], positions, rows, max(starts) + step)[:, 0])
+                    record_indices(layer)
+            decoded = torch.stack(decoded)
+            for row in range(2):
+                outputs[row].append(decoded[:, row:row + 1])
+        return [torch.cat(parts, 1) for parts in outputs], selected, layers
+
+    baseline, baseline_indices, layers = run("none")
+    weights = [layer.state_dict() for layer in layers]
+    for row, start in enumerate(starts):
+        expected = _reference(layers, inputs[row, :, :start + steps])
+        torch.testing.assert_close(baseline[row], expected, rtol=.02, atol=.002)
+    packed, packed_indices, _ = run("fp8-fp4", weights)
+    for actual, expected in zip(packed, baseline):
+        torch.testing.assert_close(actual, expected, rtol=0, atol=0)
+    assert len(packed_indices) == len(baseline_indices)
+    for actual, expected in zip(packed_indices, baseline_indices):
+        assert actual.keys() == expected.keys()
+        for key in expected:
+            torch.testing.assert_close(actual[key], expected[key], rtol=0, atol=0)
+
+
+def test_packed_attention_quantizes_raw_values_at_all_three_write_points(monkeypatch):
+    import freetoken.models.deepseek_v41.attention as attention_module
+
+    _, backend, pool, layers = _stack(kv_quant="fp8-fp4")
+    layer = layers[1]
+    torch.manual_seed(71)
+    x = torch.randn(3, 64).bfloat16()
+    positions = torch.tensor([0, 2, 4])
+    raw_window = apply_rope(layer.kv_norm(layer.wkv(x)), positions, layer.inv_freq)
+    original_fp8, original_fp4 = attention_module.pack_fp8, attention_module.pack_fp4
+    calls = []
+
+    def pack_window(value, block_size):
+        torch.testing.assert_close(value, raw_window, rtol=0, atol=0)
+        calls.append("window")
+        return original_fp8(value, block_size)
+
+    monkeypatch.setattr(attention_module, "pack_fp8", pack_window)
+    _, _, window = layer._project(x, positions)
+    backend.store_window(window, layer.layer_id, torch.tensor([1, 3, 5]))
+    torch.testing.assert_close(unpack_fp8(pool.window_pool[1][[1, 3, 5]]), fp8_roundtrip(raw_window), rtol=0, atol=0)
+    latent = torch.randn(3, 32).bfloat16()
+    raw_keys = apply_rope(layer.indexer.k_norm(layer.indexer.wk(latent)), positions, layer.inv_freq)
+    raw_compressed = apply_rope(latent.clone(), positions, layer.inv_freq)
+
+    def pack_compressed(value, block_size, scale_format):
+        expected = raw_keys if scale_format == "e8m0" else raw_compressed
+        torch.testing.assert_close(value, expected, rtol=0, atol=0)
+        calls.append(scale_format)
+        return original_fp4(value, block_size, scale_format)
+
+    monkeypatch.setattr(attention_module, "pack_fp4", pack_compressed)
+    layer._publish(latent, positions, torch.tensor([1, 3, 5]), torch.tensor([1, 3, 5]))
+    assert calls == ["window", "e8m0", "e4m3"]
+    torch.testing.assert_close(unpack_fp4(pool.idx_pool[1][[1, 3, 5]], 32, "e8m0"),
+                               fp4_roundtrip(raw_keys, 32, "e8m0"), rtol=0, atol=0)
+    torch.testing.assert_close(unpack_fp4(pool.cmp_pool[1][[1, 3, 5]], 16, "e4m3"),
+                               fp4_roundtrip(raw_compressed, 16, "e4m3"), rtol=0, atol=0)
diff --git a/tests/models/test_glm5_next_config.py b/tests/models/test_glm5_next_config.py
index ad56f7e14..388693beb 100644
--- a/tests/models/test_glm5_next_config.py
+++ b/tests/models/test_glm5_next_config.py
@@ -23,6 +23,28 @@
 _KDA_IDS = tuple(i for i in range(_NUM_LAYERS) if i not in _DSA_IDS)
 
 
+def test_nvfp4_pool_factory_preserves_glm5_hybrid_geometry():
+    import torch
+    from freetoken.kvcache import create_kvcache_pool
+    from freetoken.kvcache.dsa_pool import KpoolDSAKVCache
+
+    cfg = parse_config(_hf_config())
+    pool = create_kvcache_pool(
+        cfg, num_pages=2, page_size=64, dtype=torch.bfloat16,
+        device=torch.device("cpu"), num_req_slots=3, kv_quant="nvfp4",
+    )
+    assert isinstance(pool, KpoolDSAKVCache)
+    spec, = cfg.kv_cache_group_specs()
+    assert pool.num_layers == len(_DSA_IDS)
+    for layer in _DSA_IDS:
+        assert pool.latent_rows(layer).shape == (128, spec.head_dim // 2)
+        assert pool.latent_block_scale(layer).shape == (128, spec.head_dim // 16)
+    with pytest.raises(KeyError):
+        pool.latent_rows(_KDA_IDS[0])
+    assert pool.index_k_cache(0).dtype == torch.bfloat16
+    assert pool.tail_gate(0).dtype == torch.bfloat16
+
+
 def _layer_types() -> list[str]:
     return [
         "deepseek_sparse_attention" if i in _DSA_IDS else "linear_attention"
@@ -134,11 +156,15 @@ def _hf_config(quantization_config: dict | None = None) -> RawConfigShim:
         "group_0": {
             "targets": ["re:.*\\.layers\\.(?:[3-9]|[1-3][0-9]|4[0-4])\\.mlp\\.experts\\..*(gate|up|down)_proj$"],
             "weights": {"num_bits": 4, "type": "float", "group_size": 16, "strategy": "tensor_group"},
+            "input_activations": {"num_bits": 4, "type": "float", "group_size": 16,
+                                  "strategy": "tensor_group", "dynamic": "local"},
             "format": "nvfp4-pack-quantized",
         },
         "group_1": {
             "targets": ["re:.*\\.layers\\.45\\.mlp\\.experts\\.\\d+\\.(gate_proj|up_proj|down_proj)$"],
-            "weights": {"num_bits": 8, "type": "float", "strategy": "block"},
+            "weights": {"num_bits": 8, "type": "float", "strategy": "block", "block_structure": [128, 128]},
+            "input_activations": {"num_bits": 8, "type": "float", "group_size": 128,
+                                  "strategy": "group", "dynamic": True},
             "format": "float-quantized",
         },
     },
@@ -287,6 +313,22 @@ def test_compressed_tensors_mixed_precision_detected():
     assert cfg.expert_quant == "nvfp4"
 
 
+def test_published_nvfp4_config_binds_only_routed_experts(tmp_path):
+    from freetoken.layers.quantization import QuantKind
+    from freetoken.models.register import checkpoint_quant_config, get_model_spec
+
+    hf = _hf_config(_CT_MIXED_QUANT)
+    quant = checkpoint_quant_config(str(tmp_path), hf, get_model_spec(hf.architectures[0]))
+    for layer in (3, 4, 44):
+        scheme = quant.scheme_for(f"model.layers.{layer}.mlp.experts")
+        assert scheme.kind is QuantKind.NVFP4 and scheme.has("input_scale")
+        assert quant.scheme_for(f"model.layers.{layer}.mlp.shared_experts.gate_proj") is None
+        assert quant.scheme_for(f"model.layers.{layer}.mlp.gate") is None
+    for prefix in ("model.layers.0.self_attn.in_proj", "model.layers.3.self_attn.q_b_proj", "lm_head"):
+        assert quant.scheme_for(prefix) is None
+    assert quant.scheme_for("model.layers.45.mlp.experts").kind is QuantKind.FP8_BLOCK
+
+
 def test_compressed_tensors_mixed_precision_reads_the_expert_group():
     """A mixed export with nvfp4 dense layers but fp8 experts must not report nvfp4
     experts: the group that targets the experts decides."""
diff --git a/tests/models/test_glm5_next_model.py b/tests/models/test_glm5_next_model.py
index 9a499fba4..85a675bbb 100644
--- a/tests/models/test_glm5_next_model.py
+++ b/tests/models/test_glm5_next_model.py
@@ -1,7 +1,7 @@
-"""Glm5NextForCausalLM wiring smoke test (tiny random model, dense MLPs).
+"""Glm5NextForCausalLM wiring with dense BF16 or offloaded NVFP4 experts.
 
 The per-op math is covered elsewhere (KDA kernels/op, kpool backend, mHC); this
-test checks the ASSEMBLY: a 2-layer hybrid (KDA + DSA) model with mHC threading
+test checks the assembly: a hybrid KDA + DSA model with mHC threading
 runs prefill and decode through the real backends/pools, and the strongest
 cache invariant holds -- decoding token T after prefilling [0, T) produces the
 same logits as prefilling [0, T] outright (state handoff across the KDA
@@ -24,7 +24,7 @@
 DEV = "cuda"
 
 
-def _hf_config():
+def _hf_config(nvfp4=False):
     from freetoken.utils.hf import RawConfigShim
 
     text = {
@@ -52,24 +52,48 @@ def _hf_config():
         "n_group": 1, "topk_group": 1, "swiglu_limit": 10.0,
         "attention_bias": False, "model_type": "glm5_next_text",
     }
-    return RawConfigShim({
+    data = {
         "architectures": ["Glm5NextForConditionalGeneration"],
         "model_type": "glm5_next", "text_config": text,
-    })
+    }
+    if nvfp4:
+        from tests.models.test_glm5_next_config import _CT_MIXED_QUANT
+
+        # Preserve the published regex's layer IDs and feed a routed MoE output into KDA.
+        text.update(
+            num_hidden_layers=5,
+            layer_types=["linear_attention"] * 3 + ["deepseek_sparse_attention", "linear_attention"],
+            mlp_layer_types=["dense"] * 3 + ["sparse"] * 2,
+            first_k_dense_replace=3, indexer_types=["full"] * 5,
+            moe_intermediate_size=64,
+        )
+        data["quantization_config"] = _CT_MIXED_QUANT
+    return RawConfigShim(data)
 
 
-@pytest.fixture()
-def rig(monkeypatch):
+@pytest.fixture(params=["dense-bf16", "RedHatAI-nvfp4"])
+def rig(monkeypatch, request, tmp_path):
     from freetoken.attention.dsa_indexer_kpool import Glm5NextDSABackend
     from freetoken.distributed import set_tp_info, try_get_tp_info
-    from freetoken.kvcache.dsa_pool import KpoolDSAKVCache
+    from freetoken.kvcache import create_kvcache_pool
     from freetoken.kvcache.linear_state_pool import LinearStatePool
+    from freetoken.layers.quantization import QuantBackend, QuantKind, finalize_quant
     from freetoken.models.glm5_next.config import parse_config
     from freetoken.models.glm5_next.model import Glm5NextForCausalLM
+    from freetoken.models.register import checkpoint_quant_config, get_model_spec
 
     if try_get_tp_info() is None:
         set_tp_info(rank=0, size=1)
-    config = parse_config(_hf_config())
+    nvfp4 = request.param == "RedHatAI-nvfp4"
+    hf = _hf_config(nvfp4)
+    config = parse_config(hf)
+    quant = checkpoint_quant_config(str(tmp_path), hf, get_model_spec(hf.architectures[0]))
+    object.__setattr__(config, "quant", quant)
+    object.__setattr__(config, "moe_strategy", "offload" if nvfp4 else "resident")
+    monkeypatch.setattr(
+        "freetoken.layers.quantization.quant_backend._QUANT_BACKEND",
+        QuantBackend.parse("moe.nvfp4=triton"),
+    )
 
     prev_dtype = torch.get_default_dtype()
     torch.set_default_dtype(torch.bfloat16)
@@ -91,12 +115,13 @@ def rig(monkeypatch):
             t = t.abs() + 0.5
         rand[k] = t.to(v.dtype)
     model.load_state_dict(rand)
+    assert finalize_quant(model) > 0
+    model.prepare_for_runtime()
 
-    kv = KpoolDSAKVCache(
-        latent_dim=LATENT, num_layers=2, num_pages=4, page_size=64,
+    kv = create_kvcache_pool(
+        config, num_pages=4, page_size=64,
         dtype=torch.bfloat16, device=torch.device(DEV),
-        index_head_dim=IDX_D, num_index_layers=1,
-        index_ratio=4, num_req_slots=4,
+        num_req_slots=4, kv_quant="nvfp4" if nvfp4 else "none",
     )
     page_table = torch.full((2, 256), -1, dtype=torch.int32, device=DEV)
     page_table[0] = torch.arange(256, dtype=torch.int32, device=DEV)
@@ -107,7 +132,7 @@ def rig(monkeypatch):
 
     ctx = SimpleNamespace(
         kv_cache=kv, page_table=page_table, linear_state_pool=linear_pool,
-        attn_backend=None, batch=None,
+        attn_backend=None, batch=None, moe_offload_cache=None,
     )
     for mod in (
         "freetoken.attention.dsa.get_global_ctx",
@@ -115,9 +140,46 @@ def rig(monkeypatch):
         "freetoken.models.glm5_next.attention.get_global_ctx",
         "freetoken.models.glm5_next.model.get_global_ctx",
         "freetoken.layers.embedding.get_global_ctx",
+        "freetoken.layers.moe.get_global_ctx",
     ):
         monkeypatch.setattr(mod, lambda: ctx)
     ctx.attn_backend = Glm5NextDSABackend(config)
+    if nvfp4:
+        from freetoken.moe.expert_banks import build_expert_banks
+        from freetoken.moe.offload_cache import (
+            OffloadMoeCache, attach_offload_moe_cache, iter_offload_moe_layers,
+        )
+
+        experts = list(iter_offload_moe_layers(model))
+        assert len(experts) == 2
+        for expert in experts:
+            method = expert.quant_method
+            assert method.kind is QuantKind.NVFP4
+            assert method.kernel.name == "triton"
+            assert method.cfg.activation == "swiglu_clamp" and method.cfg.limit == 10.0
+            assert method.cfg.strategy == "offload" and method.cfg.decode_target == "gpu"
+        banks = build_expert_banks(
+            experts[0].quant_method, len(experts), None, device=torch.device(DEV), dummy=True,
+        )
+        cache = OffloadMoeCache(
+            num_layers=len(experts), num_experts=config.num_experts,
+            cache_size=config.num_experts, device=torch.device(DEV),
+            quant_format=banks.quant_format, layout=banks.layout, prefill_overlap=False,
+            max_slots=experts[0].quant_method.slot_limit(),
+        )
+        cache.set_bank_sources(banks.sources)
+        cache.set_alphas(banks.gate_up_alpha, banks.down_alpha)
+        cache.reset()
+        assert len(attach_offload_moe_cache(model, cache)) == len(experts)
+        ctx.moe_offload_cache = cache
+        assert kv.num_layers == 1
+        assert kv.latent_rows(3).shape == (256, LATENT // 2)
+        assert kv.latent_rows(3).dtype == torch.uint8
+        assert kv.latent_block_scale(3).shape == (256, LATENT // 16)
+        assert kv.latent_block_scale(3).dtype == torch.uint8
+        assert kv.latent_scale(3).dtype == torch.float32
+        assert kv.index_k_cache(0).dtype == torch.bfloat16
+        assert kv.tail_k(0).dtype == kv.tail_gate(0).dtype == torch.bfloat16
     return model, ctx
 
 
@@ -163,12 +225,32 @@ def _reset(ctx):
     ctx.kv_cache._index_k_buffer.zero_()
     ctx.kv_cache._tail_k.zero_()
     ctx.kv_cache._tail_gate.zero_()
+    if ctx.kv_cache.kv_quant == "nvfp4":
+        ctx.kv_cache._scale_buffer.zero_()
+        ctx.kv_cache._block_scale_buffer.zero_()
+        ctx.moe_offload_cache.reset()
+
+
+def _state(ctx):
+    pool = ctx.linear_state_pool
+    return pool.recurrent_states[:, 1].clone(), pool.conv_states[:, 1].clone()
+
+
+def _assert_state_close(actual, expected):
+    for name, got, ref in zip(("recurrent", "convolution"), actual, expected):
+        assert torch.isfinite(got).all()
+        for layer, (got_layer, ref_layer) in enumerate(zip(got, ref)):
+            scale = ref_layer.float().abs().max().item()
+            assert scale > 0
+            err = (got_layer.float() - ref_layer.float()).abs().max().item()
+            assert err / scale < 3e-2, f"{name} layer {layer} divergence: {err} (scale {scale})"
 
 
 def test_prefill_decode_consistency(rig):
     model, ctx = rig
     torch.manual_seed(1)
-    total = 24
+    nvfp4 = ctx.kv_cache.kv_quant == "nvfp4"
+    total, decode_tokens = (40, 3) if nvfp4 else (24, 1)
     ids = torch.randint(0, VOCAB, (total,)).tolist()
 
     # One-shot prefill over the full sequence: last-token logits per position
@@ -178,33 +260,51 @@ def test_prefill_decode_consistency(rig):
     full_logits = model.forward()  # [1, VOCAB] logits of the last position
     assert full_logits.shape == (1, VOCAB)
     assert torch.isfinite(full_logits.float()).all()
+    full_state = _state(ctx)
 
-    # Prefill [0, total-1) then decode the last token: must match the one-shot run.
+    # Continue through incomplete and completed kpool tails using the same request state.
     _reset(ctx)
-    _batch(ctx, ids[:-1], 0, "prefill")
+    prefix = total - decode_tokens
+    _batch(ctx, ids[:prefix], 0, "prefill")
     model.forward()
-    _batch(ctx, ids[-1:], total - 1, "decode")
-    dec_logits = model.forward()
+    if nvfp4:
+        packed_prefix = ctx.kv_cache.latent_rows(3)[:prefix].clone()
+        scale_prefix = ctx.kv_cache.latent_scale(3)[:prefix].clone()
+        block_prefix = ctx.kv_cache.latent_block_scale(3)[:prefix].clone()
+        assert packed_prefix.any() and block_prefix.any()
+        assert torch.isfinite(scale_prefix).all() and (scale_prefix > 0).all()
+    for pos in range(prefix, total):
+        _batch(ctx, ids[pos:pos + 1], pos, "decode")
+        dec_logits = model.forward()
     err = (dec_logits.float() - full_logits.float()).abs().max().item()
     scale = full_logits.float().abs().max().item() + 1e-8
     assert err / scale < 3e-2, f"decode/prefill divergence: {err} (scale {scale})"
+    _assert_state_close(_state(ctx), full_state)
+    if nvfp4:
+        assert torch.equal(ctx.kv_cache.latent_rows(3)[:prefix], packed_prefix)
+        assert torch.equal(ctx.kv_cache.latent_scale(3)[:prefix], scale_prefix)
+        assert torch.equal(ctx.kv_cache.latent_block_scale(3)[:prefix], block_prefix)
+        assert (ctx.kv_cache.latent_scale(3)[prefix:total] > 0).all()
 
 
 def test_chunked_prefill_consistency(rig):
     model, ctx = rig
     torch.manual_seed(2)
-    total = 28  # split 16 + 12; chunk boundary pool-aligned (16 % 4 == 0)
+    nvfp4 = ctx.kv_cache.kv_quant == "nvfp4"
+    total, split = (44, 17) if nvfp4 else (28, 16)
     ids = torch.randint(0, VOCAB, (total,)).tolist()
 
     _reset(ctx)
     _batch(ctx, ids, 0, "prefill")
     full_logits = model.forward()
+    full_state = _state(ctx)
 
     _reset(ctx)
-    _batch(ctx, ids[:16], 0, "prefill")
+    _batch(ctx, ids[:split], 0, "prefill")
     model.forward()
-    _batch(ctx, ids[16:], 16, "prefill")
+    _batch(ctx, ids[split:], split, "prefill")
     chunk_logits = model.forward()
     err = (chunk_logits.float() - full_logits.float()).abs().max().item()
     scale = full_logits.float().abs().max().item() + 1e-8
     assert err / scale < 3e-2, f"chunked/one-shot divergence: {err} (scale {scale})"
+    _assert_state_close(_state(ctx), full_state)
diff --git a/tests/models/test_glm_dsa.py b/tests/models/test_glm_dsa.py
index 0dceabadd..f09a73822 100644
--- a/tests/models/test_glm_dsa.py
+++ b/tests/models/test_glm_dsa.py
@@ -225,7 +225,7 @@ def test_splitk_matches_single_program():
     assert (o_single.float() - o_split.float()).abs().max().item() < 1e-2
 
 
-def _make_backend(dsa: bool, latent=80, dv=64, idx_dim=32, idx_heads=16, topk=64, pages=400):
+def _make_backend(dsa: bool, latent=80, dv=64, idx_dim=32, idx_heads=16, topk=64, pages=400, kv_quant="none"):
     """Minimal ctx + pool + DSAAttnBackend (no engine)."""
     from types import SimpleNamespace
 
@@ -239,9 +239,9 @@ def _make_backend(dsa: bool, latent=80, dv=64, idx_dim=32, idx_heads=16, topk=64
     ctx.page_table = torch.zeros(4, pages, dtype=torch.int32, device="cuda")
     if dsa:
         ctx.kv_cache = DSAKVCache(latent, 2, pages, 1, torch.bfloat16, torch.device("cuda"),
-                                  index_head_dim=idx_dim, num_index_layers=1)
+                                  index_head_dim=idx_dim, num_index_layers=1, kv_quant=kv_quant)
     else:
-        ctx.kv_cache = MLAKVCache(latent, 2, pages, 1, torch.bfloat16, torch.device("cuda"))
+        ctx.kv_cache = MLAKVCache(latent, 2, pages, 1, torch.bfloat16, torch.device("cuda"), kv_quant=kv_quant)
     set_global_ctx(ctx)
     args = SimpleNamespace(
         kv_lora_rank=dv, qk_rope_head_dim=latent - dv, qk_head_dim=latent,
@@ -258,7 +258,8 @@ def _ref_attend(q_cat, pool_rows, live_rows, scale, dv):
     return s.softmax(-1) @ k[:, :dv]
 
 
-def test_backend_ragged_prefill_identity_and_selection():
+@pytest.mark.parametrize("kv_quant", ["none", "nvfp4"])
+def test_backend_ragged_prefill_identity_and_selection(kv_quant):
     """Two-request ragged prefill through the BACKEND (page-table slicing,
     counts = positions + 1, per-request segmentation, leader/follower reuse):
     request A stays under index_topk (selection == identity == dense), request B
@@ -267,7 +268,7 @@ def test_backend_ragged_prefill_identity_and_selection():
 
     torch.manual_seed(5)
     dv, dr, h, idx_h, idx_d, topk = 64, 16, 8, 16, 32, 64
-    backend, ctx = _make_backend(dsa=True, topk=topk)
+    backend, ctx = _make_backend(dsa=True, topk=topk, kv_quant=kv_quant)
     pool = ctx.kv_cache
     scale = backend.sm_scale
 
@@ -311,6 +312,10 @@ def test_backend_ragged_prefill_identity_and_selection():
     # request A (kv <= topk): selection covers all live -> equals dense reference
     q_cat = torch.cat([q_nope, q_pe], -1)
     slab = pool.latent_rows(0)
+    if kv_quant == "nvfp4":
+        from tests.kernels.test_kv_nvfp4 import _decode_latent
+
+        slab = _decode_latent(pool)
     for j in range(8):  # A's queries, positions 32..39
         live = ctx.page_table[0, : 33 + j]
         ref = _ref_attend(q_cat[j], slab, live, scale, dv)
@@ -335,15 +340,17 @@ def test_backend_ragged_prefill_identity_and_selection():
     assert (o0.float() - o1.float()).abs().max().item() < 3e-2
 
     # identity wiring (dense ablation): same batch through an MLAKVCache backend
-    backend_d, ctx_d = _make_backend(dsa=False)
+    backend_d, ctx_d = _make_backend(dsa=False, kv_quant=kv_quant)
     ctx_d.page_table.copy_(ctx.page_table)
-    for lid in (0, 1):
-        ctx_d.kv_cache._kv_buffer.copy_(pool._kv_buffer)
+    ctx_d.kv_cache._kv_buffer.copy_(pool._kv_buffer)
+    if kv_quant == "nvfp4":
+        ctx_d.kv_cache._scale_buffer.copy_(pool._scale_buffer)
+        ctx_d.kv_cache._block_scale_buffer.copy_(pool._block_scale_buffer)
     batch_d = SimpleNamespace(reqs=reqs, positions=positions, out_loc=out_loc,
                               active_table_idx=None, attn_metadata=None)
     backend_d.prepare_metadata(batch_d)
     od = backend_d.mla_forward(q_nope, q_pe, c_kv, k_rope, 0, batch_d, indexer_inputs=None)
-    slab_d = ctx_d.kv_cache.latent_rows(0)
+    slab_d = _decode_latent(ctx_d.kv_cache) if kv_quant == "nvfp4" else ctx_d.kv_cache.latent_rows(0)
     for j in range(8):
         live = ctx_d.page_table[0, : 33 + j]
         ref = _ref_attend(q_cat[j], slab_d, live, scale, dv)
diff --git a/tests/models/test_minimax_m3.py b/tests/models/test_minimax_m3.py
index dc0ab90df..b43d88b05 100644
--- a/tests/models/test_minimax_m3.py
+++ b/tests/models/test_minimax_m3.py
@@ -9,6 +9,9 @@
 
 from __future__ import annotations
 
+from importlib import import_module
+from types import SimpleNamespace
+
 import pytest
 import torch
 from freetoken.layers.quantization import MoEConfig
@@ -299,3 +302,62 @@ def test_expert_source_spec_layer_to_bank():
     assert _NVFP4_SOURCE_SPEC.layer_to_bank(3, cfg) == 0
     assert _NVFP4_SOURCE_SPEC.layer_to_bank(59, cfg) == 56
     assert _NVFP4_SOURCE_SPEC.layer_to_bank(0, cfg) is None
+
+
+@pytest.mark.parametrize("num_tokens", [1, 3])
+@pytest.mark.parametrize("family,class_name", [
+    ("minimax_m3", "MiniMaxM3SparseMoeBlock"),
+    ("glm4_moe", "Glm4MoeSparseBlock"),
+    ("glm_moe_dsa", "GlmMoeDsaSparseBlock"),
+    ("glm5_next", "Glm5NextSparseBlock"),
+    ("deepseek_v41", "MoE"),
+    ("qwen3_5_moe", "Qwen3_5MoE"),
+    ("qwen4_exp", "Qwen4ExpMoE"),
+])
+def test_shared_experts_keep_original_input_when_routed_experts_overwrite_it(
+    monkeypatch, family, class_name, num_tokens,
+):
+    """An in-place routed result must not feed the independent shared expert or its gate."""
+    module = import_module(f"freetoken.models.{family}.moe")
+    cls = getattr(module, class_name)
+    original = torch.arange(num_tokens * 4, dtype=torch.float32).view(num_tokens, 4) / 4 - 1
+    hidden = original.clone()
+    weights = torch.ones(num_tokens, 1)
+    ids = torch.zeros(num_tokens, 1, dtype=torch.int32)
+
+    def shared(x):
+        return x.square()
+
+    def routed(hidden_states, *args, **kwargs):
+        hidden_states.mul_(2).add_(3)
+        return hidden_states
+
+    block = SimpleNamespace(
+        _route=lambda x: (weights, ids),
+        shared_experts=SimpleNamespace(forward=shared),
+        experts=SimpleNamespace(routed_forward=routed, forward=routed),
+    )
+    expected = original * 2 + 3 + original.square()
+    if family == "deepseek_v41":
+        block.dim = 4
+        block.gate = lambda x, image_mask: (weights, ids)
+        block.shared_experts = shared
+        hidden = hidden.unsqueeze(0)
+        expected = expected.unsqueeze(0)
+    elif family.startswith("qwen"):
+        gate_weight = torch.tensor([[-0.5, 0.25, 0.75, 1.0]])
+        block.gate = SimpleNamespace(forward=lambda x: x[:, :2].clone())
+        block.shared_expert = SimpleNamespace(forward=shared)
+        block.shared_expert_gate = SimpleNamespace(
+            weight=gate_weight, forward=lambda x: x @ gate_weight.t(),
+        )
+        expected = original * 2 + 3 + original.square() * torch.sigmoid(original @ gate_weight.t())
+        if family == "qwen4_exp":
+            # Exercise the model's branch ordering on CPU; the gate kernels have their own GPU tests.
+            monkeypatch.setattr(module, "shared_gate_sigmoid", lambda x, w: torch.sigmoid(x @ w))
+            monkeypatch.setattr(module, "shared_gate_mul_add", lambda routed, shared, gate: routed + shared * gate[:, None])
+
+    got = cls.forward(block, hidden)
+
+    torch.testing.assert_close(hidden.reshape_as(original), original * 2 + 3)
+    torch.testing.assert_close(got, expected)
diff --git a/tests/models/test_models_loader.py b/tests/models/test_models_loader.py
index b0ac2b13f..4ae6d232f 100644
--- a/tests/models/test_models_loader.py
+++ b/tests/models/test_models_loader.py
@@ -201,3 +201,86 @@ def test_stacked_expert_pieces_pair_each_layer_in_arrival_order():
     assert torch.equal(pieces[1][3]["gate_up"], torch.full((2, 3, 4), 2.0))
     with pytest.raises(ValueError, match="Missing MoE expert source layers"):
         list(stacked_expert_pieces(tensors[:3], config))
+
+
+@pytest.mark.parametrize("per_layer", [False, True], ids=["flat", "per-layer"])
+@pytest.mark.parametrize("quant_format", ["q4_0", "unknown-format"])
+def test_ftw_legacy_q4_0_banks_keep_native_bytes(tmp_path, monkeypatch, per_layer, quant_format):
+    from freetoken.checkpoint.ftw import FTWWriter, layer_bank_entry_name, load_ftw_banks
+
+    monkeypatch.setattr(torch.cuda, "is_available", lambda: False)
+    writer = FTWWriter(str(tmp_path), shard_limit=8192)
+    sources = {
+        "gate_up": torch.arange(2 * 2 * 72, dtype=torch.int64).to(torch.uint8).reshape(2, 2, 72),
+        "down": torch.arange(2 * 2 * 36, dtype=torch.uint8).reshape(2, 2, 36),
+    }
+    for name, value in sources.items():
+        if per_layer:
+            for layer, rows in enumerate(value):
+                writer.add_tensor(layer_bank_entry_name(name, layer), rows, kind="experts_bank")
+        else:
+            writer.add_tensor(name, value.flatten(0, 1), kind="experts_bank")
+    writer.finalize({"quant_format": quant_format, "expert_bank_num_layers": 2})
+
+    if quant_format != "q4_0":
+        with pytest.raises(KeyError, match="unknown-format"):
+            load_ftw_banks(str(tmp_path), num_layers=2, layer_residency=["pageable"] * 2)
+        return
+    banks = load_ftw_banks(str(tmp_path), num_layers=2, layer_residency=["pageable"] * 2)
+    assert banks.quant_format == "q4_0"
+    assert banks.kind is None and banks.kernel is None
+    assert banks.layer_residency == ["pageable"] * 2
+    assert set(banks.sources) == set(sources)
+    for name, value in sources.items():
+        assert len(banks.sources[name]) == 2
+        for layer, rows in enumerate(banks.sources[name]):
+            assert rows.dtype is torch.uint8
+            torch.testing.assert_close(rows, value[layer], rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("include_vision", [False, True])
+def test_ftw_text_only_filters_native_and_qwen_vision_weights(tmp_path, include_vision):
+    from freetoken.checkpoint.ftw import FTWWriter
+    from freetoken.models.weight import load_weight
+
+    vision_keys = ("vision.blocks.0.norm1.weight", "aligner.mlp.0.weight",
+                   "image_start", "image_end", "image_newline", "visual.blocks.0.norm1.weight")
+    text_keys = ("embed.weight", "layers.0.attn.wq_a.weight", "head.weight")
+    weights = {name: torch.full((2, 3), float(i), dtype=torch.bfloat16)
+               for i, name in enumerate((*vision_keys, *text_keys))}
+    writer = FTWWriter(str(tmp_path), shard_limit=8192)
+    for name, value in weights.items():
+        writer.add_tensor(name, value)
+    writer.finalize({})
+
+    actual = dict(load_weight(str(tmp_path), torch.device("cpu"), include_vision=include_vision))
+    assert set(actual) == set(weights if include_vision else text_keys)
+    for name, value in actual.items():
+        torch.testing.assert_close(value, weights[name], rtol=0, atol=0)
+
+
+@pytest.mark.parametrize("with_input_scale", [False, True])
+def test_ftw_activation_scale_follows_the_current_model_scheme(tmp_path, with_input_scale):
+    from freetoken.checkpoint.ftw import FTWWriter
+    from freetoken.engine.engine import _materialize_loaded_weight_state_dict
+    from freetoken.models.weight import load_weight
+
+    layer = torch.nn.Linear(3, 2, bias=False, dtype=torch.bfloat16)
+    layer.register_buffer("weight_scale", torch.empty((), dtype=torch.float32))
+    if with_input_scale:
+        layer.register_buffer("input_scale", torch.empty((), dtype=torch.float32))
+    model = torch.nn.ModuleDict({"linear": layer})
+    weights = {"linear.weight": torch.arange(6, dtype=torch.bfloat16).reshape(2, 3),
+               "linear.weight_scale": torch.tensor(0.25), "linear.input_scale": torch.tensor(0.5)}
+    writer = FTWWriter(str(tmp_path), shard_limit=8192)
+    for name, value in weights.items():
+        writer.add_tensor(name, value)
+    writer.finalize({})
+
+    loaded = _materialize_loaded_weight_state_dict(
+        model.state_dict(), load_weight(str(tmp_path), torch.device("cpu")), device=torch.device("cpu"),
+    )
+    model.load_state_dict(loaded, strict=True)
+    assert set(loaded) == set(model.state_dict())
+    for name, value in model.state_dict().items():
+        torch.testing.assert_close(value, weights[name], rtol=0, atol=0)
diff --git a/tests/models/test_quant_config.py b/tests/models/test_quant_config.py
index 6ac1b84b1..2986399d9 100644
--- a/tests/models/test_quant_config.py
+++ b/tests/models/test_quant_config.py
@@ -392,6 +392,42 @@ def test_probed_layers_get_the_method_their_config_says(case: Case, monkeypatch)
 # --------------------------------------------------------------------------- config without local weights
 
 
+@pytest.mark.parametrize("filename", [
+    "deepseek_v41_nvfp4_config.json",
+    "deepseek_v41_libertai_nvfp4_config.json",
+])
+def test_dsv41_factory_preserves_native_quantization(filename, tmp_path, monkeypatch):
+    from freetoken.utils import cached_load_hf_config
+
+    def unexpected_sidecar(*args, **kwargs):
+        pytest.fail("V4.1 native quantization must not use the generic sidecar dialect")
+
+    monkeypatch.setattr("freetoken.utils.hf.optional_hf_file", unexpected_sidecar)
+    raw = json.loads((Path(__file__).parent / "fixtures" / filename).read_text())
+    (tmp_path / "config.json").write_text(json.dumps(raw))
+    hf = cached_load_hf_config(str(tmp_path))
+    quant = checkpoint_quant_config(str(tmp_path), hf, get_model_spec(raw["architectures"][0]))
+    dense = quant.scheme_for("layers.2.attn.wq_a")
+    assert dense.kind is QuantKind.FP8_BLOCK
+    assert dense.weight.group == (32, 32)
+    assert dense.weight.scale == "e8m0"
+    assert quant.storage(dense)["weight_scale_inv"].name == "scale"
+    experts = quant.scheme_for("layers.2.ffn.experts")
+    assert experts.kind is QuantKind.NVFP4
+    assert experts.weight.group == (1, 16)
+    assert experts.weight.scale == "e4m3"
+    assert experts.roles == {"weight", "weight_scale", "weight_global"}
+    assert quant.storage(experts)["weight_global"].name == "weight_scale_2"
+    assert quant.scheme_for("head") is None
+    assert quant.scheme_for("layers.2.attn.compressor.wkv") is None
+
+
+def test_generic_fp8_factory_still_rejects_32_by_32_blocks(tmp_path):
+    raw = {"quantization_config": {"quant_method": "fp8", "weight_block_size": [32, 32]}}
+    with pytest.raises(NotImplementedError, match="only 128x128 blocks"):
+        checkpoint_quant_config(str(tmp_path), raw, get_model_spec("Qwen3MoeForCausalLM"))
+
+
 def test_the_modelopt_sidecar_is_folded_into_the_hf_config(tmp_path):
     """An old ModelOpt export keeps its quantization config only in hf_quant_config.json; the config loader folds it in, so every reader of the config sees it."""
     from freetoken.utils import cached_load_hf_config
@@ -406,6 +442,60 @@ def test_the_modelopt_sidecar_is_folded_into_the_hf_config(tmp_path):
     assert quant.scheme_for("lm_head") is None
 
 
+def test_hub_sidecar_reaches_all_config_readers_and_cached_copies(tmp_path, monkeypatch):
+    from freetoken.layers.quantization.configs.base import quantization_config_of
+    from freetoken.utils.hf import AutoConfig, RawConfigShim, _load_hf_config, cached_load_hf_config
+
+    path = "org/old-modelopt-no-input"
+    quant = {"quant_algo": "FP8", "with_input_scale": False, "exclude_modules": ["lm_head"]}
+    sidecar = tmp_path / "hf_quant_config.json"
+    sidecar.write_text(json.dumps({"quantization": quant}))
+    downloads = []
+
+    def download(repo_id, filename, **kwargs):
+        assert (repo_id, filename) == (path, "hf_quant_config.json")
+        downloads.append(filename)
+        return str(sidecar)
+
+    monkeypatch.setattr("freetoken.utils.hf.hf_hub_download", download)
+    monkeypatch.setattr(AutoConfig, "from_pretrained", lambda *args, **kwargs: RawConfigShim({
+        "architectures": ["Qwen3MoeForCausalLM"], "model_type": "unknown_export",
+    }))
+    _load_hf_config.cache_clear()
+    try:
+        hf = cached_load_hf_config(path)
+        assert quantization_config_of(hf) == {"quant_method": "modelopt", **quant}
+        direct = QuantConfig.from_hf(hf)
+        registered = checkpoint_quant_config(path, hf, get_model_spec(hf.architectures[0]))
+        for config in (direct, registered):
+            scheme = config.scheme_for("model.layers.0.self_attn.q_proj")
+            assert scheme.kind is QuantKind.FP8_TENSOR and not scheme.has("input_scale")
+            assert config.scheme_for("lm_head") is None
+        hf._data["quantization_config"]["with_input_scale"] = True
+        assert quantization_config_of(cached_load_hf_config(path))["with_input_scale"] is False
+        assert downloads == ["hf_quant_config.json"]
+    finally:
+        _load_hf_config.cache_clear()
+
+
+@pytest.mark.parametrize("nested", [False, True], ids=["top-level", "text-config"])
+def test_inline_quantization_takes_precedence_without_fetching_sidecar(tmp_path, monkeypatch, nested):
+    from freetoken.utils import cached_load_hf_config
+
+    quant = {"quant_method": "modelopt", "quant_algo": "W4A16_NVFP4"}
+    raw = {"model_type": "unknown_export", "architectures": ["Qwen3MoeForCausalLM"]}
+    if nested:
+        raw["text_config"] = {"quantization_config": quant}
+    else:
+        raw["quantization_config"] = quant
+    (tmp_path / "config.json").write_text(json.dumps(raw))
+    monkeypatch.setattr("freetoken.utils.hf.optional_hf_file",
+                        lambda *args, **kwargs: pytest.fail("inline quantization must win before sidecar lookup"))
+    config = QuantConfig.from_hf(cached_load_hf_config(str(tmp_path)))
+    scheme = config.scheme_for("model.layers.0.mlp.down_proj")
+    assert scheme.kind is QuantKind.NVFP4 and not scheme.has("input_scale")
+
+
 def test_compressed_tensors_ignore_names_the_module_alone():
     """llm-compressor lists every skipped module, containers included: an ignored ``linear_attn`` must not shield the quantized projections under it (unsloth/Qwen3.6-35B-A3B-NVFP4-Fast)."""
     q = {
@@ -480,7 +570,7 @@ def test_unsupported_dialects_fail_closed(tmp_path):
 # unsupported dialects: from_hf must refuse them, nothing else is checked
 _UNSUPPORTED = {"quark", "mxfp8", "fp_quant"}
 # tables the model code reads itself; RadixArk lists the fp8 PLE table in ``ignore`` yet ships it quantized
-_MODEL_READS = ("ple.ple_embedding",)
+_MODEL_READS = ("ple.ple_embedding", "engram.embed")
 # non-Linear leaves whose parameters happen to be called scale
 _NOT_LINEAR = {"router", "gate", "shared_expert_gate"}
 
@@ -588,11 +678,17 @@ def test_scheme_for_agrees_with_the_stored_tensors(ckpt: Path):
         spec = get_model_spec((cfg.get("architectures") or [""])[0])
     except Exception:
         spec = None
-    qc = QuantConfig.from_hf(cfg, unquantized=spec.unquantized_modules if spec else ())
+    qc = (checkpoint_quant_config(str(ckpt), cfg, spec) if spec and spec.quant_config is not None
+          else QuantConfig.from_hf(cfg, unquantized=spec.unquantized_modules if spec else ()))
     weight_map = _weight_map(ckpt)
     tensors = _tensor_info(ckpt, weight_map)
     mismatches, checked = [], 0
     for module, suffixes in _probes(weight_map):
+        if spec and spec.module == "freetoken.models.deepseek_v41" and (
+            module.startswith("mtp.") or module.endswith(".attn.wo_a")
+        ):
+            # MTP is not served; wo_a is a raw BF16 parameter dequantized by the reader.
+            continue
         expected = _expected_kinds(suffixes, tensors, module)
         scheme = qc.scheme_for(module)
         kind = scheme.kind if scheme else QuantKind.NONE
diff --git a/tests/models/test_qwen3_5_moe_weight.py b/tests/models/test_qwen3_5_moe_weight.py
index 158f1f459..5b2bee005 100644
--- a/tests/models/test_qwen3_5_moe_weight.py
+++ b/tests/models/test_qwen3_5_moe_weight.py
@@ -342,6 +342,25 @@ def test_emitted_keys_are_the_model_state_dict(checkpoint):
     assert not any(".mlp.experts." in k or k.startswith("mtp.") for k in loaded)
 
 
+@pytest.mark.parametrize("layout", ["modelopt_mixed_a16", "modelopt_mixed_noinput"])
+def test_sidecar_only_modelopt_uses_the_same_schemes_for_model_and_reader(tmp_path, layout):
+    moe, quant, raw = _layout(layout)
+    folder = _write(tmp_path, moe, None, raw)
+    (tmp_path / "hf_quant_config.json").write_text(json.dumps({
+        "producer": {"name": "modelopt"},
+        "quantization": {key: value for key, value in quant.items() if key != "quant_method"},
+    }))
+    loaded, state = _load(folder, vision=False), _meta_state_dict(folder)
+    assert set(loaded) == set(state)
+    for key, tensor in loaded.items():
+        assert tensor.shape == state[key].shape and tensor.dtype == state[key].dtype, key
+    assert loaded["model.layers.0.mlp.shared_expert.down_proj.weight"].dtype is torch.uint8
+    assert not any(key.startswith("model.layers.0.mlp.shared_expert.") and key.endswith("input_scale")
+                   for key in loaded)
+    fp8_scale = "model.layers.1.self_attn.qkv_proj.input_scale"
+    assert (fp8_scale in loaded) == (layout == "modelopt_mixed_a16")
+
+
 def test_expert_quant_tag_follows_the_config(checkpoint):
     name, folder, _raw = checkpoint
     config = parse_config(cached_load_hf_config(folder))
diff --git a/tests/moe/test_offload.py b/tests/moe/test_offload.py
index cd5f5dc77..633239f66 100644
--- a/tests/moe/test_offload.py
+++ b/tests/moe/test_offload.py
@@ -455,6 +455,61 @@ def test_adjust_config_converts_moe_cache_rate_to_cache_size(monkeypatch):
     assert is_offload_moe_strategy(config.moe_strategy)
 
 
+@pytest.mark.parametrize("unified_memory,requested,expert_quant,expected", [
+    (True, "auto", "none", "fused"),
+    (True, "auto", "fp8_block", "fused"),
+    (True, "auto", "nvfp4", "offload"),
+    (True, "offload", "none", "offload"),
+    (True, "offload", "fp8_block", "offload"),
+    (True, "offload", "nvfp4", "offload"),
+    (False, "auto", "none", "offload"),
+    (False, "auto", "fp8_block", "offload"),
+    (False, "auto", "nvfp4", "offload"),
+])
+def test_adjust_config_unified_memory_respects_expert_format(monkeypatch, unified_memory, requested, expert_quant, expected):
+    from types import SimpleNamespace
+
+    from freetoken.distributed import DistributedInfo
+    from freetoken.engine.config import EngineConfig
+    import freetoken.engine.engine as engine_module
+
+    profile_reads = []
+
+    def recommendation(fmt, **identity):
+        profile_reads.append((fmt, identity))
+        return "hybrid" if unified_memory else None
+
+    monkeypatch.setattr(engine_module, "_is_unified_memory_gpu", lambda index=None: unified_memory)
+    monkeypatch.setattr(engine_module, "_profile_gpu", lambda index=None: ("test-gpu", "test-uuid"))
+    monkeypatch.setattr("freetoken.moe.bench_profile.load_backend_recommendation", recommendation)
+    config = EngineConfig(
+        model_path="/tmp/freetoken-test-model",
+        tp_info=DistributedInfo(rank=0, size=1),
+        dtype=torch.bfloat16,
+        attention_backend="triton",
+        moe_strategy=requested,
+        moe_cache_auto=True,
+        quant_backend="moe.nvfp4=triton" if expert_quant == "nvfp4" else None,
+    )
+    model = SimpleNamespace(
+        has_swa_attention=False, has_linear_attention=False, is_moe=True,
+        num_layers=10, num_moe_layers=10, num_experts=8,
+        expert_quant=expert_quant, moe_strategy=requested,
+    )
+    object.__setattr__(config, "model_config", model)
+
+    engine_module._adjust_config(config)
+
+    assert config.moe_strategy == model.moe_strategy == expected
+    assert config.moe_cache_auto == (expected == "offload")
+    assert model.decode_target == "gpu"
+    if unified_memory:
+        assert profile_reads == []
+    else:
+        assert profile_reads == [("bf16" if expert_quant == "none" else expert_quant,
+                                  {"gpu_name": "test-gpu", "gpu_uuid": "test-uuid"})]
+
+
 def test_graph_capture_reuses_warm_offload_cache_before_capture(monkeypatch):
     import freetoken.core as core
     from freetoken.core import Context, Req, get_global_ctx
diff --git a/tests/scheduler/test_dsv4_generic_manager.py b/tests/scheduler/test_dsv4_generic_manager.py
index ad69b89ec..5005f49f5 100644
--- a/tests/scheduler/test_dsv4_generic_manager.py
+++ b/tests/scheduler/test_dsv4_generic_manager.py
@@ -1,4 +1,4 @@
-"""DSV4 through the GENERIC CacheManager (CPU, no model) -- the unified serving path.
+"""DSV4/V4.1 through the GENERIC CacheManager (CPU, no model) -- the unified serving path.
 
 The shared page_table is the virtual full-token coordinate (ShadowRadix); DSV4PagedKVCache plugs
 in as the swa_pool: window pages bind page-atomically behind token-face alloc_swa/free_swa, the
@@ -10,6 +10,9 @@
 
 from __future__ import annotations
 
+import sys
+from types import SimpleNamespace
+
 import pytest
 import torch
 
@@ -25,6 +28,22 @@
 MRR = 4
 
 
+@pytest.fixture(autouse=True, params=["v4", "v41"])
+def _pool_family(request, monkeypatch):
+    if request.param == "v41":
+        from freetoken.kvcache.dsv41_cost_model import dsv41_pool_sizes
+        from freetoken.kvcache.dsv41_paged_pool import DSV41PagedKVCache
+
+        module = sys.modules[__name__]
+        monkeypatch.setattr(module, "_args", lambda: SimpleNamespace(
+            n_layers=5, compress_ratios=(0, 2, 2, 1, 1), max_seq_len=8192,
+            kv_source_layers=(1, 3), index_source_layers=(1, 3, 4),
+            head_dim=32, index_head_dim=16, window_size=P,
+        ))
+        monkeypatch.setattr(module, "dsv4_pool_sizes", dsv41_pool_sizes)
+        monkeypatch.setattr(module, "DSV4PagedKVCache", DSV41PagedKVCache)
+
+
 def _args():
     return DeepseekV4Args(
         n_layers=8, compress_ratios=RATIOS, max_seq_len=8192,
diff --git a/tests/scheduler/test_mm.py b/tests/scheduler/test_mm.py
index 134476347..ef255d55f 100644
--- a/tests/scheduler/test_mm.py
+++ b/tests/scheduler/test_mm.py
@@ -2,11 +2,14 @@
 
 from __future__ import annotations
 
+from types import SimpleNamespace
+
+import pytest
 import torch
 
 from freetoken.message import MMItem
 from freetoken.mm.encoder_cache import EncoderCache
-from freetoken.scheduler.mm import cut_image_spans, mm_chunk_end, mm_rows_after, plan_mm_batch, plan_mm_chunk
+from freetoken.scheduler.mm import cut_image_spans, gather_legacy_mm_batch, mm_chunk_end, mm_rows_after, plan_mm_batch, plan_mm_chunk
 
 CPU = torch.device("cpu")
 
@@ -90,7 +93,53 @@ def test_batch_rows_follow_the_reqs_in_batch_order():
     assert [j.hash for j in jobs] == [7, 8]
     assert plan == [(1, 7, 0, 3, 3, 2), (2, 8, 1, 5, 5, 0)]
     assert rows == [2, 3, 4, 6, 7, 8, 9]  # req 2 starts at batch row 6; its image rows 1..5 land on its first four tokens
-    assert block_ends == [0, 0, 5, 5, 5, 0, 8, 8, 8, 8, 0, 0]  # every image row carries its span's end in request positions
+    assert block_ends == [0, 0, 5, 5, 5, 0, 8, 8, 8, 8, 0, 0]
+
+
+def test_legacy_soft_tokens_slice_at_chunk_boundaries():
+    embeds = torch.arange(24).reshape(3, 8)
+    prefix_req = SimpleNamespace(extend_len=3)
+    req = SimpleNamespace(input_ids=torch.tensor([9, 1, 9, 9, 2]), mm_embeds=embeds)
+    gathered = []
+    for lo, hi, expected_rows in ((0, 2, [3]), (2, 4, [3, 4]), (4, 5, [])):
+        req.cached_len, req.device_len, req.extend_len = lo, hi, hi - lo
+        parts, rows = gather_legacy_mm_batch([prefix_req, req], image_token_id=9)
+        assert rows == expected_rows
+        gathered.extend(parts)
+    assert torch.equal(torch.cat(gathered), embeds)
+
+
+@pytest.mark.parametrize("image_token_id,embeds,message", [
+    (None, torch.zeros(1, 8), "require an image_token_id"),
+    (9, torch.zeros(8), "must have shape"),
+    (9, torch.zeros(0, 8), "slots exceed"),
+])
+def test_legacy_soft_tokens_reject_incompatible_metadata(image_token_id, embeds, message):
+    req = SimpleNamespace(input_ids=torch.tensor([9]), mm_embeds=embeds,
+                          cached_len=0, device_len=1, extend_len=1)
+    with pytest.raises(ValueError, match=message):
+        gather_legacy_mm_batch([req], image_token_id)
+
+
+def test_scheduler_appends_legacy_rows_after_canonical_rows():
+    from freetoken.scheduler.scheduler import Scheduler
+
+    embeds = torch.arange(8).reshape(1, 8)
+    legacy = SimpleNamespace(uid=1, mm_items=None, input_ids=torch.tensor([9, 5]),
+                             mm_embeds=embeds, cached_len=0, device_len=2, extend_len=2)
+    canonical = SimpleNamespace(uid=2, mm_items=[_item(7, [[1, 3]])],
+                                cached_len=0, device_len=3, extend_len=3)
+    batch = SimpleNamespace(padded_reqs=[legacy, canonical], mm_embeds=None, mm_rows=None)
+    scheduler = SimpleNamespace(engine=SimpleNamespace(encoder_cache=None), device=CPU, _bidirectional_mm=False,
+                                config=SimpleNamespace(model_config=SimpleNamespace(image_token_id=9)))
+
+    Scheduler._gather_multimodal(scheduler, batch)
+
+    assert batch.mm_rows.tolist() == [3, 4, 0]
+    assert batch.mm_block_ends.tolist() == [0, 0, 0, 3, 3]
+    assert torch.equal(batch.mm_embeds, embeds)
+    assert batch.mm_encoder_jobs == canonical.mm_items
+    assert batch.mm_gather_plan == [(2, 7, 0, 2, 2, 1)]
 
 
 def test_chunk_end_never_lands_inside_an_image_span():
diff --git a/tests/scheduler/test_multimodal_chunks.py b/tests/scheduler/test_multimodal_chunks.py
new file mode 100644
index 000000000..41c4039e8
--- /dev/null
+++ b/tests/scheduler/test_multimodal_chunks.py
@@ -0,0 +1,200 @@
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from freetoken.core import SamplingParams
+from freetoken.message import MMItem, UserMsg
+from freetoken.mm import mm_pad_value
+from freetoken.mm.encoder_cache import EncoderCache
+from freetoken.scheduler.cache import CacheManager
+from freetoken.scheduler.mm import plan_mm_batch
+from freetoken.scheduler.prefill import ChunkedReq, PrefillManager
+from freetoken.scheduler.table import TableManager
+from freetoken.scheduler.utils import PendingReq
+
+
+def _manager(cache_type="radix", encoder_cache=None):
+    table = TableManager(4, torch.zeros(4, 64, dtype=torch.int32))
+    kwargs = {}
+    page_size = 1
+    if cache_type == "swa_radix":
+        from freetoken.distributed import set_tp_info, try_get_tp_info
+        from freetoken.kvcache.hybrid_swa_pool import HybridSWAKVCache
+        from freetoken.models.config import KVCacheGroupSpec
+
+        if try_get_tp_info() is None:
+            set_tp_info(rank=0, size=1)
+        page_size = 2
+        groups = (
+            KVCacheGroupSpec(name="full", layer_ids=(1,), num_kv_heads=1, head_dim=8,
+                             sliding_window=None),
+            KVCacheGroupSpec(name="swa", layer_ids=(0,), num_kv_heads=1, head_dim=8,
+                             sliding_window=4),
+        )
+        pool = HybridSWAKVCache(groups=groups, num_layers=2, num_full_pages=32,
+                                page_size=page_size, dtype=torch.bfloat16,
+                                device=torch.device("cpu"), num_swa_tokens=64)
+        kwargs = dict(swa_pool=pool, sliding_window_size=4)
+    cache = CacheManager(64 // page_size, page_size, table.page_table, cache_type, **kwargs)
+    manager = PrefillManager(cache, table, SimpleNamespace(inflight_tokens=0), encoder_cache=encoder_cache)
+    return cache, manager
+
+
+@pytest.mark.parametrize("cache_type", ["radix", "swa_radix"])
+@pytest.mark.parametrize("payload_name", ["media", "mm_embeds"])
+def test_media_survives_chunking_without_prefix_sharing(cache_type, payload_name):
+    cache, manager = _manager(cache_type)
+    ids = torch.tensor([4, 9, 9, 9, 9, 5, 6, 7, 8], dtype=torch.int32)
+    media = [dict(start=1, types=torch.tensor([0, 1, 2, 3]), patches=torch.zeros(1, 3, 2, 2), n_vit_h=1, n_vit_w=1)]
+    payload = media if payload_name == "media" else torch.arange(32).reshape(4, 8)
+    manager.add_one_req(UserMsg(1, ids, SamplingParams(max_tokens=1), **{payload_name: payload}))
+    original_free = len(cache.free_slots)
+    original_swa = cache.swa_available_size if cache.swa_paged else None
+    forwarded = 0
+    finished = None
+    while manager.runnable:
+        batch = manager.schedule_next_batch(3)
+        assert batch is not None
+        [req] = batch.reqs
+        assert getattr(req, payload_name) is payload
+        assert req.cached_len == forwarded
+        cache.free_swa_out_of_window_extend([req])
+        cache.allocate_paged([req])
+        forwarded += req.extend_len
+        req.cached_len = req.device_len
+        with cache.lazy_free_region():
+            cache.cache_req(req, finished=False)
+        assert cache.match_req(SimpleNamespace(input_ids=ids, input_len=len(ids), mm_embeds=None, media=media)).cuda_handle.cached_len == 0
+        assert cache.match_req(SimpleNamespace(input_ids=ids, input_len=len(ids), mm_embeds=None)).cuda_handle.cached_len == 0
+        if not isinstance(req, ChunkedReq):
+            finished = req
+    assert forwarded == len(ids)
+    with cache.lazy_free_region():
+        cache.cache_req(finished, finished=True)
+    assert len(cache.free_slots) == original_free
+    if cache.swa_paged:
+        assert cache.swa_available_size == original_swa
+    cache.check_integrity()
+
+
+@pytest.mark.parametrize("cache_type", ["radix", "swa_radix"])
+def test_hashed_image_chunks_preserve_encoder_claims_and_reuse_only_matching_content(cache_type):
+    encoder = EncoderCache(storage="cpu")
+    cache, manager = _manager(cache_type, encoder_cache=encoder)
+    pad = mm_pad_value(7)
+    ids = torch.tensor([4, pad, pad, pad, pad, 5, 6, 7], dtype=torch.int32)
+    item = MMItem(modality="image", hash=7, pad_value=pad, offsets=[[1, 5]], feature=torch.zeros(1))
+    items = [item]
+    positions = torch.arange(24, dtype=torch.int32).reshape(3, 8)
+    params = SamplingParams(max_tokens=1)
+    manager.add_one_req(UserMsg(1, ids, params, mm_items=items, mrope_positions=positions, mrope_delta=5))
+    embedding = torch.arange(32).reshape(4, 8)
+    gathered = []
+    encoded = 0
+    while manager.runnable:
+        batch = manager.schedule_next_batch(3)
+        assert batch is not None
+        [req] = batch.reqs
+        assert req.mm_items is items and req.mrope_positions_full is positions
+        assert req.mrope_delta == 5
+        assert req.media is None and req.mm_embeds is None
+        jobs, plan, rows, block_ends = plan_mm_batch(batch.reqs, encoder)
+        for job in jobs:
+            encoded += 1
+            encoder.put(job.hash, embedding)
+        expected_rows = list(range(max(1, req.cached_len) - req.cached_len,
+                                   min(5, req.device_len) - req.cached_len))
+        assert rows == expected_rows
+        assert block_ends == [5 if i in rows else 0 for i in range(req.extend_len)]
+        for uid, h, lo, hi, _, _ in plan:
+            gathered.append(encoder.get_slice(h, lo, hi, torch.device("cpu")))
+            encoder.consume(h, uid, hi - lo)
+        cache.free_swa_out_of_window_extend([req])
+        cache.allocate_paged([req])
+        req.cached_len = req.device_len
+        if not isinstance(req, ChunkedReq):
+            with cache.lazy_free_region():
+                cache.cache_req(req, finished=True)
+
+    assert encoded == 1 and torch.equal(torch.cat(gathered), embedding)
+    assert encoder.stats() == (0, 0)
+    same = PendingReq(2, ids, params, mm_items=items)
+    matched_len = (len(ids) - 1) // cache.page_size * cache.page_size
+    assert cache.match_req(same).cuda_handle.cached_len == matched_len
+    changed = ids.clone()
+    changed[1:5] = mm_pad_value(8)
+    other = MMItem(modality="image", hash=8, pad_value=mm_pad_value(8), offsets=[[1, 5]], feature=torch.zeros(1))
+    assert cache.match_req(PendingReq(3, changed, params, mm_items=[other])).cuda_handle.cached_len == 1 // cache.page_size
+    assert cache.match_req(PendingReq(4, ids, params, media=[{}])).cuda_handle.cached_len == 0
+
+    manager.add_one_req(UserMsg(2, ids, params, mm_items=items, mrope_positions=positions, mrope_delta=5))
+    [reused] = manager.schedule_next_batch(3).reqs
+    assert reused.cached_len == matched_len
+    assert plan_mm_batch([reused], encoder) == ([], [], [], [0] * reused.extend_len)
+    assert encoder.stats() == (0, 0)
+    cache.allocate_paged([reused])
+    reused.cached_len = reused.device_len
+    with cache.lazy_free_region():
+        cache.cache_req(reused, finished=True)
+    cache.check_integrity()
+
+
+@pytest.mark.parametrize("cache_type", ["radix", "swa_radix"])
+def test_atomic_images_defer_after_legacy_input_without_losing_encoder_claims(cache_type):
+    encoder = EncoderCache(storage="cpu")
+    cache, manager = _manager(cache_type, encoder_cache=encoder)
+    manager.keep_images_whole = True
+    params = SamplingParams(max_tokens=1)
+    legacy_embeds = torch.arange(8).reshape(1, 8)
+    manager.add_one_req(UserMsg(1, torch.tensor([9, 5]), params, mm_embeds=legacy_embeds))
+    pad = mm_pad_value(7)
+    items = [MMItem(modality="image", hash=7, pad_value=pad, offsets=[[lo, lo + 4]], feature=torch.zeros(1))
+             for lo in (0, 4)]
+    manager.add_one_req(UserMsg(2, torch.tensor([pad] * 8 + [5, 6]), params, mm_items=items))
+    initial_tables = manager.table_manager.available_size
+    first = manager.schedule_next_batch(5)
+    assert [(req.uid, req.extend_len) for req in first.reqs] == [(1, 2)]
+    assert first.reqs[0].mm_embeds is legacy_embeds
+    assert manager.table_manager.available_size == initial_tables - 1
+    assert encoder._entries == {}
+    cache.allocate_paged(first.reqs)
+    legacy = first.reqs[0]
+    legacy.cached_len = legacy.device_len
+    with cache.lazy_free_region():
+        cache.cache_req(legacy, finished=True)
+    manager.table_manager.free(legacy.table_idx)
+
+    embedding = torch.arange(32).reshape(4, 8)
+    gathered = []
+    encoded = 0
+    while manager.runnable:
+        batch = manager.schedule_next_batch(5)
+        assert batch is not None
+        [req] = batch.reqs
+        assert req.uid == 2
+        assert all(not lo < req.device_len < hi for item in items for lo, hi in item.offsets)
+        jobs, plan, rows, block_ends = plan_mm_batch(batch.reqs, encoder)
+        for job in jobs:
+            encoded += 1
+            encoder.put(job.hash, embedding)
+            job.feature = None
+        for uid, item_hash, lo, hi, _, _ in plan:
+            gathered.append(encoder.get_slice(item_hash, lo, hi, torch.device("cpu")))
+            encoder.consume(item_hash, uid, hi - lo)
+        assert all(block_ends[row] in (4, 8) for row in rows)
+        if req.device_len < 8:
+            assert encoder.has(7)
+        cache.free_swa_out_of_window_extend([req])
+        cache.allocate_paged([req])
+        req.cached_len = req.device_len
+        if not isinstance(req, ChunkedReq):
+            with cache.lazy_free_region():
+                cache.cache_req(req, finished=True)
+            manager.table_manager.free(req.table_idx)
+
+    assert encoded == 1
+    assert torch.equal(torch.cat(gathered), embedding.repeat(2, 1))
+    assert encoder._entries == {}
+    assert manager.table_manager.available_size == initial_tables
+    cache.check_integrity()
diff --git a/tests/server/test_anthropic_api.py b/tests/server/test_anthropic_api.py
index 327bfcd87..b32064ee2 100644
--- a/tests/server/test_anthropic_api.py
+++ b/tests/server/test_anthropic_api.py
@@ -433,6 +433,7 @@ def __init__(self, outputs):
         )
         self._uid = 0
         self._count_manager = None
+        self.last_sent = None
 
     def frontend_tokenizer(self):
         return self._count_manager
@@ -442,6 +443,7 @@ def new_user(self):
         return self._uid
 
     async def send_one(self, msg):
+        self.last_sent = msg
         return None
 
     async def wait_for_ack(self, uid):
@@ -741,6 +743,22 @@ def test_count_tokens_excluded_from_request_ring():
     request_ring.reset()
 
 
+def test_anthropic_preserves_images_in_messages_and_tool_results():
+    source = {"type": "base64", "media_type": "image/png", "data": "YWJj"}
+    req = AnthropicMessagesRequest.model_validate({"model": "deepseek-v41", "max_tokens": 4,
+        "messages": [
+            {"role": "user", "content": [{"type": "image", "source": source}, {"type": "text", "text": "describe"}]},
+            {"role": "assistant", "content": [{"type": "tool_use", "id": "call1", "name": "capture", "input": {}}]},
+            {"role": "user", "content": [{"type": "tool_result", "tool_use_id": "call1", "content": [
+                {"type": "text", "text": "screenshot"}, {"type": "image", "source": source},
+            ]}]},
+        ]})
+    messages, *_ = A.convert_anthropic_prompt(req, preserve_tool_images=True)
+    assert messages[0]["content"][0]["freetoken_ref"] == {"kind": "b64", "data": source["data"]}
+    assert messages[2]["role"] == "tool"
+    assert messages[2]["content"][1]["freetoken_ref"] == {"kind": "b64", "data": source["data"]}
+
+
 def test_count_tokens_image_only_message_400():
     # Message list is non-empty on the wire but empty after block filtering: the neutral
     # count_prompt_tokens raises ValueError -> 400, not a 500 from an empty chat template.
@@ -928,7 +946,8 @@ def test_tool_result_image_moves_to_the_following_user_turn():
     assert "vision" in r.json()["error"]["message"]
 
 
-def test_tool_result_image_outside_a_user_message_is_rejected():
+@pytest.mark.parametrize("preserve_tool_images", [False, True])
+def test_tool_result_image_outside_a_user_message_is_rejected(preserve_tool_images):
     # Anthropic only allows tool_result in user messages; an image there has no user turn to ride on.
     body = {
         "model": "claude-x",
@@ -948,7 +967,7 @@ def test_tool_result_image_outside_a_user_message_is_rejected():
         ],
     }
     with pytest.raises(ValueError, match="user message"):
-        A.convert_anthropic_prompt(AnthropicMessagesRequest.model_validate(body))
+        A.convert_anthropic_prompt(AnthropicMessagesRequest.model_validate(body), preserve_tool_images=preserve_tool_images)
 
 
 def test_image_only_tool_result_keeps_an_empty_tool_message():
@@ -976,3 +995,59 @@ def test_image_only_tool_result_keeps_an_empty_tool_message():
         "role": "user",
         "content": [{"type": "image", "freetoken_ref": {"kind": "b64", "data": "aGk="}}],
     }
+
+
+@pytest.mark.parametrize("model_cls", ["DeepseekV41ForCausalLM", "Qwen4ExpForCausalLM"])
+def test_tool_images_use_model_encoding_and_match_token_count(model_cls):
+    body = {
+        "model": "client-alias", "max_tokens": 16,
+        "messages": [
+            {"role": "user", "content": "compare the screenshots"},
+            {"role": "assistant", "content": [
+                {"type": "tool_use", "id": "call_1", "name": "shot", "input": {}},
+                {"type": "tool_use", "id": "call_2", "name": "shot", "input": {}},
+            ]},
+            {"role": "user", "content": [
+                {"type": "tool_result", "tool_use_id": "call_1", "content": [
+                    {"type": "text", "text": "before"},
+                    {"type": "image", "source": {"type": "base64", "media_type": "image/png", "data": "Zmlyc3Q="}},
+                    {"type": "text", "text": "after"},
+                ]},
+                {"type": "tool_result", "tool_use_id": "call_2", "content": [
+                    {"type": "image", "source": {"type": "base64", "media_type": "image/png", "data": "c2Vjb25k"}},
+                ]},
+                {"type": "text", "text": "compare now"},
+            ]},
+        ],
+    }
+    original = json.dumps(body)
+    fake = FakeState([("done", True, 5, 1)])
+    fake.config.model_spec = SimpleNamespace(model_cls=model_cls)
+    fake.config.served_modalities = frozenset({"image"})
+    fake._count_manager = _FakeTokenizeManager()
+    client = _client(fake)
+    response = client.post("/v1/messages", json=body)
+    assert response.status_code == 200, response.text
+    generated = fake.last_sent
+    if model_cls == "DeepseekV41ForCausalLM":
+        assert generated.text[-3:] == [
+            {"role": "tool", "tool_call_id": "call_1", "content": [
+                {"type": "text", "text": "before"}, {"type": "image"}, {"type": "text", "text": "after"},
+            ]},
+            {"role": "tool", "tool_call_id": "call_2", "content": [{"type": "image"}]},
+            {"role": "user", "content": "compare now"},
+        ]
+    else:
+        assert generated.text[-3:] == [
+            {"role": "tool", "tool_call_id": "call_1", "content": "beforeafter"},
+            {"role": "tool", "tool_call_id": "call_2", "content": ""},
+            {"role": "user", "content": [
+                {"type": "image"}, {"type": "image"}, {"type": "text", "text": "compare now"},
+            ]},
+        ]
+    assert generated.images == [b"first", b"second"]
+    response = client.post("/v1/messages/count_tokens", json={k: v for k, v in body.items() if k != "max_tokens"})
+    assert response.status_code == 200, response.text
+    counted = fake._count_manager.msgs[-1]
+    assert counted.text == generated.text and counted.images == generated.images
+    assert json.dumps(body) == original
diff --git a/tests/server/test_args.py b/tests/server/test_args.py
new file mode 100644
index 000000000..f1c332e2f
--- /dev/null
+++ b/tests/server/test_args.py
@@ -0,0 +1,26 @@
+"""CLI cache sizing is validated before loading the checkpoint."""
+
+import pytest
+
+from freetoken.server.args import parse_args
+
+
+@pytest.mark.parametrize("ratio", ["0.01", "0.02", "0.2", "1"])
+def test_window_pool_ratio_reaches_serving_config(ratio):
+    config, _ = parse_args([
+        "--model", "/models/local", "--dtype", "bfloat16",
+        "--tool-call-parser", "llama3", "--reasoning-parser", "off",
+        "--swa-full-tokens-ratio", ratio,
+    ])
+    assert config.swa_full_tokens_ratio == float(ratio)
+
+
+@pytest.mark.parametrize("ratio", ["0", "-0.1", "1.01", "nan", "inf", "invalid"])
+def test_invalid_window_pool_ratio_fails_before_model_lookup(ratio, monkeypatch):
+    def unexpected_lookup(*args, **kwargs):
+        pytest.fail("invalid ratio reached checkpoint loading")
+
+    monkeypatch.setattr("freetoken.utils.cached_load_hf_config", unexpected_lookup)
+    with pytest.raises(SystemExit) as error:
+        parse_args(["--model", "/models/local", "--swa-full-tokens-ratio", ratio])
+    assert error.value.code == 2
diff --git a/tests/server/test_function_call_parser.py b/tests/server/test_function_call_parser.py
index 34c1c04d1..87e6a48a2 100644
--- a/tests/server/test_function_call_parser.py
+++ b/tests/server/test_function_call_parser.py
@@ -275,6 +275,25 @@ def _feed(parser, chunks):
     return texts, calls
 
 
+def test_deepseek_v41_calls_across_every_split():
+    block = (
+        '<｜DSML｜ calls><｜DSML｜ invoke name="get_weather">'
+        '<｜DSML｜ parameter name="city" string="true">Tokyo</｜DSML｜ parameter>'
+        '<｜DSML｜ parameter name="days" string="false">2</｜DSML｜ parameter>'
+        '</｜DSML｜ invoke></｜DSML｜ calls>'
+    )
+    parser = FunctionCallParser(TOOLS, tool_call_parser="deepseekv41")
+    complete = parser.parse_non_stream(block)
+    assert complete.calls[0].name == "get_weather"
+    assert json.loads(complete.calls[0].parameters) == {"city": "Tokyo", "days": 2}
+    for split in range(1, len(block)):
+        parser = FunctionCallParser(TOOLS, tool_call_parser="deepseekv41")
+        texts, calls = _feed(parser, [block[:split], block[split:], ""])
+        assert not ("".join(texts) + parser.finish_stream()).strip()
+        assert [call.name for call in calls if call.name] == ["get_weather"]
+        assert json.loads("".join(call.parameters for call in calls)) == {"city": "Tokyo", "days": 2}
+
+
 @pytest.mark.parametrize("parser_name", ["qwen25", "glm47", "gemma4", "minimax", "deepseekv32", "qwen3_coder"])
 def test_streaming_plain_text_releases_per_chunk(parser_name):
     # A pure-text response must stream out chunk by chunk, not buffer to the end.
diff --git a/tests/server/test_message_wire.py b/tests/server/test_message_wire.py
index e22866715..d339f13e0 100644
--- a/tests/server/test_message_wire.py
+++ b/tests/server/test_message_wire.py
@@ -28,6 +28,26 @@
 from freetoken.core import SamplingParams
 
 
+def test_multidimensional_cpu_tensor_wire_roundtrip_and_validation():
+    import pytest
+    import torch
+    from freetoken.message import UserMsg
+    from freetoken.message.utils import serialize_type, deserialize_type
+
+    for dtype in (torch.float32, torch.bfloat16, torch.int64):
+        source = torch.arange(24).to(dtype).reshape(2, 3, 4).transpose(0, 1)
+        payload = serialize_type(source)
+        result = deserialize_type({}, payload)
+        torch.testing.assert_close(result, source)
+        payload["shape"] = [25]
+        with pytest.raises(ValueError, match="shape"):
+            deserialize_type({}, payload)
+    legacy = {"__type__": "Tensor", "dtype": "torch.int32", "buffer": torch.tensor([4, 5], dtype=torch.int32).numpy().tobytes()}
+    assert deserialize_type({}, legacy).tolist() == [4, 5]
+    empty = torch.empty(0, 3)
+    assert deserialize_type({}, serialize_type(empty)).shape == (0, 3)
+
+
 def test_cache_rebuild_msg_roundtrip():
     msg = CacheRebuildMsg(request_id="abc", moe_cache_size=8, num_pages=1024, mode="if_idle")
     out = BaseTokenizerMsg.decoder(BaseTokenizerMsg.encoder(msg))
diff --git a/tests/server/test_openai_api.py b/tests/server/test_openai_api.py
index e33018e19..b1c2ea7cc 100644
--- a/tests/server/test_openai_api.py
+++ b/tests/server/test_openai_api.py
@@ -51,6 +51,23 @@ async def wait_for_ack(self, uid: int):
             yield reply
 
 
+def test_openai_image_request_preserves_content_and_numeric_effort():
+    state = FakeState([UserReply(uid=42, incremental_output="red", finished=True)])
+    state.config.served_modalities = frozenset({"image"})
+    request = ChatCompletionRequest.model_validate({
+        "model": "deepseek-v41", "reasoning_effort": 63,
+        "messages": [{"role": "user", "content": [
+            {"type": "text", "text": "what color?"},
+            {"type": "image_url", "image_url": {"url": "data:image/png;base64,YWJj"}},
+        ]}],
+    })
+    result = run(handle_chat_completion(request, request=None, state=state, model_sampling={}))
+    assert result["choices"][0]["message"]["content"] == "red"
+    assert state.sent.text[0]["content"][1] == {"type": "image"}
+    assert state.sent.images == [b"abc"]
+    assert state.sent.chat_template_kwargs["reasoning_effort"] == 63
+
+
 def tool_schema():
     return [
         {
diff --git a/tests/server/test_parser_auto_selection.py b/tests/server/test_parser_auto_selection.py
index 78b1663e8..c432eac94 100644
--- a/tests/server/test_parser_auto_selection.py
+++ b/tests/server/test_parser_auto_selection.py
@@ -94,3 +94,7 @@ def test_an_explicit_choice_beats_inference():
         pinned, _ = parse_args(["--model", ANON_PATH, "--reasoning-parser", "qwen3"])
     assert off.reasoning_parser is None
     assert pinned.reasoning_parser == "qwen3"
+
+
+def test_deepseek_v41_uses_space_delimited_dsml():
+    assert _inferred("DeepseekV41ForCausalLM") == ("deepseekv41", "deepseekv32")
diff --git a/tests/server/test_reasoning_parser_dsv4.py b/tests/server/test_reasoning_parser_dsv4.py
index a624870b4..1097e2ccc 100644
--- a/tests/server/test_reasoning_parser_dsv4.py
+++ b/tests/server/test_reasoning_parser_dsv4.py
@@ -47,6 +47,16 @@ def test_thinking_with_end_token():
     assert content == "The answer is 42."
 
 
+def test_v41_space_delimited_calls_end_reasoning():
+    block = '<｜DSML｜ calls><｜DSML｜ invoke name="get_weather"></｜DSML｜ invoke></｜DSML｜ calls>'
+    text = "Let me check." + block
+    for split in range(1, len(text)):
+        parser = ReasoningParser("deepseekv32", force_reasoning=True)
+        reasoning, content = _stream(parser, [text[:split], text[split:]])
+        assert reasoning == "Let me check."
+        assert content == block
+
+
 def test_thinking_with_tool_block_after_end_token():
     parser = ReasoningParser("deepseekv32", force_reasoning=True)
     text = f"Let me check the weather.</think>Sure!\n\n{TOOL_BLOCK}"
diff --git a/tests/server/test_responses_api.py b/tests/server/test_responses_api.py
index 887e78e57..96538f22f 100644
--- a/tests/server/test_responses_api.py
+++ b/tests/server/test_responses_api.py
@@ -15,6 +15,8 @@
 import sys
 from types import SimpleNamespace
 
+import pytest
+
 _ROOT = os.path.dirname(os.path.dirname(os.path.dirname(os.path.abspath(__file__))))
 _PY = os.path.join(_ROOT, "python")
 if _PY not in sys.path:
@@ -62,6 +64,19 @@ def test_convert_string_input_and_instructions():
     assert spec.sampling_params.max_tokens == 50
 
 
+def test_responses_preserves_image_between_text_parts():
+    req = ResponsesRequest.model_validate({"model": "deepseek-v41", "input": [
+        {"role": "user", "content": [
+            {"type": "input_text", "text": "before"},
+            {"type": "input_image", "image_url": "https://example.com/p.png"},
+            {"type": "input_text", "text": "after"},
+        ]},
+    ]})
+    content = RP.convert_responses_to_genspec(req, {}).messages[0]["content"]
+    assert [p["type"] for p in content] == ["text", "image", "text"]
+    assert content[1]["freetoken_ref"] == {"kind": "url", "data": "https://example.com/p.png"}
+
+
 def test_convert_defaults_max_output_tokens_when_omitted():
     # codex omits max_output_tokens; must NOT fall to the 16-token floor (bug b1).
     req = ResponsesRequest.model_validate({"model": "gpt-x", "input": "hi"})
@@ -98,6 +113,68 @@ def test_convert_list_input_with_tool_roundtrip_and_tools():
     assert spec.parse_tools
 
 
+@pytest.mark.parametrize("model_cls", ["DeepseekV41ForCausalLM", "Qwen4ExpForCausalLM"])
+def test_function_output_preserves_ordered_images_and_text(model_cls):
+    output = [
+        {"type": "input_text", "text": "first image"},
+        {"type": "input_image", "image_url": "data:image/png;base64,Zmlyc3Q="},
+        {"type": "input_text", "text": "second image"},
+        {"type": "input_image", "image_url": "data:image/png;base64,aW1hZ2U="},
+        {"type": "input_text", "text": "compare them"},
+    ]
+    fake = FakeState([("same", True, 5, 1)])
+    fake.config.served_modalities = frozenset({"image"})
+    fake.config.model_spec = SimpleNamespace(model_cls=model_cls)
+    response = _client(fake).post("/v1/responses", json={"model": "deepseek-v41", "input": [
+        {"type": "function_call", "call_id": "call_image", "name": "inspect", "arguments": "{}"},
+        {"type": "function_call_output", "call_id": "call_image", "output": output},
+    ]})
+    assert response.status_code == 200, response.text
+    if model_cls == "DeepseekV41ForCausalLM":
+        assert fake.last_sent.text[-1] == {
+            "role": "tool", "tool_call_id": "call_image", "content": [
+                {"type": "text", "text": "first image"},
+                {"type": "image"},
+                {"type": "text", "text": "second image"},
+                {"type": "image"},
+                {"type": "text", "text": "compare them"},
+            ],
+        }
+    else:
+        assert fake.last_sent.text[-2:] == [
+            {"role": "tool", "tool_call_id": "call_image", "content": "first imagesecond imagecompare them"},
+            {"role": "user", "content": [{"type": "image"}, {"type": "image"}]},
+        ]
+    assert fake.last_sent.images == [b"first", b"image"]
+
+
+@pytest.mark.parametrize("output", [
+    {"result": [1, 2], "text": "raw object"},
+    [1, {"value": 2}, "three"],
+    [{"type": "input_text", "text": "data"}, {"result": 2}],
+    {"type": "input_image", "image_url": "https://example.com/data-field.png"},
+    [],
+])
+@pytest.mark.parametrize("preserve_tool_images", [False, True])
+def test_function_output_preserves_non_media_json(output, preserve_tool_images):
+    req = ResponsesRequest(model="m", input=[
+        {"type": "function_call_output", "call_id": "call_1", "output": output},
+    ])
+    message = RP.convert_responses_to_genspec(req, {}, preserve_tool_images=preserve_tool_images).messages[0]
+    assert message["content"] == json.dumps(output)
+
+
+def test_function_output_rejects_uploaded_image_file_ids():
+    fake = FakeState([])
+    response = _client(fake).post("/v1/responses", json={"model": "deepseek-v41", "input": [
+        {"type": "function_call_output", "call_id": "call_1", "output": [
+            {"type": "input_image", "file_id": "file_1"},
+        ]},
+    ]})
+    assert response.status_code == 400
+    assert "image_url" in response.json()["error"]["message"]
+
+
 def test_convert_function_call_output_image_moves_to_a_user_turn():
     # codex's view_image returns the image as the tool output; templates render tool messages as text.
     req = ResponsesRequest.model_validate(
@@ -125,6 +202,21 @@ def test_convert_function_call_output_image_moves_to_a_user_turn():
     ]
 
 
+@pytest.mark.parametrize("preserve_tool_images", [False, True])
+def test_image_only_function_output_keeps_call_id_and_image(preserve_tool_images):
+    req = ResponsesRequest(model="alias", input=[{
+        "type": "function_call_output", "call_id": "call_image", "output": [
+            {"type": "input_image", "image_url": "data:image/png;base64,aGk="},
+        ],
+    }])
+    messages = RP.convert_responses_to_genspec(req, {}, preserve_tool_images=preserve_tool_images).messages
+    image = {"type": "image", "freetoken_ref": {"kind": "url", "data": "data:image/png;base64,aGk="}}
+    assert messages[0] == {
+        "role": "tool", "tool_call_id": "call_image", "content": [image] if preserve_tool_images else "",
+    }
+    assert messages[1:] == ([] if preserve_tool_images else [{"role": "user", "content": [image]}])
+
+
 def test_convert_reasoning_item_merges_into_assistant_turn():
     # One turn's reasoning/message/function_call items fold into ONE assistant message.
     req = ResponsesRequest.model_validate(
@@ -949,7 +1041,8 @@ def test_convert_reasoning_toggle_broadcasts_every_spelling():
         }, parser
 
 
-def test_convert_function_call_output_text_list_stays_a_plain_tool_message():
+@pytest.mark.parametrize("preserve_tool_images", [False, True])
+def test_convert_function_call_output_text_list_stays_a_plain_tool_message(preserve_tool_images):
     req = ResponsesRequest.model_validate(
         {
             "model": "gpt-x",
@@ -964,6 +1057,6 @@ def test_convert_function_call_output_text_list_stays_a_plain_tool_message():
             ],
         }
     )
-    spec = RP.convert_responses_to_genspec(req, {})
+    spec = RP.convert_responses_to_genspec(req, {}, preserve_tool_images=preserve_tool_images)
     assert [m["role"] for m in spec.messages] == ["user", "assistant", "tool"]
     assert spec.messages[2]["content"] == "ab"
diff --git a/tests/tokenizer/test_tokenize.py b/tests/tokenizer/test_tokenize.py
index 369576a18..a50e5dff9 100644
--- a/tests/tokenizer/test_tokenize.py
+++ b/tests/tokenizer/test_tokenize.py
@@ -27,6 +27,151 @@ def encode(self, prompt, return_tensors=None, add_special_tokens=True):
         return torch.tensor([[1, 2, 3]], dtype=torch.long)
 
 
+class FakeDsv41Tokenizer:
+    chat_template = "must use the V4.1 encoder instead"
+    unk_token_id = -1
+
+    def __init__(self, folder):
+        self.name_or_path = str(folder)
+        self.prompt = None
+        (folder / "config.json").write_text(json.dumps({
+            "model_type": "deepseek_v41", "image_token_id": 9,
+            "vision_config": {"num_hidden_layers": 1, "patch_size": 2,
+                              "downsample_ratio": 2, "min_pixels": 16, "max_image_tokens": 24},
+        }))
+
+    def convert_tokens_to_ids(self, token):
+        return 9
+
+    def encode(self, prompt, return_tensors=None, add_special_tokens=True):
+        from freetoken.models.deepseek_v41.encoding import IMAGE_PLACEHOLDER
+        assert not add_special_tokens
+        self.prompt = prompt
+        parts = prompt.split(IMAGE_PLACEHOLDER)
+        ids = []
+        for i, part in enumerate(parts):
+            if i:
+                ids.append(9)
+            ids.extend(100 + ord(c) for c in part)
+        return ids if return_tensors is None else torch.tensor([ids])
+
+
+def test_dsv41_numeric_effort_and_image_span_survive_tokenization(tmp_path):
+    import base64
+    import io
+    from PIL import Image
+    from freetoken.message import BaseBackendMsg, UserMsg
+
+    buffer = io.BytesIO()
+    Image.new("RGB", (8, 4), "red").save(buffer, format="PNG")
+    tokenizer = FakeDsv41Tokenizer(tmp_path)
+    manager = TokenizeManager(tokenizer)
+    msg = TokenizeMsg(uid=7, text=[{"role": "user", "content": [
+        {"type": "text", "text": "What is this?"},
+        {"type": "image_url", "image_url": {"url": "data:image/png;base64," + base64.b64encode(buffer.getvalue()).decode()}},
+    ]}], sampling_params=SamplingParams(), chat_template_kwargs={"enable_thinking": True, "reasoning_effort": 63})
+    [wire] = manager.tokenize([msg])
+    ids = wire.input_ids
+    assert "Reasoning Effort: 63" in tokenizer.prompt
+    assert len(msg.media) == 1
+    item = msg.media[0]
+    assert item["types"].tolist() == [0, 1, 1, 2, 3]
+    assert ids[item["start"]:item["start"] + 5].tolist() == [9] * 5
+    received = BaseBackendMsg.decoder(wire.encoder())
+    torch.testing.assert_close(received.media[0]["patches"], item["patches"])
+    assert received.media[0]["patches"].shape == (8, 3, 2, 2)
+    assert received.media[0]["start"] == item["start"]
+
+
+def test_dsv41_main_mm_wire_keeps_reordered_tool_images_aligned(tmp_path):
+    import base64
+    import io
+    from PIL import Image
+    from freetoken.message import BaseBackendMsg, BaseTokenizerMsg
+    from freetoken.mm.config import MultimodalConfig
+    from freetoken.mm.media import collect_image_refs
+    from freetoken.models.deepseek_v41.mm_processor import DeepseekV41MMProcessor
+    from freetoken.server.generation import render_messages
+    from freetoken.utils.hf import RawConfigShim
+
+    def png(color):
+        buffer = io.BytesIO()
+        Image.new("RGB", (8, 4), color).save(buffer, format="PNG")
+        return buffer.getvalue()
+
+    red, blue = png("red"), png("blue")
+    tokenizer = FakeDsv41Tokenizer(tmp_path)
+    hf = RawConfigShim(json.loads((tmp_path / "config.json").read_text()))
+    processor = DeepseekV41MMProcessor(hf, str(tmp_path), MultimodalConfig())
+    manager = TokenizeManager(tokenizer, processor)
+    messages = [{"role": "assistant", "tool_calls": [
+        {"id": name, "type": "function", "function": {"name": "capture", "arguments": "{}"}}
+        for name in ("first", "second")
+    ]}]
+    for name, image in (("second", blue), ("first", red)):
+        messages.append({"role": "tool", "tool_call_id": name, "content": [
+            {"type": "image_url", "image_url": {"url": "data:image/png;base64," + base64.b64encode(image).decode()}},
+        ]})
+    messages = render_messages(messages)
+    msg = TokenizeMsg(7, messages, SamplingParams(), {"enable_thinking": True, "reasoning_effort": 63})
+    assert "Reasoning Effort: 63" in manager.render_prompt(msg)
+    refs = collect_image_refs(messages)
+    assert len(refs) == 2
+    msg.images = [blue, red]
+    decoded = BaseTokenizerMsg.decoder(BaseTokenizerMsg.encoder(msg))
+    result, = manager.tokenize([decoded])
+    received = BaseBackendMsg.decoder(result.encoder())
+    assert received.media is None and received.mm_embeds is None
+    assert received.mrope_positions is None and received.mrope_delta == 0
+    assert len(received.mm_items) == 2
+    first, second = received.mm_items
+    assert first.hash != second.hash and first.pad_value != second.pad_value
+    assert first.feature[:, 0].mean() > second.feature[:, 0].mean()
+    for item in received.mm_items:
+        assert item.feature.dtype == torch.bfloat16 and item.feature.device.type == "cpu"
+        assert item.types == [0, 1, 1, 2, 3]
+        assert item.feature.shape == (8, 3, 2, 2)
+        start, end = item.offsets[0]
+        assert received.input_ids[start:end].tolist() == [item.pad_value] * 5
+    assert decoded.text == messages
+    with pytest.raises(ValueError, match="image parts"):
+        manager.tokenize([TokenizeMsg(8, messages, SamplingParams(), images=[red])])
+
+
+def test_dsv41_effort_range_is_validated(tmp_path):
+    manager = TokenizeManager(FakeDsv41Tokenizer(tmp_path))
+    for effort in (0, 101):
+        with pytest.raises(ValueError, match="between 1 and 100"):
+            manager.render_prompt(TokenizeMsg(1, [{"role": "user", "content": "hi"}], SamplingParams(),
+                                              {"enable_thinking": True, "reasoning_effort": effort}))
+
+
+def test_text_model_rejects_images_before_jinja():
+    manager = TokenizeManager(FakeTokenizer())
+    with pytest.raises(ValueError, match="does not support image"):
+        manager.render_prompt(TokenizeMsg(1, [{"role": "user", "content": [
+            {"type": "image_url", "image_url": {"url": "https://example.com/p.png"}},
+        ]}], SamplingParams()))
+
+
+def test_dsv41_images_follow_sorted_tool_result_order():
+    from freetoken.models.deepseek_v41.encoding import encode_messages
+
+    tool_calls = [
+        {"id": name, "type": "function", "function": {"name": "capture", "arguments": "{}"}}
+        for name in ("first", "second")
+    ]
+    messages = [{"role": "assistant", "tool_calls": tool_calls}]
+    for name in ("second", "first"):
+        messages.append({"role": "tool", "tool_call_id": name, "content": [
+            {"type": "image_url", "image_url": {"url": f"https://example.com/{name}.png"}},
+        ]})
+    _, payload = encode_messages(messages, "thinking", return_multi_modal_data=True)
+    assert [item["url"] for item in payload["images"]] == [
+        "https://example.com/first.png", "https://example.com/second.png",
+    ]
+
+
 def test_tokenize_manager_passes_chat_template_kwargs():
     tokenizer = FakeTokenizer()
     manager = TokenizeManager(tokenizer)