feat(qwen35): add transactional speculative verifier

docs/index.md

@@ -51,6 +51,7 @@ Organized by domain (model line / subsystem / playbook / lesson) instead of by l
 | `models/qwen35/kv-admission.md` | Issue #254 complete: Qwen3.5 now uses full-lifetime KV admission, deferred pressure handling, impossible-request rejection, explicit error semantics, direct rejection-event coverage, RTX 5090 e2e, and real HTTP pressure/post-pressure validation. |
 | `models/qwen35/optimization.md` | Hybrid 24 linear + 8 full attn optimization ledger. Decode-tuning refresh fuses MLP gate/up and tunes decode cublasLt buckets, improving direct TPOT by 2-3%; vLLM still leads 1024/256 HTTP decode. |
 | `models/qwen35/accuracy.md` | Qwen3.5 HF bf16 logits goldens, size-keyed (0.8b/2b/4b/9b/27b all committed), through `past_key_values`: short replay covers sequential graph, bucket-straddling batched graph, and slot-compaction; long replay covers 4097/8192-token prompts; full GSM8K 8-shot now matches the HF baseline within 0.15 percentage points. |
+| `models/qwen35/speculative-verifier.md` | Extracts the target-only Qwen3.5 speculative verifier from PR #626 with transactional hybrid-state handling; DFlash drafter and serving integration stay out of this slice. |
 | `models/qwen35/model-crate.md` | `pegainfer-qwen35` owns Qwen3.5 model/scheduler/recurrent ops/tests/benches; feature-gated behind `qwen35` (Triton AOT is the only Python build dependency); root loads it through `EngineHandle`. Build/check/clippy, root bench sanity check, historical Qwen3.5 e2e, and scheduler e2e records live here. |
 | `models/qwen35/batched-step-tail.md` | Qwen3.5 issue #353 implementation record: final prefill tail is batched, decode/unified sample from batched logits, host full-vocab copies are logprobs-only, HF + scheduler e2e pass, and final serving A/B supports only the first-token/short-output TTFT claim. |
 | `models/qwen35/tp-design.md` | Qwen3.5 TP design: Phase 1 is eager dense TP on Qwen3's controller/worker runtime; validate TP2 first, fail closed for indivisible degrees and TP+CUDA Graph, shard dense full-attention/MLP, and leave sharded linear/GDR state to follow-up. |

docs/models/qwen35/speculative-verifier.md

@@ -0,0 +1,34 @@
+# Qwen3.5 Speculative Verifier
+
+> **TL;DR:** PR #667 extracts the target-only verifier from #626; for the C16, prompt-1024, span-5 test, reusing existing Q/K prep and paged K/V scatter cut aggregate prep-kernel GPU time by 21.1%, while drafter and serving wiring remain follow-ups.
+>
+> **Last touched:** 2026-07
+
+## Contract
+
+- Target verification only; no draft model, scheduler/server wiring, sampling, or serving-performance claim.
+- Verifier reuses the existing batched Q/K prep and paged K/V scatter; normal single-request prefill keeps its fused path.
+- Each request supplies a non-empty `[current token, draft tokens...]` span. A one-token span is valid when one output token remains.
+- Greedy acceptance commits the matching draft prefix plus one target token.
+- Full acceptance keeps verified KV and recurrent state. Partial acceptance truncates KV, restores recurrent/convolution state, and replays only the accepted span.
+- Backup, verify, commit, and rollback use the context stream. Stream overrides are rejected before mutation.
+- Any error after mutation restores every canonical state component; rollback failure is executor-fatal.
+- Verifier logits and sampling use `selection_vocab`, so checkpoint padding rows cannot produce token ids the frontend cannot decode.
+
+## Verified
+
+- Passed: Qwen3.5 release check and Clippy; RTX 5090 verifier tests 11/11. Earlier gates also passed HF golden 2/2, scheduler E2E 1/1, page-pool 4/4, and KV-pool 6/6.
+
+| C16, prompt 1024, span 5 | Aggregate prep-kernel GPU time, 3 runs | Median | Whole-test median, 5 runs |
+| --- | --- | --- | --- |
+| Fused verifier kernels (`698ccbd`) | 52.256 / 52.448 / 52.351 us | 52.351 us | 7.1696 s |
+| Shared Q/K prep + paged K/V scatter | 41.312 / 41.312 / 41.600 us | 41.312 us (-21.1%) | 7.1202 s (-0.7%) |
+
+The common attention kernel stayed within 0.2%, and both profiles launched the same number of kernels. Keep the shared path: the prep-kernel reduction was consistent, while five whole-test runs did not establish a meaningful improvement. This result is specific to this RTX 5090 verifier test and does not establish serving performance.
+
+- Claim boundary: sampling, calibrated logprob parity, and serving integration remain unverified.
+
+## Next
+
+- Add the DFlash drafter and its independent forward oracle.
+- Before serving integration, move verifier scratch and recurrent backups to an executor-owned persistent workspace; then wire opt-in fallback/admission rules and collect same-host benchmark evidence.

pegainfer-core/src/kv_pool.rs

@@ -212,6 +212,20 @@ impl KvState {
         self.seq_len += count;
     }
 
+    /// Roll this request's logical KV length back to `token_count`, returning
+    /// any now-unused tail pages to the pool.
+    pub fn truncate_to(&mut self, token_count: usize) -> Result<()> {
+        anyhow::ensure!(
+            token_count <= self.seq_len,
+            "KvState cannot truncate from {} up to {token_count}",
+            self.seq_len
+        );
+        let needed = pages_needed(token_count, self.pool.inner.layout.page_size);
+        self.permit.truncate(needed);
+        self.seq_len = token_count;
+        Ok(())
+    }
+
     /// Build kernel-facing metadata for this request's KV.
     pub fn desc(&self) -> KvDesc<'_> {
         KvDesc {
@@ -329,12 +343,39 @@ mod tests {
         assert_eq!(desc.last_page_len(), 1);
         assert_eq!(pool.available_pages(), 2);
 
+        // Truncate back into the first page: tail page returns immediately.
+        kv.truncate_to(15).unwrap();
+        assert_eq!(kv.seq_len(), 15);
+        let desc = kv.desc();
+        assert_eq!(desc.num_pages(), 1);
+        assert_eq!(desc.last_page_len(), 15);
+        assert_eq!(pool.available_pages(), 3);
+
+        // Truncate to zero releases all request pages.
+        kv.truncate_to(0).unwrap();
+        assert_eq!(kv.seq_len(), 0);
+        assert_eq!(kv.desc().num_pages(), 0);
+        assert_eq!(pool.available_pages(), 4);
+
         // Reset returns all pages
+        kv.ensure_capacity(17).unwrap();
+        kv.advance(17);
         kv.reset();
         assert_eq!(kv.seq_len(), 0);
         assert_eq!(pool.available_pages(), 4);
     }
 
+    #[test]
+    fn kv_state_rejects_truncate_forward() {
+        let pool = test_pool(16, 3);
+        let mut kv = pool.alloc();
+        kv.ensure_capacity(4).unwrap();
+        kv.advance(4);
+
+        let err = kv.truncate_to(5).unwrap_err().to_string();
+        assert!(err.contains("cannot truncate from 4 up to 5"));
+    }
+
     #[test]
     fn kv_state_out_of_pages() {
         // 3 pages total: 1 padding, 2 available → 32 tokens max

pegainfer-core/src/ops/attention.rs

@@ -196,10 +196,8 @@ pub fn paged_attention_batch_decode_via_prefill_hd256_into(
     layout: &KvLayout,
     layer: usize,
     plan: &PrefillPagedPlan,
-    positions_d: &CudaSlice<i32>,
     output: &mut HiddenStates,
     num_qo_heads: usize,
-    batch_size: usize,
 ) -> Result<()> {
     pegainfer_kernels::ops::paged_attention_batch_decode_via_prefill_hd256_into(
         ctx,
@@ -210,9 +208,7 @@ pub fn paged_attention_batch_decode_via_prefill_hd256_into(
         &layout.kernel_layout(),
         layer,
         plan,
-        positions_d,
         output,
         num_qo_heads,
-        batch_size,
     )
 }

pegainfer-core/src/page_pool.rs

@@ -105,6 +105,24 @@ impl OwnedPagePermit {
         }
         true
     }
+
+    /// Return tail pages until the permit holds exactly `new_len` pages.
+    ///
+    /// Prefix page order is preserved. Pages beyond `new_len` are returned to
+    /// the same pool immediately, matching the drop-time LIFO reuse order.
+    pub(crate) fn truncate(&mut self, new_len: usize) {
+        assert!(
+            new_len <= self.pages.len(),
+            "cannot grow an OwnedPagePermit via truncate"
+        );
+        if new_len == self.pages.len() {
+            return;
+        }
+
+        let returned = self.pages.split_off(new_len);
+        let mut free_list = self.inner.free_list.lock();
+        free_list.extend(returned.into_iter().rev());
+    }
 }
 
 impl Drop for OwnedPagePermit {
@@ -181,4 +199,27 @@ mod tests {
         // all 4 pages back after drop
         assert_eq!(pool.available_pages(), 4);
     }
+
+    #[test]
+    fn truncate_returns_tail_pages_and_preserves_prefix() {
+        let pool = PagePool::new(5);
+
+        {
+            let mut permit = pool.try_acquire_many(4).expect("initial acquire");
+            assert_eq!(
+                permit.pages(),
+                &[PageId(0), PageId(1), PageId(2), PageId(3)]
+            );
+            assert_eq!(pool.available_pages(), 1);
+
+            permit.truncate(2);
+            assert_eq!(permit.pages(), &[PageId(0), PageId(1)]);
+            assert_eq!(pool.available_pages(), 3);
+
+            let next = pool.try_acquire_many(2).expect("tail pages reusable");
+            assert_eq!(next.pages(), &[PageId(2), PageId(3)]);
+        }
+
+        assert_eq!(pool.available_pages(), 5);
+    }
 }

pegainfer-kernels/csrc/qwen35/prefill_attention_hd256.cu

@@ -291,7 +291,7 @@ __global__ void qk_norm_partial_rope_batched_decode_hd256_kernel(
 
 extern "C" {
 
-void qk_norm_partial_rope_batched_decode_hd256_cuda(
+int qk_norm_partial_rope_batched_decode_hd256_cuda(
     const __nv_bfloat16* q_full_batch,
     __nv_bfloat16* k_batch,
     const __nv_bfloat16* q_norm_weight,
@@ -323,6 +323,7 @@ void qk_norm_partial_rope_batched_decode_hd256_cuda(
         rotary_dim,
         rms_eps
     );
+    return static_cast<int>(cudaGetLastError());
 }
 
 void prefill_attention_hd256_prep_paged_cuda(

pegainfer-kernels/src/ffi/qwen35.rs

@@ -57,7 +57,7 @@ unsafe extern "C" {
         rotary_dim: i32,
         rms_eps: f32,
         stream: CUstream,
-    );
+    ) -> i32;
 
     // Gated delta rule recurrent decode (single step)
     pub fn gated_delta_rule_decode_cuda(

pegainfer-kernels/src/ops.rs

@@ -123,6 +123,7 @@ pub use linear::gemm_lt_pin_tune;
 pub use linear::gemm_lt_pin_warmup;
 pub use linear::gemm_lt_tune;
 pub use linear::gemm_per_token;
+pub use linear::gemm_per_token_into_checked;
 pub use linear::gemm_rows_into;
 pub use linear::gemm_rows_into_checked;
 pub use linear::gemm_strided_batched_bf16;
@@ -134,6 +135,7 @@ pub use linear::per_token_served;
 pub use linear::pin_served;
 pub use linear::reset_numeric_policy_counters;
 pub use linear::set_numeric_policy;
+pub use linear::with_gemm_lt_disabled;
 pub use lora::LoraDecodeGroupedProjection;
 pub use lora::lora_decode_fused_delta_group3_into;
 pub use lora::lora_decode_fused_delta_into;