perf(dflash): split-chain fast rollback with F32 SSM checkpoints (+29% decode)

.github/workflows/ci.yml

@@ -279,7 +279,7 @@ jobs:
             -o "$RUNNER_TEMP/hip_smoke" .github/ci/hip_smoke.cpp
           "$RUNNER_TEMP/hip_smoke"
 
-      - name: Build + test ROCmFP formats
+      - name: Build + test ROCm formats and inference core
         run: |
           cmake -S server -B "$RUNNER_TEMP/rocmfp-build" \
             -DDFLASH27B_GPU_BACKEND=hip \
@@ -290,11 +290,11 @@ jobs:
             -DCMAKE_HIP_FLAGS=-DDFLASH_WAVE_SIZE=32
           cmake --build "$RUNNER_TEMP/rocmfp-build" \
             --target test_rocmfp4 test_rocmfpx test_rocmfp4_hip_tail test_rocmfpx_mmq \
-              test_deepseek4_mmid_grouped_cuda \
+              test_deepseek4_mmid_grouped_cuda test_recurrent_snapshot test_server_unit \
             --parallel 8
           ctest --test-dir "$RUNNER_TEMP/rocmfp-build" \
             --output-on-failure \
-            -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail|rocmfpx_mmq|deepseek4_mmid_grouped_cuda'
+            -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail|rocmfpx_mmq|deepseek4_mmid_grouped_cuda|recurrent_snapshot|ChainRollbackPolicy'
 
   build-windows:
     name: Build Windows (MSVC + CUDA, library + server targets)

README.md

@@ -346,6 +346,8 @@ When compression is on, the request path picks one of three modes automatically,
 | `DFLASH27B_KV_TQ3=1` | (default) | Preset TQ3_0 K+V (3.5 bpv, fits 256K @ 24 GB) |
 | `DFLASH27B_KV_Q4=1` | off | Q4_0 K+V (4.5 bpv, legacy, ~128K ceiling) |
 | `--prefix-cache-slots N` | — | Live prefix-cache slot count |
+| `DFLASH_PREFIX_CACHE_SLOTS=N` | `32` | Container-entrypoint equivalent of `--prefix-cache-slots`; the native binary itself uses the CLI flag. |
+| `DFLASH_PREFILL_CACHE_SLOTS=N` | `0` | Container-entrypoint equivalent of `--prefill-cache-slots`; the native binary itself uses the CLI flag. |
 | `--kv-cache-dir <path>` | — | Persist prefix cache to disk |
 | `--kv-cache-budget N` | — | On-disk cache size cap |
 
@@ -379,6 +381,8 @@ Pages the attention KV cache through a fixed pool of GPU slots; cold 64-token ch
 | `--target-gpu N` | `0` | Target GPU index |
 | `--draft-gpu N` | same as target | Draft GPU index; offload draft to a second GPU |
 | `--target-devices <list>` / `--target-layer-split` | single GPU | Layer-split target across GPUs |
+| `--target-split-fast-rollback` | off | Qwen35 local layer-split only: enable exact F32 per-token checkpoints and skip accepted-token replay. Adds checkpoint VRAM (~1.65 GiB for the measured Qwen3.6-27B q=16 split). |
+| `DFLASH_SPLIT_FAST_ROLLBACK=1` | off | Environment equivalent of `--target-split-fast-rollback`. |
 | `--draft-ipc-bin <path>` | — | Out-of-process draft binary (mixed CUDA/HIP) |
 | `--peer-access` | off | Enable P2P between target GPUs |
 | `--chunk N` | backend default | Prefill ubatch size |
@@ -394,6 +398,7 @@ For MoE targets (`laguna`, `qwen35`/`qwen36`) whose experts don't fit in VRAM. `
 | Flag / env | Default | Effect |
 |---|---|---|
 | `--spark` | off | One-flag autotune: enable the bounded expert cache, size it from the VRAM target, auto-load and keep persisting a placement profile (`<model>.gguf.spark.csv`). |
+| `--spark-slots <N>` | auto | Explicit expert-cache slots per layer; overrides Spark auto-sizing. |
 | `--spark-vram <GiB>` | whole card | Total VRAM Spark may use; it sizes the hot tier + cache + KV under this cap. |
 | `DFLASH_SPARK=1` | off | Env equivalent of `--spark`. |
 | `DFLASH_SPARK_VRAM_MB=N` | — | Env equivalent of `--spark-vram` (in MB). |

server/CMakeLists.txt

@@ -1026,6 +1026,22 @@ if(DFLASH27B_TESTS)
                 "${CMAKE_CURRENT_SOURCE_DIR}/scripts/entrypoint.sh")
     endif()
 
+    # CPU-only contract test for the fail-closed layer-split tree boundary.
+    add_executable(test_qwen35_split_tree_guard
+        test/test_qwen35_split_tree_guard.cpp)
+    target_include_directories(test_qwen35_split_tree_guard PRIVATE
+        ${CMAKE_CURRENT_SOURCE_DIR}/src/qwen35)
+    list(APPEND _raw_unit_test_targets test_qwen35_split_tree_guard)
+
+    if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_recurrent_snapshot.cpp")
+        add_executable(test_recurrent_snapshot test/test_recurrent_snapshot.cpp)
+        target_include_directories(test_recurrent_snapshot PRIVATE
+            ${DFLASH27B_SRC_INCLUDE_DIRS})
+        target_link_libraries(test_recurrent_snapshot PRIVATE
+            dflash_common ggml ${DFLASH27B_GGML_BACKEND_TARGET})
+        list(APPEND _raw_unit_test_targets test_recurrent_snapshot)
+    endif()
+
     if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_server_unit.cpp")
         set(_server_unit_sources
             test/test_unit_main.cpp
@@ -1170,6 +1186,10 @@ if(DFLASH27B_TESTS)
                     set(_unit_ctest_name draft_topk_cuda)
                 elseif(_unit_target STREQUAL "test_deepseek4_unit")
                     set(_unit_ctest_name deepseek4_unit)
+                elseif(_unit_target STREQUAL "test_qwen35_split_tree_guard")
+                    set(_unit_ctest_name qwen35_split_tree_guard)
+                elseif(_unit_target STREQUAL "test_recurrent_snapshot")
+                    set(_unit_ctest_name recurrent_snapshot)
                 endif()
                 add_test(NAME "${_unit_ctest_name}" COMMAND ${_unit_target})
                 if(_unit_target STREQUAL "test_deepseek4_mmid_grouped_cuda")

server/docs/ENVIRONMENT.md

@@ -33,6 +33,12 @@ consolidation of this list into CLI flags is tracked as follow-up work.
 | `DFLASH_MOE_TP_*` / `DFLASH_MOE_HYBRID_PREFILL_EAGER` | unset | BURN-IN: model-neutral names for common heterogeneous-MoE scheduling and kernel policy. Existing `DFLASH_DS4_*` names remain compatibility aliases. |
 | `DFLASH_MMID_TELEMETRY` | unset | DEBUG: report MUL_MAT_ID dispatch, MMVQ variant, and per-node graph compatibility. |
 | `DFLASH_KVFLASH` | unset | Prefer the CLI: `--kvflash` (token count or `auto`). |
+| `DFLASH_PREFIX_CACHE_SLOTS` | 32 | Container-entrypoint equivalent of `--prefix-cache-slots`; not read directly by the native binary. |
+| `DFLASH_PREFILL_CACHE_SLOTS` | 0 | Container-entrypoint equivalent of `--prefill-cache-slots`; not read directly by the native binary. |
+| `DFLASH_SPLIT_FAST_ROLLBACK` | unset | OPT-IN: exact F32 checkpoints and replay-free rollback for local qwen35 target layer splits. Prefer `--target-split-fast-rollback`; adds checkpoint VRAM (~1.65 GiB for the measured Qwen3.6-27B q=16 split). |
+| `DFLASH_STALL_TOOL_PREFIX` | unset | OPT-IN: recover a stalled tool call by injecting the prepared tool prefix when generation stops after an action suffix. |
+| `DFLASH_DS4_SPEC` / `DFLASH_DS4_DRAFT` / `DFLASH_DS4_DRAFT_GPU` | unset | OPT-IN: enable DeepSeek4 DSpark, select its draft GGUF, and optionally select the local drafter GPU. See `DS4.md`. |
+| `DFLASH_DS4_CUDA_LAYERS` | auto | Override the DeepSeek4 heterogeneous layer-split heuristic. See `DS4.md`. |
 
 ## Full inventory (generated)
 
@@ -53,6 +59,8 @@ consolidation of this list into CLI flags is tracked as follow-up work.
 - `DFLASH_ADAPTIVE_WIDTH_MIN` - adaptive_verify_width.h
 - `DFLASH_ADAPTIVE_WIDTH_THETA` - adaptive_verify_width.h
 - `DFLASH_COLD_THREADS` - moe_expert_compute_cpu.cpp
+- `DFLASH_CUDA_MMVQ_MOE_ALIGN_SHARED_IDS` - moe_hybrid_ffn_eval.cpp
+- `DFLASH_CUDA_MMVQ_MOE_KERNEL` - moe_hybrid_ffn_eval.cpp
 - `DFLASH_DISABLE_DRAFT_ATTN` - draft_graph.cpp
 - `DFLASH_DISABLE_DRAFT_ATTN_GATE` - draft_graph.cpp
 - `DFLASH_DISABLE_DRAFT_AUX_NORMS` - draft_graph.cpp
@@ -64,14 +72,29 @@ consolidation of this list into CLI flags is tracked as follow-up work.
 - `DFLASH_DRAFT_KV` - laguna_backend.cpp, qwen35_backend.cpp
 - `DFLASH_DRAFT_PERSIST` - laguna_backend.cpp
 - `DFLASH_DROP_COLD` - qwen35moe_backend.cpp, qwen35moe_pipelined_decode.cpp
+- `DFLASH_DS4_ADAPTIVE_WIDTH` - deepseek4_dspark_spec.cpp
+- `DFLASH_DS4_CUDA_LAYERS` - deepseek4_layer_split_adapter.cpp
+- `DFLASH_DS4_DENSE_TP_MASK` - deepseek4_loader.cpp
+- `DFLASH_DS4_DENSE_TP_STRIX_FRACTION` - deepseek4_loader.cpp
+- `DFLASH_DS4_DRAFT` - deepseek4_backend.cpp
 - `DFLASH_DS4_DRAFT_GPU` - deepseek4_backend.cpp
+- `DFLASH_DS4_DSPARK_DEBUG` - deepseek4_graph.cpp
+- `DFLASH_DS4_FUSED_VERIFY` - deepseek4_dspark_spec.cpp, deepseek4_loader.cpp
 - `DFLASH_DS4_HOTNESS_CSV` - deepseek4_backend.cpp
 - `DFLASH_DS4_MOE_TP` - deepseek4_backend.cpp
 - `DFLASH_DS4_MOE_TP_GPU` - deepseek4_backend.cpp
 - `DFLASH_DS4_MOE_TP_INPROC` - deepseek4_backend.cpp
+- `DFLASH_DS4_ROUTING_STATS_OUT` - deepseek4_backend.cpp
+- `DFLASH_DS4_SEQ_VERIFY` - deepseek4_dspark_spec.cpp
+- `DFLASH_DS4_SPEC` - deepseek4_backend.cpp
+- `DFLASH_DS4_SPEC_Q` - deepseek4_dspark_spec.cpp
 - `DFLASH_DS4_TIMING` - deepseek4_backend.cpp, deepseek4_target_shard_ipc_daemon.cpp
+- `DFLASH_DS4_TP_CAPTURE_CACHE_SLOTS` - deepseek4_fused_verify.inc
+- `DFLASH_DS4_TP_FUSED_CACHE_SLOTS` - deepseek4_fused_verify.inc
+- `DFLASH_DS4_TOPK` - deepseek4_graph.cpp
 - `DFLASH_EXPERT_BUDGET_MB` - deepseek4_backend.cpp, laguna_backend.cpp, qwen35moe_backend.cpp
 - `DFLASH_EXPERT_BUDGET_PCT` - laguna_backend.cpp
+- `DFLASH_FAST_ROLLBACK_THRESHOLD` - chain_rollback_policy.h
 - `DFLASH_FEATURE_DTYPE` - dflash_feature_ring.cpp
 - `DFLASH_FP_ALPHA` - http_server.cpp, qwen3_graph.cpp, server_main.cpp
 - `DFLASH_FP_CHUNK_S` - qwen3_graph.cpp
@@ -133,22 +156,34 @@ consolidation of this list into CLI flags is tracked as follow-up work.
 - `DFLASH_MMQ_SUB_BATCH` - moe_hybrid_ffn_eval.cpp
 - `DFLASH_MODEL_CARDS_DIR` - model_card.cpp
 - `DFLASH_MOE_COLD_BACKEND` - deepseek4_loader.cpp
+- `DFLASH_MOE_COMPACT_MATERIALIZED` - moe_hybrid_ffn_eval.cpp
+- `DFLASH_MOE_DUPLICATE_HOT_ON_COLD` - moe_hybrid_storage.cpp
+- `DFLASH_MOE_EXPERT_COMPUTE_DAEMON_TOKEN_LOOP` - moe_expert_compute_ipc.cpp
 - `DFLASH_MOE_EXPERT_COMPUTE_IPC_BATCH_CAPACITY` - moe_expert_compute_ipc.cpp
 - `DFLASH_MOE_EXPERT_COMPUTE_IPC_DTYPE` - moe_expert_compute_ipc.cpp
+- `DFLASH_MOE_EXPERT_COMPUTE_IPC_GPU` - deepseek4_backend.cpp
 - `DFLASH_MOE_EXPERT_COMPUTE_IPC_MODE` - moe_hybrid_ffn_eval.cpp
 - `DFLASH_MOE_EXPERT_COMPUTE_IPC_PROFILE` - moe_expert_compute_ipc.cpp
 - `DFLASH_MOE_EXPERT_COMPUTE_IPC_SHARED_BYTES` - moe_expert_compute_ipc.cpp
 - `DFLASH_MOE_EXPERT_COMPUTE_IPC_TRANSPORT` - moe_expert_compute_ipc.cpp
 - `DFLASH_MOE_EXPERT_COMPUTE_THREADS` - moe_expert_compute_cpu.cpp
+- `DFLASH_MOE_EXPERT_MAJOR_GPU_REDUCE` - moe_hybrid_ffn_eval.cpp
+- `DFLASH_MOE_EXPERT_MAJOR_PREFILL` - moe_hybrid_ffn_eval.cpp
 - `DFLASH_MOE_FIXED_SLOT_GRAPHS` - moe_hybrid_ffn_eval.cpp
 - `DFLASH_MOE_FIXED_SLOT_MAX` - moe_hybrid_ffn_eval.cpp
+- `DFLASH_MOE_FULL_COLD_PARALLEL` - moe_hybrid_ffn_eval.cpp
+- `DFLASH_MOE_FUSED_COMBINE` - moe_hybrid_ffn_eval.cpp
+- `DFLASH_MOE_PREFILL_DEVICE_INPUT` - deepseek4_graph.cpp
 - `DFLASH_MOE_PREFILL_HOT_SUB_BATCH` - moe_hybrid_ffn_eval.cpp
+- `DFLASH_MOE_PREFILL_MASKED_COLD` - moe_hybrid_ffn_eval.cpp
 - `DFLASH_MOE_PREFILL_PERSISTENT_OWNER_ALLOC` - deepseek4_graph.cpp
 - `DFLASH_NO_MASK` - laguna_backend.cpp
 - `DFLASH_NO_MOE_ROUTER_FUSE` - qwen35moe_ffn.cpp
 - `DFLASH_NO_MOE_SWIGLU_FUSE` - qwen35moe_ffn.cpp
 - `DFLASH_NO_PREAD` - deepseek4_loader.cpp
 - `DFLASH_PROF` - prof_env.h
+- `DFLASH_PREFILL_CACHE_SLOTS` - scripts/entrypoint.sh (maps to `--prefill-cache-slots`)
+- `DFLASH_PREFIX_CACHE_SLOTS` - scripts/entrypoint.sh (maps to `--prefix-cache-slots`)
 - `DFLASH_QWEN35MOE_CACHE_SLOTS` - qwen35moe_backend.cpp
 - `DFLASH_QWEN35MOE_HOTNESS` - qwen35moe_backend.cpp
 - `DFLASH_QWEN35MOE_NEXT_PLACEMENT_OUT` - qwen35moe_backend.cpp
@@ -161,8 +196,14 @@ consolidation of this list into CLI flags is tracked as follow-up work.
 - `DFLASH_QWEN35_NO_KVPAD` - graph_builders.cpp
 - `DFLASH_SAMPLED_VERIFY` - laguna_backend.cpp, qwen35_backend.cpp
 - `DFLASH_SHARE_DIR` - http_server.cpp
+- `DFLASH_SINGLE_CHAIN_CHECKPOINT_F32` - chain_rollback_policy.h
+- `DFLASH_SINGLE_CHAIN_ROLLBACK_DIAG` - chain_rollback_policy.h
 - `DFLASH_SPARK` - laguna_backend.cpp, qwen35moe_backend.cpp
 - `DFLASH_SPARK_VRAM_MB` - laguna_backend.cpp, qwen35moe_backend.cpp
+- `DFLASH_SPLIT_CAPTURE_SELFTEST` - qwen35_layer_split_dflash_target.cpp
+- `DFLASH_SPLIT_CHAIN_ROLLBACK_DIAG` - qwen35_layer_split_dflash_target.cpp, qwen35_target_graph.cpp
+- `DFLASH_SPLIT_FAST_ROLLBACK` - chain_rollback_policy.h
+- `DFLASH_STALL_TOOL_PREFIX` - http_server.cpp
 - `DFLASH_SV_DEBUG` - qwen35_backend.cpp
 - `DFLASH_TARGET_SHARD_IPC_SHARED_BYTES` - target_shard_ipc.cpp
 - `DFLASH_TARGET_SHARD_IPC_TRANSPORT` - target_shard_ipc.cpp
@@ -171,6 +212,8 @@ consolidation of this list into CLI flags is tracked as follow-up work.
 - `DFLASH_VERIFY_WIDTH` - qwen35moe_backend.cpp
 - `FAST_ROLLBACK_DIAG` - qwen35_dflash_target.cpp
 - `HOME` - spark_corpus.cpp
+- `LUCE_CUDA_I32_REPEAT` - moe_hybrid_ffn_eval.cpp
+- `LUCE_MMVQ_MAX_NCOLS` - deepseek4_backend.cpp
 - `LUCE_QK_FUSE_LAYERS` - laguna_target_graph.cpp
 - `LUCE_QK_FUSE_MODE` - laguna_target_graph.cpp
 - `PFLASH_DRAFTER_EARLY_EXIT_N` - qwen3_graph.cpp

server/docs/TOOL_PREFIX_CACHE.md

@@ -40,8 +40,9 @@ never restores past the stable prefix.
 
 No tool-specific flag is required. The native server default enables the
 in-memory prefix cache with 32 slots. Direct container launches inherit that
-default; set `DFLASH_PREFIX_CACHE_SLOTS=N` to override it or explicitly set it
-to `0` to disable prefix reuse.
+default. Pass `--prefix-cache-slots N` to the native binary, or set
+`DFLASH_PREFIX_CACHE_SLOTS=N` through `server/scripts/entrypoint.sh`; use `0`
+to disable prefix reuse.
 
 ## Reproducible benchmark
 

server/hip_compat/cuda_runtime.h

@@ -29,6 +29,9 @@ using cudaDeviceProp        = hipDeviceProp_t;
 // Error codes
 #define cudaSuccess                 hipSuccess
 #define cudaErrorInvalidValue       hipErrorInvalidValue
+#define cudaErrorIllegalAddress     hipErrorIllegalAddress
+#define cudaErrorAssert             hipErrorAssert
+#define cudaErrorLaunchFailure      hipErrorLaunchFailure
 
 // Memory functions
 #define cudaMalloc                  hipMalloc
@@ -49,6 +52,7 @@ using cudaDeviceProp        = hipDeviceProp_t;
 #define cudaStreamSynchronize       hipStreamSynchronize
 #define cudaStreamDefault           hipStreamDefault
 #define cudaStreamNonBlocking       hipStreamNonBlocking
+#define cudaStreamPerThread         hipStreamPerThread
 
 // Device functions
 #define cudaGetDevice               hipGetDevice
@@ -73,6 +77,8 @@ using cudaDeviceProp        = hipDeviceProp_t;
 // Error checking
 #define cudaGetLastError            hipGetLastError
 #define cudaGetErrorString          hipGetErrorString
+#define cudaGetErrorName            hipGetErrorName
+#define cudaPeekAtLastError         hipPeekAtLastError
 
 // Launch bounds
 #define __launch_bounds__           __launch_bounds__

server/src/common/chain_rollback_policy.h

@@ -18,6 +18,13 @@ inline bool env_flag_enabled(const char * name) {
     return value != nullptr && value[0] != '\0' && std::strcmp(value, "0") != 0;
 }
 
+// Layer-split fast rollback needs exact F32 per-token checkpoints. Keep the
+// feature, its capture work, and its additional memory behind one shared
+// opt-in so allocation and runtime dispatch cannot drift apart.
+inline bool split_chain_fast_rollback_enabled() {
+    return env_flag_enabled("DFLASH_SPLIT_FAST_ROLLBACK");
+}
+
 inline ChainRollbackPolicy resolve_chain_rollback_policy() {
     ChainRollbackPolicy policy;
     policy.checkpoint_f32 = env_flag_enabled("DFLASH_SINGLE_CHAIN_CHECKPOINT_F32");

server/src/common/gpu_runtime_compat.h

@@ -14,11 +14,16 @@
 #define cudaDeviceSynchronize hipDeviceSynchronize
 #define cudaErrorPeerAccessAlreadyEnabled hipErrorPeerAccessAlreadyEnabled
 #define cudaErrorPeerAccessNotEnabled hipErrorPeerAccessNotEnabled
+#define cudaErrorIllegalAddress hipErrorIllegalAddress
+#define cudaErrorAssert hipErrorAssert
+#define cudaErrorLaunchFailure hipErrorLaunchFailure
 #define cudaError_t hipError_t
 #define cudaFree hipFree
 #define cudaGetDeviceCount hipGetDeviceCount
 #define cudaGetErrorString hipGetErrorString
+#define cudaGetErrorName hipGetErrorName
 #define cudaGetLastError hipGetLastError
+#define cudaPeekAtLastError hipPeekAtLastError
 #define cudaMalloc hipMalloc
 #define cudaMemcpy2D hipMemcpy2D
 #define cudaMemcpy2DAsync hipMemcpy2DAsync
@@ -33,6 +38,7 @@
 #define cudaSetDevice hipSetDevice
 #define cudaStreamSynchronize hipStreamSynchronize
 #define cudaStream_t hipStream_t
+#define cudaStreamPerThread hipStreamPerThread
 #define cudaSuccess hipSuccess
 #define cudaDeviceProp hipDeviceProp_t
 #define cudaDeviceReset hipDeviceReset

server/src/internal.h

@@ -392,7 +392,9 @@ struct TargetCache {
     // persistent cache memory (not tracked by the per-call gallocr), matching
     // SGLang's mamba_caches.intermediate_ssm / intermediate_conv_window pattern.
     //
-    //   ssm_intermediate: [S_v, S_v, H_v, max_q_len] f32, one per delta layer.
+    //   ssm_intermediate: [S_v, S_v, H_v, max_q_len], checkpoint dtype
+    //     (Q8_0 for direct caches, F16 for migrated single-target caches, or
+    //     F32 for opt-in exact rollback), one per delta layer.
     //     Element t on axis 3 holds the DeltaNet recurrent state after
     //     processing verify token t. Spec decode commits t = commit_n - 1.
     //   conv_input_cache: [(kernel-1) + max_q_len, conv_channels] f32, one per
@@ -420,9 +422,9 @@ struct TargetCache {
 };
 
 // Snapshot the current SSM+conv state into TargetCache::*_snap tensors.
-void snapshot_ssm_state(TargetCache & c);
+bool snapshot_ssm_state(TargetCache & c, ggml_backend_t backend);
 // Restore the SSM+conv state from the snapshot.
-void restore_ssm_state(TargetCache & c);
+bool restore_ssm_state(TargetCache & c, ggml_backend_t backend);
 // Allocate rollback snapshot tensors mirroring live ssm/conv state (MoE path).
 bool ensure_ssm_snapshot(TargetCache & c, ggml_backend_t backend);
 
@@ -536,6 +538,9 @@ bool create_target_cache(const TargetWeights & w,
                          bool prefill_only = false,
                          int ctx_alloc = 0);
 
+// `f32_ssm_intermediates` enables exact per-token checkpoints for the opt-in
+// layer-split fast rollback path. The default preserves the established Q8_0
+// allocation and avoids its ~1.65 GiB incremental memory cost.
 bool create_target_cache_partial(const TargetWeights & w,
                                  int max_ctx,
                                  int max_verify_tokens,
@@ -545,7 +550,8 @@ bool create_target_cache_partial(const TargetWeights & w,
                                  int layer_begin,
                                  int layer_end,
                                  bool allocate_target_feat,
-                                 int ctx_alloc = 0);
+                                 int ctx_alloc = 0,
+                                 bool f32_ssm_intermediates = false);
 
 void free_target_cache(TargetCache & c);
 

server/src/qwen35/layer_split_daemon_loop.cpp

@@ -6,6 +6,7 @@
 #include "layer_split_forward.h" // free_qwen35_layer_split_shards
 #include "dflash_feature_ring.h"
 #include "common/io_utils.h"
+#include "common/chain_rollback_policy.h"
 #include "common/sampler.h"
 #include "common/layer_split_utils.h"
 #include "common/gguf_inspect.h"
@@ -55,7 +56,11 @@ int run_layer_split_daemon(const LayerSplitDaemonConfig & cfg) {
                                          shard.backend, shard.cache,
                                          /*prefill_only=*/!cfg.run_dflash,
                                          shard.layer_begin, shard.layer_end,
-                                         /*allocate_target_feat=*/false)) {
+                                         /*allocate_target_feat=*/false,
+                                         /*ctx_alloc=*/0,
+                                         /*f32_ssm_intermediates=*/
+                                             cfg.run_dflash &&
+                                             split_chain_fast_rollback_enabled())) {
             std::fprintf(stderr, "target-split load/cache gpu=%d: %s\n",
                          shard.gpu, dflash27b_last_error());
             free_qwen35_layer_split_shards(shards);