psv666 · Oct 28, 2026 · Oct 29, 2026 · Oct 31, 2026 · Oct 31, 2026 · Nov 4, 2026
diff --git a/Cargo.lock b/Cargo.lock
diff --git a/docs/architecture.md b/docs/architecture.md
@@ -12,12 +12,19 @@ to separately running workers. Cua-S1, Open-Jev, Decider, and StartLux Instant
 [Qwen3.5/3.8 executor](../src/models/qwen3_5/native/), which accepts
 single prompts and bounded packed prefill. StartLux Instant uses a
 bidirectional full-attention call on each prompt; its GDN stays causal, and
-it does not pack prompts or cache prefixes. Cua-S1 and Open-Jev-9B use
-single-prompt calls; Open-Jev-27B-v1.1 packs candidates within one request for
+it does not pack prompts or cache prefixes. Cua-S1 uses single-prompt calls. Open-Jev-9B and Open-Jev-27B-v1.1
+pack candidates within one request for
 input and gate/up GEMMs. Its opt-in output/down packing retains each original
 matrix algorithm and checks the combined shape; incompatible groups keep separate
 calls. Its batched linear-attention kernels preserve independent sequence state.
-Long candidates from one question can share an exact prefix when every retained
+For compatible 9B short prefixes, pointwise rows can be compacted while mixers
+retain the complete candidate layout. The model owns expansion/gather maps and
+stream lifetime; the executor retains request and candidate ordering. Exact9B
+prefix capture explicitly requests state-only output; after final-layer K/V is
+captured, unconsumed prefix hidden outputs are omitted. Hidden-row readouts
+continue to request the complete window computation.
+On 27B (and 9B with `OPEN_JEV_EXACT_PREFIX=1`), long candidates from one question
+can share an exact prefix when every retained
 projection algorithm supports the prefix and suffix shapes. This Open-Jev path
 preserves the original complete-prompt arithmetic choices; it is separate from
 the fixed-algorithm `forward_shared` interface described below.
@@ -33,8 +40,10 @@ coordinate independent processors and executors through
 [native runtime](../src/runtime/README.md) for FIFO admission and blocking dispatch
 per loaded executor. The [Decider native worker](../recipe/decider/README.md)
 also uses this runtime and can opt into outstanding-request and processed-token
-budgets plus HTTP request deadlines. Shared processing orchestration,
-cross-request batch assembly, compatibility grouping and dynamic batching are planned.
+budgets plus HTTP request deadlines. Shared processing orchestration
+and cross-worker coordination remain planned. Open-Jev-9B default execution uses
+the runtime's weighted FIFO `BatchScheduler` to combine original candidate batches;
+exact-prefix requests remain exclusive execution items.
 The pplx-decider reference worker runs the checkpoint's own loader instead of
 this shared Qwen executor: the checkpoint's full-attention layers are
 noncausal, and the language kernels in `src/backends/cuda/qwen3_5/attention.cu`
@@ -131,7 +140,8 @@ The Cua-S1 text worker uses it
 ## Native worker boundaries
 
 The native workers separate `processing.rs` from `executor.rs`; their worker
-assembly owns a `SerialScheduler` per loaded executor, and the HTTP handler
+assembly owns a scheduler per loaded executor (`BatchScheduler` for 9B default
+execution, `SerialScheduler` for the other paths), and the HTTP handler
 coordinates the three stages. Preparation validates the entire request before
 any forward call and returns executor inputs plus a response context. The context retains question
 and candidate identity, usage, and response metadata outside the executor.
@@ -146,13 +156,16 @@ and candidate identity, usage, and response metadata outside the executor.
 | StartLux Instant 4B | One unpadded token-ID prompt per question, or per knockout group and then a final prompt, in request order. | FP32 letter logits from a CPU dot of the tied embedding rows at the last position. | Per-type temperature, none-check split, wide-choice residual, and ordered `choice`, `noul`, and `score` answers with token usage. |
 | Clef | The raw `systemone` body. Preparation stays in the pinned Python encoder. | One joint-schema forward on the executor process main thread. The Rust closure only exchanges a frame. | The checkpoint's `systemone_answer` object, forwarded with its HTTP status unchanged. |
 
-Cua-S1 and Open-Jev-9B input collections are serial work. For Open-Jev-27B-v1.1,
+Cua-S1 input collections are serial work. For both Open-Jev checkpoints,
 Open-Jev's model-specific batch adapter packs up to 16 independent candidates and
 4096 tokens per group; longer prompts
-execute alone. It restores question/candidate grouping before normalization.
-Laya batches questions within one request. Shared runtime
+execute alone. For 9B default numerics, the runtime can combine these original
+batches across requests up to 16 candidates and 8192 tokens. Input/gate-up GEMMs
+retain original batch row counts; output/down GEMMs retain candidate lengths.
+Request-local prefixes run as exclusive items. It restores question/candidate
+grouping before normalization. Laya batches questions within one request. Shared runtime
 admission precedes blocking dispatch: Cua-S1 admits one question forward at a
-time; Open-Jev, Decider, Laya, StartLux Instant, and Clef admit one complete request. Clef's admitted work is the frame exchange with its executor process. StartLux runs each prompt inside that admission one at a time. Cua-S1's CPU letter projection
+time; serial Open-Jev modes, Decider, Laya, StartLux Instant, and Clef admit one complete request. Clef's admitted work is the frame exchange with its executor process. StartLux runs each prompt inside that admission one at a time. Cua-S1's CPU letter projection
 stays outside admission; Open-Jev's scalar heads and Laya's GPU heads and
 synchronized readback remain inside their request unit. Decider includes every
 backbone row and CUDA head projection, synchronizes both streams and retires
@@ -239,6 +252,13 @@ cost; shared runtime primitives own capacity accounting, cancellation/deadline
 checks and FIFO execution admission. These request budgets do not add
 cross-request batching or replace the scheduler's single execution permit.
 
+The implemented weighted batch scheduler similarly owns one FIFO execution queue
+per loaded executor. It combines already queued atomic items without waiting for
+more arrivals, enforces model-declared costs and exclusive items, and routes outputs
+back to the requesting futures. Queued cancellation skips a request before its
+first dispatch; afterward its complete work and resources are retained. A callback
+must retire device work on failure or panic before releasing that ownership.
+
 - Batch only work accepted by the same loaded executor, with compatible
   checkpoint/adapter identity, device, dtype, and input layout. Models declare
   additional constraints and limits; the scheduler owns queue and admission

diff --git a/recipe/open_jev/native.md b/recipe/open_jev/native.md
@@ -69,7 +69,7 @@ target below is Ada (`89`); pass your GPU's compute capability explicitly.
 The CUDA shared library and Rust workers must be rebuilt together for ABI
 version 9, which includes shared vision, CUDA Graph, canonical prefix continuation,
 fixed-GEMM selection and device-feature validation.
-The 27B worker also requires the additive packed-GEMM and batched-GDN
+Both checkpoints also require the additive packed-GEMM and batched-GDN
 symbols from this checkout; rebuild the worker and library together.
 
 ```sh
@@ -128,7 +128,7 @@ and CUDA kernel tests are opt-in; the latter require a GPU reservation:
 ```sh
 # Tokenizer cases against either merged export, on CPU:
 OPEN_JEV_MODEL=weights/open-jev-9b-merged \
-  cargo test --locked -p omni-open-jev-native --lib -- --ignored
+  cargo test --locked -p omni-open-jev-native --lib tokenization_tests -- --ignored
 # Inside a GPU reservation, after building the library:
 CUA_S1_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \
   cargo test --release --locked -p omni-qwen3-5-native --test kernels -- --ignored
@@ -159,7 +159,7 @@ uses the general kernel. MLP SiLU uses
 both BF16 rounding points; other layouts use the scalar path.
 
 This recipe leaves `CUA_S1_GRAPH` unset and uses eager prefill. For
-Open-Jev-27B-v1.1, candidates within a request are packed in prepared order, up
+both Open-Jev checkpoints, candidates within a request are packed in prepared order, up
 to 16 sequences and 4096 total tokens per group; longer prompts execute alone
 without truncation. Input and gate/up
 projections share GEMMs. Output/down projections combine consecutive candidates
@@ -175,9 +175,10 @@ records latency, exact output checks and frozen controls.
 That historical comparison covers the earlier packing implementation on H200
 (sm_90). The newer projection/GDN paths have exact-output regression coverage on
 an SM90 device; their execution on other CUDA architectures remains unverified.
-Open-Jev-9B runs one forward pass per candidate: packing is not validated for it.
+Open-Jev-9B packing keeps prefix reuse disabled by default. Explicit
+`OPEN_JEV_PREFIX` modes retain their existing separate execution paths.
 
-For one question with 2–16 candidates all longer than 2048 tokens, the worker can
+On 27B, for one question with 2–16 candidates all longer than 2048 tokens, the worker can
 share an exact token prefix within compatible candidate groups. The prefix must
 be at least 512 tokens, aligned to 64 tokens, with a nonempty suffix for every
 candidate. Compatibility checks retain the complete prompt's algorithms for every
@@ -242,7 +243,7 @@ the largest shared prefix encountered. Include both allocations in memory
 measurements and distinguish first traversal from warmed latency.
 
 Set `CUA_S1_GRAPH=1` on the worker to enable CUDA Graph replay. The
-27B packed text path retains at most 256 graphs; other modes retain 64. Keys are
+packed Open-Jev text path retains at most 256 graphs; other modes retain 64. Keys are
 ordered sequence token lengths. Graph-enabled packed text initially reserves
 scratch for at least 4096 rows, rounded up in 1024-row increments. Longer supported
 inputs can still grow the buffer; growth synchronizes work and clears graphs
@@ -358,3 +359,152 @@ OPEN_JEV_MODEL=$PWD/weights/open-jev-9b-merged \
 OPEN_JEV_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \
   cargo test --release --locked -p omni-open-jev-native --test prefix -- --ignored
 ```
+
+### Exact request-local prefix reuse on 9B
+
+Set `OPEN_JEV_EXACT_PREFIX=1` with `OPEN_JEV_PREFIX=0` to opt into the
+retained-algorithm prefix path already used by 27B. This computes an identical
+long prefix once within a single question, then evaluates independent candidate
+suffixes. It preserves the complete prompts' GEMM algorithms and falls back to
+ordinary packing when shapes are incompatible. All numeric prefix state is
+recomputed for every request; no historical request or answer is cached.
+
+The current eligibility rule requires 2–16 candidates, each longer than 2048
+tokens, and a common prefix of at least 512 tokens rounded down to a 64-token
+boundary. Input token usage still counts the complete candidate prompts.
+`metadata.prefix_cache` reports actual reuse and saved computation. The switch
+cannot be combined with fixed/shared/auto modes, which select a different
+algorithm policy. `OPEN_JEV_EXACT_PREFIX=0` explicitly disables this exact path;
+unset retains each checkpoint's default (9B off, 27B on).
+
+The explicit `Engine::load_with_prefix` API retains checkpoint defaults and does
+not read this additional environment override. With the 9B export and the new
+switch enabled, run the `text_prefix_buffers_grow_and_shrink_without_reusing_prior_contents`
+GPU test in `prefill_batch` for exact output equality and prefix-state lifetime.
+
+Both pinned Open-Jev tokenizers use bounded batch encoding for independent
+candidates (up to 16 at a time). This avoids computing unused character offsets
+and allows the tokenizer's worker pool to encode candidates in parallel. Padded
+tokenizers retain serial encoding, and any batch error is retried in original
+order so validation reports the same first error. The exact-token regression
+includes Unicode, special tokens, batch boundaries and oversized prompts for 9B.
+
+On the measured device profile (CUDA-reported `NVIDIA L20X`, SM 9.0,
+132 multiprocessors, `cublasLtGetVersion() == 130101`), the 9B worker can combine 512–4096-row
+projections using a tested engine-66 profile (stage 35, no split-K or reduction,
+custom options 0/1/2/4/5/6/7, clusters 2×1×1, 4×1×1 or 1×2×1, undefined inner shape). It
+requires the listed configuration; other options retain the original comparison.
+Arithmetic equivalence between these options is empirically tested, not guaranteed
+by the cuBLAS API.
+This avoids redundant full-prefix computation for compatible long candidates.
+The optional CUDA capability defaults off for other consumers; older libraries,
+other hardware/library versions and unsupported shapes retain the conservative
+algorithm comparison. It neither changes ordinary GEMM algorithm selection nor
+shares numeric state between requests. Full-model output parity is required in
+addition to standalone tile-equivalence checks.
+
+For 9B exact-prefix mode, candidates of at least 512 tokens can also reuse a
+prefix when the original packed input/gate-up GEMMs are compatible with the
+candidate projections. The check includes each original batch's total row count.
+Any mismatch declines the entire sharing attempt before computation and retains
+the original candidate packing. Other prefix modes and the 27B policy are unchanged.
+
+### 9B packing compatibility and rollback
+
+9B packing is enabled only on the validated NVIDIA L20X / SM90 / 132-SM /
+cuBLASLt 130101 profile. Other devices/library versions, including SM89, and
+older CUDA libraries without the optional profile check automatically use the
+serial eager path described below. This guard covers candidate packing itself,
+cross-request batching, packed kernels and CUDA Graphs, not only relaxed
+projection tiles. The worker logs its effective mode at startup;
+`engine.executor.packing_enabled()` exposes it to embedded callers.
+
+This is a conservative compatibility restriction: packing changed full-model
+9B outputs on SM89 even with relaxed tiles disabled. The underlying numerical
+cause is not isolated, and a passing kernel microtest is insufficient to add a
+device to the validated profile. Automatic fallback rejects explicitly enabled
+fixed/shared/auto and exact-prefix modes rather than silently ignoring them.
+27B retains its existing execution policy.
+
+`OPEN_JEV_PACK=0` is a startup escape hatch for Open-Jev-9B. It restores serial
+tokenization and request scheduling, then runs each candidate through the eager
+`Model::forward` path. It disables candidate packing, cross-request batching,
+prefix reuse, packed GDN/projection capabilities and equivalent projection tiles.
+It forces CUDA Graph off even if `CUA_S1_GRAPH=1` is inherited. Shared kernel
+implementations still come from the loaded CUDA library; this is not a promise
+of an identical historical binary or historical performance.
+
+```sh
+OPEN_JEV_PACK=0 OPEN_JEV_PREFIX=0 OPEN_JEV_FIXED=0 OPEN_JEV_EXACT_PREFIX=0 \
+OPEN_JEV_MODEL=weights/open-jev-9b-merged \
+OPEN_JEV_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \
+  target/release/omni-open-jev-native
+```
+
+The default is `OPEN_JEV_PACK=1`, which permits packing only after the device
+check; it does not force packing on an unsupported profile. Values other than `0`/`1`, a 27B rollback,
+or combining rollback with fixed/shared/auto or enabled exact-prefix reuse fail
+at startup. Prefix options may also be unset. `Engine::load` reads this switch;
+embedded callers can use `Engine::load_with_packing(..., PrefixMode::Off, false)`.
+The explicit `load_with_prefix` API applies the same device check.
+
+The frozen original-native oracle also covers the rollback route. Inside a GPU
+reservation, with the pinned 9B export (8192-token capacity) and matching library,
+run the following for both `CUA_S1_GRAPH=0` and `1`; both must take the eager route:
+
+```sh
+OPEN_JEV_PACK=0 OPEN_JEV_PREFIX=0 OPEN_JEV_FIXED=0 OPEN_JEV_EXACT_PREFIX=0 \
+OPEN_JEV_MODEL=/absolute/path/to/open-jev-9b-merged \
+OPEN_JEV_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so CUA_S1_GRAPH=1 \
+  cargo test --release --locked -p omni-open-jev-native --lib \
+  rollback_uses_serial_candidates_and_matches_original_logits \
+  -- --ignored --test-threads=1
+```
+
+To compare the default mode with explicit rollback on the **same device**, run
+`default_matches_explicit_serial_on_same_device` from the same `--lib` test
+target with `CUA_S1_GRAPH=0` and `1`. It checks all 18 cases / 94 logits bit for
+bit without assuming that another architecture reproduces the frozen SM90
+oracle. It reports the effective packing mode; a fallback pass does not validate
+packing on that device.
+
+CPU tests check configuration rejection, tokenization equivalence, device-profile
+gating and worker failure health responses. They do not substitute for this GPU
+oracle check. An unavailable batch executor or poisoned model mutex makes
+`/health` return HTTP 503 with `{"status":"failed","model":"..."}`; ordinary
+request errors and a busy executor retain HTTP 200 / `ready`. Restart the worker
+after permanent failure. This endpoint does not probe CUDA on each request.
+
+### Cross-request batching
+
+On a validated profile, the 9B default execution path batches already queued requests through the shared
+runtime, up to 16 candidates and 8192 tokens per GPU dispatch, without an added
+batching timer. Original candidate batches remain atomic: input/gate-up projection
+algorithms are selected at their original batch row counts, while output/down
+projections retain per-candidate row counts. Request-local prefix work runs alone.
+Question calibration waits for the complete request and retains candidate order.
+Fixed/shared/auto prefix modes and 27B retain serial scheduling. CUDA Graph keys
+include both original batch groups and candidate lengths. Queued cancellation,
+post-dispatch resource retention and errors follow the [runtime contract](../../src/runtime/README.md).
+
+With 9B exact-prefix mode, a bounded single-question batch can compact an exact
+16–511-token prefix for pointwise projections, residuals and the MLP. Before each
+mixer, projected rows are expanded into the original complete candidate layout;
+attention and GDN retain their original lengths, positions and sequence boundaries.
+Their outputs are gathered back into the compact layout. All projection algorithms
+are checked against the original packed input sizes and individual output sizes
+before execution; unsupported layouts keep the complete forward. Short retained
+checks use tested 64–511-row profiles, including custom option 3 and cluster 4×2×1,
+under the same device-profile/cuBLASLt scope. Ordinary packed GEMM grouping is unchanged.
+
+This path stores only reusable row-index buffers and recomputes every request.
+It does not count projection-only reuse as full-prefix `saved_tokens`; reported
+input usage still includes every original token. The original full-prefix path
+remains available for longer prefixes.
+
+A 9B exact-prefix capture requests state only. Once the final full-attention
+layer has stored its K/V rows, it omits that prefix's attention output, output
+projection, MLP and final normalization: no caller reads those hidden rows.
+The suffix starts from fresh embeddings and consumes only the captured state.
+Other window/readout modes retain complete hidden outputs. State capture is
+synchronized before it becomes available to a suffix.