cua-s1: reuse screenshot prefixes within a request

docs/architecture.md

@@ -68,7 +68,12 @@ FLA chunk and either FlashInfer dispatch. CuteDSL and CUDA graphs are
 not selected.
 
 Cua-S1 and Open-Jev compute their decision heads on the CPU after downloading
-the final hidden state. Decider projects its selected LM-head rows on CUDA.
+the final hidden state. Cua-S1 screenshots can opt into request-local
+aligned token/position prefix reuse with fixed-M64 or full-prompt-matched GEMMs,
+host features and language Graph off. Captured contents belong to one request; one engine-owned storage
+allocation may retain capacity for later captures. Suffixes execute independently,
+and the existing letter head and token accounting are preserved. Decider projects
+its selected LM-head rows on CUDA.
 JEMM uses the shared Qwen executor with separate unmerged
 FP32 adapters and framework-reference arithmetic. It computes selected BF16
 label-head logits on CUDA, then calibrates typed answers on the CPU. JEMM admits

recipe/cua_s1/native_screenshot.md

@@ -94,3 +94,148 @@ Recorded RTX 4090 evidence covers native screenshot parity and Graph regressions
 the resident-feature comparison covers host/device output equality on its fixed
 real-screenshot corpus. These scopes do not establish general GUI task accuracy
 or native equivalence to every reference precision and environment.
+
+## Request-local language prefix reuse
+
+`CUA_S1_MM_PREFIX=shared` opts into sharing the aligned language prefix of all
+questions in one screenshot request. Unset or `off` keeps the existing independent
+forwards. `fixed` runs complete prompts independently with the same fixed GEMM
+algorithms as `shared`, providing the numerical control. The `matched` mode
+below instead checks ordinary full-prompt algorithms. Other values fail before
+checkpoint or CUDA loading.
+
+```sh
+export CUA_S1_MM_PREFIX=shared
+export CUA_S1_GRAPH=0 CUA_S1_DEVICE_FEATURES=0
+# Start omni-cua-s1-vision using the pinned artifacts above.
+```
+
+The planner compares prepared token IDs and all three position axes, rounds the
+common prefix down to 64 tokens, and leaves a nonempty suffix for each question.
+It captures attention K/V, GDN recurrent state and convolution history once, then
+continues each suffix independently in request order. The engine retains one
+snapshot allocation, growing it when a larger shared prefix needs capacity and
+freeing it with the engine. Every shared request captures fresh contents before
+reading them; a repeated or changed screenshot is always recomputed. One question,
+or less than 64 shared tokens, uses complete fixed passes and never reads a
+previous snapshot. Retained capacity is additional device memory, separate from
+the vision-cache budget.
+
+All active prefix modes require host features and language Graph off; vision Graph
+remains independently selectable. They keep the existing prompt, CPU letter head,
+answer order and usage accounting. Fixed GEMMs can round differently from default
+execution: compare `shared` with `fixed` for continuation parity, and separately
+report probability drift/decision changes against `off`. No automatic mode or
+cross-request content cache is included. Fixed algorithms still use reference
+M=64; the experimental M=256 profile is not part of this implementation.
+
+On one RTX 4090 (SM8.9), a matched eight-payload synthetic HTTP comparison with
+vision cache 8 entries/1024 MiB measured shared M64 p50/p95 149.213/223.530 ms
+before allocation reuse and 138.404/214.262 ms after it. All response bytes
+matched; sampled process maxima were 10282/10288 MiB. Each side excluded an
+8-request growth pass, an 8-request all-shape stabilization pass and 5 warmups,
+then measured 3 rounds of 8 requests, concurrency 1. The clock includes full
+HTTP response parsing and schema validation. The ordinary off control was
+118.804/275.748 ms, so this result does not justify enabling shared mode by default.
+Source pins, raw samples, exclusions and reproduction inputs are in the
+[follow-up qualification](https://github.com/Levius-Fubuki/system1-omni/releases/download/cua-mm-qualification-20261025/cua-mm-followup-qualification-20261025.tar.gz).
+These historical measurements apply to the fixed/shared implementation at
+`d4908d95c`, before the matched mode below. Other GPUs and general GUI task
+success remain unverified.
+
+CPU control/planner tests run in the ordinary Cua suite. On a reserved GPU with
+the four artifact variables above, run the exact model and recovery checks:
+
+```sh
+cargo test --release --locked -p omni-cua-s1-native --lib prefix_tests \
+  -- --ignored --test-threads=1 --nocapture
+```
+
+The screenshot corpus is a numerical and execution regression; it does not
+establish general GUI task accuracy or a speedup for every request.
+
+### Full-prompt-matched algorithms
+
+`CUA_S1_MM_PREFIX=matched` shares the same request-local aligned prefix using
+algorithms selected at the first question's complete prompt length, with the
+ordinary BF16 reduction policy. Before any prefix execution, every layer's
+projection must have identical algorithm descriptors for every complete question;
+the selected algorithms must also support every prefix and suffix window within
+the workspace limit. A mismatch or unsupported window falls back to ordinary
+independent full passes. A single question or no aligned prefix takes that path
+directly. Allocation/setup and execution errors remain errors.
+
+```sh
+export CUA_S1_MM_PREFIX=matched
+export CUA_S1_GRAPH=0 CUA_S1_DEVICE_FEATURES=0
+# Rebuild the CUDA library with src/backends/cuda/qwen3_5/build.sh.
+```
+
+This mode requires the optional additive ABI 9 symbols `cs1_gemm_create_matched`
+and `cs1_gemm_compatible`. Both are checked during the single-question startup
+warmup, before the worker binds its health endpoint; an incomplete library fails
+with a rebuild error. Older libraries still serve the existing modes.
+The engine retains at most 16 matched GEMM workspaces, 32 MiB each. On a new
+reference length beyond that cap, it drains the language stream and releases the
+previous generation before allocating another. This is separate from retained
+prefix-state capacity and vision-cache limits. No screenshot contents survive
+as reusable results between requests. Errors and panics drain submitted work
+and restore the ordinary/fixed-M64 handle selection before returning/unwinding.
+
+Descriptor equality and algorithm support do not establish arithmetic invariance
+across row counts on every GPU. Qualify the exact GPU/cuBLAS/model combination
+before relying on ordinary-output parity. Keep `off` as the ordinary numerical
+control and `fixed`/`shared` as the separate M64 controls. No automatic policy,
+language Graph integration, changed prompt or new precision profile is enabled.
+
+```sh
+cargo test --release --locked -p omni-qwen3-5-native --test kernels \
+  matched_gemm_rows_equal_complete_ordinary_projection \
+  -- --ignored --test-threads=1 --nocapture
+# Also run the prefix_tests command above with vision Graph 0 and 1.
+```
+
+At runtime source `02d75b3b4`, on RTX 4090/SM8.9, driver 595.71.05,
+CUDA 13.0/cuBLASLt 13.1.1.3, a comparison against `d4908d95c` used concurrency 1
+and a direct native-worker HTTP
+POST through full JSON/schema validation. Each configuration ran in forward and
+reverse order, four timed traversals per visit, after excluding growth,
+all-shape stabilization and five warmups. The frozen real suite uses 12 screenshots
+(desktop/mobile/web, four each), four correlated click-choice permutations each.
+
+| Workload / mode | timed requests | mean ms | p50 ms | p95 ms | sampled max MiB |
+| --- | ---: | ---: | ---: | ---: | ---: |
+| synthetic / off | 64 | 133.30 | 119.31 | 277.13 | 10138 |
+| synthetic / shared M64 | 64 | 134.73 | 139.35 | 215.54 | 10288 |
+| synthetic / matched | 64 | 107.17 | 103.24 | 173.69 | 10448 |
+| real / off | 96 | 256.36 | 265.33 | 271.98 | 10210 |
+| real / shared M64 | 96 | 205.53 | 209.87 | 215.85 | 10318 |
+| real / matched | 96 | 152.83 | 152.94 | 160.08 | 10606 |
+
+On this real suite, matched reduces mean latency 25.6% versus shared M64 and
+40.4% versus off; p95 reductions are 25.8% and 41.1%. The paired-workload-round
+bootstrap 95% intervals for mean reduction are 25.1–26.1% and 39.9–40.7%; these
+describe repeated timings, not general task performance. Synthetic mean
+reductions are 20.5% and 19.6%. One small two-question synthetic case regresses
+4.8% versus off. There is no claim of a speedup for every shape. Real-suite sampled
+maximum increases 288 MiB versus M64 and 396 MiB versus off. Memory is sampled
+at about 100 ms, not an exact allocator peak. Vision-cache controls are 8 entries/
+1024 MiB for synthetic and 16 entries/2048 MiB for real, not new defaults.
+
+Five profiles each completed 80 real/stress layouts twice and 32 single-question
+extractions (960 calls total), plus malformed-input/recovery probes. Matched and
+candidate off have exactly the ordinary-control FP32 logit bits, complete
+responses, token IDs and all three position axes; candidate shared retains the
+original M64 results. All single-question extractions match their parent answers.
+The 320 projection-window comparisons and prefix model regressions with vision
+Graph 0/1 passed on this GPU. Existing default-versus-reference limits remain:
+the adapted 48-screenshot/192-correlated-choice set scores 119 for off/matched
+and 120 for M64. This is not official ScreenSpot accuracy or evidence of improved
+reference quality. Defaults remain off; other GPUs, libraries and concurrency
+need qualification.
+
+Raw timings, exclusions, logits, model/source/binary hashes, reproduction inputs
+and acquisition scripts, failures and final-source mapping are in the
+[2026-10-26 qualification](https://github.com/Levius-Fubuki/system1-omni/releases/download/cua-mm-qualification-20261026/cua-mm-matched-qualification-20261026.tar.gz). Original real screenshot pixels are omitted.
+These measurements precede the startup-symbol guard. That guard has host-only
+missing-symbol regression coverage; no new GPU numerical or timing run is claimed.

src/backends/cuda/qwen3_5/README.md

@@ -176,3 +176,22 @@ does not establish a fresh device pass or whole-model 27B qualification.
 
 Shared preprocessing and vision adaptation notices accompany the Qwen crate in
 [THIRD_PARTY_NOTICES.md](../../../models/qwen3_5/native/THIRD_PARTY_NOTICES.md).
+
+## Optional full-prompt-matched GEMMs
+
+`cs1_gemm_create_matched(workspace_bytes, reference_m)` and
+`cs1_gemm_compatible(ordinary, matched, M, N, K)` are additive ABI 9 symbols.
+Existing consumers do not resolve them. The matched handle retains the ordinary
+BF16 reduction mask and chooses one algorithm per weight shape at `reference_m`;
+legacy `cs1_gemm_create_fixed` still excludes split-K reductions. Compatibility
+returns 1 only for equal complete algorithm descriptors and supported layouts
+within the workspace budget. Passing a null ordinary handle checks window
+support without requiring another heuristic to choose the same algorithm.
+Neither compatibility check executes a GEMM.
+
+Rust owns admission, request scope and the bounded workspace cache. Descriptor
+matching is not a guarantee of row-invariant arithmetic on arbitrary hardware;
+run the matched projection regression in `tests/qwen3_5/kernels.rs` and the
+[Cua screenshot qualification](../../../../recipe/cua_s1/native_screenshot.md#full-prompt-matched-algorithms)
+for each supported GPU/library/model combination. Other numerical profiles,
+vision GEMMs and the existing default/fixed policies stay on their own handles.

src/backends/cuda/qwen3_5/gemm.cu

@@ -63,6 +63,7 @@ struct Gemm {
     void* workspace = nullptr;
     size_t workspace_bytes = 0;
     bool reference = false;
+    bool matched = false;
     int reference_m = 0;  // > 0: one algorithm per weight shape, chosen at this M
     std::map<Key, Plan> plans;
     std::map<std::tuple<int, int, int, bool, bool>, cublasLtMatmulAlgo_t> fixed;  // N, K, ldy, FP32, bias
@@ -126,9 +127,12 @@ int fixed_choice(Gemm& g, int N, int K, int ldy, bool fp32, bool bias, Plan& p)
     if (it == g.fixed.end()) {
         Plan r;
         int rc = describe(g.reference_m, N, K, ldy, r, fp32, bias);
-        if (rc == 0) rc = first_choice(g, r, CUBLASLT_REDUCTION_SCHEME_NONE);
+        const uint32_t schemes = g.matched
+            ? CUBLASLT_REDUCTION_SCHEME_MASK & ~CUBLASLT_REDUCTION_SCHEME_INPLACE
+            : CUBLASLT_REDUCTION_SCHEME_NONE;
+        if (rc == 0) rc = first_choice(g, r, schemes);
         if (rc == 0) {
-            // Check that the choice really has no split-K reduction.
+            // Legacy fixed excludes split-K; matched retains ordinary deterministic reductions.
             int32_t splits = 0;
             uint32_t scheme = 0;
             size_t written = 0;
@@ -137,7 +141,8 @@ int fixed_choice(Gemm& g, int N, int K, int ldy, bool fp32, bool bias, Plan& p)
                                                      sizeof(splits), &written) == CUBLAS_STATUS_SUCCESS &&
                 cublasLtMatmulAlgoConfigGetAttribute(&r.algo, CUBLASLT_ALGO_CONFIG_REDUCTION_SCHEME, &scheme,
                                                      sizeof(scheme), &written) == CUBLAS_STATUS_SUCCESS;
-            if (!read || splits != 1 || scheme != CUBLASLT_REDUCTION_SCHEME_NONE)
+            if (!read || (!g.matched && (splits != 1 || scheme != CUBLASLT_REDUCTION_SCHEME_NONE)) ||
+                scheme == CUBLASLT_REDUCTION_SCHEME_INPLACE)
                 rc = status(CUBLAS_STATUS_NOT_SUPPORTED);
         }
         if (rc == 0) it = g.fixed.emplace(key, r.algo).first;
@@ -148,6 +153,7 @@ int fixed_choice(Gemm& g, int N, int K, int ldy, bool fp32, bool bias, Plan& p)
     cublasLtMatmulHeuristicResult_t check{};
     const cublasStatus_t s = cublasLtMatmulAlgoCheck(g.handle, p.op, p.a, p.b, p.c, p.c, &p.algo, &check);
     if (s != CUBLAS_STATUS_SUCCESS) return status(s);
+    if (g.matched && check.state != CUBLAS_STATUS_SUCCESS) return status(check.state);
     if (check.workspaceSize > g.workspace_bytes) return status(CUBLAS_STATUS_NOT_SUPPORTED);
     return 0;
 }
@@ -432,3 +438,20 @@ extern "C" void* cs1_reference_gemm_create(size_t workspace_bytes) {
     if(g) g->reference=true;
     return g;
 }
+
+// Additive ABI 9 entry points. Ordinary and fixed-M64 handles keep their policy.
+extern "C" void* cs1_gemm_create_matched(size_t bytes, int reference_m) {
+    auto* g = static_cast<Gemm*>(cs1_gemm_create_fixed(bytes, reference_m));
+    if (g) g->matched = true;
+    return g;
+}
+
+// ordinary == nullptr checks that the selected algorithm supports a subwindow.
+extern "C" int cs1_gemm_compatible(void* ordinary, void* matched, int M, int N, int K) {
+    if (!matched || M <= 0 || N <= 0 || K <= 0) return 0;
+    Plan *a = nullptr, *b = nullptr;
+    if (plan_for(*static_cast<Gemm*>(matched), M, N, K, N, b)) return 0;
+    if (!ordinary) return 1;
+    if (plan_for(*static_cast<Gemm*>(ordinary), M, N, K, N, a)) return 0;
+    return std::memcmp(&a->algo, &b->algo, sizeof(a->algo)) == 0;
+}

src/backends/cuda/qwen3_5/ops.h

@@ -165,6 +165,11 @@ void* cs1_gemm_create(size_t workspace_bytes);
 // runs on). Null if reference_m <= 0 or setup fails. cs1_gemm with this handle returns a
 // cuBLAS status for an M the algorithm cannot serve, rather than switching algorithms.
 void* cs1_gemm_create_fixed(size_t workspace_bytes, int reference_m);
+// Optional additive ABI 9: choose the ordinary reduction mask at reference_m.
+// Compatibility returns 1 for equal complete descriptors, else 0. A null ordinary
+// handle tests support only (used for prefix/suffix windows); no GEMM executes.
+void* cs1_gemm_create_matched(size_t workspace_bytes, int reference_m);
+int cs1_gemm_compatible(void* ordinary, void* matched, int M, int N, int K);
 void cs1_gemm_destroy(void* gemm);
 int cs1_gemm(void* gemm, const void* x, const void* w, void* y, int M, int N, int K, int ldy,
              void* stream);

src/models/cua_s1/README.md

@@ -159,3 +159,24 @@ The historical cumulative A800 investigation and frozen-corpus results are recor
 - More than 26 options per question.
 - The Metal backend.
 - The 0.1 adapters, which are a separate artifact.
+
+## Native screenshot prefix modes
+
+The screenshot worker optionally shares a request's exact token/3D-position
+prefix with `CUA_S1_MM_PREFIX=shared`. The `fixed` control uses independent full
+prompts with the same fixed GEMMs; unset/`off` preserves default execution.
+The `matched` opt-in instead checks the ordinary full-prompt GEMM algorithms
+for all questions and supports each prefix/suffix shape before sharing. If plans
+differ or a window is unsupported, it runs ordinary independent full prompts;
+a single question also uses the ordinary path. All active prefix modes reject
+language Graph and resident features at startup. Prefix contents
+are request-local and branches execute serially. One engine-owned snapshot storage
+slot grows as needed; each shared request recaptures it before use. Requests with
+no shared prefix ignore the retained contents. Prompt construction, the CPU
+letter head, answer order and token accounting are unchanged. Fixed versus default
+rounding is a separate comparison from shared-versus-fixed continuation parity.
+The matched handle cache retains at most 16 workspaces of 32 MiB, in addition
+to snapshot storage and vision-cache budgets. Descriptor matching is admission,
+not a portable numerical guarantee; GPU/library/model combinations need explicit
+row-invariance and model qualification. The original M64 controls retain their
+own arithmetic. See the [recipe](../../../recipe/cua_s1/native_screenshot.md#request-local-language-prefix-reuse).