[psv666]

Summary

Reduce Open-Jev-9B prefill latency and increase throughput while preserving original native outputs on the validated execution profile. Candidate packing, request-local prefix reuse, cross-request batching, attention/GEMM optimizations and CUDA Graphs share one compatibility gate.

load pinned 9B export
  explicit rollback or unvalidated device/library
    serial preparation → serial admission → eager per-candidate forward
  validated device/library
    candidate packing → FIFO batch scheduler → packed kernels / optional prefix reuse
return scores in the original question/candidate order

The validated profile is NVIDIA L20X, CUDA-reported SM90, 132 SMs and cuBLASLt 130101. Unknown profiles, including the reviewed SM89 system, and older libraries without the optional check automatically use the serial/eager path. OPEN_JEV_PACK=1 permits packing after validation; it cannot bypass the gate. Explicit fixed/shared/auto or exact-prefix modes are rejected when fallback is active. OPEN_JEV_PACK=0 remains the explicit 9B rollback. 27B retains its existing policy. Configuration and regression commands.

This addresses the SM89 review: the previous head gated only relaxed projection tiles, leaving packing active on unsupported hardware. The reviewer observed 88/94 differing logits and two HTTP decision flips; disabling packing restored same-device base parity. The fix prevents entry into that unvalidated path. It does not claim to have isolated or repaired the underlying SM89 numerical operator.

Weights remain BF16; all input tokens and candidates are evaluated. Prefix reuse stays within the current request; graphs retain execution state, not responses. Permanent executor failure returns HTTP 503 from the health endpoint.

Current main aec1540 is integrated. Its window Graphs retain owned prefix buffers; the merge preserves that lifetime management together with this PR's tree-prefix and state-only paths.

Review size against current main: 53 files, 7,814 total changed lines (+7,278/−536); 5,608 authored-code changed lines (+5,107/−501), counting Rust/CUDA/C++/headers/Python/shell including tests. This exceeds the 3,000-line review threshold. Split rationale: the runtime, Open-Jev adapter and numerical optimizations share one original-native baseline and cross-component regressions; independent kernel/runtime extraction remains a maintainer discussion. Review map: runtime batching → Open-Jev startup and execution → Qwen/CUDA numerical and graph paths → tests/oracles and recipe. The current fix adds startup-path regression coverage, effective-mode assertions and a same-device oracle comparison.

Evidence

Benchmark setup and metric definitions

Model: ZefanCai/Open-Jev-9B, BF16, revision 47e966881e489511c0c7f5633a9e1960a676a551. Device reports NVIDIA L20X, approximately 140 GiB; CUDA measures SM90. CUDA 13.0, driver 570.133.20.

The public-task benchmark contains 231 authored evaluation tasks: 139 Choice, 74 Noul, and 18 Score. Each concurrency uses three rounds with five warmup requests per round: 693 measured requests. Both backends use context capacity 4096; startup and warmup are excluded.

• End-to-end request latency (E2EL) runs from the client's HTTP POST to complete response-body receipt, JSON parsing, and schema validation. It includes preparation, queueing, execution, and return overhead.
• Mean E2EL: the arithmetic mean of all measured request latencies; it captures the average cost and is sensitive to very slow requests.
• Median E2EL (P50): the middle latency after sorting all measured requests; approximately 50% finish within this time.
• 95th-percentile E2EL (P95): 95% of measured requests finish within this time; it describes the slower end of the distribution.
• Throughput: completed HTTP requests divided by measurement wall time, in requests/second.
• Concurrent HTTP requests: client requests in flight, not the GPU batch size. Mean and percentiles are computed from all request samples pooled across the three rounds; round percentiles are not averaged.

These metrics jointly describe average latency (mean), the middle of the distribution (P50), slower requests (P95), and processing capacity (throughput). No production latency objective was supplied, so no single metric is used to declare an overall winner. Input-length/task composition, request order, concurrency, repetitions, and aggregation must be specified alongside the metrics. This is a fixed-concurrency closed-loop test: each completed request triggers the next, so throughput and mean latency are related. All reported requests succeeded; output parity is checked separately below.

Native configuration: OPEN_JEV_PREFIX=0 OPEN_JEV_EXACT_PREFIX=1 CUA_S1_GRAPH=1. The vllm-jev reference uses vLLM 0.29.0, BF16, default compilation/graphs, prefix caching with a fresh salt per request, a 128-sequence limit, and an 8192-token batch budget. Full commands and pinned dependencies are in the replay instructions.

Original native baseline, this PR, and vllm-jev

Before is original system1-omni at 22b7717, with prefix reuse disabled. After is the optimized inference implementation before the review-guard follow-up (production sources preserved through ecea1a3). These are historical measurements, not a benchmark of current head e42fa25. Bold marks the best observed value at each concurrency; ↓ lower is better, ↑ higher is better.

Concurrent HTTP requests	Implementation	Mean E2EL (ms) ↓	Median E2EL (P50, ms) ↓	95th-percentile E2EL (P95, ms) ↓	Throughput (requests/s) ↑
1	system1-omni — before	75.42	36.69	335.49	13.17
1	system1-omni — after	25.72	13.79	83.55	38.23
1	vllm-jev	61.87	61.09	115.37	15.97
8	system1-omni — before	550.30	274.57	1832.90	13.88
8	system1-omni — after	183.89	98.04	517.25	41.63
8	vllm-jev	211.10	110.95	591.35	36.62
16	system1-omni — before	1082.79	518.69	2726.15	13.87
16	system1-omni — after	355.96	206.21	830.02	42.51
16	vllm-jev	383.95	169.21	1011.87	39.30
32	system1-omni — before	2092.47	968.06	5307.22	13.86
32	system1-omni — after	682.80	345.98	1615.56	42.41
32	vllm-jev	719.96	380.60	1837.53	40.49

Historical cross-device change from original native to optimized native, computed from unrounded measurements (not a code-only speedup estimate):

Concurrent HTTP requests	Mean E2EL reduction	Median E2EL reduction (P50)	95th-percentile E2EL reduction (P95)	Throughput increase
1	65.89%	62.41%	75.10%	+190.27%
8	66.58%	64.29%	71.78%	+199.96%
16	67.13%	60.25%	69.55%	+206.57%
32	67.37%	64.26%	69.56%	+205.97%

Comparison limits: original-native and optimized measurements used different physical GPUs of the same reported model. These percentages are historical end-to-end changes, not a controlled estimate of code-only speedup. Optimized native and vllm-jev used the same physical GPU. The table records a fixed-order workload; median-latency rankings vary across repetitions and request permutations. It does not establish a consistent win at every concurrency.

A separate synthetic matrix covers candidate lengths 128/512/2048/4096 and concurrency 1/8/16/32, with two service launches and at least 30 measured seconds per cell. Both engines use context capacity 8192 with unchanged inputs up to 4096 tokens. In the recorded matrix, this PR wins all three metrics in 15/16 cells. At 128 tokens/concurrency 32, median latency is 383.85 vs 378.67 ms and throughput is 82.68 vs 83.96 requests/s, favoring vllm-jev. No matching unoptimized synthetic baseline is available.

Correctness and reproducibility

The GPU and HTTP results in this subsection are historical. They establish the recorded implementation's behavior; current-head numerical and graph regressions are reported separately below. No performance or full HTTP campaign was rerun for the current fix.

The performance target is not yet met across all workloads. Correctness is checked against original native; existing original-native/vllm-jev output differences are not treated as optimization regressions.

• All 40,884 formal optimized-native responses matched original-native decision/probability/Score values, with zero failed requests and zero observed output drift.
• Mixed candidate packing matches original native bit for bit on all 188 FP32 learned-head logits across eager and graph modes, before response normalization. The 18-case fixture covers mixed lengths, changed inputs, permutations, packing boundaries, and multiple questions.
• Fresh HTTP checks cover 231 public tasks and 192 synthetic cases, with zero decision flips or returned probability/Score drift.

Evidence index · Replay commands · Mixed-packing regression and oracle regeneration

The evidence archive contains fixed inputs, per-request timings/status/complete parsed answers, test logs, source/binary hashes, and an audit script. Run python3 audit.py in its checkout to verify checksums, recompute recorded metrics, and check response parity. It is a separate evidence branch, outside this PR's code diff. Full HTTP wire dumps and profiler artifacts are outside the published bundle.

Test Plan and Results

Validated head: e42fa256ffc2a6efed35428de2e9a85cb87e2efa; source tree: c64eb255ff9ba67a2eac652e219c4e60e4feb7b2; base: aec15406909ccbf95a5478d39a72f44be7fc159d.

Before / after regression: unvalidated_packing_uses_serial_eager_execution drives real executor startup with a host-only backend. Before the fix, an unsupported-profile response still produced pack=true and the test failed. After the fix it passes all six subprocess cases: unsupported profile, absent optional symbol, supported profile, setup error, explicit rollback and 27B. It also checks effective scheduling, disabled graphs and prefix-bypass rejection.

Local validation of this source tree:

• 416 workspace CPU tests passed, 117 ignored; 2 pinned-tokenizer CPU tests passed when explicitly selected. Formatting, Clippy with warnings denied, workspace release build and strict documentation build passed.
• 27 CUDA backend CPU tests passed. CUDA contract check: zero errors, one existing Laya ABI-declaration warning.
• 29 GPU test executions passed on one reserved NVIDIA L20X. CUDA runtime reports capability 9.0 / 132 SMs; cuBLASLt 130101, driver 570.133.20. NVML reports capability 8.9 on this host, so the gate and qualification use the CUDA-runtime values, not that NVML label. The unchanged CUDA backend sources use the previously built matching sm_90 library, SHA256 0a930affafd31c1ea4ab13d619d645ac75f45d017e88fd4248af8cc11e437362.
• Both eager and graph settings pass the 18-case / 94-logit frozen-original packing comparison, explicit rollback comparison, and new same-device default-versus-serial bitwise comparison. Effective packing is asserted, so automatic fallback cannot masquerade as positive packing coverage.
• All 7 Open-Jev prefill/isolation/prefix tests pass in each graph setting. Fixed/shared prefix, state-only capture, ordinary/tree Graph and compact Graph regressions also pass.
• Automatic fallback on real GPU: with only the optional profile symbol hidden in a copied library (simulating an older ABI-9 library), both graph settings choose serial/eager execution and match explicit rollback on all 94 logits. This is a guard-negative fixture on L20X, not an SM89 hardware retest.
• Three window Graph mechanism tests pass using the pinned 9B weights with a synthetic image-token placeholder and image rows. They exercise the integrated allocation ownership, replay and retirement paths; they do not establish Cua-S1 checkpoint accuracy.

Reproduction entry points:

cargo fmt --all --check
cargo clippy --workspace --locked --all-targets --features omni-laya/serve -- -D warnings
cargo test --workspace --locked --features omni-laya/serve
cargo build --workspace --release --locked --features omni-laya/serve
OPEN_JEV_MODEL=/path/to/pinned-export RAYON_NUM_THREADS=4 \
  cargo test --locked -p omni-open-jev-native --lib tokenization_tests -- --ignored
python3 -m unittest discover -s tests/cuda -p 'test_*.py' -v
python3 src/backends/cuda/check_contract.py --repo-root .
mkdocs build --strict

# Inside a GPU reservation, with the pinned export and matching CUDA library.
# Run each with CUA_S1_GRAPH=0 and 1.
cargo test --release --locked -p omni-open-jev-native --lib \
  default_matches_explicit_serial_on_same_device -- --ignored --nocapture --test-threads=1
cargo test --release --locked -p omni-open-jev-native --test packing_parity \
  mixed_packing_matches_frozen_original_native -- --ignored --exact --nocapture --test-threads=1

Exact-head remote checks: CI passed (Rust, benchmark helpers, CUDA contract and CUDA compilation); Docs passed. These CI jobs do not run the ignored GPU tests; GPU results above were run locally against the identical source tree.

Limits: the reviewer's SM89 hardware has not been rerun with this head. Its unsafe packing path is now gated off; re-enabling it requires full-model same-device qualification. No new throughput/latency claim or full HTTP response-matrix result is made. Local current-head logs and binary hashes are retained, but are not yet in the public historical evidence archive. Ignored tests outside the listed GPU selections remain unexecuted.

Required human contributor answers

Contributor request: compare and optimize Open-Jev-9B against vllm-jev while preserving native results, and submit a Draft PR. Implementation choices were delegated to the agent.

1. Outcome: lower request latency and increase throughput with unchanged native decisions and scores.
2. Highest-leverage variable: repeated candidate computation and execution overhead; selected within that delegated implementation scope.
3. How this work moves it: reuse compatible candidate work and improve batching/kernel execution, with original-native parity checks.
4. User and need: the contributor evaluating these two inference backends; fixed benchmark workloads were requested without a specific production traffic distribution.

Is a PR needed? Yes; the contributor requested the original performance PR and the current #196 repair. Bug-fix evidence: the linked SM89 review supplies the same-device failure/rollback comparison; the new startup regression fails before the fix and passes afterward. Current-head local validation and remaining hardware limits are stated above.

Merge Danger

Door: Two-way; the code can be reverted without a data migration.

Blast Radius: Inference. Shared batching and Qwen3.5 CUDA paths are affected. Review numerical compatibility, request-state isolation, cancellation, and graph/buffer lifetimes when reviewing this integration.

Self-review

• I have reviewed the full diff and addressed the issues I found.
• I have checked that the change follows the project's architecture and stays focused on the stated purpose.
• I have run the checks appropriate to this change and reported commands, results, and anything I could not verify above.
• I have checked that the PR description, documentation, and any accuracy or performance claims match the implementation and available evidence.

[psv666]

@hsliuustc0106 @Levius-Fubuki PTAL

[Levius-Fubuki]

The 9B packing test uses 2112-token candidates, so each runs separately under the 4096-token budget. Could you add shorter, mixed-length candidates that actually get packed and compare their outputs against the original native baseline?

Please also share the GPU validation records, frozen manifests, and replay commands so the numerical parity and performance results can be independently checked.

[psv666]

@Levius-Fubuki You're right—the old 2112-token case did not cover shorter candidates sharing one original packing group. I've added that coverage in ecea1a3 and published the evidence.

The new regression uses 18 fixed cases, including 128/257/511-token candidates packed together, permutations, changed inputs followed by replay of the original inputs, 4095/4096/4097-token totals, 16/17 candidates, and multiple questions. It checks the production grouping function against explicit expected groups and asserts that 9B packing is enabled with prefix reuse off. The old long-candidate test remains with a more specific name.

The reference was generated in a separate checkout of original native 22b7717, using its original executor and CUDA library, where 9B candidates execute individually. All 94 FP32 learned-head logits per mode matched the original bit for bit in both eager and graph modes (188/188 comparisons, max error 0). This comparison is before response normalization. The reference is scoped to the pinned BF16 9B export on NVIDIA L20X / measured SM90, CUDA 13.0.

The seven existing GPU regressions also passed in each mode. Fresh HTTP parity runs through the published replay tool covered 231 public + 192 synthetic cases, with zero decision flips and zero probability/Score drift against original native.

The evidence is available at immutable links:

• GPU validation and provenance: original output, optimized eager, optimized graph.
• Frozen manifests: packing cases, original oracle, Public231, synthetic shapes.
• Oracle regeneration / regression commands, main benchmark replay commands, and archive index and checksum inventory.

Running python3 audit.py in the evidence checkout verifies the checksums, recomputes metrics from 116,106 historical request records, and checks all 40,884 formal optimized-native answers against original-native references. I also ran this audit on the archive fetched back from GitHub.

The archive preserves timing/status and complete parsed answers without rounding; full HTTP wire dumps and profiler traces are outside its stated scope. It also retains the initial harness failure caused by sharing a Cargo target directory across checkouts; the corrected probes were rebuilt separately, and the replay instructions require separate target directories. Production inference code is unchanged by this follow-up, and these additional correctness checks make no new performance claim.

[MrDongsls]

The CPU-verifiable surface checks out — scheduler routing and order restoration, request-scoped prefix reuse, the unchanged 27B path, and the frozen packing oracle are all solid, and thanks for publishing the evidence. Remaining points, none of which overlap the thread above:

• [P2] src/backends/cuda/qwen3_5/gemm.cu:518 — the equivalent-tiles relaxation enables on SM90 + cuBLASLt 13.1.1 alone, but the bitwise-parity evidence is from a single L20X. On a different SM90 card (H100/H200) the relaxed profile enables automatically, and any accumulation-order difference there drifts outputs silently. A device allowlist or a short startup self-check (the check_fixed_rows precedent) would gate it per device.
• [P2] src/models/open_jev/native/src/contract.rs:45 — with no environment variables set, the 9B default path now packs candidates, batches across requests, and tokenizes in parallel, and no switch combination restores the old per-candidate serial execution; on unverified architectures the new kernels enable on symbol presence alone. An OPEN_JEV_PACK=0-style escape hatch (or a capability self-check with fallback at the enable points) would give a runtime rollback path.
• [P3] src/runtime/src/batching.rs:186 — one panicked batch permanently closes the executor until process restart. Deliberate and tested, but a net availability regression vs SerialScheduler; exposing liveness on /health (or rebuilding the executor) would make it orchestratable.
• [P3] REPLAY.md and the evidence archive live on the f5ed2d47 branch, so merged main would have no in-repo replay instructions — worth landing them in recipe/open_jev/ so performance comparisons on main are self-contained.

[hsliuustc0106]

Draft PR review: #196 "perf(open-jev): optimize 9B prefill while preserving native outputs"

Reviewer pass: first review, pass102 (2026-10-31). Verdict: COMMENTED (draft PR — no approval or readiness implied). ## Status block

• Base: main @ 22b7717c9 (PR-declared and merge-base-confirmed). Head: ecea1a31e287890fee5a595504203b53734ffe2a (matches origin/pr/196 and PR API headRefOid; rechecked this pass).
• Diff (two-dot git diff 22b7717c9..ecea1a31e, confirmed identical to the GitHub API numbers): 40 files, +6,461/−360. The briefing's "55 files +6497/−1057" does not match the current PR state at this head; the API and local git agree on 40/+6,461/−360.
• Draft state: DRAFT (API isDraft: true). Stay in draft until the items below are addressed.
• Standing sync note: merge-base 22b7717c9 is an older main; current tip is e573260cc (several commits behind). git merge-tree shows no content conflicts against e573260cc, but a main sync/refresh is required before ready so CI and any semantic drift are revalidated at the new base.
• Author: psv666; head hosted on InferMatrixLab (codex/open-jev-9b-performance).

Skill / evidence provenance

• Skill: .agents/skills/system1-omni-review/SKILL.md (canonical repo path), bytes read this pass, SHA-256 58be3bc60806c4c5af7038a0c6a80fadb94f3360d714556a3e965c8903dde108; last commit touching it: ef9279638 ("skills: require extra contributor review for large code changes"). References also read: docs/architecture.md (SHA-256 526e0f7a...), CONTRIBUTING.md (8040ab6f...), references/repository-contracts.md (d6ae9650...).
• Separation of sources used below: [inspected] = I read the diff/tree at the pinned SHAs; [run] = commands I executed locally; [CI] = GitHub Actions API results at the exact head SHA; [author] = claims in the PR body I did not independently verify.

Template / body check

• Four post-#138 contributor answers: present (Outcome / highest-leverage variable / how the work moves it / user and need), plus explicit "Is a PR needed? Yes; the contributor explicitly requested a Draft PR" and a bug-fix-evidence N/A statement. [inspected]
• Self-review checklist: present, all four boxes checked by the author. The underlying work items (component map with review order, per-area validation, limits) are substantively present; one required number is missing (see F-1).
• Template headings deviate from .github/PULL_REQUEST_TEMPLATE.md ("Summary"/"Evidence"/"Test Plan and Results" instead of Purpose/Test Plan/Test Result/Demo-evidence; base/head given inline instead of under "System1-Omni Version / Commit:"; no separate "Demo / evidence" heading). Content-equivalent information is present, including cold/warm separation and N/A-style scoping. Informational only.
• Performance claims are honestly hedged in the body: "The performance target is not yet met across all workloads" (loses the 128-token/c32 synthetic cell to vllm-jev), and "Comparison limits" discloses the cross-GPU baseline. Good.

Commit-by-commit review [inspected]

1. 43ccbbcf "Optimize Open-Jev-9B prefill while preserving native outputs" — 35 files, +4,427/−360. Snapshot message pins the measured local head (95c0e70f2) and exact source tree (044617e2b). Contents: new shared src/runtime/src/batching.rs (FIFO batch scheduler), new prefix_tree.rs/prefix_compaction.rs in qwen3_5, open_jev executor/batched.rs, batched/shared attention + shared-state GDN + equivalent-tile GEMM + vectorized residual-RMSNorm CUDA kernels with matching Rust cuda.rs bindings, contract defaults for 9B, docs, and GPU-ignored regression tests.
2. ecea1a31 "test(open-jev): compare mixed packing with frozen original native" — 6 files, +2,035/−1: parity test tests/open_jev/packing_parity.rs, frozen oracle fixtures (packing_cases.json 588 lines, packing_original.json 1,200 lines), packing_parity.md replay protocol, and a test registration in src/models/open_jev/native/Cargo.toml (test-only; no production inference source changes). Matches the body's claim that the last commit adds tests/fixtures/docs only.

No merge commits; linear history; clean hygiene.

Contract review

Perf-contract completeness (frozen-before-compare)

• Baseline/candidate frozen and declared before comparing: baseline = original native at 22b7717c9 (the PR merge-base itself — good), candidate = ecea1a3. [inspected body; SHAs verified]
• Hardware/toolchain/checkpoint/precision declared: NVIDIA L20X (~140 GiB, SM90), CUDA 13.0, driver 570.133.20; ZefanCai/Open-Jev-9B BF16 revision 47e96688... (matches contract.rs pinned checkpoint_revision) [inspected src/models/open_jev/native/src/contract.rs:30-52].
• Input manifest declared: 231 authored public tasks (139 Choice / 74 Noul / 18 Score), 3 rounds × 5 warmups, 693 measured requests; synthetic matrix (128/512/2048/4096 × c1/8/16/32, ≥30 s/cell) separately declared.
• Metrics and aggregation declared (mean/P50/P95/throughput, pooled rounds, closed-loop relation, no single declared winner without a supplied SLO). Cold/startup/warmup excluded from measurements; capture cost separated conceptually ("startup and warmup are excluded") and a dedicated tests/qwen3_5/capture_only.rs exists.
• F-1 (P2, process/readiness): authored-code line count not reported separately from total diff. CONTRIBUTING "Large code changes" (>3,000 authored code lines) requires the authored count reported separately from total; the body gives only totals (6,461/360). Even excluding ~1,789 lines of static fixture JSON and docs, the code diff is plausibly above 3,000 → the extra-attention requirements apply. The substantive requirements (full self-review, component map + review order, per-area validation) are present, so this is a disclosure gap, not a correctness finding. Trigger: maintainer cannot route review size without the count. Consequence: ready-marking would lack a required disclosure.
• F-2 (P2, perf evidence): baseline and candidate measured on different physical GPUs (both "L20X-labeled" but distinct devices; optimized-vs-vllm-jev used one shared device). The headline +190–207% throughput / ~65–67% latency-reduction figures are therefore cross-device historical changes, not controlled code-only speedups. The body discloses this under "Comparison limits" — but the numbers remain quotable without the caveat. Trigger: readers cite the percentages as pure code speedup. Consequence: overclaim risk; a same-device baseline rerun would close it. [author-reported measurements; disclosure inspected]

"Preserving native outputs" — bit-identical or tolerance-gated?

• Bit-identical, strictly gated in code: packing_parity.rs asserts returned FP32 learned-head logits match the frozen original bit for bit (188 logits, 18-case fixture), eager and graph modes, before response normalization. The oracle is frozen from original native at the merge-base revision and the replay doc tests/open_jev/packing_parity.md specifies independent oracle regeneration with a strict "mismatch is a failure; do not overwrite the committed oracle" policy, separate Cargo target dirs, and pinned arch 90 build. [inspected]
• HTTP-level correctness is observed-zero-drift ("40,884 responses matched, zero drift", "zero decision flips") — observational, not a tolerance gate, and the body says so. Existing original-vs-vllm differences are explicitly not treated as regressions. [author-reported for the numbers; gate mechanics inspected]
• Scope honestly stated: bitwise agreement is scoped to the pinned export + SM90 + CUDA 13.0 toolchain; other architectures/toolchains unverified.

CUDA kernel/ABI, graph policy, buffer lifetimes [inspected]

• ABI lockstep: verified by inspection. Four new optional symbols (cs1_attention_gated_batch, cs1_attention_gated_shared_batch, cs1_gdn_prefill_batch_shared, cs1_gemm_enable_equivalent_tiles) match ops.h signatures parameter-by-parameter in src/models/qwen3_5/native/src/cuda.rs (types/order/count). All resolved as Option<> with graceful degradation; legacy cs1_gdn_prefill_batch now delegates to the shared variant with state=nullptr, which issues the identical pre-existing kernel call — legacy arithmetic unchanged.
• F-3 (P2→confirm at GPU validation, behavior change): graph capture policy changed for the pinned 9B geometry (src/models/qwen3_5/native/src/model.rs, graph-capture path): the pre-capture cuda::synchronize is skipped when (hidden, intermediate, layers) == (4096, 12288, 32) and no reference/lora, so stream capture records while the eager warm run is still in flight ("recording can overlap eager GPU work. Readback still retires that eager work."). Rationale is defensible (capture does not execute; plans are warm after the eager submission; buffers are shared but untouched during recording; readback forces retirement), and the guard is conservative; failure mode should be loud (capture error), not silent. Still: this is new capture-over-pending-work behavior on the serve path whose only evidence is author-run GPU tests (graph replay with changed inputs, capture-only test). Trigger: driver/toolchain where capture with in-flight prior work behaves differently (e.g., implicit sync inside a captured op). Consequence: capture failure or, worst case, nondeterministic replay if any op under capture synchronizes implicitly. Requires explicit GPU replay evidence at the final head before ready.
• Equivalent-tiles GEMM is empirically validated, not guaranteed (src/backends/cuda/qwen3_5/gemm.cu): validated_long_profile() pins exact cuBLASLt algo config attributes (ID 66, split-K 1, reduction 0, swizzle 0, stages 35, restricted custom options/cluster shapes); cs1_gemm_enable_equivalent_tiles hard-gates on SM90 and cublasLtGetVersion() == 130101, refuses after plans exist, and returns 0 (unsupported → unchanged fallback) off-profile. The file's own comment concedes cuBLAS "does not promise arithmetic equivalence between custom options". Default path (opt-out) is byte-identical to before (compatible_algorithms short-circuits to memcmp). The bit-exact claim is measured-then-declared and scoped; acceptable for draft, keep the version/device gate. Note the packing group predicate now uses lengths[begin]/lengths[end] for profile row thresholds — semantics documented in ops.h ("tested512..4096", "64..511" — also fix the missing spaces in those comments).
• Norm kernel: new add_rms_norm_vector_4096_kernel claims to retain the scalar kernel's FP32 reduction order via shared-memory transpose; per-thread strided accumulation order is preserved by construction and block_sum is reused; dedicated kernel/reference alignment tests are included (GPU-ignored). Plausible; GPU-verified only by author report.
• Buffer lifetimes / batching: src/runtime/src/batching.rs handles cancellation before dispatch (response-closed check), retains inputs after first dispatch, contains executor panics (worker exits, callers get errors, no new work on a panicked executor), routes outputs positionally with count validation, and avoids self-join on drop. Good. Architecture ownership is respected: shared runtime owns batching; model-local code owns numerics.

27B compatibility [inspected]

Open-Jev-27B-v1.1 checkpoint values are unchanged at baseline (already pack_candidates: true, parallel_tokenization: true; now also share_packed_prefix: true, which preserves its existing long-prefix reuse path independently of packing). The changed defaults (9B packing/parallel tokenization, batched scheduler for alias open-jev-9b with PrefixMode::Off) are exactly the evidence scope of this PR.

CI (observed) [CI]

Queried both repos' Actions APIs filtered to the exact head SHA ecea1a31e287890fee5a595504203b53734ffe2a, event pull_request:

• ThinkFlowLab/system1-omni: CI run 40807444862 success; Docs run 40807444863 success.
• InferMatrixLab/system1-omni: same runs reported success (the two repos expose identical run IDs for this head).
• Evidence links in the body resolve: InferMatrixLab commit f5ed2d47 exists ("evidence(pr196): publish audited Open-Jev-9B experiments and replay inputs") with REPLAY.md present. The evidence archive itself (checksums, raw timings) was not downloaded or audited in this pass [author-reported content].

Locally executed [run]

Environment: macOS arm64, no cargo, no GPU; /Users/hsliu/miniconda3/bin/python 3.12.8 (pytest 8.4.1, mkdocs 1.6.1). Worktree at ecea1a31e, removed after.

• PYTHONPATH=. python -m pytest tests/benchmarks -q → 133 passed.
• mkdocs build --strict → exit 0.
• Rust checks (cargo fmt/clippy/test/build): not executable locally (no cargo); covered by observed green CI plus author-reported results (407 CPU tests, 106 ignored). GPU tests are #[ignore]-gated and were not executed in this pass [author-reported GPU evidence only].

Material unverified checks

• All GPU execution results (perf tables, 40,884-response parity, 7+7 eager/graph tests, packing bit-parity runs, graph replay with changed inputs) are author-reported; CI compiles but does not execute GPU tests.
• The published evidence archive contents and audit.py recomputation were not independently verified here.
• The capture-policy change (F-3) and the norm/tile bit-exactness on non-pinned toolchains remain open until a GPU validation pass at the final head after a main sync.

[psv666]

@MrDongsls I've implemented the first three changes in 9f782c6. GPU execution of this follow-up is still pending, so I am keeping the PR in draft.

1. Device gate: equivalent projection tiles now require the measured CUDA-reported profile: NVIDIA L20X, SM 9.0, 132 multiprocessors, and cublasLtGetVersion() == 130101. H100/H200, unknown names, partition sizes, or other library versions retain the conservative algorithm comparison. The existing algorithm/shape checks still apply. A CPU regression compiles the production predicate and checks the allow/deny cases; this is a profile allowlist, not a claim of cross-device numerical validation.

2. Rollback: OPEN_JEV_PACK=0 restores serial tokenization, complete-request serial admission and eager per-candidate Model::forward execution for 9B. It disables candidate/cross-request packing, prefix reuse, packed GDN/projection capabilities and equivalent tiles, and forces graphs off even if CUA_S1_GRAPH=1 is inherited. Conflicting prefix modes fail at startup. The recipe gives the launch command and oracle check. This restores the execution strategy; shared kernel implementations still come from the loaded library, so it does not claim to recreate the historical binary.

3. Failure visibility: batch-worker availability is shared across scheduler clones and becomes false before the panicked request receives its error. Thread exit also clears it. Open-Jev checks this state and model-mutex poisoning in /health, returning HTTP 503 / status: failed after permanent failure. Ordinary request errors and busy workers remain healthy. Recovery requires process restart; we do not reuse potentially corrupted model state. One nuance: serial admission itself can recover after panic, but an Open-Jev model mutex can already be poisoned, so the health check covers that case too.

Validation: 410 workspace CPU tests passed (107 ignored), plus two pinned-tokenizer tests and 27 CUDA backend CPU tests. Format, Clippy, release build, modified CUDA library compilation, contract checking and strict docs passed. The liveness/health regression exercises an ordinary error followed by successful work, then panic, HTTP 503 and rejection of later work. The new rollback test compiles and is designed to assert serial/eager routing plus all 94 frozen original logits.

Still unverified: current-head GPU execution for rollback, default packing/parity and graph replay. All four test-host GPUs are reserved by another user. Previous GPU results are explicitly labelled historical in the updated PR body; they are not presented as validation of this commit.

For point 4, I agree that replay instructions need to be discoverable. The maintained mixed-packing regression and its oracle-regeneration instructions already live in the PR. I am keeping the full one-off benchmark harness and results in the immutable evidence archive, linked from the PR. Promoting that harness into main should be a separate decision with an agreed scope, maintenance owner and regression expectations; moving it alone would create an ongoing maintenance obligation.

[psv666]

Follow-up at 9152ee4: main is now refreshed to 5fa5179. The preceding CI merge preview exposed two missing WindowReadout arguments in main's new capture_fixed / continue_fixed calls. Both now explicitly use Hidden, retaining their original complete computation/readout semantics.

The integrated tree passes 414 workspace CPU tests (109 ignored), the two pinned-tokenizer tests, 27 CUDA backend CPU tests, formatting, Clippy, release build and strict docs. Both CI and Docs passed at this exact head.

I queued one GPU for the authorized maximum of 30 minutes (2026-10-31, 01:30:04–02:00:04 UTC). The reservation ended with wait timeout exceeded; no GPU test process started, and no queue entry remains. The frozen current-head suite is ready for rollback with graph settings 0/1, mixed packing against the original oracle, the seven Open-Jev regressions in each mode, and the fixed-prefix cache integration check. These GPU checks remain unverified. The PR body records the timeout and keeps earlier GPU/performance results clearly historical. The PR remains draft.

[Levius-Fubuki]

2026-11-02 实测复核：PR #196 — 通用路径通过，原目标配置仍需验收

PR: #196
Head: 9152ee46b2f8a2589d64f8f193919d8916b96618
Base: 5fa5179d57f9267bb0d20162004d8ce26352aa9d

结论：当前SM120配置未观察到新增数值或服务回归，可以继续推进review；合并前应补齐原目标profile的优化路径验证。本轮没有确认需要修改某一行实现的通用路径缺陷，也没有批准尚未执行的专用优化或历史性能结论。

已完成实测

• 独立Cargo target：workspace414 passed /0 failed /109 ignored；Python27 passed；fmt、严格Clippy、release build通过。ignored不算通过，下面列出另行执行的GPU项目。
• 固定Qwen/Qwen3.5-9B@c202236235762e1c871ad0ccb60c8ee5ba337b9a、ZefanCai/Open-Jev-9B@47e966881e489511c0c7f5633a9e1960a676a551，全部输入校验hash后真实CPU合并导出；同一导出用于所有模式。max_length在推理前固定为8192，以覆盖4224token增长案例。真实tokenizer两项检查通过。
• RTX6000D /SM120 /156SM，driver595.91.07、NVCC13.2.86、系统cuBLASLtGetVersion=130100。匹配该硬件的新建CUDA库：通用内核30 passed /0 failed；4项过滤分别需要另一个多模态checkpoint或原目标profile。
• 不修改提交的18case/94logit清单及原始oracle，捕获同GPU当前main、head关闭packing的Graph0/1、head开启packing的Graph0/1五份结果。四个head模式各94/94 logits与当前main逐位相同；独立严格核对了全部case ID、数组形状与数量。
• 原SM90 oracle在新机上有60/94位模式差异，最大原始score差0.03560042381286621；当前main和所有head模式的差异完全相同。原oracle测试Graph0/1均exit101，失败日志保留，未更换golden。这不足以归因为PR新增回归，也不能证明差异仅由硬件造成：历史基线、工具链和导出也可能影响结果。
• 隔离/顺序/长候选/Graph shape三项真实模型测试各Graph0/1通过，共6项；Graph测试确认实际replay发生。附加并发驱动在两种模式各执行5对请求，覆盖3/17候选、长度变化及返回旧输入，全部learned logits与顺序控制相同。
• 四个真实HTTP服务模式各231条Public请求，共924条：当前main、head packed eager、head packed Graph、head显式exact-prefix+Graph。均HTTP200，原bench响应合同/质量检查通过；相对main概率漂移0、Score漂移0、Choice/Noul翻转0、usage差异0。题目和候选wire顺序相同；响应差异仅为metadata.inference_seconds及prefix_cache.enabled/saved_tokens/scope。显式exact-prefix模式有71条实际复用。
• 观察到真实复用后，补跑原有未修改测试：缓存增长/缩小、改变输入后恢复、packed-prefix跨budget、并发packed/prefix路由，分别Graph0/1，共6项通过。完整hidden-vector前缀树测试另1项通过，8个场景实际执行、2个安全回退。这7项补充结果更新了早期196-full-v1中not_run的阶段性记录。

采样窗口内最大显存18015MiB，采样间隔约0.5秒；这是本次负载观测，不是最低显存或全部容量边界认证。CPU导出环境缺少FLA的提示发生于准备阶段，不是Rust/CUDA服务回退。

推进前需要补齐

专用等价tile优化的guard在src/backends/cuda/qwen3_5/gemm_profile.h:10-13：设备名必须为NVIDIA L20X、capability9.0、132SM、cuBLASLtGetVersion=130101。单凭SM90不能推定设备合格。本机不匹配，未伪造设备信息或放宽guard。

因此，relaxed equivalent-tile、actual-row heuristic、要求guard为真的pointwise compaction及相关tree/compact Graph正向路径，仍需在原目标配置对当前head补测并保留原始结果。若计划支持新profile，应先补足相应验证；本轮SM120通用与精确匹配路径通过不能代替这些项目。本轮也没有重现历史3倍性能结论。

与#197共同修改Qwen模型/CUDA代码；两个分支独立通过的部分不能代替冲突解决后的集成验收。

证据与 review 范围

以上是2026-11-02本轮独立运行的结果，不是转述作者历史 benchmark。发出前已复核本条绑定的 head/base 未变。所有 Rust 计数均来自各 PR 隔离 Cargo target 后的重建；初版共享 target 的 Rust 结果已作废，ignored 未计为通过。

沿用可信 base 的 .agents/skills/system1-omni-review/SKILL.md，源提交 5fa5179d57f9267bb0d20162004d8ce26352aa9d，文件 SHA256 3a4b876162da14eee8db56af86664576aa5755b8f3e80b3072e944018647c3bc；已读取 contributor/architecture、repository contracts、artifact hygiene 及 CI/Docs 规范。本条补充当前提交的实测证据，不代表逐条关闭旧 head 上的全部 review 事项。

下列证据包保留原始成功/失败结果及环境、命令和输入 hash，已本地归档并核验 SHA256，尚未上传为公开附件；文件名与 hash 用于标识本次记录，不是公开下载链接。仅发布 COMMENT review；不代填人工清单、不变更 Draft、不执行合并。

归档：196-full-evidence.tar.gz，420580bytes/96成员，SHA256 71283d3d14266f7d6ab5751dd54d3867c5917c1cb0ea8752b6983ae660a6065b。包含原oracle失败、五模式raw logits、四模式HTTP响应、并发、13项真实模型测试、导出hash、驱动源码/Cargo.lock、环境及GPU采样；196-final-audit.json和补充日志更新了初始campaign的阶段性not_run记录。离线重新比较原始logit位模式、HTTP字段/顺序，结论一致。

[twu3202]

I ran the current head (9152ee46b2f8a2589d64f8f193919d8916b96618) on an sm_89 card against its base (5fa5179d57f9267bb0d20162004d8ce26352aa9d) on the same GPU, as a third device next to the L20X/SM90 evidence and Levius's SM120 run.

Setup: device NVIDIA RTX 6000 Ada Generation, compute capability 8.9, driver 595.91.07, CUDA 13.2, CUDA library built from each tree with build.sh <dir> 89. One export from Qwen/Qwen3.5-9B@c202236 and ZefanCai/Open-Jev-9B@47e9668 for every run, one server at a time on loopback.

Outputs: the default packed path is not bit-identical to serial native on sm_89

• I followed tests/open_jev/packing_parity.md, but recorded the oracle on this GPU from 5fa5179, copying packing_parity.rs into the base tree the way the doc does for 22b7717. Base itself differs from the committed SM90 oracle in all 94 logits, so, as on SM120, the committed oracle can't be used here.
• Against that sm_89 base oracle, the head with packing on differs in 88 of 94 logits, max |Δ| 0.028, with graph 0 and 1 alike. The 6 that match are the candidates that run alone: singleton-after-pack, the three in long-split, and the last candidate of budget-4097 and of count-17.
• packed_candidates_preserve_isolation_order_and_graph_shapes fails its 0.01 check with graph 0 and 1: candidate 0: packed=-9.731991 independent=-9.777791.
• It doesn't look like leakage between sequences. Reordering a pack (mixed-reverse, mixed-permute, questions-reverse) gives each candidate the same bits, lengthening one candidate (budget-4095 vs budget-4096) leaves the other two unchanged, and the *-return* cases reproduce their first run bit for bit.
• Public231 over HTTP: 126 of 231 responses differ from base, all with more than one candidate (111 choice, 15 score, no noul). Max probability difference is 0.025, with 2 decision flips, both close calls (0.521/0.478 → 0.495/0.504 and 0.403/0.390 → 0.397/0.400). Base and current main give identical responses.
• OPEN_JEV_PACK=0 brings it back: the rollback test passes against the sm_89 base oracle, and Public231 responses are identical to base (0 drift, 0 flips).

The kernel tests that check the new paths bit for bit against the old ones all pass here: both batched_attention tests, both batched_gdn tests, packed_gemm::packed_projections_match_individual_shapes_and_changed_input_replay, fixed_gemm_rows_do_not_depend_on_m and the vector RMSNorm tests. That is 30 of 34 in --test kernels; the other 4 need the L20X profile or QWEN3_5_CHECKPOINT. So the drift comes from somewhere those tests don't reach (the packed GEMM test, for one, only uses 27B shapes), and I didn't narrow it down further.

Since the 9B default now packs on every device, "preserving native outputs" is backed by the earlier L20X runs and Levius's SM120 run at this head, but doesn't hold on this sm_89 card. Gating 9B packing with the same kind of profile check as the equivalent tiles, or a load-time check that packed and serial rows match (like check_fixed_rows), would cover it. Another option is for the docs to say outputs can move on other GPUs. Your call which fits.

Same-device latency, Public231 C1

bench.py from main, warmup 5, concurrency 1, three rounds per runtime, clock post_to_body_json_and_schema_validation, no thinking field in the requests. The 231 requests are the ones in workloads/public231.jsonl at f5ed2d4 except for the question key (decision instead of q), the order, and prompt_tokens, which here counts all candidates' tokens (that file uses the longest candidate), so 2048+ is 52 requests rather than 36; manifest sha256 d13f6dd95ada14763aa5ff216bb1a56e0f78b73a5d09b4adfdbe15f3dcaa660f. Every round had 231/231 successes and evidence_complete: true.

runtime	pooled p50	pooled p95	2048+ p95
base 5fa5179	112.1 ms	1787.1 ms	2737.3 ms
head 9152ee4	74.4 ms	1842.6 ms	2679.7 ms
base 5fa5179, rerun after head	113.4 ms	1851.8 ms	2721.9 ms

p50 drops by 34%, mostly on short multi-candidate requests: the median goes from 115 to 75 ms with 5 candidates and from 142 to 90 ms with 6. p95 is inside the spread between base's two runs. The 2048+ p95 is 42–58 ms (about 2%) below both, within base's round-to-round range (2637–2763 ms). The equivalent projection tiles (gated to the L20X profile) never enable on this card, so none of this gain depends on them.

Per-round p50/p95 (ms): base 111.7/1735.2, 114.3/1850.6, 115.0/1861.1; head 73.1/1830.3, 75.1/1861.6, 74.1/1842.6; base rerun 113.0/1817.5, 114.7/1868.6, 114.3/1851.8.

Other checks at this head

• CUA_S1_GRAPH=1: Public231 responses are identical to eager.
• Concurrency 8, two rounds: responses are identical to concurrency 1. Throughput was 2.84–2.89 req/s, against 2.65 for base at concurrency 8 and 2.84 for the head at concurrency 1.
• With OPEN_JEV_PREFIX=0 OPEN_JEV_EXACT_PREFIX=1:
  • packed_prefix_matches_original_batches_across_budget_boundaries and text_prefix_buffers_grow_and_shrink_without_reusing_prior_contents pass.
  • concurrent_packed_and_prefix_requests_preserve_routing_and_state fails its check that some request actually shared a prefix (line 328).
  • pointwise_prefix_compaction_* needs the L20X profile.