decider: reuse request prefixes with conservative auto selection

Cargo.toml

@@ -1,3 +1,3 @@
 [workspace]
-members = ["src/frontend", "src/runtime", "src/models/clm", "src/models/cua_s1/native", "src/models/qwen3_5/native", "src/models/open_jev/native", "src/models/laya", "src/backends/cuda"]
+members = ["src/frontend", "src/runtime", "src/models/clm", "src/models/cua_s1/native", "src/models/qwen3_5/native", "src/models/open_jev/native", "src/models/laya", "src/backends/cuda", "src/models/decider/native"]
 resolver = "3"

recipe/decider/download_weights.py

@@ -0,0 +1,58 @@
+#!/usr/bin/env python3
+"""Download and SHA-256 verify the immutable Decider-2B v11 artifacts."""
+import argparse
+import hashlib
+import os
+from pathlib import Path
+import tempfile
+import urllib.request
+
+REVISION = "533964dae8be954c5b5e19fa4948e48408094c1e"
+FILES = {
+    "config.json": (1790, "6cb8daca9fb653c61485ff7452fc068bacd5c27cbee659ecd24b47186b0d1b52"),
+    "decider_config.json": (1240, "6e4891f2754a1c18a10f8dadb0c04e439e7f79fab0333d56641491bd4a05e722"),
+    "tokenizer.json": (19989325, "06b9509352d2af50381ab2247e083b80d32d5c0aba91c272ca9ff729b6a0e523"),
+    "tokenizer_config.json": (1127, "171ecbe7ddae98d11840698f7df2b8d5b4722139db0f0620d3bbf429bd656250"),
+    "model.safetensors": (3763692048, "acaef2228b134dcdc20cad4ee79219482c927ec819aa3687b9b8a575c338817f"),
+}
+
+
+def verify(path, expected):
+    size, digest = expected
+    if path.stat().st_size != size:
+        raise ValueError(f"{path}: size mismatch")
+    actual = hashlib.sha256()
+    with path.open("rb") as source:
+        for block in iter(lambda: source.read(1024 * 1024), b""):
+            actual.update(block)
+    if actual.hexdigest() != digest:
+        raise ValueError(f"{path}: checksum mismatch")
+
+
+def main():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("directory", type=Path)
+    parser.add_argument("--endpoint", default="https://huggingface.co")
+    args = parser.parse_args()
+    args.directory.mkdir(parents=True, exist_ok=True)
+    for name, expected in FILES.items():
+        destination = args.directory / name
+        if destination.exists():
+            verify(destination, expected)
+        else:
+            fd, temporary = tempfile.mkstemp(prefix=name + ".", suffix=".part", dir=args.directory)
+            temporary = Path(temporary)
+            try:
+                url = f"{args.endpoint.rstrip('/')}/Mapika/decider-2b/resolve/{REVISION}/{name}"
+                with os.fdopen(fd, "wb") as output, urllib.request.urlopen(url, timeout=120) as response:
+                    for block in iter(lambda: response.read(1024 * 1024), b""):
+                        output.write(block)
+                verify(temporary, expected)
+                temporary.replace(destination)
+            finally:
+                temporary.unlink(missing_ok=True)
+        print(f"verified {name}", flush=True)
+
+
+if __name__ == "__main__":
+    main()

src/backends/cuda/qwen3_5/attention.cu

@@ -5,7 +5,9 @@
 // four warps of 16 rows each, and walks the keys up to its last query in tiles of
 // 32, keeping the output and the online softmax in registers. The probabilities are
 // rounded to bfloat16 for the P*V product, as in flash attention; the running sums
-// stay float32.
+// stay float32. The queries can be the last Tq of Tk positions (cached keys before
+// them); key tiles always start at position 0, so a query sees the same tiles in the
+// same order wherever the queries start, and tiles past its own position change nothing.
 #include "common.cuh"
 #include "mma.cuh"
 #include "ops.h"
@@ -81,7 +83,7 @@ constexpr int SMEM_BYTES = (BM + 2 * BN) * LDS * 2;
 template <bool Gated>
 __global__ void __launch_bounds__(THREADS)
     flash_kernel(const bf16* __restrict__ q, const bf16* __restrict__ k, const bf16* __restrict__ v, int ldv,
-                 const bf16* __restrict__ gate, bf16* __restrict__ out, int T, int Hq, int Hk,
+                 const bf16* __restrict__ gate, bf16* __restrict__ out, int Tq, int Tk, int Hq, int Hk,
                  float scale_log2) {
     extern __shared__ __align__(16) unsigned char smem[];
     bf16* qs = reinterpret_cast<bf16*>(smem);
@@ -92,10 +94,11 @@ __global__ void __launch_bounds__(THREADS)
     const int tid = threadIdx.x, warp = tid / 32, lane = tid % 32;
     const int g = lane / 4, t = lane % 4;
     const int row0 = q0 + warp * 16;  // this warp's first query
+    const int offset = Tk - Tq;       // the position of query 0
 
     for (int c = tid; c < BM * (D / 8); c += THREADS) {
         const int r = c / (D / 8), col = (c % (D / 8)) * 8, row = q0 + r;
-        cp_async16(qs + r * LDS + col, q + ((size_t)min(row, T - 1) * Hq + h) * D + col, row < T);
+        cp_async16(qs + r * LDS + col, q + ((size_t)min(row, Tq - 1) * Hq + h) * D + col, row < Tq);
     }
     cp_async_commit();
 
@@ -104,23 +107,24 @@ __global__ void __launch_bounds__(THREADS)
     for (int n = 0; n < D / 8; n++) o[n][0] = o[n][1] = o[n][2] = o[n][3] = 0.f;
     float m[2] = {-INFINITY, -INFINITY}, l[2] = {0.f, 0.f};
 
-    const int kv_end = min(T, q0 + BM);
+    const int kv_end = min(Tk, offset + q0 + BM);
     for (int k0 = 0; k0 < kv_end; k0 += BN) {
         for (int c = tid; c < BN * (D / 8); c += THREADS) {
             const int r = c / (D / 8), col = (c % (D / 8)) * 8, s = k0 + r;
-            cp_async16(ks + r * LDS + col, k + ((size_t)min(s, T - 1) * Hk + hk) * D + col, s < T);
+            cp_async16(ks + r * LDS + col, k + ((size_t)min(s, Tk - 1) * Hk + hk) * D + col, s < Tk);
         }
         cp_async_commit();
         for (int c = tid; c < BN * (D / 8); c += THREADS) {
             const int r = c / (D / 8), col = (c % (D / 8)) * 8, s = k0 + r;
-            cp_async16(vs + r * LDS + col, v + (size_t)min(s, T - 1) * ldv + (size_t)hk * D + col, s < T);
+            cp_async16(vs + r * LDS + col, v + (size_t)min(s, Tk - 1) * ldv + (size_t)hk * D + col, s < Tk);
         }
         cp_async_commit();
         cp_async_wait<1>();  // Q and K
         __syncthreads();
 
-        // keys past every query of this warp contribute nothing
-        const bool active = k0 <= row0 + 15;
+        // keys past every query of this warp contribute nothing, and rows past the
+        // last query are never stored
+        const bool active = row0 < Tq && k0 <= offset + row0 + 15;
         float sc[BN / 8][4];
 #pragma unroll
         for (int n = 0; n < BN / 8; n++) sc[n][0] = sc[n][1] = sc[n][2] = sc[n][3] = 0.f;
@@ -146,8 +150,8 @@ __global__ void __launch_bounds__(THREADS)
             for (int n = 0; n < BN / 8; n++) {
 #pragma unroll
                 for (int e = 0; e < 4; e++) {
-                    const int key = k0 + n * 8 + 2 * t + (e & 1), row = row0 + g + (e >> 1) * 8;
-                    sc[n][e] = (key <= row && key < T) ? sc[n][e] * scale_log2 : -INFINITY;
+                    const int key = k0 + n * 8 + 2 * t + (e & 1), pos = offset + row0 + g + (e >> 1) * 8;
+                    sc[n][e] = (key <= pos && key < Tk) ? sc[n][e] * scale_log2 : -INFINITY;
                     mx[e >> 1] = fmaxf(mx[e >> 1], sc[n][e]);
                 }
             }
@@ -213,7 +217,7 @@ __global__ void __launch_bounds__(THREADS)
 #pragma unroll
     for (int r = 0; r < 2; r++) {
         const int row = row0 + g + r * 8;
-        if (row >= T) continue;
+        if (row >= Tq) continue;
         bf16* dst = out + ((size_t)row * Hq + h) * D + 2 * t;
 #pragma unroll
         for (int n = 0; n < D / 8; n++) {
@@ -232,20 +236,20 @@ __global__ void __launch_bounds__(THREADS)
 
 template <bool Gated>
 int launch(const void* q, const void* k, const void* v, int ldv, const void* gate, void* out,
-           int T, int Hq, int Hk, int Dh, float scale, void* stream) {
-    if (Dh != D || Hk <= 0 || Hq <= 0 || Hq % Hk != 0 || ldv % 8 != 0 || ldv < Hk * Dh || T < 0)
+           int Tq, int Tk, int Hq, int Hk, int Dh, float scale, void* stream) {
+    if (Dh != D || Hk <= 0 || Hq <= 0 || Hq % Hk != 0 || ldv % 8 != 0 || ldv < Hk * Dh || Tq < 0 || Tk < Tq)
         return cudaErrorInvalidValue;
-    if (T == 0) return cudaSuccess;
+    if (Tq == 0) return cudaSuccess;
     if (Gated && gate == nullptr) return cudaErrorInvalidValue;
     // Once per specialization (for the device current at the first call).
     static const cudaError_t configured = cudaFuncSetAttribute(
         flash_kernel<Gated>, cudaFuncAttributeMaxDynamicSharedMemorySize, SMEM_BYTES);
     if (configured != cudaSuccess) return configured;
     constexpr float LOG2E = 1.4426950408889634f;
-    flash_kernel<Gated><<<dim3((T + BM - 1) / BM, Hq), THREADS, SMEM_BYTES,
+    flash_kernel<Gated><<<dim3((Tq + BM - 1) / BM, Hq), THREADS, SMEM_BYTES,
                          static_cast<cudaStream_t>(stream)>>>(
         static_cast<const bf16*>(q), static_cast<const bf16*>(k), static_cast<const bf16*>(v), ldv,
-        static_cast<const bf16*>(gate), static_cast<bf16*>(out), T, Hq, Hk, scale * LOG2E);
+        static_cast<const bf16*>(gate), static_cast<bf16*>(out), Tq, Tk, Hq, Hk, scale * LOG2E);
     return cudaGetLastError();
 }
 
@@ -274,10 +278,16 @@ extern "C" int cs1_attn_prep(const void* qg, const void* kr, int ld, const void*
 
 extern "C" int cs1_attention(const void* q, const void* k, const void* v, int ldv, void* out, int T, int Hq, int Hk,
                              int Dh, float scale, void* stream) {
-    return flash::launch<false>(q, k, v, ldv, nullptr, out, T, Hq, Hk, Dh, scale, stream);
+    return flash::launch<false>(q, k, v, ldv, nullptr, out, T, T, Hq, Hk, Dh, scale, stream);
 }
 
 extern "C" int cs1_attention_gated(const void* q, const void* k, const void* v, int ldv, const void* gate,
                                    void* out, int T, int Hq, int Hk, int Dh, float scale, void* stream) {
-    return flash::launch<true>(q, k, v, ldv, gate, out, T, Hq, Hk, Dh, scale, stream);
+    return flash::launch<true>(q, k, v, ldv, gate, out, T, T, Hq, Hk, Dh, scale, stream);
+}
+
+extern "C" int cs1_attention_gated_cached(const void* q, const void* k, const void* v, int ldv, const void* gate,
+                                          void* out, int Tq, int Tk, int Hq, int Hk, int Dh, float scale,
+                                          void* stream) {
+    return flash::launch<true>(q, k, v, ldv, gate, out, Tq, Tk, Hq, Hk, Dh, scale, stream);
 }

src/backends/cuda/qwen3_5/elementwise.cu

@@ -16,10 +16,13 @@ __global__ void embed_kernel(const int32_t* __restrict__ ids, const Pack8* __res
 }
 
 // F.conv1d in bfloat16 (float32 accumulation, rounded), then SiLU (rounded again),
-// written to three contiguous outputs.
+// written to three contiguous outputs. Positions before the first row come from
+// history [3, channels] when there is one; without it they are skipped, as the zero
+// padding of a sequence's start adds nothing.
 __global__ void gdn_conv_kernel(const bf16* __restrict__ qkv, int ld, const bf16* __restrict__ w,
-                                bf16* __restrict__ q, bf16* __restrict__ k, bf16* __restrict__ v, int T,
-                                int key_dim, int value_dim) {
+                                const bf16* __restrict__ history, bf16* __restrict__ q,
+                                bf16* __restrict__ k, bf16* __restrict__ v, int T, int key_dim,
+                                int value_dim) {
     const int channels = 2 * key_dim + value_dim;
     const size_t idx = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
     if (idx >= (size_t)T * channels) return;
@@ -28,7 +31,10 @@ __global__ void gdn_conv_kernel(const bf16* __restrict__ qkv, int ld, const bf16
 #pragma unroll
     for (int j = 0; j < 4; j++) {
         const int s = t - 3 + j;
-        if (s >= 0) acc = fmaf(f32(w[c * 4 + j]), f32(qkv[(size_t)s * ld + c]), acc);
+        if (s >= 0)
+            acc = fmaf(f32(w[c * 4 + j]), f32(qkv[(size_t)s * ld + c]), acc);
+        else if (history)
+            acc = fmaf(f32(w[c * 4 + j]), f32(history[(size_t)(3 + s) * channels + c]), acc);
     }
     const bf16 y = to_bf16(silu(round_bf16(acc)));
     if (c < key_dim)
@@ -39,6 +45,16 @@ __global__ void gdn_conv_kernel(const bf16* __restrict__ qkv, int ld, const bf16
         v[(size_t)t * value_dim + c - 2 * key_dim] = y;
 }
 
+// The conv inputs of the last three positions, from qkv or, before its first row, from
+// history (zeros without one).
+__global__ void gdn_conv_history_kernel(const bf16* __restrict__ qkv, int ld, const bf16* __restrict__ history,
+                                        bf16* __restrict__ out, int T, int channels) {
+    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= 3 * channels) return;
+    const int r = idx / channels, c = idx % channels, s = T - 3 + r;
+    out[idx] = s >= 0 ? qkv[(size_t)s * ld + c] : history ? history[(3 + s) * channels + c] : to_bf16(0.f);
+}
+
 // beta = sigmoid(b) in bfloat16; g = -exp(A_log) * softplus(a + dt_bias) in float32
 // (F.softplus with threshold 20).
 __global__ void gdn_gates_kernel(const bf16* __restrict__ b, const bf16* __restrict__ a, int ld,
@@ -101,17 +117,30 @@ extern "C" int cs1_embed(const int32_t* ids, const void* table, void* out, int T
     return cudaGetLastError();
 }
 
-extern "C" int cs1_gdn_conv(const void* qkv, int ld, const void* w, void* q, void* k, void* v, int T, int key_dim,
-                            int value_dim, void* stream) {
+extern "C" int cs1_gdn_conv_history(const void* qkv, int ld, const void* w, const void* history, void* history_out,
+                                    void* q, void* k, void* v, int T, int key_dim, int value_dim, void* stream) {
     if (T < 0 || key_dim < 0 || value_dim < 0 || ld < 2 * key_dim + value_dim) return cudaErrorInvalidValue;
-    const size_t n = (size_t)T * (2 * key_dim + value_dim);
-    if (n == 0) return cudaSuccess;
-    gdn_conv_kernel<<<blocks(n), THREADS, 0, static_cast<cudaStream_t>(stream)>>>(
-        static_cast<const bf16*>(qkv), ld, static_cast<const bf16*>(w), static_cast<bf16*>(q),
-        static_cast<bf16*>(k), static_cast<bf16*>(v), T, key_dim, value_dim);
+    const int channels = 2 * key_dim + value_dim;
+    const size_t n = (size_t)T * channels;
+    const cudaStream_t st = static_cast<cudaStream_t>(stream);
+    if (n > 0) {
+        gdn_conv_kernel<<<blocks(n), THREADS, 0, st>>>(
+            static_cast<const bf16*>(qkv), ld, static_cast<const bf16*>(w), static_cast<const bf16*>(history),
+            static_cast<bf16*>(q), static_cast<bf16*>(k), static_cast<bf16*>(v), T, key_dim, value_dim);
+    }
+    if (history_out && channels > 0) {
+        gdn_conv_history_kernel<<<blocks((size_t)3 * channels), THREADS, 0, st>>>(
+            static_cast<const bf16*>(qkv), ld, static_cast<const bf16*>(history), static_cast<bf16*>(history_out),
+            T, channels);
+    }
     return cudaGetLastError();
 }
 
+extern "C" int cs1_gdn_conv(const void* qkv, int ld, const void* w, void* q, void* k, void* v, int T, int key_dim,
+                            int value_dim, void* stream) {
+    return cs1_gdn_conv_history(qkv, ld, w, nullptr, nullptr, q, k, v, T, key_dim, value_dim, stream);
+}
+
 extern "C" int cs1_gdn_gates(const void* b, const void* a, int ld, const void* A_log, const void* dt_bias,
                              void* beta, float* g, int T, int H, void* stream) {
     if (T < 0 || H < 0 || ld < H) return cudaErrorInvalidValue;