From c896840da764fcba99762cb5966f44aa9f266466 Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Tue, 13 Oct 2026 17:58:32 +0800 Subject: [PATCH 1/7] Fix Cua-S1 text parity with the pinned no-FLA reference --- README.md | 2 +- docs/architecture.md | 5 +- .../cua-s1-text-parity-20261013/README.md | 128 +++++++ mkdocs.yml | 1 + recipe/README.md | 3 +- recipe/cua_s1/export_text_unmerged.py | 84 +++++ recipe/cua_s1/native.md | 36 +- src/backends/cuda/qwen3_5/README.md | 21 ++ src/backends/cuda/qwen3_5/build.sh | 8 +- src/backends/cuda/qwen3_5/gdn_fp32.cu | 352 ++++++++++++++++++ src/backends/cuda/qwen3_5/ops.h | 8 + .../cuda/qwen3_5/qwen3_5.backend.json | 1 + src/backends/cuda/qwen3_5/reference_norm.cu | 4 +- src/models/cua_s1/README.md | 24 +- src/models/cua_s1/native/src/engine.rs | 15 +- src/models/cua_s1/native/src/executor.rs | 26 +- src/models/cua_s1/native/src/provenance.rs | 31 +- src/models/qwen3_5/native/src/cuda.rs | 110 ++++++ src/models/qwen3_5/native/src/lora.rs | 1 + src/models/qwen3_5/native/src/model.rs | 96 ++++- tests/cua_s1/head.rs | 16 +- tests/cua_s1/text_export.rs | 62 +++ tests/qwen3_5/lora.rs | 23 ++ tests/qwen3_5/test_gdn_fp32.py | 129 +++++++ 24 files changed, 1138 insertions(+), 48 deletions(-) create mode 100644 docs/benchmarks/cua-s1-text-parity-20261013/README.md create mode 100644 recipe/cua_s1/export_text_unmerged.py create mode 100644 src/backends/cuda/qwen3_5/gdn_fp32.cu create mode 100644 tests/cua_s1/text_export.rs create mode 100644 tests/qwen3_5/test_gdn_fp32.py diff --git a/README.md b/README.md index ddd896fe..f1994af2 100644 --- a/README.md +++ b/README.md @@ -167,7 +167,7 @@ v1.1 has its checkpoint export and CPU contract, with no worker yet: | Model | Status | | --- | --- | | LAYA | [External worker](recipe/laya/README.md); [Python worker on Apple Silicon (MPS) and CPU](recipe/laya/apple-silicon.md); [CPU checkpoint reader](src/models/laya/README.md); [native Rust/CUDA worker on Hopper](recipe/laya/native/README.md) | -| Cua-S1 4B 0.2 (`text` adapter) | [Python worker](recipe/cua_s1/text.md); [native worker](recipe/cua_s1/native.md), CUDA, run on sm_89 | +| Cua-S1 4B 0.2 (`text` adapter) | [Python worker](recipe/cua_s1/text.md); [native worker](recipe/cua_s1/native.md), CUDA; unmerged profile validated on the cumulative A800 sm_80 build, earlier merged path on sm_89 | | Cua-S1 4B 0.2 (`multimodal` adapter) | [Python CUDA worker](src/frontend/cua_s1.py); one PNG/JPEG screenshot, `choice`; native screenshot execution remains in progress | | JEMM (Qwen3.8-27B) | [Native Rust/CUDA worker](recipe/jemm/README.md); [text and 1–4-image frozen-corpus parity on A800](src/models/jemm/README.md#validation) | | Valen-Preview-0923 | [Python reference worker](recipe/valen/reference.md); text state or one inline PNG/JPEG image, `choice`; author-run on RTX 5060 Laptop GPU (sm_120, WSL2), owner: community | diff --git a/docs/architecture.md b/docs/architecture.md index 9fb69e11..433c5ddd 100644 --- a/docs/architecture.md +++ b/docs/architecture.md @@ -56,8 +56,9 @@ dispatch retains the scheduler permit until execution completes. Laya synchronizes and disposes a failed model before returning an inference error and reports unavailable health thereafter. -Qwen executors own loaded models and CPU head weights, preserving FP64 accumulation and -the existing FP32 rounding and bias order. Finishing checks output cardinality +Qwen executors own loaded models and CPU head weights with FP64 accumulation. +Cua-S1 rounds answer-letter logits to the tied `lm_head` BF16 output dtype before +FP32 softmax; learned scalar heads retain their FP32 rounding and bias order. Finishing checks output cardinality before reconstruction. HTTP validation, error status/body conventions, and real warmup before readiness remain model-specific and unchanged. diff --git a/docs/benchmarks/cua-s1-text-parity-20261013/README.md b/docs/benchmarks/cua-s1-text-parity-20261013/README.md new file mode 100644 index 00000000..9baf163d --- /dev/null +++ b/docs/benchmarks/cua-s1-text-parity-20261013/README.md @@ -0,0 +1,128 @@ +# Cua-S1 text native/reference parity on A800 + +This investigation follows the five decision flips reported in +[PR #141](https://github.com/ThinkFlowLab/system1-omni/pull/141). +The original L20X run used a different GPU. On A800, the frozen 139-request corpus +produced two stable flips, not the original five. Two independent HTTP passes +per worker reproduced both flips. This report qualifies the A800 environment; +it does not establish that all five original L20X cases have been reproduced. + +## PR dependencies and evidence provenance + +This independent change is stacked on [#118](https://github.com/ThinkFlowLab/system1-omni/pull/118) +(`codex/jemm-native`, `ed9f9d6eb566d226e2441c5e9d84187abe1e4d63`) for the +unmerged FP32 language adapter loader and optional reference norm/GEMM/attention +operations. #118 is itself stacked on [#117](https://github.com/ThinkFlowLab/system1-omni/pull/117). +The frozen corpus comes from #141; merging #141 is not a runtime dependency. +Decider execution modes and Cua screenshot cache/resident-feature PRs are outside +this change's dependency chain and diff. + +The GPU results below were measured on cumulative candidate +`76af07996f60905082b8d48199d4a7003f33f4f6`, whose parent is +`18e1a5a6f68b9d11d3e11c7ea7b939af87272f53`. That integrated build used ABI 8. +The independently rebased branch retains #118's ABI 6 and adds only optional +FP32 GDN symbols. Rebuild the library with the worker; no cuDNN is needed for +this text path. Results from the cumulative build are historical evidence for +the pinned implementation, not a fresh GPU qualification of the rebased PR head. +The GPU server was shut down after the campaign; no new device run is implied +by packaging this PR. CPU/build/docs checks are run again on the independent head. + +## Frozen inputs and environment + +- Baseline source: `3e9cc33a61ddf110d97512fb6f296cf540377c61`. +- Corpus: PR #141 `869260649f7dd5b56cac4d91eaae9087f2e985ee`, + `benchmarks/e2e/choice.jsonl`; 139 requests/questions, 172–3988 prompt tokens. +- Historical GPU candidate parent: cumulative ABI-8/unmerged-LoRA shared executor, + `18e1a5a6f68b9d11d3e11c7ea7b939af87272f53`. +- GPU: A800 PCIe 80 GB, sm_80; CUDA toolkit 13.0.88; NVIDIA driver 595.71.05. +- Python 3.12, Torch 2.14.0+cu130, Transformers 5.17.0, PEFT 0.21.0; + TF32 disabled in the independent reference; no FLA or causal-conv1d installed. +- Base `Qwen/Qwen3.5-4B@851bf6e806efd8d0a36b00ddf55e13ccb7b8cd0a`; + text adapter `cua-ai/cua-s1-4b-0.2@16818868b0cc7813808aae4e87b417657046ab79`. + Pinned files were SHA-256 checked. The new export copies the original adapter + unchanged, with hash `9b59c5aed96171a50b26526613766bbf44347a5c7af70f81efe6bcc6e9dbfb0e`. + +The reference worker uses the repository's independent Transformers/PEFT text +implementation. No reference logits, labels or probabilities are modified. +Requests are sent through the Rust frontend to a real worker. The missing model +field in corpus rows is filled with the required `cua-s1-4b-0.2`, as the benchmark +wire adapter does. Raw request IDs, bodies, statuses, usage and model identities +are retained with source/build/launch manifests in the local investigation archive. + +## Confirmed numerical causes + +Token IDs matched on all 139 requests. The differences begin in execution: + +1. The old export merges FP32 LoRA into BF16 base weights. PEFT instead computes + the unmerged LoRA branch in FP32 and casts the combined projection to BF16. +2. The old GEMM reduction policy and normalization reductions differ from the + pinned framework. With identical first-layer MLP down-projection input, + 244,845 BF16 base outputs differed; the existing reference GEMM handle matched. + With identical normalization input, two BF16 values differed; reference norm + matched after adding the 2560-hidden shape. +3. The old GDN uses TF32 pair products and BF16 chunk intermediates/state; + the no-FLA Torch implementation uses FP32 chunked triangular solves and state. + On identical first-layer input, old relative L2 error was 0.00217711 and + 326,687 of 954,368 BF16 outputs differed. Keeping recurrence FP32 alone did not + match the reference's chunk arithmetic. Matching the chunked computation and + Torch's cumulative-decay scan association produced identical BF16 outputs. +4. Native letter projection previously emitted FP32 logits. Reference `lm_head` + emits BF16 before casting to FP32 softmax. Restoring that BF16 boundary matters + for tied options and removes the remaining readout differences. + +The fix selects the precision path using `cua-s1-text-unmerged/1`. Legacy merged +exports remain loadable. Other model adapters retain their existing numerical +mode. Prefix continuation is rejected in the new mode; eager is the default. +The probability gate of 0.02 was not relaxed. + +## Results + +| Real HTTP mode | Passes × requests | Flips vs reference | Maximum probability difference | Correct | +| --- | --- | --- | --- | --- | +| Baseline native | 2 × 139 | 2 | 0.04872608184814453 | 97/139 | +| Fixed eager | 2 × 139 | 0 | 5.960464477539063e-8 | 97/139 | +| Fixed Graph | 2 × 139 | 0 | 5.960464477539063e-8 | 97/139 | + +Every fixed request returned HTTP 200. Each mode's two passes produced identical +response bytes; model identity, usage and option/question order matched reference. +Graph trace independently records 207 captures and 72 replays, with no fallback. +The two baseline flips were `original-routing-06-1` (tools/general) and +`hard-sol-b-temporal_numeric-03` (9_credits/12_credits); both reference maxima +are ties. Accuracy staying at 97/139 is expected: this repairs implementation +parity and does not retrain or improve the model. + +A controlled full-model probe additionally matched all 139 final BF16 hidden +vectors and selected BF16 logit rows exactly. The slight HTTP probability drift +comes after that BF16 boundary in FP32 softmax/response finishing. + +Fourteen maintained GPU cases compare the new GDN with the pinned no-FLA Torch +implementation: 1, 31, 32, 33, 63, 64, 65, 127, 128, 129, 233, 513 and 2049 tokens +with 32 value/16 key heads, plus 129 tokens with 4 value/2 key heads. All BF16 +outputs matched exactly. Repeated changed-input execution resets state, invalid +shapes/scales fail, and repeated context creation/destruction releases scratch. + +## Reproduction and limits + +Use the [native text recipe](../../../recipe/cua_s1/native.md) to build sm_80, +export original FP32 LoRA separately and start the worker. Use the pinned +[reference text recipe](../../../recipe/cua_s1/text.md) in an environment without +FLA/causal-conv1d, then replay the frozen corpus twice through each worker/frontend. +Run the maintained operator regression in that reference environment: + +```sh +.venv/bin/python tests/qwen3_5/test_gdn_fp32.py \ + --library target/release/libqwen3_5_cuda.so +``` + +The historical integrated candidate passed Cargo fmt, full-workspace Clippy +with warnings denied, full-workspace tests and +release build with `--locked --features omni-laya/serve` where applicable. +Backend contract checker/tests, 21 pinned Python text contract/server tests, CUDA +operator regressions and strict MkDocs build also passed. The checkpoint-dependent +prefix test used canonical merged text weights with the original pinned multimodal +config for a zero-image contract fixture; it is not numerical-reference evidence. + +Independent-PR-head GPU execution, other GPUs, framework versions, FLA execution +and screenshot-model numerical +parity require separate qualification. This corpus does not establish general +GUI accuracy or a performance improvement. diff --git a/mkdocs.yml b/mkdocs.yml index e8488667..69226bac 100644 --- a/mkdocs.yml +++ b/mkdocs.yml @@ -65,6 +65,7 @@ nav: - Laya on Apple Silicon: recipe/laya/apple-silicon.md - Cua-S1 text worker: recipe/cua_s1/text.md - Cua-S1 native text worker: recipe/cua_s1/native.md + - Cua-S1 text numerical parity: docs/benchmarks/cua-s1-text-parity-20261013/README.md - Valen reference worker: recipe/valen/reference.md - Cua-S1 vision CUDA Graph: recipe/cua_s1/native_vision_graph.md - Open-Jev native text worker: recipe/open_jev/native.md diff --git a/recipe/README.md b/recipe/README.md index c2a31891..e85b7d6b 100644 --- a/recipe/README.md +++ b/recipe/README.md @@ -14,7 +14,8 @@ For a first real decision, follow the [complete CPU walkthrough](../docs/getting - [Valen-Preview-0923 reference worker](valen/reference.md): prepare the pinned Preview checkpoint and Qwen base, start the text-or-image `choice` worker and verify the frontend. - [Cua-S1 4B 0.2 native text worker](cua_s1/native.md): build the CUDA library and - the Rust worker, export the merged weights and start the worker. + the Rust worker, export BF16 weights with the original FP32 adapter and start + the worker. - [Open-Jev native text worker](open_jev/native.md): export the merged text backbone and trained decision head of Open-Jev-27B-v1.1 or Open-Jev-9B, then serve with Rust and CUDA. [Open-Jev-9B validation](open_jev/validation-9b.md) diff --git a/recipe/cua_s1/export_text_unmerged.py b/recipe/cua_s1/export_text_unmerged.py new file mode 100644 index 00000000..d78e1657 --- /dev/null +++ b/recipe/cua_s1/export_text_unmerged.py @@ -0,0 +1,84 @@ +"""Export BF16 language weights with the original FP32 text adapter kept separate.""" + +import argparse +import hashlib +import json +import shutil +from pathlib import Path + + +def validate_adapter(config): + targets = { + "q_proj", + "k_proj", + "v_proj", + "o_proj", + "gate_proj", + "up_proj", + "down_proj", + } + if config.get("r") != 16 or config.get("lora_alpha") != 32: + raise ValueError("Cua-S1 0.2 text requires rank 16 and alpha 32") + if set(config.get("target_modules", [])) != targets: + raise ValueError("expected the Cua-S1 text adapter, not a multimodal adapter") + if ( + config.get("bias", "none") != "none" + or config.get("lora_bias") + or config.get("use_dora") + or config.get("use_rslora") + or config.get("fan_in_fan_out") + or config.get("modules_to_save") + or config.get("rank_pattern") + or config.get("alpha_pattern") + ): + raise ValueError("unsupported text adapter scaling or projection policy") + + +def export(base: Path, adapter: Path, out: Path): + validate_adapter(json.loads((adapter / "adapter_config.json").read_text())) + if out.exists() and any(out.iterdir()): + raise ValueError("output directory must be empty") + import torch + from safetensors import safe_open + from transformers import AutoModelForCausalLM, AutoTokenizer + + adapter_file = adapter / "adapter_model.safetensors" + with safe_open(adapter_file, framework="pt", device="cpu") as tensors: + if len(tensors.keys()) != 256 or any( + tensors.get_slice(name).get_dtype() != "F32" for name in tensors.keys() + ): + raise ValueError("expected the complete 256-tensor FP32 text adapter") + model = AutoModelForCausalLM.from_pretrained( + base, dtype=torch.bfloat16, device_map="cpu" + ) + if model.config.hidden_size != 2560 or model.config.num_hidden_layers != 32: + raise ValueError("expected the Qwen3.5-4B language model") + out.mkdir(parents=True, exist_ok=True) + model.save_pretrained(out, max_shard_size="5GB") + AutoTokenizer.from_pretrained(base).save_pretrained(out) + shutil.copyfile(adapter_file, out / "adapter.safetensors") + with (out / "adapter.safetensors").open("rb") as file: + digest = hashlib.file_digest(file, "sha256").hexdigest() + (out / "cua_s1_export.json").write_text( + json.dumps( + { + "format": "cua-s1-text-unmerged/1", + "lora_scale": 2.0, + "adapter_sha256": digest, + }, + indent=2, + ) + ) + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--base", required=True, type=Path) + parser.add_argument("--adapter", required=True, type=Path) + parser.add_argument("--out", required=True, type=Path) + args = parser.parse_args() + export(args.base, args.adapter, args.out) + + +if __name__ == "__main__": + main() diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index 78d14b1a..2a23a35f 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -1,26 +1,26 @@ # Cua-S1 4B 0.2 native text worker -The native worker ([`src/models/cua_s1/native/`](../../src/models/cua_s1/native/)) serves the `text` adapter like the reference worker in [`text.md`](text.md), with the Qwen3.5-4B forward pass on the CUDA kernels in [`src/backends/cuda/qwen3_5/`](../../src/backends/cuda/qwen3_5/) and no Python or PyTorch. It needs an NVIDIA GPU with compute capability 8.0 or newer; only an RTX 6000 Ada (sm_89) with CUDA 13.2 has been run. +The native worker ([`src/models/cua_s1/native/`](../../src/models/cua_s1/native/)) serves the `text` adapter like the reference worker in [`text.md`](text.md), with the Qwen3.5-4B forward pass on the CUDA kernels in [`src/backends/cuda/qwen3_5/`](../../src/backends/cuda/qwen3_5/) and no Python or PyTorch. It needs an NVIDIA GPU with compute capability 8.0 or newer; the unmerged precision profile was checked on the cumulative A800 80 GB (sm_80) build with CUDA 13.0; device qualification of this independent #118-based branch remains pending. The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. Run the commands from the repository root. Pass your GPU's compute capability to `build.sh` (89 for Ada, 80 for A100, 90 for H100); the worker finds `libqwen3_5_cuda.so` next to its executable, or at `CUA_S1_CUDA_LIB`: ```sh -src/backends/cuda/qwen3_5/build.sh target/release 89 # needs nvcc and cuBLASLt +src/backends/cuda/qwen3_5/build.sh target/release 89 # needs nvcc, cuBLAS and cuBLASLt cargo build --release --locked -p omni-cua-s1-native ``` -The worker loads the weights with the `text` adapter merged in. With the reference worker's environment and weights from `text.md`, export them once (about 8.5 GB): +With the pinned reference environment and weights from `text.md`, export the BF16 language weights and retain the original FP32 text adapter separately (about 8.5 GB): ```sh -PYTHONPATH=src .venv/bin/python recipe/cua_s1/export_text_merged.py \ +PYTHONPATH=src .venv/bin/python recipe/cua_s1/export_text_unmerged.py \ --base weights/Qwen3.5-4B --adapter weights/cua-s1-4b-0.2/text \ - --out weights/cua-s1-4b-0.2-text-merged + --out weights/cua-s1-4b-0.2-text-unmerged ``` Start the worker (`CUA_S1_HOST` and `CUA_S1_PORT` default to `127.0.0.1` and `8000`), then the frontend and requests as in `text.md`: ```sh -CUA_S1_MODEL=weights/cua-s1-4b-0.2-text-merged target/release/omni-cua-s1-native +CUA_S1_MODEL=weights/cua-s1-4b-0.2-text-unmerged target/release/omni-cua-s1-native ``` For the local CUDA Graph experiment, also set `CUA_S1_GRAPH=1`. The first use of @@ -31,12 +31,34 @@ the eager control. Rebuild both the worker and CUDA library together (ABI 6). If capture fails, the worker returns the completed eager result and disables Graph capture/replay for its remaining lifetime, logging the failure to stderr. -Each question runs one forward pass over its prompt, eagerly by default or through exact-length CUDA Graph replay when enabled; the final hidden state at the last position times the 26 letter rows of the output projection gives the option probabilities. The probabilities are not bitwise identical to the reference worker's, since the adapter is merged and the kernels differ; they are held to the tolerance in [`src/models/cua_s1/README.md`](../../src/models/cua_s1/README.md#validation). Error messages are worded differently, and bodies nested more than 127 levels deep are refused. +Each question runs one forward pass over its prompt. The `cua-s1-text-unmerged/1` +marker selects framework-compatible BF16 normalization and cuBLASLt GEMM, +FP32 LoRA branches, and FP32 chunked Gated DeltaNet. The GDN cumulative decay +uses the pinned Torch scan order; state is reset for every prompt. The letter +projection rounds its result to BF16 before FP32 softmax, matching `lm_head`. +The marker records scale 2 and the SHA-256 of the unchanged adapter; incompatible +markers or altered adapter bytes fail before CUDA loading. Prefix continuation +is unsupported in this precision mode. + +Existing `cua-s1-text-merged/1` exports remain loadable with the earlier merged +execution path. Merging the adapter, TF32/BF16 GDN intermediates, and different +GEMM reduction policies can change decisions on close or tied options. Re-export +with the command above to use the precision profile. Rebuild worker and library +together: this branch retains #118's ABI 6, but the precision profile requires the new optional +FP32 GDN symbols. No cuDNN or Python is needed at serving time. + +The validation scope and tolerances are in +[`src/models/cua_s1/README.md`](../../src/models/cua_s1/README.md#validation). +Error messages are worded differently, and bodies nested more than 127 levels +deep are refused. The request tests need no GPU; the kernel tests compare attention and the chunked Gated DeltaNet prefill with float64 references: ```sh cargo test -p omni-cua-s1-native +# In the pinned Python reference environment, on the same GPU: +.venv/bin/python tests/qwen3_5/test_gdn_fp32.py \ + --library target/release/libqwen3_5_cuda.so CUA_S1_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \ cargo test --release -p omni-qwen3-5-native --test kernels -- --ignored ``` diff --git a/src/backends/cuda/qwen3_5/README.md b/src/backends/cuda/qwen3_5/README.md index 5e230dd4..b2a710cf 100644 --- a/src/backends/cuda/qwen3_5/README.md +++ b/src/backends/cuda/qwen3_5/README.md @@ -49,3 +49,24 @@ on this opt-in path. The shared workers' default arithmetic remains unchanged. The [JEMM numerical record](../../../../docs/benchmarks/jemm-reference-20261011/README.md) contains a frozen 1–4-image corpus and reproducible operator/HTTP checks. + +## Optional Cua-S1 no-FLA text arithmetic + +The unmerged Cua-S1 text export reuses the reference norm, attention and GEMM +operations without requiring cuDNN. Norm additionally supports hidden size 2560. +The additive `cs1_gdn_fp32_*` group owns a cuBLAS handle and 32 MiB internal +workspace; the caller supplies shape-dependent FP32 buffers. It implements +64-token chunked GDN, FP32 normalization, triangular solves and recurrent state, +Torch 2.14 cumulative-decay association, and a final BF16 output cast. Each call +starts with zero state; prefix continuation is rejected by the model adapter. +The context drains its last stream before destruction, and is retired before +that stream. Rebuild the library for this optional group; ABI remains 6 on this #118-based branch; the historical cumulative validation +build used ABI 8. +Builds now link both cuBLAS and cuBLASLt. + +`tests/qwen3_5/test_gdn_fp32.py` compares against the pinned Transformers 5.17 +no-FLA implementation. On the cumulative A800 sm_80/CUDA 13.0/Torch 2.14 build, all 14 cases passed +with identical BF16 outputs, covering chunk boundaries, grouped heads, +changed-input replay/state reset, invalid dimensions/scales and context scratch +release. The fixed relative-L2 gate is 1e-4. This does not qualify other GPUs or +Torch versions; the original approximate GDN entry points remain available. diff --git a/src/backends/cuda/qwen3_5/build.sh b/src/backends/cuda/qwen3_5/build.sh index bdf323ef..b17e33a0 100755 --- a/src/backends/cuda/qwen3_5/build.sh +++ b/src/backends/cuda/qwen3_5/build.sh @@ -3,10 +3,10 @@ # # src/backends/cuda/qwen3_5/build.sh [compute capability, e.g. 89] # -# Needs nvcc (from NVCC, CUDA_HOME/bin or PATH) and cuBLASLt; the compute capability +# Needs nvcc (from NVCC, CUDA_HOME/bin or PATH) and cuBLAS/cuBLASLt; the compute capability # defaults to CUDA_COMPUTE_CAP, else 89. The library holds machine code for that # compute capability and PTX that newer GPUs can compile at load time. The CUDA -# runtime is linked statically and cuBLASLt dynamically, with an rpath to the +# runtime is linked statically and cuBLAS/cuBLASLt dynamically, with an rpath to the # toolkit's library directory when there is one next to nvcc. set -euo pipefail here=$(cd "$(dirname "$0")" && pwd) @@ -16,9 +16,9 @@ nvcc=${NVCC:-} if [ -z "$nvcc" ]; then if [ -n "${CUDA_HOME:-}" ]; then nvcc=$CUDA_HOME/bin/nvcc; else nvcc=$(command -v nvcc); fi fi -link=(-lcublasLt) +link=(-lcublasLt -lcublas) if lib=$(cd "$(dirname "$nvcc")/../lib64" 2>/dev/null && pwd); then - link=(-L"$lib" -lcublasLt -Xlinker -rpath -Xlinker "$lib") + link=(-L"$lib" -lcublasLt -lcublas -Xlinker -rpath -Xlinker "$lib") fi extra=() # Optional JEMM framework-compatible Conv3D. Legacy builds do not require cuDNN. diff --git a/src/backends/cuda/qwen3_5/gdn_fp32.cu b/src/backends/cuda/qwen3_5/gdn_fp32.cu new file mode 100644 index 00000000..ba7baa41 --- /dev/null +++ b/src/backends/cuda/qwen3_5/gdn_fp32.cu @@ -0,0 +1,352 @@ +// FP32 chunked Gated DeltaNet for Transformers' no-FLA text reference. +// Normalize/scale before dot products; solve the unit lower-triangular systems +// in FP32; keep every recurrent-state operation FP32 until the BF16 output. +#include "common.cuh" +#include "ops.h" +#include +#include +#include + +namespace { +constexpr int C = 64, D = 128; +struct Context { + cublasHandle_t blas = nullptr; + void *scratch = nullptr; + size_t bytes = 0; + cudaStream_t stream = nullptr; + int device = 0; +}; +struct Layout { + size_t q, k, kbeta, qd, kd, u, w, a, p, cum, state, vnew, temp, out, plus, + next, pa, pu, pw, total; + Layout(int tokens, int heads) { + size_t nc = (tokens + C - 1) / C, ncc = nc * C, at = 0; + auto take = [&](size_t bytes) { + size_t offset = at; + at = (at + bytes + 255) / 256 * 256; + return offset; + }; + size_t vector = (size_t)heads * ncc * D * 4, + matrix = (size_t)heads * nc * C * C * 4; + q = take(vector); + k = take(vector); + kbeta = take(vector); + qd = take(vector); + kd = take(vector); + u = take(vector); + w = take(vector); + a = take(matrix); + p = take(matrix); + cum = take((size_t)heads * ncc * 4); + state = take((size_t)heads * D * D * 4); + vnew = take((size_t)heads * C * D * 4); + temp = take((size_t)heads * C * D * 4); + out = take((size_t)heads * C * D * 4); + plus = take((size_t)heads * C * D * 4); + next = take((size_t)heads * D * D * 4); + pa = take((size_t)heads * nc * sizeof(void *)); + pu = take((size_t)heads * nc * sizeof(void *)); + pw = take((size_t)heads * nc * sizeof(void *)); + total = at; + } +}; +struct Work { + float *q, *k, *kbeta, *qd, *kd, *u, *w, *a, *p, *cum, *state, *vnew, *temp, + *out, *plus, *next; + const float **pa; + float **pu; + float **pw; +}; +Work bind(float *workspace, const Layout &l) { + auto p = [&](size_t offset) { return (float *)((char *)workspace + offset); }; + return {p(l.q), + p(l.k), + p(l.kbeta), + p(l.qd), + p(l.kd), + p(l.u), + p(l.w), + p(l.a), + p(l.p), + p(l.cum), + p(l.state), + p(l.vnew), + p(l.temp), + p(l.out), + p(l.plus), + p(l.next), + (const float **)p(l.pa), + (float **)p(l.pu), + (float **)p(l.pw)}; +} +// Association order: Torch v2.14.0 aten/native/cuda/ScanUtils.cuh, +// scan_innermost_dim (not a sequential recurrence). +// Match the reference's row-dependent Sklansky scan, including the carry +// inserted into the first element when 64 values are scanned in two blocks. +__global__ void cumulative(const float *g, float *out, int tokens, int heads, + int nc, int width) { + __shared__ float values[C]; + int b = blockIdx.x, head = b / nc, chunk = b % nc, lane = threadIdx.x; + int count = width == 16 ? 32 : 64; + float carry = 0.f; + for (int base = 0; base < C; base += count) { + for (int i = lane; i < count; i += 32) { + int token = chunk * C + base + i; + values[i] = token < tokens ? g[(size_t)token * heads + head] : 0.f; + } + __syncthreads(); + if (lane == 0) + values[0] = __fadd_rn(values[0], carry); + __syncthreads(); + for (int step = 1; step < count; step *= 2) { + if (lane < count / 2) { + int a = (lane / step) * (2 * step) + step; + int ti = a + lane % step; + values[ti] = __fadd_rn(values[ti], values[a - 1]); + } + __syncthreads(); + } + for (int i = lane; i < count; i += 32) + out[(size_t)b * C + base + i] = values[i]; + carry = values[count - 1]; + __syncthreads(); + } +} +__global__ void prepare(const cs1::bf16 *q, const cs1::bf16 *k, + const cs1::bf16 *v, const cs1::bf16 *beta, Work ws, + int tokens, int heads, int key_heads, int nc, + float scale) { + int head = blockIdx.x / (nc * C), row = blockIdx.x % (nc * C), + chunk = row / C, r = row % C, lane = threadIdx.x; + int kh = head / (heads / key_heads), token = chunk * C + r; + float qv[4], kv[4], qs[4] = {}, ks[4] = {}; +#pragma unroll + for (int j = 0; j < 4; j++) { + int d = lane * 4 + j; + qv[j] = token < tokens + ? cs1::f32(q[((size_t)token * key_heads + kh) * D + d]) + : 0.f; + kv[j] = token < tokens + ? cs1::f32(k[((size_t)token * key_heads + kh) * D + d]) + : 0.f; + qs[j] = qv[j] * qv[j]; + ks[j] = kv[j] * kv[j]; + } + float qsum = ((qs[0] + qs[1]) + qs[2]) + qs[3], + ksum = ((ks[0] + ks[1]) + ks[2]) + ks[3]; + for (int step = 16; step > 0; step /= 2) { + qsum += __shfl_down_sync(0xffffffff, qsum, step); + ksum += __shfl_down_sync(0xffffffff, ksum, step); + } + float qi = rsqrtf(__shfl_sync(0xffffffff, qsum, 0) + 1e-6f), + ki = rsqrtf(__shfl_sync(0xffffffff, ksum, 0) + 1e-6f); + float cum = ws.cum[(size_t)head * nc * C + row], + last = ws.cum[((size_t)head * nc + chunk) * C + C - 1]; + float b = token < tokens ? cs1::f32(beta[(size_t)token * heads + head]) : 0.f; + float ec = expf(cum), er = expf(last - cum); +#pragma unroll + for (int j = 0; j < 4; j++) { + int d = lane * 4 + j; + size_t index = ((size_t)head * nc * C + row) * D + d; + float qn = __fmul_rn(__fmul_rn(qv[j], qi), scale), + kn = __fmul_rn(kv[j], ki), kb = __fmul_rn(kn, b); + ws.q[index] = qn; + ws.k[index] = kn; + ws.kbeta[index] = kb; + ws.qd[index] = __fmul_rn(qn, ec); + ws.kd[index] = __fmul_rn(kn, er); + ws.u[index] = + token < tokens + ? __fmul_rn(cs1::f32(v[((size_t)token * heads + head) * D + d]), b) + : 0.f; + ws.w[index] = __fmul_rn(kb, ec); + } +} +__global__ void mask(Work ws, int batches) { + size_t index = (size_t)blockIdx.x * blockDim.x + threadIdx.x, + total = (size_t)batches * C * C; + if (index >= total) + return; + int i = (index % (C * C)) / C, j = index % C; + size_t batch = index / (C * C); + float decay = + j <= i ? expf(ws.cum[batch * C + i] - ws.cum[batch * C + j]) : 0.f; + ws.a[index] = __fmul_rn(ws.a[index], decay); + ws.p[index] = __fmul_rn(ws.p[index], decay); +} +__global__ void pointers(Work ws, int batches) { + int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= batches) + return; + ws.pa[b] = ws.a + (size_t)b * C * C; + ws.pu[b] = ws.u + (size_t)b * C * D; + ws.pw[b] = ws.w + (size_t)b * C * D; +} +__global__ void subtract(Work ws, int heads, int nc, int chunk) { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x, + total = (size_t)heads * C * D; + if (i >= total) + return; + size_t h = i / (C * D), r = i % (C * D); + ws.vnew[i] = __fsub_rn(ws.u[(h * nc + chunk) * C * D + r], ws.temp[i]); +} +__global__ void output(Work ws, cs1::bf16 *out, int tokens, int heads, + int chunk) { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x, + total = (size_t)heads * C * D; + if (i >= total) + return; + int head = i / (C * D), row = (i % (C * D)) / D, d = i % D, + token = chunk * C + row; + if (token < tokens) + out[((size_t)token * heads + head) * D + d] = + cs1::to_bf16(__fadd_rn(ws.out[i], ws.plus[i])); +} +__global__ void update(Work ws, int heads, int nc, int chunk) { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x, + total = (size_t)heads * D * D; + if (i >= total) + return; + size_t h = i / (D * D); + float decay = expf(ws.cum[(h * nc + chunk) * C + C - 1]); + ws.state[i] = __fadd_rn(__fmul_rn(ws.state[i], decay), ws.next[i]); +} +int status(cublasStatus_t s) { + return s == CUBLAS_STATUS_SUCCESS ? 0 : 1000 + (int)s; +} +} // namespace +extern "C" void *cs1_gdn_fp32_create(size_t bytes) { + Context *c = new (std::nothrow) Context; + if (!c) + return nullptr; + c->bytes = bytes; + if (cudaGetDevice(&c->device) != cudaSuccess || + cublasCreate(&c->blas) != CUBLAS_STATUS_SUCCESS) { + delete c; + return nullptr; + } + if (bytes && cudaMalloc(&c->scratch, bytes) != cudaSuccess) { + cublasDestroy(c->blas); + delete c; + return nullptr; + } + if (cublasSetMathMode(c->blas, CUBLAS_DEFAULT_MATH) != + CUBLAS_STATUS_SUCCESS) { + if (c->scratch) + cudaFree(c->scratch); + cublasDestroy(c->blas); + delete c; + return nullptr; + } + return c; +} +extern "C" void cs1_gdn_fp32_destroy(void *opaque) { + auto *c = (Context *)opaque; + if (!c) + return; + int previous = 0; + cudaGetDevice(&previous); + cudaSetDevice(c->device); + cudaStreamSynchronize(c->stream); + cublasDestroy(c->blas); + if (c->scratch) + cudaFree(c->scratch); + delete c; + cudaSetDevice(previous); +} +extern "C" size_t cs1_gdn_fp32_workspace_floats(int tokens, int heads) { + if (tokens <= 0 || tokens > 65536 || heads <= 0 || heads > 256) + return 0; + return Layout(tokens, heads).total / 4; +} +extern "C" int cs1_gdn_prefill_fp32(void *opaque, const void *q, const void *k, + const void *v, const float *g, + const void *beta, void *out, + float *workspace, int tokens, int heads, + int key_heads, float scale, void *stream) { + auto *c = (Context *)opaque; + if (!c || !workspace || !q || !k || !v || !g || !beta || !out || + !cs1_gdn_fp32_workspace_floats(tokens, heads) || key_heads <= 0 || + heads % key_heads != 0 || !std::isfinite(scale) || scale <= 0.f) + return cudaErrorInvalidValue; + int nc = (tokens + C - 1) / C, batches = heads * nc; + c->stream = (cudaStream_t)stream; + Work ws = bind(workspace, Layout(tokens, heads)); + int rc = status(cublasSetStream(c->blas, c->stream)); + if (rc) + return rc; + if (c->bytes) { + rc = status(cublasSetWorkspace(c->blas, c->scratch, c->bytes)); + if (rc) + return rc; + } + int log_rows = 0; + while ((1 << log_rows) < batches) + ++log_rows; + int log_width = (9 + 6 - log_rows) / 2; + int width = 1 << max(4, min(9, log_width)); + cumulative<<stream>>>(g, ws.cum, tokens, heads, nc, + width); + prepare<<stream>>>( + (const cs1::bf16 *)q, (const cs1::bf16 *)k, (const cs1::bf16 *)v, + (const cs1::bf16 *)beta, ws, tokens, heads, key_heads, nc, scale); + rc = cudaGetLastError(); + if (rc) + return rc; + const float one = 1.f, zero = 0.f; +#define BLAS(expr) \ + do { \ + rc = status(expr); \ + if (rc) \ + return rc; \ + } while (0) + BLAS(cublasSgemmStridedBatched(c->blas, CUBLAS_OP_T, CUBLAS_OP_N, C, C, D, + &one, ws.k, D, C * D, ws.kbeta, D, C * D, + &zero, ws.a, C, C * C, batches)); + BLAS(cublasSgemmStridedBatched(c->blas, CUBLAS_OP_T, CUBLAS_OP_N, C, C, D, + &one, ws.k, D, C * D, ws.q, D, C * D, &zero, + ws.p, C, C * C, batches)); + mask<<<((size_t)batches * C * C + 255) / 256, 256, 0, c->stream>>>(ws, + batches); + pointers<<<(batches + 127) / 128, 128, 0, c->stream>>>(ws, batches); + rc = cudaGetLastError(); + if (rc) + return rc; + BLAS(cublasStrsmBatched(c->blas, CUBLAS_SIDE_RIGHT, CUBLAS_FILL_MODE_UPPER, + CUBLAS_OP_N, CUBLAS_DIAG_UNIT, D, C, &one, ws.pa, C, + ws.pu, D, batches)); + BLAS(cublasStrsmBatched(c->blas, CUBLAS_SIDE_RIGHT, CUBLAS_FILL_MODE_UPPER, + CUBLAS_OP_N, CUBLAS_DIAG_UNIT, D, C, &one, ws.pa, C, + ws.pw, D, batches)); + rc = cudaMemsetAsync(ws.state, 0, (size_t)heads * D * D * 4, c->stream); + if (rc) + return rc; + for (int chunk = 0; chunk < nc; chunk++) { + size_t offset = (size_t)chunk * C * D; + long long stride = (long long)nc * C * D; + BLAS(cublasSgemmStridedBatched(c->blas, CUBLAS_OP_N, CUBLAS_OP_N, D, C, D, + &one, ws.state, D, D * D, ws.w + offset, D, + stride, &zero, ws.temp, D, C * D, heads)); + subtract<<<((size_t)heads * C * D + 255) / 256, 256, 0, c->stream>>>( + ws, heads, nc, chunk); + BLAS(cublasSgemmStridedBatched(c->blas, CUBLAS_OP_N, CUBLAS_OP_N, D, C, D, + &one, ws.state, D, D * D, ws.qd + offset, D, + stride, &zero, ws.out, D, C * D, heads)); + BLAS(cublasSgemmStridedBatched( + c->blas, CUBLAS_OP_N, CUBLAS_OP_N, D, C, C, &one, ws.vnew, D, C * D, + ws.p + (size_t)chunk * C * C, C, (long long)nc * C * C, &zero, ws.plus, + D, C * D, heads)); + output<<<((size_t)heads * C * D + 255) / 256, 256, 0, c->stream>>>( + ws, (cs1::bf16 *)out, tokens, heads, chunk); + BLAS(cublasSgemmStridedBatched(c->blas, CUBLAS_OP_N, CUBLAS_OP_T, D, D, C, + &one, ws.vnew, D, C * D, ws.kd + offset, D, + stride, &zero, ws.next, D, D * D, heads)); + update<<<((size_t)heads * D * D + 255) / 256, 256, 0, c->stream>>>( + ws, heads, nc, chunk); + rc = cudaGetLastError(); + if (rc) + return rc; + } +#undef BLAS + return cudaSuccess; +} diff --git a/src/backends/cuda/qwen3_5/ops.h b/src/backends/cuda/qwen3_5/ops.h index 487aab3a..04c9e655 100644 --- a/src/backends/cuda/qwen3_5/ops.h +++ b/src/backends/cuda/qwen3_5/ops.h @@ -72,6 +72,14 @@ int cs1_gdn_gates(const void* b, const void* a, int ld, const void* A_log, const // Chunked gated delta rule, q and k L2-normalized inside, q scaled by `scale`. // q, k [T, HK, 128], v [T, H, 128], g float [T, H], beta [T, H], o [T, H, 128]. size_t cs1_gdn_workspace_floats(int T, int H); +// Optional no-FLA text precision mode: FP32 normalized Q/K and recurrent state, +// BF16 output. Context and scratch remain model-owned; prefix continuation is unsupported. +void* cs1_gdn_fp32_create(size_t workspace_bytes); +void cs1_gdn_fp32_destroy(void* context); +size_t cs1_gdn_fp32_workspace_floats(int T,int H); +int cs1_gdn_prefill_fp32(void* context,const void* q, const void* k, const void* v, const float* g, + const void* beta, void* o, float* workspace, + int T, int H, int HK, float scale, void* stream); int cs1_gdn_prefill(const void* q, const void* k, const void* v, const float* g, const void* beta, void* o, float* workspace, int T, int H, int HK, float scale, void* stream); diff --git a/src/backends/cuda/qwen3_5/qwen3_5.backend.json b/src/backends/cuda/qwen3_5/qwen3_5.backend.json index 66ed2bce..40f3ef37 100644 --- a/src/backends/cuda/qwen3_5/qwen3_5.backend.json +++ b/src/backends/cuda/qwen3_5/qwen3_5.backend.json @@ -10,6 +10,7 @@ "elementwise.cu", "attention.cu", "gdn_prefill.cu", + "gdn_fp32.cu", "gemm.cu", "runtime.cu", "vision.cu", diff --git a/src/backends/cuda/qwen3_5/reference_norm.cu b/src/backends/cuda/qwen3_5/reference_norm.cu index 61f7cdeb..c4f13891 100644 --- a/src/backends/cuda/qwen3_5/reference_norm.cu +++ b/src/backends/cuda/qwen3_5/reference_norm.cu @@ -25,10 +25,10 @@ extern "C" int reference_norm_run(const void* x,const void* delta,const void* w, } extern "C" int cs1_reference_rms_norm(const void* x,const void* w,void* out,int rows,int d,float eps,void* stream) { - if(rows<=0 || d!=5120) return cudaErrorInvalidValue; + if(rows<=0 || (d!=5120 && d!=2560)) return cudaErrorInvalidValue; return reference_norm_run(x,nullptr,w,out,rows,d,eps,stream); } extern "C" int cs1_reference_add_rms_norm(void* x,const void* delta,const void* w,void* out,int rows,int d,float eps,void* stream) { - if(rows<=0 || d!=5120) return cudaErrorInvalidValue; + if(rows<=0 || (d!=5120 && d!=2560)) return cudaErrorInvalidValue; return reference_norm_run(x,delta,w,out,rows,d,eps,stream); } diff --git a/src/models/cua_s1/README.md b/src/models/cua_s1/README.md index 36530206..c24f3ee0 100644 --- a/src/models/cua_s1/README.md +++ b/src/models/cua_s1/README.md @@ -10,7 +10,8 @@ The native worker follows the [architecture contracts](../../../docs/architectur with independent [processing](native/src/processing.rs) and [executor](native/src/executor.rs) modules. Preparation returns token IDs and an option count per question, plus a response context for identity/order and usage. -The executor returns FP32 letter logits; response finishing owns per-question +The executor projects the letter rows, rounds to the BF16 `lm_head` output dtype, +then returns FP32 letter logits; response finishing owns per-question softmax and the choice/confidence interpretation. The engine owns a [shared serial scheduler](../../runtime/README.md) that admits each question's forward before blocking dispatch. Execution remains one prompt per question @@ -117,6 +118,27 @@ The bfloat16 worker's own difference from the fp32 worker is reported next to ea **Performance.** Performance results will state the hardware, revisions, commands and raw results. Load and warmup are reported separately from warm latency, and warm latency is broken down by prompt length and option count. +## Native text precision profile + +The unmerged text export retains the original 256 FP32 LoRA tensors (rank 16, +scale 2). It selects the shared executor's framework-compatible BF16 norm, +GEMM and attention operations, and a model-owned FP32 chunked GDN context. +Q/K normalization, triangular solves and recurrent state stay FP32 until the +GDN output is cast to BF16. Cumulative decay follows Torch 2.14's scan association +order. State resets on each prompt; prefix continuation is rejected. Context, +workspace and streams retire after outstanding work and captures complete. +The BF16 letter-logit boundary is applied before response softmax. + +This profile targets the pinned no-FLA reference described above. Installing +FLA changes Transformers' GDN implementation; comparisons must record which +implementation actually ran. A different Torch version or GPU may choose other +reduction orders, so qualification remains tied to the recorded environment. +The old merged export remains an approximate execution mode. The precision +profile does not change the screenshot adapter's execution mode. + +The historical cumulative A800 investigation and frozen-corpus results are recorded in +[the numerical report](../../../docs/benchmarks/cua-s1-text-parity-20261013/README.md). + ## Not covered yet - Native execution of the `multimodal` adapter: image preprocessing, the vision tower and the vision LoRA. This is tracked in [#10](https://github.com/ThinkFlowLab/system1-omni/issues/10). diff --git a/src/models/cua_s1/native/src/engine.rs b/src/models/cua_s1/native/src/engine.rs index 21e2fd18..5c0b434e 100644 --- a/src/models/cua_s1/native/src/engine.rs +++ b/src/models/cua_s1/native/src/engine.rs @@ -2,7 +2,7 @@ use std::path::Path; -use anyhow::{Context, Result, ensure}; +use anyhow::{Context, Result}; use omni_runtime::SerialScheduler; use crate::contract::LETTERS; @@ -17,12 +17,7 @@ pub struct Engine { impl Engine { pub async fn load(dir: &Path, library: &Path) -> Result { - // Written by export_text_merged.py; without it `dir` may hold the base model alone. - ensure!( - dir.join("cua_s1_export.json").exists(), - "{} is not a merged text checkpoint; see recipe/cua_s1/native.md", - dir.display() - ); + let export = crate::provenance::text_export(dir)?; let processor = Processor::load(dir)?; let ids = LETTERS .chars() @@ -33,7 +28,7 @@ impl Engine { .context("letter token") }) .collect::>>()?; - let executor = Executor::load(dir, library, &ids).await?; + let executor = Executor::load(dir, library, &ids, export).await?; Ok(Self { processor, scheduler: SerialScheduler::default(), @@ -41,3 +36,7 @@ impl Engine { }) } } + +#[cfg(test)] +#[path = "../../../../../tests/cua_s1/text_export.rs"] +mod tests; diff --git a/src/models/cua_s1/native/src/executor.rs b/src/models/cua_s1/native/src/executor.rs index 3db4c425..6bcbdba2 100644 --- a/src/models/cua_s1/native/src/executor.rs +++ b/src/models/cua_s1/native/src/executor.rs @@ -8,6 +8,7 @@ use omni_runtime::SerialScheduler; use serde_json::Value as Json; use crate::model::Model; +use crate::provenance::TextExport; /// Tied with the output projection in Qwen3.5-4B. const EMBEDDING: &str = "model.language_model.embed_tokens.weight"; @@ -24,9 +25,23 @@ pub struct Executor { } impl Executor { - pub async fn load(dir: &Path, library: &Path, letter_ids: &[u32]) -> Result { + pub(crate) async fn load( + dir: &Path, + library: &Path, + letter_ids: &[u32], + export: TextExport, + ) -> Result { let (d, lib) = (dir.to_path_buf(), library.to_path_buf()); - let model = tokio::task::spawn_blocking(move || Model::load(&d, &lib)).await??; + let model = tokio::task::spawn_blocking(move || match export { + TextExport::Merged => Model::load(&d, &lib), + TextExport::Unmerged => { + let mut model = + Model::load_with_lora(&d, &lib, &d.join("adapter.safetensors"), 2.0)?; + model.enable_reference_text_numerics()?; + Ok(model) + } + }) + .await??; let letters = letter_rows(dir, letter_ids, model.cfg.hidden)?; Ok(Self { model: Arc::new(Mutex::new(model)), @@ -63,10 +78,13 @@ fn letter_logits(letters: &[f32], last: &[f32], n_options: usize) -> Result() as f32 + .sum::() as f32; + // The tied lm_head emits BF16 before reference option selection/softmax. + half::bf16::from_f32(dot).to_f32() }) .collect(); // Abort before submitting a later prompt, as the former per-question softmax did. diff --git a/src/models/cua_s1/native/src/provenance.rs b/src/models/cua_s1/native/src/provenance.rs index 4ea295bb..352b3ab2 100644 --- a/src/models/cua_s1/native/src/provenance.rs +++ b/src/models/cua_s1/native/src/provenance.rs @@ -1,4 +1,4 @@ -//! Hash checks before assigning the pinned multimodal model identity. +//! Validate text export policy and multimodal artifact identity before CUDA loading. use anyhow::{Context, Result, ensure}; use serde_json::Value; use sha2::{Digest, Sha256}; @@ -9,6 +9,35 @@ use std::{ }; const LOCK_SHA: &str = "9820bd232c5762f114e19680c0f8203d7e1faaf8a60c196cfe01964d6d8a6c09"; +#[derive(Clone, Copy)] +pub(crate) enum TextExport { + Merged, + Unmerged, +} + +pub(crate) fn text_export(dir: &Path) -> Result { + let marker: Value = serde_json::from_slice(&fs::read(dir.join("cua_s1_export.json"))?) + .context("invalid text export marker")?; + match marker["format"].as_str() { + Some("cua-s1-text-merged/1") => Ok(TextExport::Merged), + Some("cua-s1-text-unmerged/1") => { + ensure!( + marker["lora_scale"].as_f64() == Some(2.0), + "Cua-S1 text LoRA scale must be 2" + ); + let expected = marker["adapter_sha256"] + .as_str() + .context("missing adapter checksum")?; + ensure!( + hash_file(&dir.join("adapter.safetensors"))? == expected, + "adapter checksum mismatch" + ); + Ok(TextExport::Unmerged) + } + _ => anyhow::bail!("unsupported text export; see recipe/cua_s1/native.md"), + } +} + fn hash_file(path: &Path) -> Result { let mut file = File::open(path)?; let mut hash = Sha256::new(); diff --git a/src/models/qwen3_5/native/src/cuda.rs b/src/models/qwen3_5/native/src/cuda.rs index 277f565e..73ab545b 100644 --- a/src/models/qwen3_5/native/src/cuda.rs +++ b/src/models/qwen3_5/native/src/cuda.rs @@ -81,6 +81,10 @@ macro_rules! api { } impl Api { + /// Language arithmetic does not require cuDNN patch convolution. + pub fn language_reference(&self) -> Result<&Reference> { + self.reference.as_ref().context("missing framework language CUDA symbols; rebuild the library") + } pub fn reference(&self) -> Result<&Reference> { let reference=self.reference.as_ref().context("JEMM requires the optional reference CUDA symbols; rebuild libqwen3_5_cuda.so")?; // SAFETY: this capability query takes no arguments. @@ -196,6 +200,112 @@ api! { static API: OnceLock = OnceLock::new(); +type Fp32GdnRun = unsafe extern "C" fn( + *mut c_void, + *const c_void, + *const c_void, + *const c_void, + *const f32, + *const c_void, + *mut c_void, + *mut f32, + c_int, + c_int, + c_int, + f32, + Stream, +) -> c_int; +/// Model-owned FP32 chunk solver. Destruction drains its last producer stream. +pub struct Fp32Gdn { + handle: *mut c_void, + run: Fp32GdnRun, + destroy: unsafe extern "C" fn(*mut c_void), + size: unsafe extern "C" fn(c_int, c_int) -> usize, +} +impl Fp32Gdn { + pub fn workspace_floats(&self, tokens: usize, heads: usize) -> Result { + ensure!(tokens <= 65536 && heads <= 256, "FP32 GDN shape limit"); + // SAFETY: this getter does not dereference memory. + let size = unsafe { (self.size)(tokens as c_int, heads as c_int) }; + ensure!(size > 0, "cannot size FP32 GDN workspace"); + Ok(size) + } + /// # Safety + /// All pointers hold the documented complete tensors; workspace has at least + /// workspace_floats(tokens,heads) entries. Serialize this context and retain + /// every allocation until stream completion. + #[allow(clippy::too_many_arguments)] // Exact CUDA operation boundary. + pub unsafe fn run( + &self, + q: *const c_void, + k: *const c_void, + v: *const c_void, + g: *const f32, + beta: *const c_void, + out: *mut c_void, + workspace: *mut f32, + tokens: c_int, + heads: c_int, + key_heads: c_int, + scale: f32, + stream: Stream, + ) -> c_int { + unsafe { + (self.run)( + self.handle, + q, + k, + v, + g, + beta, + out, + workspace, + tokens, + heads, + key_heads, + scale, + stream, + ) + } + } +} +impl Drop for Fp32Gdn { + fn drop(&mut self) { + // SAFETY: this object exclusively owns the context; the backend drains before freeing it. + unsafe { (self.destroy)(self.handle) }; + } +} +impl Api { + pub fn fp32_gdn(&self) -> Result { + // SAFETY: these optional signatures match ops.h; the library lives in static API. + let (create, destroy, run, size) = unsafe { + ( + *self + ._lib + .get:: *mut c_void>(b"cs1_gdn_fp32_create\0")?, + *self + ._lib + .get::(b"cs1_gdn_fp32_destroy\0")?, + *self._lib.get::(b"cs1_gdn_prefill_fp32\0")?, + *self + ._lib + .get:: usize>( + b"cs1_gdn_fp32_workspace_floats\0", + )?, + ) + }; + // SAFETY: creation owns an independent cuBLAS handle and workspace. + let handle = unsafe { create(32 << 20) }; + ensure!(!handle.is_null(), "FP32 GDN cuBLAS initialization failed"); + Ok(Fp32Gdn { + handle, + run, + destroy, + size, + }) + } +} + /// The library next to the running executable. pub fn default_library() -> Result { let exe = std::env::current_exe()?; diff --git a/src/models/qwen3_5/native/src/lora.rs b/src/models/qwen3_5/native/src/lora.rs index cb838a49..aa390d2d 100644 --- a/src/models/qwen3_5/native/src/lora.rs +++ b/src/models/qwen3_5/native/src/lora.rs @@ -187,6 +187,7 @@ fn parse(cfg: &Config, bytes: &[u8], scale: f32) -> Result { } let native = name .strip_prefix(PREFIX) + .or_else(|| name.strip_prefix("base_model.model.model.")) .context("unsupported LoRA tensor prefix")?; let (target, letter) = if let Some(target) = native.strip_suffix(".lora_A.weight") { (target, 0) diff --git a/src/models/qwen3_5/native/src/model.rs b/src/models/qwen3_5/native/src/model.rs index 1ad7e5aa..eb6f0220 100644 --- a/src/models/qwen3_5/native/src/model.rs +++ b/src/models/qwen3_5/native/src/model.rs @@ -450,6 +450,7 @@ impl Scratch { stream: Stream, lora: Option<&lora::Lora>, reference: Option<&cuda::Reference>, + fp32_gdn: Option<&cuda::Fp32Gdn>, ) -> Result { let (h, kd, vd, hv) = (cfg.hidden, cfg.key_dim(), cfg.value_dim(), cfg.lin_v_heads); let (hq, hk, hd) = (cfg.heads, cfg.kv_heads, cfg.head_dim); @@ -473,6 +474,11 @@ impl Scratch { } else { unsafe { (cuda::api().cs1_gdn_workspace_floats)(cap as i32, hv as i32) } }; + let ws_floats = if let Some(gdn) = fp32_gdn { + ws_floats.max(gdn.workspace_floats(cap, hv)?) + } else { + ws_floats + }; let offsets = [ take(cap * 4), take(cap * h * BF16), @@ -620,6 +626,7 @@ pub struct Model { stream: Stream, gemm: *mut c_void, reference: Option<&'static cuda::Reference>, + fp32_gdn: Option, /// Buffers for the largest packed token count so far; grows as needed. scratch: Option, /// Opt-in replay with at most 64 captures per mode, keyed by sequence lengths. @@ -653,6 +660,8 @@ impl Drop for Model { if self.lora.is_some() { self.scratch = None; } + // Retire the FP32 cuBLAS context while its last stream is still valid. + self.fp32_gdn = None; // SAFETY: created by cs1_gemm_create and not destroyed before. unsafe { (cuda::api().cs1_gemm_destroy)(self.gemm) }; } @@ -761,6 +770,7 @@ impl Model { stream, gemm, reference: None, + fp32_gdn: None, scratch: None, graph_enabled: std::env::var("CUA_S1_GRAPH").as_deref() == Ok("1"), graphs: VecDeque::new(), @@ -770,9 +780,37 @@ impl Model { Ok(model) } + /// Cua-S1's unmerged text path: framework norm/GEMM/attention plus FP32 GDN. + pub fn enable_reference_text_numerics(&mut self) -> Result<()> { + ensure!( + self.cfg.hidden == 2560 && self.lora.is_some(), + "reference text numerics require an unmerged 4B language model" + ); + ensure!( + self.scratch.is_none() && self.reference.is_none() && self.fp32_gdn.is_none(), + "select reference text numerics before the first forward" + ); + cuda::set_device(0)?; + let gdn = cuda::api().fp32_gdn()?; + let reference = cuda::api().language_reference()?; + cuda::synchronize(self.stream)?; + // SAFETY: both handles are model-owned; no queued work or captures exist. + let gemm = unsafe { (reference.gemm_create)(GEMM_WORKSPACE) }; + ensure!(!gemm.is_null(), "reference text cuBLASLt setup failed"); + unsafe { (cuda::api().cs1_gemm_destroy)(self.gemm) }; + self.gemm = gemm; + self.reference = Some(reference); + self.fp32_gdn = Some(gdn); + Ok(()) + } + /// Select JEMM's pinned framework arithmetic before the first forward. /// Prefix caching is outside this execution contract. pub fn enable_reference_numerics(&mut self) -> Result<()> { + ensure!( + self.fp32_gdn.is_none(), + "cannot combine JEMM reference numerics with FP32 GDN" + ); ensure!( self.cfg.hidden == 5120 && self.cfg.head_dim == 256 && self.lora.is_some(), "reference numerics require the unmerged JEMM 27B model" @@ -924,6 +962,7 @@ impl Model { self.stream, self.lora.as_ref(), self.reference, + self.fp32_gdn.as_ref(), )?); } Ok(()) @@ -1122,6 +1161,10 @@ impl Model { /// Allocate an uninitialized prefix of `len` tokens (a multiple of 64). /// Capture must finish successfully on this model before continuation. pub fn alloc_prefix(&self, len: usize) -> Result { + ensure!( + self.fp32_gdn.is_none(), + "FP32 GDN does not support prefix caching" + ); ensure!( self.reference.is_none(), "JEMM reference numerics do not support prefix caching" @@ -1383,22 +1426,39 @@ impl Model { )?; for &(offset, length) in &sequences { check( - (self - .reference - .map_or(cuda::api().cs1_gdn_prefill, |r| r.gdn_prefill))( - p(s.lq + offset * kd * BF16), - p(s.lk + offset * kd * BF16), - p(s.lv + offset * vd * BF16), - p(s.g + offset * hv * F32).cast(), - p(s.beta + offset * hv * BF16), - p(s.lo + offset * vd * BF16), - p(s.workspace).cast(), - length, - hv as i32, - cfg.lin_k_heads as i32, - (cfg.lin_k_dim as f32).powf(-0.5), - st, - ), + if let Some(gdn) = &self.fp32_gdn { + gdn.run( + p(s.lq + offset * kd * BF16), + p(s.lk + offset * kd * BF16), + p(s.lv + offset * vd * BF16), + p(s.g + offset * hv * F32).cast(), + p(s.beta + offset * hv * BF16), + p(s.lo + offset * vd * BF16), + p(s.workspace).cast(), + length, + hv as i32, + cfg.lin_k_heads as i32, + (cfg.lin_k_dim as f32).powf(-0.5), + st, + ) + } else { + (self + .reference + .map_or(cuda::api().cs1_gdn_prefill, |r| r.gdn_prefill))( + p(s.lq + offset * kd * BF16), + p(s.lk + offset * kd * BF16), + p(s.lv + offset * vd * BF16), + p(s.g + offset * hv * F32).cast(), + p(s.beta + offset * hv * BF16), + p(s.lo + offset * vd * BF16), + p(s.workspace).cast(), + length, + hv as i32, + cfg.lin_k_heads as i32, + (cfg.lin_k_dim as f32).powf(-0.5), + st, + ) + }, "gdn prefill", )?; } @@ -1581,6 +1641,10 @@ impl Model { capture: Option<&mut PrefixState>, prefix: Option<&PrefixState>, ) -> Result<()> { + ensure!( + self.fp32_gdn.is_none(), + "FP32 GDN does not support prefix continuation" + ); ensure!( qb <= tend && (qb == 0) == prefix.is_none() && (qb == 0 || capture.is_none()), "prefix window shape" diff --git a/tests/cua_s1/head.rs b/tests/cua_s1/head.rs index 19dac505..37e03680 100644 --- a/tests/cua_s1/head.rs +++ b/tests/cua_s1/head.rs @@ -1,7 +1,21 @@ use super::*; #[test] -fn letter_projection_preserves_fp64_accumulation_and_fp32_readout() { +fn letter_projection_matches_the_reference_bf16_output_boundary() { + assert_eq!( + letter_logits(&[1.0078125], &[1.0078125], 1).unwrap(), + [1.015625] + ); + // Both exact dot products belong to the same BF16 bin. Retaining FP32 + // invents a winning second option where the reference has an ordered tie. + assert_eq!( + letter_logits(&[20.0, 0.03125, 20.0, 0.046875], &[1.0, 1.0], 2).unwrap(), + [20.0, 20.0] + ); +} + +#[test] +fn letter_projection_preserves_fp64_accumulation_and_bf16_readout() { let letters = [ 16_777_216.0, 1.0, diff --git a/tests/cua_s1/text_export.rs b/tests/cua_s1/text_export.rs new file mode 100644 index 00000000..de813edb --- /dev/null +++ b/tests/cua_s1/text_export.rs @@ -0,0 +1,62 @@ +use super::*; +use serde_json::json; +use std::{ + fs, + path::PathBuf, + time::{SystemTime, UNIX_EPOCH}, +}; + +struct Directory(PathBuf); +impl Directory { + fn new(marker: serde_json::Value) -> Self { + let nonce = SystemTime::now() + .duration_since(UNIX_EPOCH) + .unwrap() + .as_nanos(); + let path = + std::env::temp_dir().join(format!("cua-text-export-{}-{nonce}", std::process::id())); + fs::create_dir(&path).unwrap(); + fs::write(path.join("cua_s1_export.json"), marker.to_string()).unwrap(); + Self(path) + } +} +impl Drop for Directory { + fn drop(&mut self) { + fs::remove_dir_all(&self.0).unwrap(); + } +} + +#[tokio::test] +async fn unknown_export_is_rejected_before_tokenizer_or_cuda_loading() { + let dir = Directory::new(json!({"format":"unknown"})); + let error = Engine::load(&dir.0, Path::new("unused-cuda-library")) + .await + .err() + .expect("must reject unknown export"); + assert!( + error.to_string().contains("unsupported text export"), + "{error}" + ); +} + +#[tokio::test] +async fn reference_export_rejects_incorrect_adapter_scale_before_loading() { + let dir = Directory::new(json!({"format":"cua-s1-text-unmerged/1", "lora_scale":2.1})); + let error = Engine::load(&dir.0, Path::new("unused-cuda-library")) + .await + .err() + .expect("must reject incorrect scale"); + assert!(error.to_string().contains("LoRA scale"), "{error}"); +} + +#[tokio::test] +async fn reference_export_rejects_changed_adapter_bytes_before_cuda_loading() { + let dir = Directory::new(json!({"format":"cua-s1-text-unmerged/1", "lora_scale":2.0, + "adapter_sha256":"0".repeat(64)})); + fs::write(dir.0.join("adapter.safetensors"), b"changed checkpoint").unwrap(); + let error = Engine::load(&dir.0, Path::new("unused-cuda-library")) + .await + .err() + .expect("must reject changed adapter"); + assert!(error.to_string().contains("adapter checksum"), "{error}"); +} diff --git a/tests/qwen3_5/lora.rs b/tests/qwen3_5/lora.rs index f32b78c7..55b7496d 100644 --- a/tests/qwen3_5/lora.rs +++ b/tests/qwen3_5/lora.rs @@ -52,6 +52,29 @@ fn pair(target: &str, rows: usize, cols: usize, rank: usize) -> Vec { ]) } +#[test] +fn text_adapter_prefix_preserves_the_same_fp32_projection() { + let cfg = config(); + let bytes = tensors(&[ + ( + "base_model.model.model.layers.0.mlp.down_proj.lora_A.weight", + vec![16, cfg.intermediate], + Dtype::F32, + 0.125, + ), + ( + "base_model.model.model.layers.0.mlp.down_proj.lora_B.weight", + vec![cfg.hidden, 16], + Dtype::F32, + -0.25, + ), + ]); + let parsed = parse(&cfg, &bytes, 2.0).unwrap(); + assert_eq!(parsed.pairs.len(), 1); + assert_eq!(parsed.pairs[0].target.group, Group::Down); + assert_eq!(parsed.pairs[0].rank, 16); +} + #[test] fn layouts_cover_both_mixers_and_the_small_gate_projections() { for (hidden, intermediate, heads, values) in [(5120, 17408, 24, 48), (2560, 9216, 16, 32)] { diff --git a/tests/qwen3_5/test_gdn_fp32.py b/tests/qwen3_5/test_gdn_fp32.py new file mode 100644 index 00000000..fb235409 --- /dev/null +++ b/tests/qwen3_5/test_gdn_fp32.py @@ -0,0 +1,129 @@ +"""GPU regression against the pinned Transformers 5.17 no-FLA implementation. + +Run with Python 3.12, torch 2.14.0+cu130, transformers 5.17.0: + python tests/qwen3_5/test_gdn_fp32.py --library target/release/libqwen3_5_cuda.so +The relative-L2 gate is fixed at 1e-4; repeatability/state isolation is exact. +""" + +import argparse +import ctypes +import inspect +import json +import math + +import torch +import transformers +from transformers.models.qwen3_5.modeling_qwen3_5 import torch_chunk_gated_delta_rule + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--library", required=True) + args = parser.parse_args() + assert torch.__version__.split("+")[0] == "2.14.0", torch.__version__ + assert transformers.__version__ == "5.17.0", transformers.__version__ + torch.backends.cuda.matmul.allow_tf32 = False + torch.manual_seed(141) + lib = ctypes.CDLL(args.library) + create = lib.cs1_gdn_fp32_create + create.argtypes, create.restype = [ctypes.c_size_t], ctypes.c_void_p + destroy = lib.cs1_gdn_fp32_destroy + destroy.argtypes, destroy.restype = [ctypes.c_void_p], None + size = lib.cs1_gdn_fp32_workspace_floats + size.argtypes, size.restype = [ctypes.c_int, ctypes.c_int], ctypes.c_size_t + run = lib.cs1_gdn_prefill_fp32 + run.argtypes = ( + [ctypes.c_void_p] * 8 + [ctypes.c_int] * 3 + [ctypes.c_float, ctypes.c_void_p] + ) + run.restype = ctypes.c_int + # Initialize cuBLAS before checking repeated context ownership. + warm = create(32 << 20) + assert warm + destroy(warm) + torch.cuda.synchronize() + free_before = torch.cuda.mem_get_info()[0] + for _ in range(3): + context = create(32 << 20) + assert context + destroy(context) + torch.cuda.synchronize() + assert torch.cuda.mem_get_info()[0] == free_before, "context scratch leaked" + for tokens, heads in [(0, 32), (-1, 32), (65537, 32), (1, 0), (1, 257)]: + assert size(tokens, heads) == 0 + reference = inspect.unwrap(torch_chunk_gated_delta_rule) + context = create(32 << 20) + assert context + try: + for tokens, heads, key_heads in [ + (t, 32, 16) + for t in [1, 31, 32, 33, 63, 64, 65, 127, 128, 129, 233, 513, 2049] + ] + [(129, 4, 2)]: + q = torch.randn( + 1, tokens, key_heads, 128, device="cuda", dtype=torch.bfloat16 + ) + k = torch.randn_like(q) + v = torch.randn(1, tokens, heads, 128, device="cuda", dtype=torch.bfloat16) + g = -torch.rand(1, tokens, heads, device="cuda") * 0.25 + beta = torch.rand(1, tokens, heads, device="cuda", dtype=torch.bfloat16) + work = torch.empty(size(tokens, heads), device="cuda", dtype=torch.float32) + out = torch.empty_like(v) + stream = torch.cuda.current_stream().cuda_stream + pointers = [x.data_ptr() for x in (q, k, v, g, beta, out, work)] + call = [ + context, + *pointers, + tokens, + heads, + key_heads, + 1 / math.sqrt(128), + stream, + ] + expected, _ = reference( + q.repeat_interleave(heads // key_heads, dim=2), + k.repeat_interleave(heads // key_heads, dim=2), + v, + g, + beta, + use_qk_l2norm_in_kernel=True, + ) + assert run(*call) == 0 + torch.cuda.synchronize() + saved = out.clone() + relative = ( + torch.linalg.vector_norm(out.float() - expected.float()) + / torch.linalg.vector_norm(expected.float()) + ).item() + assert relative <= 1e-4, (tokens, heads, relative) + # Poison scratch and run unrelated data before replaying the original input. + work.fill_(float("nan")) + changed = torch.randn_like(v) + changed_call = call.copy() + changed_call[3] = changed.data_ptr() + assert run(*changed_call) == 0 + assert run(*call) == 0 + torch.cuda.synchronize() + assert torch.equal(out, saved), "previous request leaked recurrent state" + # Invalid shape/scale must fail before a launch or device dereference. + for index, invalid in [(8, 0), (10, 3), (11, float("nan")), (11, 0.0)]: + invalid_call = call.copy() + invalid_call[index] = invalid + assert run(*invalid_call) != 0 + print( + json.dumps( + { + "tokens": tokens, + "heads": heads, + "key_heads": key_heads, + "relative_l2": relative, + "bf16_mismatches": int((saved != expected).sum().item()), + "state_reset": True, + } + ), + flush=True, + ) + finally: + destroy(context) + + +if __name__ == "__main__": + main() From 27f77e49868d84b091aa804f02ccb4d3474a9a00 Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Wed, 14 Oct 2026 09:05:25 +0800 Subject: [PATCH 2/7] docs(cua-s1): track independent-head qualification separately --- .../cua-s1-text-parity-20261013/README.md | 14 ++++++++------ recipe/cua_s1/native.md | 2 +- 2 files changed, 9 insertions(+), 7 deletions(-) diff --git a/docs/benchmarks/cua-s1-text-parity-20261013/README.md b/docs/benchmarks/cua-s1-text-parity-20261013/README.md index 9baf163d..dd485921 100644 --- a/docs/benchmarks/cua-s1-text-parity-20261013/README.md +++ b/docs/benchmarks/cua-s1-text-parity-20261013/README.md @@ -23,9 +23,11 @@ The GPU results below were measured on cumulative candidate The independently rebased branch retains #118's ABI 6 and adds only optional FP32 GDN symbols. Rebuild the library with the worker; no cuDNN is needed for this text path. Results from the cumulative build are historical evidence for -the pinned implementation, not a fresh GPU qualification of the rebased PR head. -The GPU server was shut down after the campaign; no new device run is implied -by packaging this PR. CPU/build/docs checks are run again on the independent head. +the pinned implementation. Fresh independent-branch qualification, including +exact source/library/binary hashes, repeated HTTP results and self-review, is +recorded separately in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). +Use each run's source pin; do not substitute the historical integrated result +for a newer branch's validation. ## Frozen inputs and environment @@ -122,7 +124,7 @@ operator regressions and strict MkDocs build also passed. The checkpoint-depende prefix test used canonical merged text weights with the original pinned multimodal config for a zero-image contract fixture; it is not numerical-reference evidence. -Independent-PR-head GPU execution, other GPUs, framework versions, FLA execution -and screenshot-model numerical -parity require separate qualification. This corpus does not establish general +Other GPUs, framework versions, FLA execution and screenshot-model numerical +parity require separate qualification. Independent-head evidence remains +source-pinned in #143 rather than inferred from these historical results. This corpus does not establish general GUI accuracy or a performance improvement. diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index 2a23a35f..1da0660e 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -1,6 +1,6 @@ # Cua-S1 4B 0.2 native text worker -The native worker ([`src/models/cua_s1/native/`](../../src/models/cua_s1/native/)) serves the `text` adapter like the reference worker in [`text.md`](text.md), with the Qwen3.5-4B forward pass on the CUDA kernels in [`src/backends/cuda/qwen3_5/`](../../src/backends/cuda/qwen3_5/) and no Python or PyTorch. It needs an NVIDIA GPU with compute capability 8.0 or newer; the unmerged precision profile was checked on the cumulative A800 80 GB (sm_80) build with CUDA 13.0; device qualification of this independent #118-based branch remains pending. The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. +The native worker ([`src/models/cua_s1/native/`](../../src/models/cua_s1/native/)) serves the `text` adapter like the reference worker in [`text.md`](text.md), with the Qwen3.5-4B forward pass on the CUDA kernels in [`src/backends/cuda/qwen3_5/`](../../src/backends/cuda/qwen3_5/) and no Python or PyTorch. It needs an NVIDIA GPU with compute capability 8.0 or newer; the unmerged precision profile was checked on the cumulative A800 80 GB (sm_80) build with CUDA 13.0; source-pinned qualification of the independent #118-based branch is tracked in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. Run the commands from the repository root. Pass your GPU's compute capability to `build.sh` (89 for Ada, 80 for A100, 90 for H100); the worker finds `libqwen3_5_cuda.so` next to its executable, or at `CUA_S1_CUDA_LIB`: From 47fe6b2c6ff34ead2b12d0469d4b382a1e3061fa Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Wed, 14 Oct 2026 11:11:48 +0800 Subject: [PATCH 3/7] docs(cua-s1): clarify merged readout and cuBLAS upgrade impact --- recipe/cua_s1/native.md | 6 +++++- src/backends/cuda/qwen3_5/README.md | 5 ++++- src/models/cua_s1/README.md | 6 +++++- 3 files changed, 14 insertions(+), 3 deletions(-) diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index 1da0660e..959a6755 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -41,7 +41,11 @@ markers or altered adapter bytes fail before CUDA loading. Prefix continuation is unsupported in this precision mode. Existing `cua-s1-text-merged/1` exports remain loadable with the earlier merged -execution path. Merging the adapter, TF32/BF16 GDN intermediates, and different +backbone execution path. PR #143 changes their letter-logit readout too: it now +rounds to BF16 instead of retaining FP32, so close or tied choices can change +after upgrading. No merged-path corpus A/B run qualifies that behavior change; +the A800 parity results cover the unmerged precision profile only. +Merging the adapter, TF32/BF16 GDN intermediates, and different GEMM reduction policies can change decisions on close or tied options. Re-export with the command above to use the precision profile. Rebuild worker and library together: this branch retains #118's ABI 6, but the precision profile requires the new optional diff --git a/src/backends/cuda/qwen3_5/README.md b/src/backends/cuda/qwen3_5/README.md index b2a710cf..d625faf6 100644 --- a/src/backends/cuda/qwen3_5/README.md +++ b/src/backends/cuda/qwen3_5/README.md @@ -62,7 +62,10 @@ starts with zero state; prefix continuation is rejected by the model adapter. The context drains its last stream before destruction, and is retired before that stream. Rebuild the library for this optional group; ABI remains 6 on this #118-based branch; the historical cumulative validation build used ABI 8. -Builds now link both cuBLAS and cuBLASLt. +Builds now link both cuBLAS and cuBLASLt. Every consumer of the rebuilt shared +library must provide both runtime libraries (`libcublas.so` and +`libcublasLt.so`), even when it does not enable the optional FP32 GDN group. +An existing cuBLASLt-only deployment must add cuBLAS before upgrading. `tests/qwen3_5/test_gdn_fp32.py` compares against the pinned Transformers 5.17 no-FLA implementation. On the cumulative A800 sm_80/CUDA 13.0/Torch 2.14 build, all 14 cases passed diff --git a/src/models/cua_s1/README.md b/src/models/cua_s1/README.md index c24f3ee0..f36eceed 100644 --- a/src/models/cua_s1/README.md +++ b/src/models/cua_s1/README.md @@ -133,7 +133,11 @@ This profile targets the pinned no-FLA reference described above. Installing FLA changes Transformers' GDN implementation; comparisons must record which implementation actually ran. A different Torch version or GPU may choose other reduction orders, so qualification remains tied to the recorded environment. -The old merged export remains an approximate execution mode. The precision +The old merged export remains an approximate execution mode. PR #143 also +changes that mode's letter-logit readout from FP32 to BF16 rounding, which can +change close or tied choices after upgrading. The A800 unmerged-profile parity +results do not qualify this merged-path behavior change; no merged-path corpus +A/B run was performed. The precision profile does not change the screenshot adapter's execution mode. The historical cumulative A800 investigation and frozen-corpus results are recorded in From 22b45e8889382f3ed366f71cfc7afae0112773a5 Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Thu, 15 Oct 2026 09:40:47 +0800 Subject: [PATCH 4/7] docs(cua-s1): separate current ABI 8 from historical GPU evidence --- .../cua-s1-text-parity-20261013/README.md | 34 +++++++++++-------- recipe/cua_s1/native.md | 21 ++++++++---- src/backends/cuda/qwen3_5/README.md | 6 ++-- 3 files changed, 38 insertions(+), 23 deletions(-) diff --git a/docs/benchmarks/cua-s1-text-parity-20261013/README.md b/docs/benchmarks/cua-s1-text-parity-20261013/README.md index dd485921..62153d46 100644 --- a/docs/benchmarks/cua-s1-text-parity-20261013/README.md +++ b/docs/benchmarks/cua-s1-text-parity-20261013/README.md @@ -9,25 +9,29 @@ it does not establish that all five original L20X cases have been reproduced. ## PR dependencies and evidence provenance -This independent change is stacked on [#118](https://github.com/ThinkFlowLab/system1-omni/pull/118) -(`codex/jemm-native`, `ed9f9d6eb566d226e2441c5e9d84187abe1e4d63`) for the -unmerged FP32 language adapter loader and optional reference norm/GEMM/attention -operations. #118 is itself stacked on [#117](https://github.com/ThinkFlowLab/system1-omni/pull/117). -The frozen corpus comes from #141; merging #141 is not a runtime dependency. -Decider execution modes and Cua screenshot cache/resident-feature PRs are outside -this change's dependency chain and diff. +The independently qualified source `27f77e49868d84b091aa804f02ccb4d3474a9a00` +was stacked on [#118](https://github.com/ThinkFlowLab/system1-omni/pull/118) +at `ed9f9d6eb566d226e2441c5e9d84187abe1e4d63` for the unmerged FP32 language +adapter loader and optional reference norm/GEMM/attention operations. +The 2026-10-15 parent synchronization used #118 at +`3be46fa270cb74c33d4db4ffd761405d0522018a`, which inherits +[#117](https://github.com/ThinkFlowLab/system1-omni/pull/117) and its merged +resident-feature increment #135. Rebuild the library and all consumers together +for ABI 8. This parent synchronization has CPU/CI coverage; current integrated +GPU qualification remains unverified. Decider modes and the separate #127 cache +increment are not included. The frozen corpus comes from #141; merging #141 is +not a runtime dependency. The GPU results below were measured on cumulative candidate `76af07996f60905082b8d48199d4a7003f33f4f6`, whose parent is `18e1a5a6f68b9d11d3e11c7ea7b939af87272f53`. That integrated build used ABI 8. -The independently rebased branch retains #118's ABI 6 and adds only optional -FP32 GDN symbols. Rebuild the library with the worker; no cuDNN is needed for -this text path. Results from the cumulative build are historical evidence for -the pinned implementation. Fresh independent-branch qualification, including -exact source/library/binary hashes, repeated HTTP results and self-review, is -recorded separately in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). -Use each run's source pin; do not substitute the historical integrated result -for a newer branch's validation. +The separate independent qualification at `27f77e49` used ABI 6 with additive +optional FP32 GDN symbols. No cuDNN is needed for this text path. Its exact +source/library/binary hashes, repeated HTTP results and self-review are recorded +in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). +Both campaigns remain historical, source-pinned evidence. ABI 8 in the cumulative +campaign does not establish qualification of the current ABI-8 branch; use each +run's source pin rather than substituting either result for a newer build. ## Frozen inputs and environment diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index f2b3e908..3fae7add 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -1,6 +1,13 @@ # Cua-S1 4B 0.2 native text worker -The native worker ([`src/models/cua_s1/native/`](../../src/models/cua_s1/native/)) serves the `text` adapter like the reference worker in [`text.md`](text.md), with the Qwen3.5-4B forward pass on the CUDA kernels in [`src/backends/cuda/qwen3_5/`](../../src/backends/cuda/qwen3_5/) and no Python or PyTorch. It needs an NVIDIA GPU with compute capability 8.0 or newer; the unmerged precision profile was checked on the cumulative A800 80 GB (sm_80) build with CUDA 13.0; source-pinned qualification of the independent #118-based branch is tracked in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. +The native worker ([`src/models/cua_s1/native/`](../../src/models/cua_s1/native/)) serves the `text` adapter like the reference worker in [`text.md`](text.md), with the Qwen3.5-4B forward pass on the CUDA kernels in [`src/backends/cuda/qwen3_5/`](../../src/backends/cuda/qwen3_5/) and no Python or PyTorch. It needs an NVIDIA GPU with compute capability 8.0 or newer. +The unmerged precision profile has historical A800 80 GB (sm_80), CUDA 13.0 +evidence for a cumulative ABI-8 build and the independent ABI-6 source +`27f77e49868d84b091aa804f02ccb4d3474a9a00`, linked in +[PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). +The current branch uses ABI 8 after synchronizing with its updated parent; +neither historical campaign qualifies this integrated source. +The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. Run the commands from the repository root. Pass your GPU's compute capability to `build.sh` (89 for Ada, 80 for A100, 90 for H100); the worker finds `libqwen3_5_cuda.so` next to its executable, or at `CUA_S1_CUDA_LIB`: @@ -71,8 +78,10 @@ For the screenshot worker, [the vision Graph recipe](native_vision_graph.md) documents `CUA_S1_VISION_GRAPH`, retained four-grid FIFO scratch, memory/lifetime limits and the real CUDA capture-failure wrapper build and test commands. -The ABI-8 parent synchronization preserves the unmerged precision profile, but -the recorded independent-head GPU campaign below belongs to its pinned earlier -source. It is historical evidence, not a fresh GPU qualification of this -integrated head. Rebuild consumers together and rerun the pinned device checks -before extending qualification to this build or another device. +The ABI-8 parent synchronization preserves the unmerged precision profile. +The historical independent-head GPU campaign is linked in +[PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143); the +[numerical investigation](../../docs/benchmarks/cua-s1-text-parity-20261013/README.md) +records the separate cumulative campaign. Rebuild consumers together and rerun +the pinned device checks before extending qualification to this build or +another device. diff --git a/src/backends/cuda/qwen3_5/README.md b/src/backends/cuda/qwen3_5/README.md index 817dfd33..13d48921 100644 --- a/src/backends/cuda/qwen3_5/README.md +++ b/src/backends/cuda/qwen3_5/README.md @@ -65,8 +65,10 @@ workspace; the caller supplies shape-dependent FP32 buffers. It implements Torch 2.14 cumulative-decay association, and a final BF16 output cast. Each call starts with zero state; prefix continuation is rejected by the model adapter. The context drains its last stream before destruction, and is retired before -that stream. Rebuild the library for this optional group; ABI remains 6 on this #118-based branch; the historical cumulative validation -build used ABI 8. +that stream. Rebuild the library and all consuming workers together for the +current ABI 8, including this optional group. The independent A800 qualification +at `27f77e49868d84b091aa804f02ccb4d3474a9a00` used ABI 6; the earlier cumulative +campaign used ABI 8. Neither campaign qualifies the current integrated source. Builds now link both cuBLAS and cuBLASLt. Every consumer of the rebuilt shared library must provide both runtime libraries (`libcublas.so` and `libcublasLt.so`), even when it does not enable the optional FP32 GDN group. From f212c56f23ac750ebfa38f4c72e9fcdc914bf1cd Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Thu, 15 Oct 2026 10:58:36 +0800 Subject: [PATCH 5/7] fix(cua-s1): preserve legacy merged readout on upgrade --- docs/architecture.md | 6 +- .../cua-s1-text-parity-20261013/README.md | 50 +++- recipe/cua_s1/native.md | 47 +++- src/models/cua_s1/README.md | 22 +- src/models/cua_s1/native/src/engine.rs | 4 + src/models/cua_s1/native/src/executor.rs | 24 +- tests/cua_s1/fixtures/LICENSE.jevbench | 21 ++ tests/cua_s1/fixtures/README.md | 31 +++ .../fixtures/legacy_merged_readout.json | 228 ++++++++++++++++++ tests/cua_s1/head.rs | 121 +++++++++- tests/cua_s1/legacy_merged_gpu.rs | 110 +++++++++ 11 files changed, 632 insertions(+), 32 deletions(-) create mode 100644 tests/cua_s1/fixtures/LICENSE.jevbench create mode 100644 tests/cua_s1/fixtures/README.md create mode 100644 tests/cua_s1/fixtures/legacy_merged_readout.json create mode 100644 tests/cua_s1/legacy_merged_gpu.rs diff --git a/docs/architecture.md b/docs/architecture.md index 433c5ddd..5677b13a 100644 --- a/docs/architecture.md +++ b/docs/architecture.md @@ -57,8 +57,10 @@ synchronizes and disposes a failed model before returning an inference error and reports unavailable health thereafter. Qwen executors own loaded models and CPU head weights with FP64 accumulation. -Cua-S1 rounds answer-letter logits to the tied `lm_head` BF16 output dtype before -FP32 softmax; learned scalar heads retain their FP32 rounding and bias order. Finishing checks output cardinality +Cua-S1's unmerged text profile rounds answer-letter logits to the tied `lm_head` +BF16 output dtype before FP32 softmax. Legacy merged text and screenshot +readouts retain FP32 logits; learned scalar heads retain their FP32 rounding +and bias order. Finishing checks output cardinality before reconstruction. HTTP validation, error status/body conventions, and real warmup before readiness remain model-specific and unchanged. diff --git a/docs/benchmarks/cua-s1-text-parity-20261013/README.md b/docs/benchmarks/cua-s1-text-parity-20261013/README.md index 62153d46..21a8b3e6 100644 --- a/docs/benchmarks/cua-s1-text-parity-20261013/README.md +++ b/docs/benchmarks/cua-s1-text-parity-20261013/README.md @@ -77,10 +77,58 @@ Token IDs matched on all 139 requests. The differences begin in execution: for tied options and removes the remaining readout differences. The fix selects the precision path using `cua-s1-text-unmerged/1`. Legacy merged -exports remain loadable. Other model adapters retain their existing numerical +exports remain loadable with their original FP32 readout; BF16 letter-logit +rounding is restricted to the unmerged text profile. Other model adapters retain their existing numerical mode. Prefix continuation is rejected in the new mode; eager is the default. The probability gate of 0.02 was not relaxed. +## Legacy merged upgrade regression + +The [2026-10-15 A800 ABI-8 evidence](https://github.com/Levius-Fubuki/system1-omni/releases/tag/cua-a800-abi8-20261015) +compares parent `5bdd625b4c27efb02f43a1f71150e1a7d7c0f071` with candidate +`7293c71cfb738504a30ff033ae3567f3368333c1`, using the same legacy merged artifact +and frozen 139-request corpus for two passes each. The candidate changed four +decisions, with maximum probability change 0.02441754937171936; correctness fell +from 97/139 to 96/139. Each changed decision collapsed the two leading options +into equal probabilities, then selected the earlier option in request order: + +| Request | Legacy choice | Rounded merged choice | Frozen task label | +| --- | --- | --- | --- | +| `original-routing-06-1` | `tools` | `general` | `general` | +| `hard-opus-a-temporal_numeric-07` | `sep_28` | `sep_24` | `sep_26` | +| `hard-sol-b-long_policy-05` | `delete_profile_omar_april_and_table` | `delete_everything` | `delete_profile_omar_april_and_table` | +| `hard-sol-b-temporal_numeric-03` | `9_credits` | `12_credits` | `9_credits` | + +The cause is the text executor's unconditional BF16 conversion of its FP32 dot +products. The model's merged loader still selects the old backbone arithmetic; +the unmerged reference-profile operations are opt-in. The compatibility fix +keeps the validated `TextExport` on each executor and applies BF16 readout only +to `Unmerged`. `Merged` retains the parent's FP64 accumulation followed by FP32 +readout. Softmax, option order, tie handling and screenshot execution are unchanged. +Restoring parent behavior reverses the candidate's one task-label improvement +and two regressions; the fourth case is wrong under both choices. Compatibility +is measured against the parent, independently of task labels. + +A CPU regression first failed with `[20.0, 20.0]` instead of +`[20.03125, 20.046875]`, then passed with the profile-specific readout. Tests also +cover both profiles' FP64 accumulation, nonfinite rejection, finite FP32 values +outside BF16 range, and the four requests' option order through real response +finishing with explicitly synthetic BF16 projection inputs. These synthetic +vectors are not captured checkpoint activations. The maintained four-request +device test uses the actual frozen inputs and parent responses; run it as in the +[native recipe](../../../recipe/cua_s1/native.md). + +Fresh GPU execution of the compatibility fix is **pending**: the earlier server +was shut down after preserving the published evidence. In particular, restored +97/139 correctness and zero upgrade flips have not been measured for this new +source. Requalify the full corpus on pinned A800: two passes each for legacy +parent and repaired merged eager/Graph, requiring zero parent/candidate decision +changes and probability error at most 0.02; retain exact-byte differences too. +Repeat unmerged eager/Graph/reference checks at the original 0.02/zero-flip gates. +The previous candidate's unmerged results were 97/139 in all three modes, maximum +probability difference 5.960464477539063e-8 and zero flips. They remain tied to +`7293c71c`; they are not a new-device pass for this follow-up. + ## Results | Real HTTP mode | Passes × requests | Flips vs reference | Maximum probability difference | Correct | diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index 3fae7add..e4421d42 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -1,12 +1,12 @@ # Cua-S1 4B 0.2 native text worker The native worker ([`src/models/cua_s1/native/`](../../src/models/cua_s1/native/)) serves the `text` adapter like the reference worker in [`text.md`](text.md), with the Qwen3.5-4B forward pass on the CUDA kernels in [`src/backends/cuda/qwen3_5/`](../../src/backends/cuda/qwen3_5/) and no Python or PyTorch. It needs an NVIDIA GPU with compute capability 8.0 or newer. -The unmerged precision profile has historical A800 80 GB (sm_80), CUDA 13.0 -evidence for a cumulative ABI-8 build and the independent ABI-6 source -`27f77e49868d84b091aa804f02ccb4d3474a9a00`, linked in +The unmerged precision profile has source-pinned A800 80 GB (sm_80), CUDA 13.0 +evidence at ABI-8 source `7293c71cfb738504a30ff033ae3567f3368333c1`, plus earlier +cumulative ABI-8 and independent ABI-6 campaigns linked in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). -The current branch uses ABI 8 after synchronizing with its updated parent; -neither historical campaign qualifies this integrated source. +The current branch uses ABI 8. The later legacy-readout compatibility fix has +CPU regression coverage; fresh GPU validation of that follow-up is pending. The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. Run the commands from the repository root. Pass your GPU's compute capability to `build.sh` (89 for Ada, 80 for A100, 90 for H100); the worker finds `libqwen3_5_cuda.so` next to its executable, or at `CUA_S1_CUDA_LIB`: @@ -48,10 +48,13 @@ markers or altered adapter bytes fail before CUDA loading. Prefix continuation is unsupported in this precision mode. Existing `cua-s1-text-merged/1` exports remain loadable with the earlier merged -backbone execution path. PR #143 changes their letter-logit readout too: it now -rounds to BF16 instead of retaining FP32, so close or tied choices can change -after upgrading. No merged-path corpus A/B run qualifies that behavior change; -the A800 parity results cover the unmerged precision profile only. +backbone execution path and retain FP32 letter logits after FP64 accumulation. +The BF16 readout boundary is restricted to the validated unmerged export profile, +not a request flag or process-global setting. An earlier #143 build applied it to +merged exports too, creating four decision changes and reducing frozen-corpus +correctness from 97/139 to 96/139. The follow-up restores legacy readout policy; +that restoration still needs fresh device validation. The A800 unmerged parity +results are scoped to their recorded source and artifact. Merging the adapter, TF32/BF16 GDN intermediates, and different GEMM reduction policies can change decisions on close or tied options. Re-export with the command above to use the precision profile. Rebuild worker and library @@ -74,14 +77,36 @@ CUA_S1_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \ cargo test --release -p omni-qwen3-5-native --test kernels -- --ignored ``` +The maintained legacy-upgrade regression uses the four frozen requests and +parent-baseline responses in `tests/cua_s1/fixtures/legacy_merged_readout.json`. +Those responses are compatibility oracles, not task ground truth. With the same +pinned A800 merged artifact, run each Graph setting in a separate process: + +```sh +for graph in 0 1; do + CUA_S1_GRAPH=$graph NVIDIA_TF32_OVERRIDE=0 \ + CUA_S1_MERGED_TEST_MODEL=$PWD/weights/cua-s1-4b-0.2-text-merged \ + CUA_S1_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \ + cargo test --release --locked -p omni-cua-s1-native --lib \ + engine::legacy_merged_tests::legacy_merged_upgrade_preserves_four_decisions \ + -- --ignored --exact --test-threads=1 --nocapture +done +``` + +This selects one device test per process, running all four requests twice through +prepare/execute/finish. It requires unchanged baseline decisions, probability +error at most 0.02, identity/usage/option order and exact repeat bytes. It does +not substitute for the full 139-request legacy upgrade and unmerged-reference +comparisons. See the [four-case diagnosis and retest scope](../../docs/benchmarks/cua-s1-text-parity-20261013/README.md#legacy-merged-upgrade-regression). + For the screenshot worker, [the vision Graph recipe](native_vision_graph.md) documents `CUA_S1_VISION_GRAPH`, retained four-grid FIFO scratch, memory/lifetime limits and the real CUDA capture-failure wrapper build and test commands. The ABI-8 parent synchronization preserves the unmerged precision profile. -The historical independent-head GPU campaign is linked in +The source-pinned GPU campaigns are linked in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143); the [numerical investigation](../../docs/benchmarks/cua-s1-text-parity-20261013/README.md) -records the separate cumulative campaign. Rebuild consumers together and rerun +records their scopes and the legacy-readout follow-up. Rebuild consumers together and rerun the pinned device checks before extending qualification to this build or another device. diff --git a/src/models/cua_s1/README.md b/src/models/cua_s1/README.md index bb41afdf..08e835bd 100644 --- a/src/models/cua_s1/README.md +++ b/src/models/cua_s1/README.md @@ -10,8 +10,10 @@ The native worker follows the [architecture contracts](../../../docs/architectur with independent [processing](native/src/processing.rs) and [executor](native/src/executor.rs) modules. Preparation returns token IDs and an option count per question, plus a response context for identity/order and usage. -The executor projects the letter rows, rounds to the BF16 `lm_head` output dtype, -then returns FP32 letter logits; response finishing owns per-question +The executor projects the letter rows with FP64 accumulation and returns FP32 +letter logits. Only the unmerged text profile first rounds to the BF16 `lm_head` +output dtype; legacy merged text and screenshot readouts retain FP32. The text +export policy is validated once and owned by each executor. Response finishing owns per-question softmax and the choice/confidence interpretation. The engine owns a [shared serial scheduler](../../runtime/README.md) that admits each question's forward before blocking dispatch. Execution remains one prompt per question @@ -133,12 +135,16 @@ This profile targets the pinned no-FLA reference described above. Installing FLA changes Transformers' GDN implementation; comparisons must record which implementation actually ran. A different Torch version or GPU may choose other reduction orders, so qualification remains tied to the recorded environment. -The old merged export remains an approximate execution mode. PR #143 also -changes that mode's letter-logit readout from FP32 to BF16 rounding, which can -change close or tied choices after upgrading. The A800 unmerged-profile parity -results do not qualify this merged-path behavior change; no merged-path corpus -A/B run was performed. The precision -profile does not change the screenshot adapter's execution mode. +The old merged export remains an approximate execution mode. Its legacy +FP64-accumulated, FP32 letter-logit readout is preserved; BF16 rounding is selected +only by `cua-s1-text-unmerged/1`. Applying that rounding to merged exports at +`7293c71c` created four ordered-tie decision changes in a same-artifact A800 +upgrade comparison. The compatibility fix is covered by CPU red/green projection +and response tests; a frozen four-request GPU regression is also provided. +See the [numerical report](../../../docs/benchmarks/cua-s1-text-parity-20261013/README.md#legacy-merged-upgrade-regression) +for the four cases, source pins and device-validation limits. Restoring merged +compatibility does not make merged weights equivalent to the unmerged reference. +The precision profile does not change the screenshot adapter's execution mode. The historical cumulative A800 investigation and frozen-corpus results are recorded in [the numerical report](../../../docs/benchmarks/cua-s1-text-parity-20261013/README.md). diff --git a/src/models/cua_s1/native/src/engine.rs b/src/models/cua_s1/native/src/engine.rs index 5c0b434e..86dcbcec 100644 --- a/src/models/cua_s1/native/src/engine.rs +++ b/src/models/cua_s1/native/src/engine.rs @@ -40,3 +40,7 @@ impl Engine { #[cfg(test)] #[path = "../../../../../tests/cua_s1/text_export.rs"] mod tests; + +#[cfg(test)] +#[path = "../../../../../tests/cua_s1/legacy_merged_gpu.rs"] +mod legacy_merged_tests; diff --git a/src/models/cua_s1/native/src/executor.rs b/src/models/cua_s1/native/src/executor.rs index 6bcbdba2..2acba636 100644 --- a/src/models/cua_s1/native/src/executor.rs +++ b/src/models/cua_s1/native/src/executor.rs @@ -22,6 +22,7 @@ pub struct Input { pub struct Executor { model: Arc>, letters: Vec, + export: TextExport, } impl Executor { @@ -46,6 +47,7 @@ impl Executor { Ok(Self { model: Arc::new(Mutex::new(model)), letters, + export, }) } @@ -67,13 +69,23 @@ impl Executor { .forward(&input.ids) }) .await?; - rows.push(letter_logits(&self.letters, &last, input.n_options)?); + rows.push(letter_logits( + &self.letters, + &last, + input.n_options, + self.export, + )?); } Ok(rows) } } -fn letter_logits(letters: &[f32], last: &[f32], n_options: usize) -> Result> { +fn letter_logits( + letters: &[f32], + last: &[f32], + n_options: usize, + export: TextExport, +) -> Result> { let logits: Vec = letters .chunks_exact(last.len()) .take(n_options) @@ -83,8 +95,12 @@ fn letter_logits(letters: &[f32], last: &[f32], n_options: usize) -> Result() as f32; - // The tied lm_head emits BF16 before reference option selection/softmax. - half::bf16::from_f32(dot).to_f32() + match export { + // Preserve the original merged worker's FP32 readout on upgrade. + TextExport::Merged => dot, + // The unmerged reference profile's tied lm_head emits BF16. + TextExport::Unmerged => half::bf16::from_f32(dot).to_f32(), + } }) .collect(); // Abort before submitting a later prompt, as the former per-question softmax did. diff --git a/tests/cua_s1/fixtures/LICENSE.jevbench b/tests/cua_s1/fixtures/LICENSE.jevbench new file mode 100644 index 00000000..ebdd5738 --- /dev/null +++ b/tests/cua_s1/fixtures/LICENSE.jevbench @@ -0,0 +1,21 @@ +MIT License + +Copyright (c) 2026 Florian Standhartinger and contributors + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. diff --git a/tests/cua_s1/fixtures/README.md b/tests/cua_s1/fixtures/README.md new file mode 100644 index 00000000..a22a9f1f --- /dev/null +++ b/tests/cua_s1/fixtures/README.md @@ -0,0 +1,31 @@ +# Legacy merged readout regression oracle + +`legacy_merged_readout.json` keeps four requests from +[`benchmarks/e2e/choice.jsonl` at `86926064`](https://github.com/ThinkFlowLab/system1-omni/blob/869260649f7dd5b56cac4d91eaae9087f2e985ee/benchmarks/e2e/choice.jsonl). +The original corpus SHA-256 is recorded in the fixture. Its JevBench-derived +request material is MIT licensed; the original copyright and full terms are +preserved in [`LICENSE.jevbench`](LICENSE.jevbench). + +The baseline responses come from A800/ABI-8 parent `5bdd625b`, using the same +legacy merged artifact as candidate `7293c71c`. The exact source, model, artifact +and environment hashes are in the +[public raw-evidence archive](https://github.com/Levius-Fubuki/system1-omni/releases/tag/cua-a800-abi8-20261015): +`evidence/pr143/merged-before/pass0.jsonl` supplies each `expected_response`, +`evidence/pr143/merged-after/pass0.jsonl` supplies `rounded_choice`, and +`evidence/pr143/native-eager/manifest.jsonl` supplies the requests. Both measured +passes produced identical response bytes. Only the four cases whose decisions +changed are retained here as a maintained compatibility oracle. Expected +responses preserve the parent behavior; they are **not task ground truth**. + +Consumers: + +- `tests/cua_s1/head.rs` uses the actual option order with explicitly synthetic + BF16 projection vectors to test the CPU readout/response boundary in both modes. +- `tests/cua_s1/legacy_merged_gpu.rs` runs the actual frozen requests twice through + the legacy model after verifying every recorded artifact file's size and + SHA-256, requiring baseline decisions, probability error at + most 0.02, identity/usage/order and exact repeat bytes. Run Graph 0/1 in separate + processes as documented in `recipe/cua_s1/native.md`. + +The GPU oracle is specific to the recorded A800 artifact/environment. Compilation +or a passing synthetic CPU test does not establish a new real-model GPU pass. diff --git a/tests/cua_s1/fixtures/legacy_merged_readout.json b/tests/cua_s1/fixtures/legacy_merged_readout.json new file mode 100644 index 00000000..3e00e482 --- /dev/null +++ b/tests/cua_s1/fixtures/legacy_merged_readout.json @@ -0,0 +1,228 @@ +{ + "schema": "cua-s1-legacy-merged-readout-regression/1", + "purpose": "Compatibility oracle for the four merged-export upgrade changes; expected responses preserve the legacy baseline, not task ground truth.", + "baseline_source": "5bdd625b4c27efb02f43a1f71150e1a7d7c0f071", + "rejected_source": "7293c71cfb738504a30ff033ae3567f3368333c1", + "evidence": "https://github.com/Levius-Fubuki/system1-omni/releases/tag/cua-a800-abi8-20261015", + "corpus_source": "869260649f7dd5b56cac4d91eaae9087f2e985ee", + "corpus_sha256": "8f4b912e5dfd8706992e90f1b87bb07a78a8339e5a9fad3ef64b3b04771f9daf", + "environment": { + "gpu": "A800 sm80", + "abi": 8, + "base_revision": "851bf6e806efd8d0a36b00ddf55e13ccb7b8cd0a", + "adapter_revision": "16818868b0cc7813808aae4e87b417657046ab79", + "torch": "2.14.0+cu130", + "transformers": "5.17.0", + "tf32_override": "0", + "export_format": "cua-s1-text-merged/1" + }, + "cases": [ + { + "id": "original-routing-06-1", + "request": { + "state": "Give me three creative pet-dragon names.", + "questions": { + "q": { + "criteria": { + "coding": "Self-contained code writing or explanation without repository operations", + "coding_agent": "Inspect/edit repository files or run tests", + "document": "Answer from a supplied document", + "general": "None of the specialist categories", + "math": "Self-contained calculation or proof", + "tools": "Carry out an external service action" + }, + "instructions": "Choose the specialist needed for the request. File edits with test execution use coding_agent, even if code-related.", + "type": "choice" + } + }, + "model": "cua-s1-4b-0.2" + }, + "expected_response": { + "model": "cua-ai/cua-s1-4b-0.2@16818868b0cc7813808aae4e87b417657046ab79:text", + "answers": { + "q": { + "type": "choice", + "choice": "tools", + "probabilities": { + "coding": 0.12521012127399445, + "coding_agent": 0.2372714877128601, + "document": 0.0032456598710268736, + "general": 0.3135629892349243, + "math": 0.0015466281911358237, + "tools": 0.31916311383247375 + }, + "confidence": 0.24194648916375006 + } + }, + "usage": { + "input_tokens": 233, + "output_tokens": 0 + } + }, + "rounded_choice": "general" + }, + { + "id": "hard-opus-a-temporal_numeric-07", + "request": { + "state": "NIMBUSGRID CLOUD — BUDGET ALERT CONFIGURATION AND DAILY COST EXPORT (ACCOUNT ANALYTICS-PROD, SEPTEMBER 2026, USD)\n\nBudget BGT-ANALYTICS-09: monthly budget USD 12,000.00.\nAlert rule A1: \"Fires once, at the end of the first day on which the month-to-date net cost is at or above 80% of the monthly budget.\"\nDefinitions (NimbusGrid billing guide): \"Net cost\" = usage charges minus credits applied, excluding taxes. Taxes are shown for information and are never counted toward budgets. Credits are counted on the day they are applied.\nNote: a promotional credit of USD 600.00 was applied on 18 September (migration incentive).\n\nDaily export:\nSep 01 | usage 402.10 | credits 0.00 | tax (8%) 32.17\nSep 02 | usage 389.55 | credits 0.00 | tax (8%) 31.16\nSep 03 | usage 415.00 | credits 0.00 | tax (8%) 33.20\nSep 04 | usage 398.40 | credits 0.00 | tax (8%) 31.87\nSep 05 | usage 377.25 | credits 0.00 | tax (8%) 30.18\nSep 06 | usage 361.80 | credits 0.00 | tax (8%) 28.94\nSep 07 | usage 420.65 | credits 0.00 | tax (8%) 33.65\nSep 08 | usage 433.10 | credits 0.00 | tax (8%) 34.65\nSep 09 | usage 409.95 | credits 0.00 | tax (8%) 32.80\nSep 10 | usage 395.00 | credits 0.00 | tax (8%) 31.60\nSep 11 | usage 388.70 | credits 0.00 | tax (8%) 31.10\nSep 12 | usage 402.35 | credits 0.00 | tax (8%) 32.19\nSep 13 | usage 371.20 | credits 0.00 | tax (8%) 29.70\nSep 14 | usage 366.45 | credits 0.00 | tax (8%) 29.32\nSep 15 | usage 418.90 | credits 0.00 | tax (8%) 33.51\nSep 16 | usage 427.30 | credits 0.00 | tax (8%) 34.18\nSep 17 | usage 411.05 | credits 0.00 | tax (8%) 32.88\nSep 18 | usage 399.60 | credits 600.00 | tax (8%) 31.97\nSep 19 | usage 404.25 | credits 0.00 | tax (8%) 32.34\nSep 20 | usage 392.80 | credits 0.00 | tax (8%) 31.42\nSep 21 | usage 386.15 | credits 0.00 | tax (8%) 30.89\nSep 22 | usage 375.40 | credits 0.00 | tax (8%) 30.03\nSep 23 | usage 410.00 | credits 0.00 | tax (8%) 32.80\nSep 24 | usage 421.75 | credits 0.00 | tax (8%) 33.74\nSep 25 | usage 398.00 | credits 0.00 | tax (8%) 31.84\nSep 26 | usage 223.30 | credits 0.00 | tax (8%) 17.86\nSep 27 | usage 380.00 | credits 0.00 | tax (8%) 30.40\nSep 28 | usage 412.60 | credits 0.00 | tax (8%) 33.01\nSep 29 | usage 405.10 | credits 0.00 | tax (8%) 32.41\nSep 30 | usage 399.85 | credits 0.00 | tax (8%) 31.99\n\nFinOps channel, 24 September: \"With tax we're already over 9,600 — the alert should have fired today. Something's broken?\"\n", + "questions": { + "q": { + "criteria": { + "sep_24": "At the end of 24 September.", + "sep_25": "At the end of 25 September.", + "sep_26": "At the end of 26 September.", + "sep_27": "At the end of 27 September.", + "sep_28": "At the end of 28 September." + }, + "instructions": "On which day does alert rule A1 fire?", + "type": "choice" + } + }, + "model": "cua-s1-4b-0.2" + }, + "expected_response": { + "model": "cua-ai/cua-s1-4b-0.2@16818868b0cc7813808aae4e87b417657046ab79:text", + "answers": { + "q": { + "type": "choice", + "choice": "sep_28", + "probabilities": { + "sep_24": 0.3386138677597046, + "sep_25": 0.07059328258037567, + "sep_26": 0.07448268681764603, + "sep_27": 0.16069446504116058, + "sep_28": 0.3556157052516937 + }, + "confidence": 0.12471221294469215 + } + }, + "usage": { + "input_tokens": 1381, + "output_tokens": 0 + } + }, + "rounded_choice": "sep_24" + }, + { + "id": "hard-sol-b-long_policy-05", + "request": { + "state": "RECORDS DELETION AND LEGAL HOLD HANDBOOK\n1. Rule hierarchy. Legal holds override ordinary deletion schedules. A later written release controls the named scope only. The privacy deletion process does not authorize destruction barred by a hold. The clarification memorandum at the end is binding.\n\n2. Definitions. \"Responsive record\" means a record within the custodians, systems, subjects, and time range stated by a hold. Exact duplicates are each responsive when stored in separate systems unless the hold notice permits deduplication. \"Derived analytics\" are records produced from source data; they are not source records. \"Verified deletion\" means deletion from active systems and expiration from ordinary backups under the stated schedule, with a logged confirmation. Cryptographic erasure counts only if the relevant keys exclusively protect the named records.\n\n3. Customer requests. A verified identity request is due within 30 calendar days. Data needed for contract performance may be retained while necessary. Data under legal hold is retained even after the ordinary purpose ends, and the customer receives the standard limitation notice. Non-held data in the same account remains subject to deletion.\n\n4. Backup process. Active copies are deleted by the due date. Immutable backups expire on a 35-day rotation; restoration procedures must reapply deletion markers. The completion record may therefore distinguish active deletion from final backup expiry.\n\n5. Holds. Hold H-17 covers support messages by custodians June and Omar concerning Project Cedar from 1 January through 31 March 2026, including attachments and exports. Hold H-22 covers all billing records for account 448 from 1 February onward. Neither permits deduplication. A legal email may release a hold only if it identifies the hold number and scope.\n\n6. Case file. Account 448 submitted a verified deletion request on 6 April. Contract service ended 5 April. The account contains profile data, invoices dated February and March, support messages with June about Project Cedar dated 20 March, unrelated support messages with Omar dated 2 April, and a dashboard table derived by counting Cedar messages by week. The table contains only week and count. On 10 April Legal wrote: \"Release H-17 for Omar's records only; June remains held. H-22 remains fully active.\" Privacy proposes deleting the profile, every support message, the derived table, and all invoices from active storage on 20 April, leaving deletion markers for backups.\n\n7. Scope application. A release for one custodian does not release another. Records outside a hold's date range are not held merely because a named custodian created them. Derived analytics are covered only if expressly included or if they independently meet the hold's stated record description.\n\n8. Decision options. The deletion team must choose the option that can be executed on 20 April without violating a hold. Ordinary backup expiry may follow later and does not prevent timely active-system action.\n\nCLARIFICATION MEMORANDUM, 12 April. H-22's term \"billing records\" includes invoices and payment ledger entries but not profile contact fields used only for login. H-17's term \"support messages\" does not include aggregate derived tables that contain no message text, attachment, customer identifier, or custodian identifier. This memorandum clarifies existing scope and does not release either hold.\n\nADMINISTRATIVE AND REFERENCE MATERIAL (part of the document):\n\nAppendix 1. identity verification. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-01 is an index locator, not a priority ranking or outcome code.\n\nAppendix 2. backup inventories. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-02 is an index locator, not a priority ranking or outcome code.\n\nAppendix 3. custodian interviews. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-03 is an index locator, not a priority ranking or outcome code.\n\nAppendix 4. deletion markers. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-04 is an index locator, not a priority ranking or outcome code.\n\nAppendix 5. request correspondence. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-05 is an index locator, not a priority ranking or outcome code.\n\nAppendix 6. analytics catalog. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-06 is an index locator, not a priority ranking or outcome code.\n\nAppendix 7. restoration drills. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-07 is an index locator, not a priority ranking or outcome code.\n\nAppendix 8. hold training. The organization keeps a contemporaneous record for this topic. Routine records identify the owner, the date recorded, the system of record, and any later correction. A draft note is not an approval, silence is not consent, and a copied recipient is not treated as the decision maker unless a controlling section expressly says otherwise. Staff should preserve ordinary supporting material and use the defined business calendar. This appendix supplies administration only: it does not enlarge eligibility, erase a deadline, reverse a precedence rule, or create an exception to a requirement stated in the operative sections. Where examples use names or dates, they illustrate recordkeeping and do not amend the operative case. Historical versions remain searchable for audit purposes, but a superseded version has no prospective effect. Questions are logged with enough context for a reviewer to reproduce the decision without relying on memory. The records team checks completeness monthly and reports aggregate defects; it does not re-decide the underlying case. Reference code REC-08 is an index locator, not a priority ranking or outcome code.", + "questions": { + "q": { + "criteria": { + "delete_everything": "The case results in delete everything under all controlling provisions.", + "delete_nothing": "The case results in delete nothing under all controlling provisions.", + "delete_only_profile": "The case results in delete only profile under all controlling provisions.", + "delete_profile_omar_april_and_table": "The case results in delete profile omar april and table under all controlling provisions." + }, + "instructions": "Apply the controlling document and select the single authorized outcome for the case file.", + "type": "choice" + } + }, + "model": "cua-s1-4b-0.2" + }, + "expected_response": { + "model": "cua-ai/cua-s1-4b-0.2@16818868b0cc7813808aae4e87b417657046ab79:text", + "answers": { + "q": { + "type": "choice", + "choice": "delete_profile_omar_april_and_table", + "probabilities": { + "delete_everything": 0.3959157466888428, + "delete_nothing": 0.1504596620798111, + "delete_only_profile": 0.014968502335250378, + "delete_profile_omar_april_and_table": 0.4386560916900635 + }, + "confidence": 0.22369902796665087 + } + }, + "usage": { + "input_tokens": 2630, + "output_tokens": 0 + } + }, + "rounded_choice": "delete_everything" + }, + { + "id": "hard-sol-b-temporal_numeric-03", + "request": { + "state": { + "event": "Classify the earned service credit.", + "rules": "An annual service credit is prorated by active calendar days in a leap year, including both activation and cancellation dates. The annual amount is 1,098 credits for 2028, which has 366 days. The service activated February 28 and cancelled March 2. A suspension covering February 29 removes that day from active days, but March 1 remains active. Compute exactly before classifying; no rounding is needed because the daily amount is integral." + }, + "questions": { + "q": { + "criteria": { + "12_credits": "The correctly calculated outcome is 12 credits.", + "15_credits": "The correctly calculated outcome is 15 credits.", + "6_credits": "The correctly calculated outcome is 6 credits.", + "9_credits": "The correctly calculated outcome is 9 credits." + }, + "instructions": "Apply all stated timing, unit, inclusion, correction, and threshold rules; select the exact band.", + "type": "choice" + } + }, + "model": "cua-s1-4b-0.2" + }, + "expected_response": { + "model": "cua-ai/cua-s1-4b-0.2@16818868b0cc7813808aae4e87b417657046ab79:text", + "answers": { + "q": { + "type": "choice", + "choice": "9_credits", + "probabilities": { + "12_credits": 0.3566144108772278, + "15_credits": 0.15715421736240387, + "6_credits": 0.1183929517865181, + "9_credits": 0.36783841252326965 + }, + "confidence": 0.07737968315894639 + } + }, + "usage": { + "input_tokens": 318, + "output_tokens": 0 + } + }, + "rounded_choice": "12_credits" + } + ], + "artifact_files": { + "chat_template.jinja": { + "size": 7756, + "sha256": "a4aee8afcf2e0711942cf848899be66016f8d14a889ff9ede07bca099c28f715" + }, + "config.json": { + "size": 1979, + "sha256": "63f47812d0f11118e4d252d2b3ad488707eb9287a11589f4fd382a1d31182724" + }, + "cua_s1_export.json": { + "size": 34, + "sha256": "472b609a176f714e2dc06db60a88659c30d87c6af487ae4986ba6cbf70deb101" + }, + "generation_config.json": { + "size": 116, + "sha256": "62153eb6c69f2e1f426beaa8002b7186437e949c7588167085df14e10e9c0a73" + }, + "model-00001-of-00002.safetensors": { + "size": 4972947968, + "sha256": "d2544c2400e723881244e260ce905eec23b3dca12e949cdcdc491a1457f34004" + }, + "model-00002-of-00002.safetensors": { + "size": 3438610328, + "sha256": "31bf465a185fd605099a11d155582432499c13e2257e4c4fa13d1024630b4349" + }, + "model.safetensors.index.json": { + "size": 41993, + "sha256": "83620b40d724f8b286580b5c7c504222b0bdbdd7dc9f068e565132409c7eec9d" + }, + "tokenizer.json": { + "size": 19989325, + "sha256": "06b9509352d2af50381ab2247e083b80d32d5c0aba91c272ca9ff729b6a0e523" + }, + "tokenizer_config.json": { + "size": 1124, + "sha256": "66e427c470fe580fe8c7b5725d857af23d8417e37fae62667ec698306a19987b" + } + } +} diff --git a/tests/cua_s1/head.rs b/tests/cua_s1/head.rs index 37e03680..bcb34cc5 100644 --- a/tests/cua_s1/head.rs +++ b/tests/cua_s1/head.rs @@ -1,15 +1,38 @@ use super::*; +#[test] +fn legacy_merged_projection_preserves_fp32_close_scores() { + // Real BF16-representable weights and hidden values, with two distinct + // FP32 dot products in the same BF16 bin. Legacy merged exports must not + // acquire the new reference profile's ordered tie when the worker upgrades. + assert_eq!( + letter_logits( + &[20.0, 0.03125, 20.0, 0.046875], + &[1.0, 1.0], + 2, + TextExport::Merged + ) + .unwrap(), + [20.03125, 20.046875] + ); +} + #[test] fn letter_projection_matches_the_reference_bf16_output_boundary() { assert_eq!( - letter_logits(&[1.0078125], &[1.0078125], 1).unwrap(), + letter_logits(&[1.0078125], &[1.0078125], 1, TextExport::Unmerged).unwrap(), [1.015625] ); // Both exact dot products belong to the same BF16 bin. Retaining FP32 // invents a winning second option where the reference has an ordered tie. assert_eq!( - letter_logits(&[20.0, 0.03125, 20.0, 0.046875], &[1.0, 1.0], 2).unwrap(), + letter_logits( + &[20.0, 0.03125, 20.0, 0.046875], + &[1.0, 1.0], + 2, + TextExport::Unmerged + ) + .unwrap(), [20.0, 20.0] ); } @@ -30,14 +53,100 @@ fn letter_projection_preserves_fp64_accumulation_and_bf16_readout() { // The first dot product would be zero with FP32 accumulation. The second // must round before postprocessing; unused letter rows must not be read. assert_eq!( - letter_logits(&letters, &[1.0, 1.0, 1.0], 2).unwrap(), + letter_logits(&letters, &[1.0, 1.0, 1.0], 2, TextExport::Unmerged).unwrap(), + [1.0, 16_777_216.0] + ); + assert_eq!( + letter_logits(&letters, &[1.0, 1.0, 1.0], 1, TextExport::Unmerged).unwrap(), + [1.0] + ); + assert_eq!( + letter_logits(&letters, &[1.0, 1.0, 1.0], 2, TextExport::Merged).unwrap(), [1.0, 16_777_216.0] ); - assert_eq!(letter_logits(&letters, &[1.0, 1.0, 1.0], 1).unwrap(), [1.0]); +} + +#[test] +fn export_profiles_keep_the_four_observed_near_tie_orders_separate() { + use serde_json::{Value, json}; + use tokenizers::{Tokenizer, models::wordlevel::WordLevel}; + + let fixture: Value = + serde_json::from_str(include_str!("fixtures/legacy_merged_readout.json")).unwrap(); + // Tokenization is incidental here: run the real response finishing on the + // frozen request's option order, but use synthetic BF16 projection inputs. + // These vectors reproduce the readout boundary, not the model's hidden state. + let tokenizer = Tokenizer::new( + WordLevel::builder() + .vocab([("[UNK]".into(), 0)].into_iter().collect()) + .unk_token("[UNK]".into()) + .build() + .unwrap(), + ); + let processor = crate::processing::Processor { tokenizer }; + let cases = fixture["cases"].as_array().unwrap(); + assert_eq!(cases.len(), 4); + for case in cases { + let keys: Vec<_> = case["request"]["questions"]["q"]["criteria"] + .as_object() + .unwrap() + .keys() + .collect(); + let legacy = case["expected_response"]["answers"]["q"]["choice"] + .as_str() + .unwrap(); + let rounded = case["rounded_choice"].as_str().unwrap(); + let legacy_index = keys.iter().position(|key| key.as_str() == legacy).unwrap(); + let rounded_index = keys.iter().position(|key| key.as_str() == rounded).unwrap(); + assert!(rounded_index < legacy_index); + let mut letters = [10.0, 0.0].repeat(keys.len()); + letters[rounded_index * 2..rounded_index * 2 + 2].copy_from_slice(&[20.0, 0.03125]); + letters[legacy_index * 2..legacy_index * 2 + 2].copy_from_slice(&[20.0, 0.046875]); + // Alternating profiles must not leak a readout policy between exports. + for export in [TextExport::Merged, TextExport::Unmerged, TextExport::Merged] { + let prepared = processor + .prepare(&serde_json::to_vec(&case["request"]).unwrap()) + .ok() + .unwrap(); + let logits = letter_logits(&letters, &[1.0, 1.0], keys.len(), export).unwrap(); + let response = prepared.context.finish(vec![logits]).unwrap(); + let choice = match export { + TextExport::Merged => legacy, + TextExport::Unmerged => rounded, + }; + assert_eq!( + response["answers"]["q"]["choice"], + json!(choice), + "{}", + case["id"] + ); + let probabilities = response["answers"]["q"]["probabilities"] + .as_object() + .unwrap(); + assert_eq!(probabilities.keys().collect::>(), keys); + match export { + TextExport::Merged => { + assert!(probabilities[legacy].as_f64() > probabilities[rounded].as_f64()) + } + TextExport::Unmerged => assert_eq!(probabilities[legacy], probabilities[rounded]), + } + } + } +} + +#[test] +fn finite_fp32_legacy_scores_are_not_overflowed_by_bf16_conversion() { + assert_eq!( + letter_logits(&[f32::MAX], &[1.0], 1, TextExport::Merged).unwrap(), + [f32::MAX] + ); + assert!(letter_logits(&[f32::MAX], &[1.0], 1, TextExport::Unmerged).is_err()); } #[test] fn letter_projection_rejects_nonfinite_scores_before_later_prompts() { - assert!(letter_logits(&[f32::MAX], &[2.0], 1).is_err()); - assert!(letter_logits(&[f32::NAN], &[1.0], 1).is_err()); + for export in [TextExport::Merged, TextExport::Unmerged] { + assert!(letter_logits(&[f32::MAX], &[2.0], 1, export).is_err()); + assert!(letter_logits(&[f32::NAN], &[1.0], 1, export).is_err()); + } } diff --git a/tests/cua_s1/legacy_merged_gpu.rs b/tests/cua_s1/legacy_merged_gpu.rs new file mode 100644 index 00000000..833170c1 --- /dev/null +++ b/tests/cua_s1/legacy_merged_gpu.rs @@ -0,0 +1,110 @@ +//! Frozen A800 compatibility oracle, separate from unmerged/reference parity. +use super::*; +use serde_json::Value; +use sha2::{Digest, Sha256}; +use std::io::Read; + +#[tokio::test] +#[ignore = "requires CUDA and the frozen A800 legacy merged artifact; run serially at Graph 0/1"] +async fn legacy_merged_upgrade_preserves_four_decisions() { + let fixture: Value = + serde_json::from_str(include_str!("fixtures/legacy_merged_readout.json")).unwrap(); + let model = std::path::PathBuf::from( + std::env::var_os("CUA_S1_MERGED_TEST_MODEL").expect("CUA_S1_MERGED_TEST_MODEL"), + ); + let library = + std::path::PathBuf::from(std::env::var_os("CUA_S1_CUDA_LIB").expect("CUA_S1_CUDA_LIB")); + assert!(matches!( + crate::provenance::text_export(&model).unwrap(), + crate::provenance::TextExport::Merged + )); + // Freeze weights/tokenizer as well as requests; another merged checkpoint + // must not accidentally qualify against this A800 compatibility oracle. + for (name, expected) in fixture["artifact_files"].as_object().unwrap() { + let path = model.join(name); + assert_eq!( + std::fs::metadata(&path).unwrap().len(), + expected["size"].as_u64().unwrap(), + "{name}" + ); + let mut file = std::fs::File::open(&path).unwrap(); + let mut hash = Sha256::new(); + let mut buffer = vec![0; 64 * 1024]; + loop { + let n = file.read(&mut buffer).unwrap(); + if n == 0 { + break; + } + hash.update(&buffer[..n]); + } + assert_eq!( + format!("{:x}", hash.finalize()), + expected["sha256"].as_str().unwrap(), + "{name}" + ); + } + let engine = Engine::load(&model, &library).await.unwrap(); + let cases = fixture["cases"].as_array().unwrap(); + assert_eq!(cases.len(), 4); + let mut first_pass = Vec::new(); + for pass in 0..2 { + for (index, case) in cases.iter().enumerate() { + let prepared = engine + .processor + .prepare(&serde_json::to_vec(&case["request"]).unwrap()) + .ok() + .unwrap(); + let rows = engine + .executor + .execute(&engine.scheduler, prepared.inputs) + .await + .unwrap(); + let response = prepared.context.finish(rows).unwrap(); + let expected = &case["expected_response"]; + assert_eq!(response["model"], expected["model"]); + assert_eq!(response["usage"], expected["usage"]); + let actual = &response["answers"]["q"]; + let baseline = &expected["answers"]["q"]; + assert_eq!(actual["type"], baseline["type"]); + assert_eq!( + actual["choice"], baseline["choice"], + "{} pass {pass}", + case["id"] + ); + let actual_probs = actual["probabilities"].as_object().unwrap(); + let baseline_probs = baseline["probabilities"].as_object().unwrap(); + assert_eq!( + actual_probs.keys().collect::>(), + baseline_probs.keys().collect::>() + ); + let max_error = baseline_probs + .iter() + .map(|(key, value)| { + let p = actual_probs[key].as_f64().unwrap(); + assert!(p.is_finite() && (0.0..=1.0).contains(&p)); + (p - value.as_f64().unwrap()).abs() + }) + .fold(0.0_f64, f64::max); + assert!( + max_error <= 0.02, + "{} pass {pass}: drift {max_error}", + case["id"] + ); + assert!( + (actual["confidence"].as_f64().unwrap() - baseline["confidence"].as_f64().unwrap()) + .abs() + <= 0.02 + ); + let bytes = serde_json::to_vec(&response).unwrap(); + if pass == 0 { + first_pass.push(bytes); + } else { + assert_eq!(bytes, first_pass[index], "{} repeat drift", case["id"]); + } + eprintln!( + "legacy merged pass={pass} id={} choice={} max_probability_error={max_error}", + case["id"], actual["choice"] + ); + } + } +} From 16bf31b802257cf17f6921210b6094f9b0e43cf3 Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Thu, 15 Oct 2026 11:10:27 +0800 Subject: [PATCH 6/7] docs(cua-s1): record merged readout repair qualification --- .../cua-s1-text-parity-20261013/README.md | 50 ++++++++++++++----- recipe/cua_s1/native.md | 12 +++-- src/models/cua_s1/README.md | 6 ++- 3 files changed, 48 insertions(+), 20 deletions(-) diff --git a/docs/benchmarks/cua-s1-text-parity-20261013/README.md b/docs/benchmarks/cua-s1-text-parity-20261013/README.md index 21a8b3e6..8a639512 100644 --- a/docs/benchmarks/cua-s1-text-parity-20261013/README.md +++ b/docs/benchmarks/cua-s1-text-parity-20261013/README.md @@ -17,12 +17,12 @@ The 2026-10-15 parent synchronization used #118 at `3be46fa270cb74c33d4db4ffd761405d0522018a`, which inherits [#117](https://github.com/ThinkFlowLab/system1-omni/pull/117) and its merged resident-feature increment #135. Rebuild the library and all consumers together -for ABI 8. This parent synchronization has CPU/CI coverage; current integrated -GPU qualification remains unverified. Decider modes and the separate #127 cache +for ABI 8. The later ABI-8 source `7293c71c` and legacy-readout repair `f212c56f` +have separate source-pinned A800 campaigns described below. Decider modes and the separate #127 cache increment are not included. The frozen corpus comes from #141; merging #141 is not a runtime dependency. -The GPU results below were measured on cumulative candidate +The historical cumulative results in the Results section were measured on candidate `76af07996f60905082b8d48199d4a7003f33f4f6`, whose parent is `18e1a5a6f68b9d11d3e11c7ea7b939af87272f53`. That integrated build used ABI 8. The separate independent qualification at `27f77e49` used ABI 6 with additive @@ -118,16 +118,40 @@ vectors are not captured checkpoint activations. The maintained four-request device test uses the actual frozen inputs and parent responses; run it as in the [native recipe](../../../recipe/cua_s1/native.md). -Fresh GPU execution of the compatibility fix is **pending**: the earlier server -was shut down after preserving the published evidence. In particular, restored -97/139 correctness and zero upgrade flips have not been measured for this new -source. Requalify the full corpus on pinned A800: two passes each for legacy -parent and repaired merged eager/Graph, requiring zero parent/candidate decision -changes and probability error at most 0.02; retain exact-byte differences too. -Repeat unmerged eager/Graph/reference checks at the original 0.02/zero-flip gates. -The previous candidate's unmerged results were 97/139 in all three modes, maximum -probability difference 5.960464477539063e-8 and zero flips. They remain tied to -`7293c71c`; they are not a new-device pass for this follow-up. +The [repair GPU campaign](https://github.com/Levius-Fubuki/system1-omni/releases/tag/cua143-merged-readout-f212c56f-20261015) +tested source `f212c56f23ac750ebfa38f4c72e9fcdc914bf1cd` on pinned A800/ABI 8. +First, two four-request passes on rejected source `7293c71c` reproduced all four +changes. The repaired four-request Engine regression then passed twice per +Graph setting, with zero probability error. Graph 1 logged 6 captures, 2 replays +and no fallback; the test verified the nine artifact files' SHA-256 before loading. + +| Fresh full-corpus mode | Passes × requests | Comparator | Maximum probability difference | Decision changes | Correct | +| --- | --- | --- | --- | --- | --- | +| Legacy parent | 2 × 139 | Parent | 0 | 0 | 97/139 | +| Repaired merged eager | 2 × 139 | Parent | 0 | 0 | 97/139 | +| Repaired merged Graph | 2 × 139 | Parent | 0 | 0 | 97/139 | +| Repaired unmerged eager | 2 × 139 | Python reference | 5.960464477539063e-8 | 0 | 97/139 | +| Repaired unmerged Graph | 2 × 139 | Python reference | 5.960464477539063e-8 | 0 | 97/139 | +| Python reference | 2 × 139 | Reference | 0 | 0 | 97/139 | + +Every repaired merged response body and content type matched the parent byte +for byte. All modes repeated identical response bytes; unmerged eager/Graph +bodies also matched each other. Each full-corpus Graph worker logged 207 +captures and 73 replays, with no fallback. The original 0.02/zero-flip gates +were retained and legacy compatibility additionally required byte identity. +Counts are 1,668 full-corpus HTTP measurements, 8 rejected-source reproductions, +7 excluded feasibility requests and 16 separate direct Engine predictions. +Initial bootstrap rejected a packaging-version mismatch before any GPU request: +Torch metadata is `2.14.0` while runtime is `2.14.0+cu130`; both fields and CUDA +runtime `13.0` were then checked explicitly. The failure is preserved. No source +or tolerance changed during the campaign. Raw evidence was verified locally +before the server was shut down. Later documentation-only updates preserve all +non-Markdown source and fixture bytes from the tested repair commit. + +Legacy compatibility remains distinct from reference parity: restoring the +original merged behavior does not remove its known approximation relative to +unmerged FP32 LoRA. These results do not establish full 27B, L20X/Hopper, +FLA-enabled, performance or general GUI task-quality qualification. ## Results diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index e4421d42..75e714f4 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -5,8 +5,9 @@ The unmerged precision profile has source-pinned A800 80 GB (sm_80), CUDA 13.0 evidence at ABI-8 source `7293c71cfb738504a30ff033ae3567f3368333c1`, plus earlier cumulative ABI-8 and independent ABI-6 campaigns linked in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). -The current branch uses ABI 8. The later legacy-readout compatibility fix has -CPU regression coverage; fresh GPU validation of that follow-up is pending. +The current branch uses ABI 8. The legacy-readout compatibility fix was separately +GPU-tested at `f212c56f23ac750ebfa38f4c72e9fcdc914bf1cd` on the same pinned A800. +See the [source-pinned repair evidence](https://github.com/Levius-Fubuki/system1-omni/releases/tag/cua143-merged-readout-f212c56f-20261015). The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. Run the commands from the repository root. Pass your GPU's compute capability to `build.sh` (89 for Ada, 80 for A100, 90 for H100); the worker finds `libqwen3_5_cuda.so` next to its executable, or at `CUA_S1_CUDA_LIB`: @@ -52,9 +53,10 @@ backbone execution path and retain FP32 letter logits after FP64 accumulation. The BF16 readout boundary is restricted to the validated unmerged export profile, not a request flag or process-global setting. An earlier #143 build applied it to merged exports too, creating four decision changes and reducing frozen-corpus -correctness from 97/139 to 96/139. The follow-up restores legacy readout policy; -that restoration still needs fresh device validation. The A800 unmerged parity -results are scoped to their recorded source and artifact. +correctness from 97/139 to 96/139. At repaired source `f212c56f`, merged eager and +Graph each matched every parent response byte in two 139-request passes, with +zero upgrade flips and 97/139 correct. Unmerged eager/Graph retained zero flips +against the pinned reference. Results remain scoped to that source and artifact. Merging the adapter, TF32/BF16 GDN intermediates, and different GEMM reduction policies can change decisions on close or tied options. Re-export with the command above to use the precision profile. Rebuild worker and library diff --git a/src/models/cua_s1/README.md b/src/models/cua_s1/README.md index 08e835bd..0f7904d3 100644 --- a/src/models/cua_s1/README.md +++ b/src/models/cua_s1/README.md @@ -139,8 +139,10 @@ The old merged export remains an approximate execution mode. Its legacy FP64-accumulated, FP32 letter-logit readout is preserved; BF16 rounding is selected only by `cua-s1-text-unmerged/1`. Applying that rounding to merged exports at `7293c71c` created four ordered-tie decision changes in a same-artifact A800 -upgrade comparison. The compatibility fix is covered by CPU red/green projection -and response tests; a frozen four-request GPU regression is also provided. +upgrade comparison. The compatibility fix has CPU red/green projection and +response coverage. At source `f212c56f`, the four-request device regression passed +in eager/Graph, and both modes matched every legacy parent response byte in two +139-request passes. Unmerged/reference parity remained within its original gates. See the [numerical report](../../../docs/benchmarks/cua-s1-text-parity-20261013/README.md#legacy-merged-upgrade-regression) for the four cases, source pins and device-validation limits. Restoring merged compatibility does not make merged weights equivalent to the unmerged reference. From 3d150b0117ff06e28eafde463cbe87754b4e52c6 Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Thu, 15 Oct 2026 11:17:18 +0800 Subject: [PATCH 7/7] docs(cua-s1): clarify native softmax accumulation precision --- docs/architecture.md | 3 ++- recipe/cua_s1/native.md | 4 +++- src/models/cua_s1/README.md | 5 +++-- 3 files changed, 8 insertions(+), 4 deletions(-) diff --git a/docs/architecture.md b/docs/architecture.md index 5677b13a..b354aebf 100644 --- a/docs/architecture.md +++ b/docs/architecture.md @@ -58,7 +58,8 @@ and reports unavailable health thereafter. Qwen executors own loaded models and CPU head weights with FP64 accumulation. Cua-S1's unmerged text profile rounds answer-letter logits to the tied `lm_head` -BF16 output dtype before FP32 softmax. Legacy merged text and screenshot +BF16 output dtype before response finishing. Cua-S1 finishing normalizes in +FP64 and returns FP32 probabilities. Legacy merged text and screenshot readouts retain FP32 logits; learned scalar heads retain their FP32 rounding and bias order. Finishing checks output cardinality before reconstruction. HTTP validation, error status/body conventions, and real diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index 75e714f4..0737bca7 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -43,7 +43,9 @@ Each question runs one forward pass over its prompt. The `cua-s1-text-unmerged/1 marker selects framework-compatible BF16 normalization and cuBLASLt GEMM, FP32 LoRA branches, and FP32 chunked Gated DeltaNet. The GDN cumulative decay uses the pinned Torch scan order; state is reset for every prompt. The letter -projection rounds its result to BF16 before FP32 softmax, matching `lm_head`. +projection rounds its result to the BF16 `lm_head` output dtype before the existing +per-question softmax. Native finishing normalizes in FP64 and returns FP32 +probabilities. The marker records scale 2 and the SHA-256 of the unchanged adapter; incompatible markers or altered adapter bytes fail before CUDA loading. Prefix continuation is unsupported in this precision mode. diff --git a/src/models/cua_s1/README.md b/src/models/cua_s1/README.md index 0f7904d3..72945c23 100644 --- a/src/models/cua_s1/README.md +++ b/src/models/cua_s1/README.md @@ -13,8 +13,9 @@ option count per question, plus a response context for identity/order and usage. The executor projects the letter rows with FP64 accumulation and returns FP32 letter logits. Only the unmerged text profile first rounds to the BF16 `lm_head` output dtype; legacy merged text and screenshot readouts retain FP32. The text -export policy is validated once and owned by each executor. Response finishing owns per-question -softmax and the choice/confidence interpretation. The engine owns a +export policy is validated once and owned by each executor. Response finishing +owns per-question softmax (FP64 normalization returning FP32 probabilities) and +the choice/confidence interpretation. The engine owns a [shared serial scheduler](../../runtime/README.md) that admits each question's forward before blocking dispatch. Execution remains one prompt per question with a model-state mutex and CPU answer-letter projection after CUDA prefill,