diff --git a/README.md b/README.md index aa5450e9..2cf66d04 100644 --- a/README.md +++ b/README.md @@ -168,7 +168,7 @@ v1.1 has its checkpoint export and CPU contract, with no worker yet: | --- | --- | | 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 (`multimodal` adapter) | [Python CUDA worker](src/frontend/cua_s1.py); one PNG/JPEG screenshot, `choice`; native screenshot execution remains in progress | +| Cua-S1 4B 0.2 (`multimodal` adapter) | [Python CUDA reference](src/frontend/cua_s1.py) and [native Rust/CUDA screenshot worker](recipe/cua_s1/native_screenshot.md); one PNG/JPEG image, `choice`; scoped RTX 4090 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 | | Open-Jev-27B-v1.1 | [Native Rust/CUDA worker](recipe/open_jev/native.md); eager independent text candidates; [H200 validation](recipe/open_jev/validation.md) | | Open-Jev-9B | The same [native Rust/CUDA worker](recipe/open_jev/native.md); eager independent text candidates; [reference comparison on sm_89](recipe/open_jev/validation-9b.md) | diff --git a/docs/supported-models.md b/docs/supported-models.md index c92e1aff..e5a03579 100644 --- a/docs/supported-models.md +++ b/docs/supported-models.md @@ -15,10 +15,11 @@ Models that are being added are also tracked in issues labeled [new model](https | Cua-S1 4B 0.2, `text` adapter | [Reference worker](../recipe/cua_s1/text.md) on Transformers and PEFT | Unverified | Validated ([#13](https://github.com/ThinkFlowLab/system1-omni/pull/13)) | Unverified | Python 3.12, the versions in `requirements-text.txt` | | Cua-S1 4B 0.2, `text` adapter | [Native Rust worker](../recipe/cua_s1/native.md) on the [Qwen3.5 CUDA kernels](../src/backends/cuda/qwen3_5/README.md) | Not supported | Validated on compute capability 8.9 ([#19](https://github.com/ThinkFlowLab/system1-omni/pull/19), [#52](https://github.com/ThinkFlowLab/system1-omni/pull/52)) | Not supported | Compute capability 8.0 or newer, the CUDA toolkit to build, weights merged with `export_text_merged.py` | | Cua-S1 4B 0.2, `multimodal` adapter | Reference worker on Transformers and PEFT, [`src/frontend/cua_s1.py`](../src/frontend/cua_s1.py); no recipe yet | Not supported | Validated ([#17](https://github.com/ThinkFlowLab/system1-omni/pull/17), [#18](https://github.com/ThinkFlowLab/system1-omni/pull/18)) | Not supported | The state is one PNG or JPEG image; upstream's `weights.lock.json` next to the base weights | +| Cua-S1 4B 0.2, `multimodal` adapter | [Native Rust/CUDA screenshot worker](../recipe/cua_s1/native_screenshot.md) | Preparation checks only; no CPU inference | RTX 4090 / sm_89, native vision and language Graph regressions; [scoped evidence and controls](../recipe/cua_s1/native_vision_graph.md) | Not supported | Verified BF16 language export, original vision checkpoint and FP32 vision adapter, rebuilt ABI 8 CUDA library; one PNG/JPEG and `choice` | | Valen-Preview-0923 | [Python reference worker](../recipe/valen/reference.md); text state or one inline PNG/JPEG image; `choice` questions | Not supported | Author-run on RTX 5060 Laptop (sm_120, WSL2): one text and one 2×2 PNG request; input and readout parity ≤1e-6 against reference (model outputs not compared). Owner: community. | Not supported | Python 3.12, pinned Valen source/checkpoint, Qwen3.5-2B base, BF16 CUDA | | Open-Jev-27B-v1.1 | [Native Rust/CUDA worker](../recipe/open_jev/native.md) on the shared Qwen3.5/3.8 executor | Not supported | Validated on H200 (sm_90) for the [74 single-candidate workload](../recipe/open_jev/validation.md) | Not supported | Compute capability 8.0 or newer, CUDA toolkit to build, exported merged weights and trained head | | Open-Jev-9B | The same [native Rust/CUDA worker](../recipe/open_jev/native.md) | Not supported | Validated on compute capability 8.9 against the reference for [253 requests](../recipe/open_jev/validation-9b.md) | Not supported | Compute capability 8.0 or newer, CUDA toolkit to build, exported merged weights and trained head | -| autotrust/JEV-27B-VL | [Experimental native worker](../recipe/jev_vl/README.md), single-question API | Contract tests only; no CPU inference | H800 frozen-corpus check; [experimental scope and limits](../recipe/jev_vl/validation.md) | Not supported | Pinned merged checkpoint, rebuilt ABI 6 CUDA library; image inputs require offline Transformers preencoding | +| autotrust/JEV-27B-VL | [Experimental native worker](../recipe/jev_vl/README.md), single-question API | Contract tests only; no CPU inference | H800 frozen-corpus check; [experimental scope and limits](../recipe/jev_vl/validation.md) | Not supported | Pinned merged checkpoint, rebuilt ABI 8 CUDA library; image inputs require offline Transformers preencoding | | CLM-v0.1-8B | [External `clm-serve` recipe](../recipe/clm/README.md) with a CPU stub embeddings server | **Stub-encoder contract checks only** ([#23](https://github.com/ThinkFlowLab/system1-omni/pull/23)); not real Qwen3-8B decisions | Real encoder unverified by the merged recipe | Unverified | Python, upstream CLM and head checkpoint; a real encoder requires a separate embeddings server | | LFM2.5-350M | [Transformers choice worker](../recipe/lfm2/README.md) | Tiny-model regression tests; full checkpoint unverified | L40S FP16 [recorded execution checks](../recipe/lfm2/RESULTS.md) | Unverified | Python 3.12, [pinned dependencies](../src/models/lfm2/requirements.lock.txt), explicit candidate batch size | | CLM-v0.1-8B | [Native Rust engine](../src/models/clm/README.md), `clm-run`, in front of a Qwen3-8B `/v1/embeddings` server | Engine-only agreement **4.5e-06** on identical vectors ([#29](https://github.com/ThinkFlowLab/system1-omni/pull/29), [validation](../recipe/clm/native/VALIDATION.md)); contract checks documented ([#23](https://github.com/ThinkFlowLab/system1-omni/pull/23)) | **Validated on compute capability 8.9** ([#29](https://github.com/ThinkFlowLab/system1-omni/pull/29), [validation](../recipe/clm/native/VALIDATION.md)) | Unverified | Rust to build and a separate Qwen3-8B `/v1/embeddings` server for the encoder; the engine itself has no device-specific code | diff --git a/mkdocs.yml b/mkdocs.yml index 7033af43..098269fa 100644 --- a/mkdocs.yml +++ b/mkdocs.yml @@ -66,6 +66,7 @@ nav: - Cua-S1 text worker: recipe/cua_s1/text.md - Cua-S1 native text worker: recipe/cua_s1/native.md - Valen reference worker: recipe/valen/reference.md + - Cua-S1 native screenshots: recipe/cua_s1/native_screenshot.md - Cua-S1 vision CUDA Graph: recipe/cua_s1/native_vision_graph.md - Open-Jev native text worker: recipe/open_jev/native.md - Open-Jev-27B H200 validation: recipe/open_jev/validation.md diff --git a/recipe/README.md b/recipe/README.md index ce860871..fdb4f9e3 100644 --- a/recipe/README.md +++ b/recipe/README.md @@ -15,6 +15,8 @@ For a first real decision, follow the [complete CPU walkthrough](../docs/getting 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. +- [Cua-S1 native screenshot worker](cua_s1/native_screenshot.md): export the + multimodal language checkpoint and serve one-image decisions with Rust/CUDA. - [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/native.md b/recipe/cua_s1/native.md index 78d14b1a..a1a4ef33 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -27,7 +27,7 @@ For the local CUDA Graph experiment, also set `CUA_S1_GRAPH=1`. The first use of each exact ordered prompt-length vector warms the GEMM plans and captures the forward pass; later requests replay it with freshly uploaded token ids. Each language mode keeps a separate 64-entry FIFO cache. Growing the scratch allocation clears the captures before freeing their buffers. Capture adds first-use latency; leave the variable unset to use -the eager control. Rebuild both the worker and CUDA library together (ABI 6). +the eager control. Rebuild both the worker and CUDA library together (ABI 8). If capture fails, the worker returns the completed eager result and disables Graph capture/replay for its remaining lifetime, logging the failure to stderr. @@ -42,5 +42,5 @@ CUA_S1_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \ ``` For the screenshot worker, [the vision Graph recipe](native_vision_graph.md) -documents `CUA_S1_VISION_GRAPH`, retained single-grid scratch, memory/lifetime +documents `CUA_S1_VISION_GRAPH`, retained four-grid FIFO scratch, memory/lifetime limits and the real CUDA capture-failure wrapper build and test commands. diff --git a/recipe/cua_s1/native_screenshot.md b/recipe/cua_s1/native_screenshot.md new file mode 100644 index 00000000..82709a6e --- /dev/null +++ b/recipe/cua_s1/native_screenshot.md @@ -0,0 +1,96 @@ +# Native Cua-S1 screenshot serving + +The `omni-cua-s1-vision` worker serves the Cua-S1 4B 0.2 `multimodal` +adapter with native Rust/CUDA vision and language execution. It accepts one +inline PNG/JPEG screenshot and 1–8 independent `choice` questions. The vision +encoder runs once per request; each question keeps its own language prompt. +See the [model contract](../../src/models/cua_s1/README.md) for pinned revisions +and the [vision Graph recipe](native_vision_graph.md) for device tests and controls. + +## Prepare and build + +Prepare the pinned base and adapter using the [reference setup](text.md), keeping +the upstream `weights.lock.json` alongside the downloaded artifacts. Use the +reference Python environment only for this one-time language export. The vision +worker uses the original BF16 vision tensors and unmerged FP32 vision LoRA: + +```sh +PYTHONPATH=src .venv/bin/python recipe/cua_s1/export_multimodal_language.py \ + --base weights/Qwen3.5-4B \ + --adapter weights/cua-s1-4b-0.2/multimodal \ + --out weights/cua-s1-multimodal-language + +src/backends/cuda/qwen3_5/build.sh target/release 89 +cargo build --release --locked -p omni-cua-s1-native -p omni-jev +``` + +The export refuses an existing output directory. Keep the export and original +weights immutable while serving. `89` targets the validated RTX 4090; choose your +device's supported compute capability when building. Rebuild the CUDA library +and every native consumer together when the ABI changes. This branch requires +ABI 8, including the default host-feature path. + +## Start the worker and frontend + +```sh +CUA_S1_BASE=weights/Qwen3.5-4B \ +CUA_S1_VISION_ADAPTER=weights/cua-s1-4b-0.2/multimodal \ +CUA_S1_MODEL=weights/cua-s1-multimodal-language \ +CUA_S1_CUDA_LIB=$PWD/target/release/libqwen3_5_cuda.so \ +CUA_S1_HOST=127.0.0.1 CUA_S1_PORT=8000 \ +CUA_S1_GRAPH=0 CUA_S1_VISION_GRAPH=0 CUA_S1_DEVICE_FEATURES=0 \ +target/release/omni-cua-s1-vision +``` + +In another terminal, start the existing Rust frontend: + +```sh +OMNI_JEV_BIND=127.0.0.1:8080 \ +OMNI_JEV_BACKEND_URL=http://127.0.0.1:8000 target/release/omni-jev +curl --fail http://127.0.0.1:8080/health +``` + +Create a request from a local screenshot, then send it through the frontend: + +```sh +python3 - screenshot.png > /tmp/cua-screenshot-request.json <<'PY' +import base64, json, pathlib, sys +image = base64.b64encode(pathlib.Path(sys.argv[1]).read_bytes()).decode() +json.dump({ + "model": "cua-s1-4b-0.2", + "state": {"image": "data:image/png;base64," + image}, + "questions": {"pick": { + "type": "choice", "instructions": "Keep the files.", + "criteria": {"delete": "Click Delete", "cancel": "Click Cancel"} + }} +}, sys.stdout) +PY +curl --fail-with-body http://127.0.0.1:8080/v1/systemone \ + -H 'Content-Type: application/json' \ + --data-binary @/tmp/cua-screenshot-request.json +``` + +Use matching `data:image/jpeg;base64,` for a JPEG. Returned option probabilities +and the selected key depend on the screenshot. The response identifies the +multimodal adapter and reports prompt-token usage with zero generated tokens. + +## Execution controls and limits + +`CUA_S1_GRAPH=1` enables host-feature language Graph replay; +`CUA_S1_VISION_GRAPH=1` independently enables vision Graph replay. +`CUA_S1_DEVICE_FEATURES=1` keeps vision features on the GPU and copies image rows +directly into language scratch. Its language path remains eager; it can be paired +with either vision Graph setting. All three controls default off. See the +[resident-feature contract](native_vision_graph.md#gpu-resident-vision-features) +for strict control parsing, synchronization and reproduction commands. + +Requests are serialized on the loaded GPU executor. HTTP concurrency includes +preparation and queueing; there is no cross-request GPU batch. Input is limited +to an 8 MiB body, 4 MiB image file, dimensions up to 2048 on each axis +and 1,048,576 pixels, one image, and 1–26 options per question. Animated images, +remote image URLs, multi-image state, `score`, `noul` and native Metal are unsupported. + +Recorded RTX 4090 evidence covers native screenshot parity and Graph regressions; +the resident-feature comparison covers host/device output equality on its fixed +real-screenshot corpus. These scopes do not establish general GUI task accuracy +or native equivalence to every reference precision and environment. diff --git a/recipe/cua_s1/native_vision_graph.md b/recipe/cua_s1/native_vision_graph.md index 0a7a944f..feb7a179 100644 --- a/recipe/cua_s1/native_vision_graph.md +++ b/recipe/cua_s1/native_vision_graph.md @@ -33,7 +33,9 @@ Use the pinned base and multimodal adapter revisions listed in the [model contract](../../src/models/cua_s1/README.md#pinned-revisions), and the verified language export produced by `export_multimodal_language.py`. Keep all checkpoint files immutable during execution. Build the worker and the current -CUDA library together; this change adds no CUDA ABI symbols. +CUDA library together for **ABI 8**. Resident feature validation adds +`cs1_get_device` and `cs1_bf16_finite`; older libraries are rejected even when +resident features are disabled. ```sh src/backends/cuda/qwen3_5/build.sh target/release 89 @@ -49,13 +51,99 @@ CUA_S1_VISION_GRAPH=1 target/release/omni-cua-s1-vision Set `CUA_S1_PORT` to choose the listening port (default 8000). Only the existing single-image native CUDA path is covered. Python workers and Metal are unchanged. +## GPU-resident vision features + +Set `CUA_S1_DEVICE_FEATURES=1` before loading `omni-cua-s1-vision` to keep the +vision merger's BF16 output on the GPU. Unset or `0` selects the host path; +every other explicit value, including an empty or non-Unicode value, fails +before checkpoint or CUDA loading. Export the control before running the launch +command above: + +```sh +export CUA_S1_DEVICE_FEATURES=1 +``` + +This switch is independent of `CUA_S1_VISION_GRAPH`. Vision can run eagerly or +replay a Graph; the resident-feature language path remains eager in both modes. +The host-feature language path retains the independent `CUA_S1_GRAPH` control. +The resident path reuses one vision output across all questions in the request +and copies contiguous image-placeholder rows directly into language scratch. +It avoids the full feature download and per-question feature upload. It still +downloads the final hidden state for the existing CPU letter readout. + +`VisionModel::forward_device` returns a borrowed, row-major BF16 matrix after +finishing the vision stream. The borrow prevents producer scratch mutation or +retirement while the features are in use. Language execution validates the +exact allocation extent, shape, device, placeholder order and positions, then +drains its stream on success or error before the borrow can end. Each question +checks finite BF16 values with a GPU scan and a four-byte result download; +flag allocation and synchronization overhead remain. The feature remains +disabled by default; reduced transfers alone do not establish a general +latency or throughput improvement. + +`DeviceBf16Matrix::download` and `VisionEngine::score_device` provide diagnostic +feature/readout comparisons. Serving uses the resident rows without downloading +the complete matrix. The root tests below exercise both diagnostic APIs. + +### Resident-feature regressions + +Weight-free CPU checks run in ordinary CI: + +```sh +cargo test --locked -p omni-cua-s1-native --test device_controls +cargo test --locked -p omni-qwen3-5-native --test device_features +cargo test --locked -p omni-qwen3-5-native --lib device_input_tests +``` + +GPU tests compile in CI but are ignored by default. On a reserved GPU, export +the four checkpoint/library variables from the launch command and run serially. +The finite-scan test needs only `CUA_S1_CUDA_LIB`; the other tests need all four: + +```sh +cargo test --release --locked -p omni-qwen3-5-native --lib device_tests \ + -- --ignored --test-threads=1 +CUA_S1_VISION_GRAPH=0 cargo test --release --locked -p omni-cua-s1-native \ + --lib device_feature_tests -- --ignored --test-threads=1 +CUA_S1_VISION_GRAPH=1 cargo test --release --locked -p omni-cua-s1-native \ + --lib device_feature_tests -- --ignored --test-threads=1 +cargo test --release --locked -p omni-qwen3-5-native --lib device_vision_graph_ \ + -- --ignored --test-threads=1 +``` + +These tests compare feature bits, hidden states, logits, probabilities and +complete responses across changed pixels/goals, repeated questions, different +grids with equal patch counts, scratch growth and short-request recovery. +They also reject invalid metadata and non-finite features. For capture failure, +build the wrapper in the next section and run in a separate process: + +```sh +CUA_S1_CUDA_LIB=/tmp/cua-fail-capture.so CUA_S1_TEST_CAPTURE_FAILURE=1 \ + cargo test --release --locked -p omni-qwen3-5-native \ + --lib device_vision_capture_failure_ -- --ignored --test-threads=1 +``` + +The queued-copy wrapper enqueues one real D2D copy, returns an injected error, +and records a successful real stream drain. The test checks that record before +releasing the vision view, then verifies request recovery across grid changes. +Use an absolute real-library path and a new evidence path for each run: + +```sh +export CUA_S1_REAL_CUDA_LIB="$CUA_S1_CUDA_LIB" +g++ -std=c++17 -shared -fPIC -Wall -Wextra tests/cua_s1/fail_device_copy.cpp \ + -Wl,--no-as-needed "$CUA_S1_REAL_CUDA_LIB" -ldl -o /tmp/cua-fail-device-copy.so +CUA_S1_CUDA_LIB=/tmp/cua-fail-device-copy.so CUA_S1_TEST_COPY_FAILURE=1 \ +CUA_S1_TEST_COPY_FAILURE_EVIDENCE=/tmp/fresh-copy-drain.txt \ + cargo test --release --locked -p omni-cua-s1-native \ + --lib device_copy_failure_tests -- --ignored --test-threads=1 --nocapture +``` + ## Regressions and native measurements Export the same `CUA_S1_BASE`, `CUA_S1_VISION_ADAPTER`, `CUA_S1_MODEL`, and `CUA_S1_CUDA_LIB` variables for these commands: ```sh -cargo test --release --locked -p omni-qwen3-5-native --lib vision_graph_ \ +cargo test --release --locked -p omni-qwen3-5-native --lib vision::model::vision_tests::vision_graph_ \ -- --ignored --test-threads=1 --nocapture CUA_S1_GRAPH=0 CUA_S1_VISION_GRAPH=0 cargo run --release --locked -p omni-cua-s1-native \ --example vision_graph_bench -- /tmp/vision-eager.json @@ -75,7 +163,7 @@ g++ -std=c++17 -shared -fPIC -Wall -Wextra tests/cua_s1/fail_graph_capture.cpp \ -o /tmp/cua-fail-capture.so CUA_S1_CUDA_LIB=/tmp/cua-fail-capture.so CUA_S1_TEST_CAPTURE_FAILURE=1 \ CUA_S1_GRAPH_TRACE=0 cargo test --release --locked -p omni-qwen3-5-native \ - --lib vision_capture_failure_ -- --ignored --test-threads=1 --nocapture + --lib vision::model::vision_tests::vision_capture_failure_ -- --ignored --test-threads=1 --nocapture ``` The root GPU tests cover changed pixels, exact-grid residency and eviction, retirement, diff --git a/recipe/jev_vl/README.md b/recipe/jev_vl/README.md index 6304c364..940dc7ab 100644 --- a/recipe/jev_vl/README.md +++ b/recipe/jev_vl/README.md @@ -39,7 +39,7 @@ The author measured the worker on one H800 80 GB (`sm_90`). A maintainer also on an earlier revision. These runs do not validate later code changes or maximum context lengths. Use an allocated GPU on scheduled hosts. Build the CUDA library and Rust workers from the same revision; -the integrated backend uses **ABI 6** and older libraries must be rebuilt. +the integrated backend uses **ABI 8** and older libraries must be rebuilt. ```sh src/backends/cuda/qwen3_5/build.sh target/release 90 @@ -141,7 +141,7 @@ The image-prefix tokenizer check requires the exported checkpoint. First [restore the frozen corpus](validation.md#restore-the-frozen-corpus). It uses the corpus's fixed `[1, 60, 60]` grid and synthetic embedding rows, so it checks token/position splitting rather than the vision encoder. -CUDA kernel tests require an allocated GPU and rebuilt ABI 6 library: +CUDA kernel tests require an allocated GPU and rebuilt ABI 8 library: ```sh JEV_VL_EXPORT=$PWD/weights/jev-vl-merged \ diff --git a/src/backends/cuda/qwen3_5/README.md b/src/backends/cuda/qwen3_5/README.md index 9056ca23..7481c3b1 100644 --- a/src/backends/cuda/qwen3_5/README.md +++ b/src/backends/cuda/qwen3_5/README.md @@ -16,8 +16,13 @@ The norm, elementwise and q/k preparation kernels round to bfloat16 where Transf `cs1_attention_gated` fuses the sigmoid gate into the attention epilogue, preserving the BF16 rounding of both attention and sigmoid before multiplication. The native workers use this entry point; the separate operations remain available for kernel -comparisons. Rebuild the library and workers together for ABI version 6, which -includes vision, CUDA Graph, gated attention and prefix-continuation entry points. +comparisons. Rebuild the library and all native Rust consumers together for +**ABI version 8**, even with resident features disabled. ABI 8 retains the vision, +CUDA Graph, gated attention and prefix-continuation entry points and adds +`cs1_get_device` and `cs1_bf16_finite` for device feature validation. Older +libraries are rejected at load. See the +[Cua-S1 resident-feature recipe](../../../../recipe/cua_s1/native_vision_graph.md#gpu-resident-vision-features) +for the opt-in path, lifetime contract and regression commands. Gated DeltaNet preparation stores converted TF32 operands in three-byte component planes, preserves the original four-term TF32 accumulation, and writes U/W fragments directly as bfloat16. Dynamic shared memory is 72 KiB per block. The [H200 comparison](../../../../benchmarks/gdn/README.md) records complete GDN call latency, numerical checks, and the small end-to-end change measured with the Open-Jev worker from PR #55. diff --git a/src/backends/cuda/qwen3_5/elementwise.cu b/src/backends/cuda/qwen3_5/elementwise.cu index 97ae9c5a..0e4c0526 100644 --- a/src/backends/cuda/qwen3_5/elementwise.cu +++ b/src/backends/cuda/qwen3_5/elementwise.cu @@ -86,6 +86,13 @@ __global__ void silu_mul_packed_kernel(const Pack8* __restrict__ gate_up, int ld out[i] = result; } +__global__ void bf16_finite_kernel(const uint16_t* x, size_t n, uint32_t* flag) { + for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < n; + i += static_cast(blockDim.x) * gridDim.x) { + if ((x[i] & 0x7f80u) == 0x7f80u) atomicExch(flag, 1u); + } +} + unsigned blocks(size_t n) { return (unsigned)((n + THREADS - 1) / THREADS); } } // namespace @@ -144,3 +151,12 @@ extern "C" int cs1_silu_mul(const void* gate_up, int ld, void* out, int T, int I } return cudaGetLastError(); } + +extern "C" int cs1_bf16_finite(const void* x, size_t n, uint32_t* flag, void* stream) { + auto st = static_cast(stream); + cudaError_t e = cudaMemsetAsync(flag, 0, sizeof(uint32_t), st); + if (e != cudaSuccess) return e; + if (n) bf16_finite_kernel<<(((n - 1) / 256 + 1) < 65535 ? ((n - 1) / 256 + 1) : 65535), 256, 0, st>>>( + static_cast(x), n, flag); + return cudaGetLastError(); +} diff --git a/src/backends/cuda/qwen3_5/ops.h b/src/backends/cuda/qwen3_5/ops.h index 984db6f6..fcb7477d 100644 --- a/src/backends/cuda/qwen3_5/ops.h +++ b/src/backends/cuda/qwen3_5/ops.h @@ -13,7 +13,7 @@ #include // Bumped whenever the required interface below changes. -#define CS1_ABI_VERSION 6 +#define CS1_ABI_VERSION 8 #ifdef __cplusplus extern "C" { @@ -24,6 +24,9 @@ extern "C" { uint32_t cs1_abi_version(void); const char* cs1_error_string(int code); int cs1_set_device(int device); +int cs1_get_device(int* device); +// Set flag to zero, then atomically set to one if any BF16 value is non-finite. +int cs1_bf16_finite(const void* x, size_t n, uint32_t* flag, void* stream); int cs1_malloc(void** ptr, size_t bytes); int cs1_free(void* ptr); int cs1_stream_create(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 f09139e7..5b019dbf 100644 --- a/src/backends/cuda/qwen3_5/qwen3_5.backend.json +++ b/src/backends/cuda/qwen3_5/qwen3_5.backend.json @@ -1,6 +1,6 @@ { "name": "qwen3_5", - "abi_version": 6, + "abi_version": 8, "status": "experimental", "sources": [ "ops.h", diff --git a/src/backends/cuda/qwen3_5/runtime.cu b/src/backends/cuda/qwen3_5/runtime.cu index 5ac61ceb..4ba8a385 100644 --- a/src/backends/cuda/qwen3_5/runtime.cu +++ b/src/backends/cuda/qwen3_5/runtime.cu @@ -12,6 +12,8 @@ const char* cs1_error_string(int code) { return cudaGetErrorString(static_cast Result> { self.0.forward(image) } + /// Borrow synchronized shared-executor output until the language consumer drains. + pub fn forward_device(&mut self, image: &ProcessedImage) -> Result> { + self.0.forward_device(image) + } pub fn forward_with_trace( &mut self, image: &ProcessedImage, diff --git a/src/models/cua_s1/native/src/vision_engine.rs b/src/models/cua_s1/native/src/vision_engine.rs index 9d621767..69ff501e 100644 --- a/src/models/cua_s1/native/src/vision_engine.rs +++ b/src/models/cua_s1/native/src/vision_engine.rs @@ -3,8 +3,9 @@ pub use crate::vision_processing::PreparedRequest; use crate::vision_processing::{VisionProcessor, probabilities}; use crate::{ contract::{LETTERS, Question}, + cuda::DeviceBf16Matrix, image_request::ImageRequest, - inputs::MultimodalInput, + inputs::{DeviceMultimodalInput, MultimodalInput}, model::Model, multimodal::ImagePrompt, vision::VisionModel, @@ -22,6 +23,7 @@ pub struct VisionEngine { pub vision: VisionModel, language: Model, letters: Vec, + device_features: bool, } pub struct Readout { @@ -32,6 +34,13 @@ pub struct Readout { impl VisionEngine { pub fn load(base: &Path, adapter: &Path, language: &Path, library: &Path) -> Result { + let device_features = match std::env::var("CUA_S1_DEVICE_FEATURES") { + Ok(value) => device_features_enabled(Some(&value))?, + Err(std::env::VarError::NotPresent) => false, + Err(std::env::VarError::NotUnicode(_)) => { + anyhow::bail!("CUA_S1_DEVICE_FEATURES must be unset, 0, or 1") + } + }; let marker: Value = serde_json::from_slice( &std::fs::read(language.join("cua_s1_language_export.json")) .context("export the multimodal language checkpoint first")?, @@ -72,6 +81,7 @@ impl VisionEngine { vision, language: model, letters, + device_features, }) } @@ -103,13 +113,17 @@ impl VisionEngine { &prompt.position_ids[2], ], })?; - let logits: Vec = self - .letters + let logits = Self::letter_logits(&self.letters, &hidden, options)?; + Ok((hidden, logits)) + } + + fn letter_logits(letters: &[f32], hidden: &[f32], options: usize) -> Result> { + let logits: Vec = letters .chunks_exact(hidden.len()) .take(options) .map(|w| { w.iter() - .zip(&hidden) + .zip(hidden) .map(|(&a, &b)| a as f64 * b as f64) .sum::() as f32 }) @@ -118,9 +132,48 @@ impl VisionEngine { logits.iter().all(|x| x.is_finite()), "non-finite candidate logits" ); + Ok(logits) + } + + fn raw_score_device( + language: &mut Model, + letters: &[f32], + prompt: &ImagePrompt, + features: &DeviceBf16Matrix<'_>, + options: usize, + ) -> Result<(Vec, Vec)> { + ensure!((1..=26).contains(&options), "expected 1 to 26 options"); + let hidden = language.forward_multimodal_device(&DeviceMultimodalInput { + token_ids: &prompt.token_ids, + image_token_indices: &prompt.image_token_indices, + image_embeddings: *features, + position_ids: [ + &prompt.position_ids[0], + &prompt.position_ids[1], + &prompt.position_ids[2], + ], + })?; + let logits = Self::letter_logits(letters, &hidden, options)?; Ok((hidden, logits)) } + /// Diagnostic parity readout from externally produced, synchronized device rows. + pub fn score_device( + &mut self, + prompt: &ImagePrompt, + features: &DeviceBf16Matrix<'_>, + options: usize, + ) -> Result { + let (hidden, logits) = + Self::raw_score_device(&mut self.language, &self.letters, prompt, features, options)?; + let probabilities = probabilities(&logits)?; + Ok(Readout { + hidden, + logits, + probabilities, + }) + } + /// Diagnostic readout, including processor normalization of raw model logits. pub fn score( &mut self, @@ -156,16 +209,37 @@ impl VisionEngine { prepared.prompts.len() == counts.len() && !counts.is_empty(), "question/prompt count mismatch" ); - let features = self.vision.forward(&prepared.image)?; - prepared - .prompts - .iter() - .zip(counts) - .map(|(p, count)| { - self.raw_score(p, &features, count) + if self.device_features { + let features = self.vision.forward_device(&prepared.image)?; + // Each consumer synchronizes before returning, including errors; + // the view remains borrowed across all questions in this request. + prepared + .prompts + .iter() + .zip(counts) + .map(|(p, count)| { + Self::raw_score_device( + &mut self.language, + &self.letters, + p, + &features, + count, + ) .map(|(_, logits)| logits) - }) - .collect() + }) + .collect() + } else { + let features = self.vision.forward(&prepared.image)?; + prepared + .prompts + .iter() + .zip(counts) + .map(|(p, count)| { + self.raw_score(p, &features, count) + .map(|(_, logits)| logits) + }) + .collect() + } })(); // Also finish queued work on an error before releasing runtime admission. let vision_sync = self.vision.synchronize(); @@ -181,6 +255,15 @@ impl VisionEngine { } } +/// Serving opt-in: unset/0 selects host; 1 selects resident features. +pub fn device_features_enabled(value: Option<&str>) -> Result { + match value { + None | Some("0") => Ok(false), + Some("1") => Ok(true), + Some(_) => anyhow::bail!("CUA_S1_DEVICE_FEATURES must be unset, 0, or 1"), + } +} + fn letter_rows(dir: &Path, ids: &[u32], hidden: usize) -> Result> { let names = [ "embed_tokens.weight", @@ -229,3 +312,11 @@ fn letter_rows(dir: &Path, ids: &[u32], hidden: usize) -> Result> { } Ok(result) } + +#[cfg(test)] +#[path = "../../../../../tests/cua_s1/device_features_gpu.rs"] +mod device_feature_tests; + +#[cfg(test)] +#[path = "../../../../../tests/cua_s1/device_copy_failure_gpu.rs"] +mod device_copy_failure_tests; diff --git a/src/models/qwen3_5/native/Cargo.toml b/src/models/qwen3_5/native/Cargo.toml index 4e9ad660..4b3c2ed0 100644 --- a/src/models/qwen3_5/native/Cargo.toml +++ b/src/models/qwen3_5/native/Cargo.toml @@ -21,3 +21,7 @@ path = "../../../../tests/qwen3_5/config.rs" [[test]] name = "kernels" path = "../../../../tests/qwen3_5/kernels.rs" + +[[test]] +name = "device_features" +path = "../../../../tests/qwen3_5/device_features.rs" diff --git a/src/models/qwen3_5/native/src/cuda.rs b/src/models/qwen3_5/native/src/cuda.rs index c3d1ac8b..81b7199a 100644 --- a/src/models/qwen3_5/native/src/cuda.rs +++ b/src/models/qwen3_5/native/src/cuda.rs @@ -9,7 +9,7 @@ use std::sync::OnceLock; use anyhow::{Context, Result, bail, ensure}; /// `CS1_ABI_VERSION` in ops.h. -const ABI_VERSION: u32 = 6; +const ABI_VERSION: u32 = 8; pub const LIBRARY: &str = "libqwen3_5_cuda.so"; /// A `cudaStream_t`. @@ -113,6 +113,8 @@ api! { cs1_abi_version() -> u32; cs1_error_string(code: c_int) -> *const c_char; cs1_set_device(device: c_int) -> c_int; + cs1_get_device(device: *mut c_int) -> c_int; + cs1_bf16_finite(x: *const c_void, n: usize, flag: *mut u32, stream: Stream) -> c_int; cs1_malloc(ptr: *mut *mut c_void, bytes: usize) -> c_int; cs1_free(ptr: *mut c_void) -> c_int; cs1_stream_create(stream: *mut Stream) -> c_int; @@ -246,6 +248,7 @@ pub fn set_device(device: i32) -> Result<()> { pub struct DeviceBuffer { ptr: *mut c_void, bytes: usize, + device: i32, } // SAFETY: the pointer is a device address; access is serialized by the owner. @@ -254,12 +257,23 @@ unsafe impl Sync for DeviceBuffer {} impl DeviceBuffer { pub fn new(bytes: usize) -> Result { + let device = current_device()?; let mut ptr = std::ptr::null_mut(); if bytes > 0 { // SAFETY: `ptr` is a valid out-pointer. check(unsafe { (api().cs1_malloc)(&mut ptr, bytes) }, "cudaMalloc")?; } - Ok(Self { ptr, bytes }) + Ok(Self { ptr, bytes, device }) + } + + /// Borrow an exact contiguous row-major BF16 matrix. No raw pointer constructor + /// is exposed; the allocation must outlive all consumers. The producer must + /// finish writes before passing the view to a different stream. + pub fn bf16_matrix(&self, rows: usize, columns: usize) -> Result> { + Ok(DeviceBf16Matrix { + buffer: self, + layout: DeviceFeatureLayout::new(rows, columns, self.bytes, self.device)?, + }) } /// The device address `offset` bytes into the buffer. @@ -269,6 +283,112 @@ impl DeviceBuffer { } } +/// Checked metadata; constructing metadata alone never creates a device view. +#[derive(Clone, Copy, Debug)] +pub struct DeviceFeatureLayout { + rows: usize, + columns: usize, + device: i32, +} + +impl DeviceFeatureLayout { + pub fn new(rows: usize, columns: usize, bytes: usize, device: i32) -> Result { + ensure!(rows > 0 && columns > 0, "empty device feature matrix"); + ensure!(device >= 0, "invalid feature device"); + ensure!( + rows.checked_mul(columns).and_then(|n| n.checked_mul(2)) == Some(bytes), + "device feature allocation/shape mismatch" + ); + Ok(Self { + rows, + columns, + device, + }) + } + + pub fn validate(&self, rows: usize, columns: usize, device: i32) -> Result<()> { + ensure!( + self.rows == rows && self.columns == columns, + "image embedding shape mismatch" + ); + ensure!( + self.device == device, + "image features belong to a different CUDA device" + ); + Ok(()) + } +} + +/// Borrowed resident image features. Mutation and reclamation of the producer +/// allocation are prevented while this view is alive. +#[derive(Clone, Copy)] +pub struct DeviceBf16Matrix<'a> { + buffer: &'a DeviceBuffer, + layout: DeviceFeatureLayout, +} + +impl DeviceBf16Matrix<'_> { + pub fn layout(&self) -> DeviceFeatureLayout { + self.layout + } + pub(crate) fn at(&self, offset: usize) -> *const c_void { + self.buffer.at(offset) + } + pub(crate) fn validate_finite(&self, stream: Stream) -> Result<()> { + let flag = DeviceBuffer::new(4)?; + // SAFETY: the checked allocation contains rows*columns BF16 elements; + // the kernel initializes its flag and download waits for all prior work. + let result = (|| unsafe { + check( + (api().cs1_bf16_finite)( + self.at(0), + self.layout.rows * self.layout.columns, + flag.at(0).cast(), + stream, + ), + "BF16 feature finite check", + )?; + let mut bytes = [0; 4]; + download(&mut bytes, flag.at(0), stream)?; + ensure!(u32::from_le_bytes(bytes) == 0, "non-finite image embedding"); + Ok(()) + })(); + let synced = synchronize(stream); + synced?; + result + } + + /// Diagnostic readback; serving does not download the full feature matrix. + /// The producer must have completed writes before calling this method. + pub fn download(&self) -> Result> { + set_device(self.layout.device)?; + let stream = new_stream()?; + let mut bytes = vec![0; self.layout.rows * self.layout.columns * 2]; + // SAFETY: the view validates the allocation's exact byte extent. + let result = unsafe { download(&mut bytes, self.at(0), stream) }; + let synced = synchronize(stream); + unsafe { (api().cs1_stream_destroy)(stream) }; + synced?; + result?; + Ok(bytes + .as_chunks::<2>() + .0 + .iter() + .map(|&b| half::bf16::from_le_bytes(b)) + .collect()) + } +} + +pub fn current_device() -> Result { + let mut device = -1; + // SAFETY: valid out-pointer, plain runtime call. + check( + unsafe { (api().cs1_get_device)(&mut device) }, + "cudaGetDevice", + )?; + Ok(device) +} + impl Drop for DeviceBuffer { fn drop(&mut self) { if !self.ptr.is_null() { @@ -402,3 +522,7 @@ impl Drop for Graph { unsafe { (api().cs1_graph_destroy)(self.exec) }; } } + +#[cfg(test)] +#[path = "../../../../../tests/qwen3_5/device_cuda.rs"] +mod device_tests; diff --git a/src/models/qwen3_5/native/src/inputs.rs b/src/models/qwen3_5/native/src/inputs.rs index bf719eca..f43b8380 100644 --- a/src/models/qwen3_5/native/src/inputs.rs +++ b/src/models/qwen3_5/native/src/inputs.rs @@ -1,8 +1,13 @@ //! Batch-one, unpadded inputs at the adapted-vision / language-model boundary. +use crate::cuda::DeviceBf16Matrix; use anyhow::{Result, ensure}; use half::bf16; +#[cfg(test)] +#[path = "../../../../../tests/qwen3_5/device_inputs.rs"] +mod device_input_tests; + /// Image rows are already adapted to the language hidden size, in placeholder /// order. Positions are the T/H/W slices of an int64 `[3, 1, sequence]` tensor. /// The caller owns preprocessing, vision execution and the position calculation. @@ -21,6 +26,28 @@ impl MultimodalInput<'_> { vocab: usize, image_token: u32, max_position: usize, + ) -> Result<()> { + self.validate_metadata( + hidden, + vocab, + image_token, + max_position, + self.image_embeddings.len(), + )?; + ensure!( + self.image_embeddings.iter().all(|x| x.is_finite()), + "non-finite image embedding" + ); + self.validate_positions(max_position) + } + + fn validate_metadata( + &self, + hidden: usize, + vocab: usize, + image_token: u32, + max_position: usize, + embedding_elements: usize, ) -> Result<()> { let t = self.token_ids.len(); ensure!(t > 0 && t <= max_position, "empty or oversized prompt"); @@ -39,13 +66,14 @@ impl MultimodalInput<'_> { "image indices must exactly match the ordered placeholders" ); ensure!( - Some(self.image_embeddings.len()) == expected.len().checked_mul(hidden), + Some(embedding_elements) == expected.len().checked_mul(hidden), "image embedding shape mismatch" ); - ensure!( - self.image_embeddings.iter().all(|x| x.is_finite()), - "non-finite image embedding" - ); + Ok(()) + } + + fn validate_positions(&self, max_position: usize) -> Result<()> { + let t = self.token_ids.len(); ensure!( self.position_ids.iter().all(|axis| axis.len() == t), "position_ids must have shape [3, 1, sequence]" @@ -61,6 +89,44 @@ impl MultimodalInput<'_> { } } +/// The same prompt/position boundary with borrowed, contiguous device BF16 rows. +/// Vision completion precedes consumption; language execution synchronizes before +/// returning, including errors, so the producer borrow can then be released. +pub struct DeviceMultimodalInput<'a> { + pub token_ids: &'a [u32], + pub image_token_indices: &'a [usize], + pub image_embeddings: DeviceBf16Matrix<'a>, + pub position_ids: [&'a [i64]; 3], +} + +impl DeviceMultimodalInput<'_> { + pub fn validate( + &self, + hidden: usize, + vocab: usize, + image_token: u32, + max_position: usize, + device: i32, + ) -> Result<()> { + self.image_embeddings + .layout() + .validate(self.image_token_indices.len(), hidden, device)?; + let elements = self + .image_token_indices + .len() + .checked_mul(hidden) + .ok_or_else(|| anyhow::anyhow!("image embedding shape overflow"))?; + let metadata = MultimodalInput { + token_ids: self.token_ids, + image_token_indices: self.image_token_indices, + image_embeddings: &[], + position_ids: self.position_ids, + }; + metadata.validate_metadata(hidden, vocab, image_token, max_position, elements)?; + metadata.validate_positions(max_position) + } +} + /// Qwen3.5's interleaved recomposition: overwrite H at 1::3 and W at 2::3 up /// to section[axis] * 3, retaining T elsewhere. The second rotary half repeats /// these frequencies, which the existing attention-prep kernel handles. diff --git a/src/models/qwen3_5/native/src/model.rs b/src/models/qwen3_5/native/src/model.rs index 1e65fdb8..1f59987a 100644 --- a/src/models/qwen3_5/native/src/model.rs +++ b/src/models/qwen3_5/native/src/model.rs @@ -18,7 +18,7 @@ use anyhow::{Context, Result, bail, ensure}; use serde_json::Value as Json; use crate::cuda::{self, DeviceBuffer, Stream, check}; -use crate::inputs::{MultimodalInput, rotary_tables}; +use crate::inputs::{DeviceMultimodalInput, MultimodalInput, rotary_tables}; const ALIGN: usize = 256; const BF16: usize = 2; @@ -896,6 +896,60 @@ impl Model { } } + /// Eager multimodal execution using borrowed device-resident image rows. + /// The producer must finish writes before entry. All queued consumer work is + /// completed on every exit path, before the external allocation can be reused. + pub fn forward_multimodal_device( + &mut self, + input: &DeviceMultimodalInput<'_>, + ) -> Result> { + let result = (|| { + cuda::set_device(0)?; + let image_token = self + .cfg + .image_token_id + .context("checkpoint has no image_token_id")?; + input.validate( + self.cfg.hidden, + self.embed.shape[0], + image_token, + self.cfg.max_positions, + 0, + )?; + input.image_embeddings.validate_finite(self.stream)?; + let t = input.token_ids.len(); + self.prepare_scratch(t)?; + let s = self.scratch.as_ref().unwrap(); + self.upload_positions(s, input.position_ids)?; + self.embed_tokens(s, input.token_ids)?; + let indices = input.image_token_indices; + let row_bytes = self.cfg.hidden * BF16; + let mut begin = 0; + while begin < indices.len() { + let mut end = begin + 1; + while end < indices.len() && indices[end] == indices[end - 1] + 1 { + end += 1; + } + // SAFETY: metadata validates source rows, device and allocation + // extent and indices exactly match placeholders inside s.res. + unsafe { + cuda::copy_dd( + s.at(s.res + indices[begin] * row_bytes), + input.image_embeddings.at(begin * row_bytes), + (end - begin) * row_bytes, + self.stream, + )?; + } + begin = end; + } + self.run(s, &[t], true)?; + self.last_hidden(s, t) + })(); + let synced = self.synchronize(); + synced?; + result + } + fn upload_positions(&self, s: &Scratch, positions: [&[i64]; 3]) -> Result<()> { let (cos, sin) = rotary_tables( positions, diff --git a/src/models/qwen3_5/native/src/vision/model.rs b/src/models/qwen3_5/native/src/vision/model.rs index 89baf9c9..b2f39e4e 100644 --- a/src/models/qwen3_5/native/src/vision/model.rs +++ b/src/models/qwen3_5/native/src/vision/model.rs @@ -1,7 +1,7 @@ //! GPU vision forward for the structurally validated, unmerged checkpoint. use super::{VisionCheckpoint, VisionConfig, geometry::VisionGeometry}; use crate::{ - cuda::{self, DeviceBuffer, Stream, api, check}, + cuda::{self, DeviceBf16Matrix, DeviceBuffer, Stream, api, check}, image_preprocess::ProcessedImage, }; use anyhow::{Result, ensure}; @@ -315,7 +315,69 @@ impl VisionModel { ) -> Result> { self.run(image, Some(&mut callback)) } + /// Synchronized resident merger output; the borrow prevents cache mutation/retirement. + pub fn forward_device(&mut self, image: &ProcessedImage) -> Result> { + let result: Result<(usize, usize)> = (|| { + let (n, index, capture) = self.enqueue_image(image, &mut None)?; + self.synchronize()?; + if capture { + self.capture_current(index); + } + Ok((n, index)) + })(); + if result.is_err() { + // Drain while every referenced allocation and Graph is still owned. + self.synchronize()?; + } + let (n, index) = result?; + self.scratch[index] + .out + .bf16_matrix(n / 4, self.config.out_hidden_size) + } fn run(&mut self, image: &ProcessedImage, mut callback: Trace<'_>) -> Result> { + let result = (|| { + let (n, index, capture) = self.enqueue_image(image, &mut callback)?; + let output = self.read( + &self.scratch[index].out, + n / 4 * self.config.out_hidden_size, + )?; + if let Some(callback) = callback.as_mut() { + callback("merger.output", &output)?; + } else if capture { + self.capture_current(index); + } + Ok(output) + })(); + if result.is_err() { + self.synchronize()?; + } + result + } + fn capture_current(&mut self, index: usize) { + let scratch = &self.scratch[index]; + // The caller completes warmup before recording; capture never reruns math. + let captured = cuda::Graph::capture(self.stream.0, || self.enqueue(scratch, &mut None)); + match captured { + Ok(graph) => { + let scratch = &mut self.scratch[index]; + scratch.graph = Some(graph); + graph_trace("captured", scratch.grid); + } + Err(error) => { + eprintln!("Vision CUDA Graph capture failed; using eager execution: {error:#}"); + self.graph_enabled = false; + for scratch in &mut self.scratch { + scratch.graph = None; + } + } + } + } + + fn enqueue_image( + &mut self, + image: &ProcessedImage, + callback: &mut Trace<'_>, + ) -> Result<(usize, usize, bool)> { cuda::set_device(0)?; let cached = self .scratch @@ -368,31 +430,10 @@ impl VisionModel { { graph.launch(self.stream.0)?; graph_trace("replayed", scratch.grid); - return self.read(&scratch.out, n / 4 * self.config.out_hidden_size); - } - self.enqueue(scratch, &mut callback)?; - let result = self.read(&scratch.out, n / 4 * self.config.out_hidden_size)?; - if let Some(callback) = callback.as_mut() { - callback("merger.output", &result)?; - } else if self.graph_enabled { - // Readback completed warmup. Capture records without executing; keep - // the eager result even when recording/instantiation fails. - let captured = cuda::Graph::capture(self.stream.0, || self.enqueue(scratch, &mut None)); - match captured { - Ok(graph) => { - self.scratch[index].graph = Some(graph); - graph_trace("captured", image.image_grid_thw); - } - Err(error) => { - eprintln!("Vision CUDA Graph capture failed; using eager execution: {error:#}"); - self.graph_enabled = false; - for scratch in &mut self.scratch { - scratch.graph = None; - } - } - } + return Ok((n, index, false)); } - Ok(result) + self.enqueue(scratch, callback)?; + Ok((n, index, callback.is_none() && self.graph_enabled)) } fn enqueue(&self, scratch: &Scratch, callback: &mut Trace<'_>) -> Result<()> { let n = scratch.grid[1] * scratch.grid[2]; diff --git a/tests/cua_s1/device_controls.rs b/tests/cua_s1/device_controls.rs new file mode 100644 index 00000000..a1806e70 --- /dev/null +++ b/tests/cua_s1/device_controls.rs @@ -0,0 +1,46 @@ +use omni_cua_s1_native::vision_engine::device_features_enabled; + +#[test] +fn device_features_accepts_only_unset_zero_and_one() { + assert!(!device_features_enabled(None).unwrap()); + assert!(!device_features_enabled(Some("0")).unwrap()); + assert!(device_features_enabled(Some("1")).unwrap()); + for value in ["", "true", " 1", "1 ", "2"] { + assert!(device_features_enabled(Some(value)).is_err()); + } +} + +fn rejects_before_loading(value: &std::ffi::OsStr) { + let missing = + std::env::temp_dir().join(format!("cua-device-control-missing-{}", std::process::id())); + assert!(!missing.exists()); + let output = std::process::Command::new(env!("CARGO_BIN_EXE_omni-cua-s1-vision")) + .env("CUA_S1_DEVICE_FEATURES", value) + .env("CUA_S1_BASE", &missing) + .env("CUA_S1_VISION_ADAPTER", &missing) + .env("CUA_S1_MODEL", &missing) + .env("CUA_S1_CUDA_LIB", &missing) + .output() + .unwrap(); + assert!(!output.status.success()); + let error = String::from_utf8(output.stderr).unwrap(); + assert!( + error.contains("CUA_S1_DEVICE_FEATURES must be unset, 0, or 1"), + "{error}" + ); + assert!(!error.contains("export the multimodal language checkpoint")); +} + +#[test] +fn invalid_feature_controls_fail_before_checkpoint_or_cuda_loading() { + for value in ["", "true", "2"] { + rejects_before_loading(value.as_ref()); + } +} + +#[cfg(unix)] +#[test] +fn non_unicode_feature_control_fails_before_checkpoint_or_cuda_loading() { + use std::os::unix::ffi::OsStrExt; + rejects_before_loading(std::ffi::OsStr::from_bytes(&[0xff])); +} diff --git a/tests/cua_s1/device_copy_failure_gpu.rs b/tests/cua_s1/device_copy_failure_gpu.rs new file mode 100644 index 00000000..e8140817 --- /dev/null +++ b/tests/cua_s1/device_copy_failure_gpu.rs @@ -0,0 +1,42 @@ +use super::device_feature_tests::{load, questions}; +use super::*; + +#[test] +#[ignore = "requires test-only queued-copy failure wrapper and CUA_S1_* weights"] +fn device_features_queued_copy_failure_drains_before_release_and_recovers() { + assert_eq!( + std::env::var("CUA_S1_TEST_COPY_FAILURE").as_deref(), + Ok("1") + ); + let evidence = + std::path::PathBuf::from(std::env::var_os("CUA_S1_TEST_COPY_FAILURE_EVIDENCE").unwrap()); + assert!(!evidence.exists(), "use a fresh copy-drain evidence path"); + let mut engine = load(); + let first = engine + .prepare(256, 256, &vec![80; 256 * 256 * 3], &questions(1)) + .unwrap(); + let changed = engine + .prepare(512, 256, &vec![120; 512 * 256 * 3], &questions(2)) + .unwrap(); + engine.device_features = false; + let expected_first = engine.predict_prepared(&first).unwrap(); + let expected_changed = engine.predict_prepared(&changed).unwrap(); + { + let features = engine.vision.forward_device(&first.image).unwrap(); + let error = VisionEngine::raw_score_device( + &mut engine.language, + &engine.letters, + &first.prompts[0], + &features, + 2, + ) + .expect_err("injected D2D failure must surface"); + assert!(error.to_string().contains("device copy"), "{error:#}"); + // The wrapper records a successful real stream drain after queuing D2D, + // before the producer's borrowed allocation is allowed to be released. + assert_eq!(std::fs::read_to_string(&evidence).unwrap(), "drained\n"); + } + engine.device_features = true; + assert_eq!(engine.predict_prepared(&changed).unwrap(), expected_changed); + assert_eq!(engine.predict_prepared(&first).unwrap(), expected_first); +} diff --git a/tests/cua_s1/device_features_gpu.rs b/tests/cua_s1/device_features_gpu.rs new file mode 100644 index 00000000..52c344ba --- /dev/null +++ b/tests/cua_s1/device_features_gpu.rs @@ -0,0 +1,202 @@ +//! GPU parity and recovery checks for the borrowed vision/language boundary. +use super::*; +use crate::image_preprocess::ProcessedImage; + +pub(super) fn load() -> VisionEngine { + let path = |name| std::path::PathBuf::from(std::env::var_os(name).expect(name)); + VisionEngine::load( + &path("CUA_S1_BASE"), + &path("CUA_S1_VISION_ADAPTER"), + &path("CUA_S1_MODEL"), + &path("CUA_S1_CUDA_LIB"), + ) + .unwrap() +} +pub(super) fn questions(seed: usize) -> Vec { + vec![ + Question { + name: "move".into(), + goal: format!("Select the next move {seed}."), + keys: vec!["left".into(), "right".into()], + labels: vec!["Move left".into(), "Move right".into()], + }, + Question { + name: "visible".into(), + goal: "Is the scene visible?".into(), + keys: vec!["yes".into(), "no".into()], + labels: vec!["Yes".into(), "No".into()], + }, + ] +} +fn bits(values: &[f32]) -> Vec { + values.iter().map(|v| v.to_bits()).collect() +} + +#[test] +#[ignore = "requires CUDA and pinned CUA_S1_* weights; run serially"] +fn device_features_host_parity_reuse_geometry_growth_and_recovery() { + let mut engine = load(); + // Run this test in separate processes with CUA_S1_VISION_GRAPH=0 and =1. + let mut previous = None; + for (width, height, seed, text_seed, count) in [ + (256, 256, 1, 1, 1), + (256, 256, 2, 1, 2), + (256, 256, 2, 2, 2), + (512, 256, 2, 2, 2), + (256, 512, 3, 3, 2), + (512, 512, 4, 4, 2), + (256, 256, 1, 1, 8), + ] { + let rgb: Vec = (0..width * height * 3) + .map(|i| ((i * 17 + seed * 31) % 251) as u8) + .collect(); + let qs: Vec<_> = (0..count) + .map(|i| { + let mut q = questions(text_seed).remove(i % 2); + q.name = format!("q{i}"); + q + }) + .collect(); + let prepared = engine.prepare(width, height, &rgb, &qs).unwrap(); + let host = engine.vision.forward(&prepared.image).unwrap(); + let expected: Vec<_> = prepared + .prompts + .iter() + .map(|p| engine.score(p, &host, 2).unwrap()) + .collect(); + if width == 256 && height == 256 && count <= 2 { + let feature_bits: Vec<_> = host.iter().map(|v| v.to_bits()).collect(); + let hidden_bits = bits(&expected[0].hidden); + if let Some((old_seed, old_text, old_features, old_hidden)) = previous { + if old_seed != seed { + assert_ne!(feature_bits, old_features, "changed pixels oracle"); + } + if old_seed == seed && old_text != text_seed { + assert_eq!(feature_bits, old_features); + assert_ne!(hidden_bits, old_hidden, "changed text oracle"); + } + } + previous = Some((seed, text_seed, feature_bits, hidden_bits)); + } + let features = engine.vision.forward_device(&prepared.image).unwrap(); + assert_eq!( + features + .download() + .unwrap() + .iter() + .map(|v| v.to_bits()) + .collect::>(), + host.iter().map(|v| v.to_bits()).collect::>() + ); + for (prompt, expected) in prepared + .prompts + .iter() + .zip(&expected) + .cycle() + .take(expected.len() + 2) + { + let (hidden, logits) = VisionEngine::raw_score_device( + &mut engine.language, + &engine.letters, + prompt, + &features, + 2, + ) + .unwrap(); + assert_eq!(bits(&hidden), bits(&expected.hidden)); + assert_eq!(bits(&logits), bits(&expected.logits)); + assert_eq!( + bits(&probabilities(&logits).unwrap()), + bits(&expected.probabilities) + ); + } + let p = &prepared.prompts[0]; + let mut invalid_prompt = ImagePrompt { + token_ids: p.token_ids.clone(), + image_token_indices: p.image_token_indices.clone(), + position_ids: p.position_ids.clone(), + }; + invalid_prompt.image_token_indices.pop(); + assert!( + VisionEngine::raw_score_device( + &mut engine.language, + &engine.letters, + &invalid_prompt, + &features, + 2 + ) + .is_err() + ); + assert!( + VisionEngine::raw_score_device( + &mut engine.language, + &engine.letters, + &prepared.prompts[0], + &features, + 2 + ) + .is_ok() + ); + engine.device_features = false; + let host_response = engine.predict_prepared(&prepared).unwrap(); + engine.device_features = true; + assert_eq!(engine.predict_prepared(&prepared).unwrap(), host_response); + let mut invalid = ProcessedImage { + image_grid_thw: prepared.image.image_grid_thw, + resized_height: prepared.image.resized_height, + resized_width: prepared.image.resized_width, + pixel_values: prepared.image.pixel_values.clone(), + }; + invalid.pixel_values[0] = f32::NAN; + assert!(engine.vision.forward_device(&invalid).is_err()); + assert_eq!(engine.predict_prepared(&prepared).unwrap(), host_response); + } +} + +#[test] +#[ignore = "requires CUDA and pinned CUA_S1_* weights; run serially"] +fn device_features_reject_nonfinite_and_recover() { + let mut engine = load(); + let rgb = vec![32; 256 * 256 * 3]; + let prepared = engine.prepare(256, 256, &rgb, &questions(1)).unwrap(); + let host = engine.vision.forward(&prepared.image).unwrap(); + let expected = engine.score(&prepared.prompts[0], &host, 2).unwrap(); + let stream = crate::cuda::new_stream().unwrap(); + let buffer = crate::cuda::DeviceBuffer::new(host.len() * 2).unwrap(); + for value in [bf16::NAN, bf16::INFINITY, bf16::NEG_INFINITY] { + let mut invalid = host.clone(); + invalid[0] = value; + let bytes: Vec<_> = invalid.iter().flat_map(|v| v.to_le_bytes()).collect(); + unsafe { + crate::cuda::upload(buffer.at(0), &bytes, stream).unwrap(); + } + let view = buffer.bf16_matrix(host.len() / 2560, 2560).unwrap(); + assert!(engine.score_device(&prepared.prompts[0], &view, 2).is_err()); + let recovered = engine.score(&prepared.prompts[0], &host, 2).unwrap(); + assert_eq!(bits(&recovered.hidden), bits(&expected.hidden)); + } + crate::cuda::synchronize(stream).unwrap(); + unsafe { + (crate::cuda::api().cs1_stream_destroy)(stream); + } +} + +#[test] +#[ignore = "requires CUDA and pinned CUA_S1_* weights; run serially"] +fn device_features_language_scratch_growth_and_short_prompt_recovery() { + let mut engine = load(); + let rgb = vec![80; 256 * 256 * 3]; + for words in [1, 1400, 1] { + let mut qs = questions(1); + qs.truncate(1); + qs[0].goal = "Continue ".repeat(words); + let request = engine.prepare(256, 256, &rgb, &qs).unwrap(); + if words > 1 { + assert!(request.prompts[0].token_ids.len() > 1024); + } + engine.device_features = false; + let expected = engine.predict_prepared(&request).unwrap(); + engine.device_features = true; + assert_eq!(engine.predict_prepared(&request).unwrap(), expected); + } +} diff --git a/tests/cua_s1/fail_device_copy.cpp b/tests/cua_s1/fail_device_copy.cpp new file mode 100644 index 00000000..a4c0376c --- /dev/null +++ b/tests/cua_s1/fail_device_copy.cpp @@ -0,0 +1,49 @@ +// Test-only wrapper: enqueue the real D2D copy, fail once, and record its drain. +// Every CUDA operation delegates to the ABI-matched real library. +#include +#include +#include +#include +#include +namespace { +std::atomic injected{false}; +std::atomic pending{false}; +void* real_library() { + const char* path = std::getenv("CUA_S1_REAL_CUDA_LIB"); + if (!path) std::abort(); + void* library = dlopen(path, RTLD_NOW | RTLD_LOCAL); + if (!library) std::abort(); + return library; +} +} +extern "C" int cs1_copy_dd(void* dst, const void* src, size_t bytes, void* stream) { + void* library = real_library(); + using Copy = int (*)(void*, const void*, size_t, void*); + auto copy = reinterpret_cast(dlsym(library, "cs1_copy_dd")); + if (!copy || copy == &cs1_copy_dd) std::abort(); + int result = copy(dst, src, bytes, stream); + if (result == 0 && !injected.exchange(true)) { + pending.store(true); + std::fputs("Injected device-copy failure after a real queued D2D copy\n", stderr); + result = 2; // cudaErrorMemoryAllocation, after the real copy was enqueued. + } + dlclose(library); + return result; +} +extern "C" int cs1_stream_sync(void* stream) { + void* library = real_library(); + using Sync = int (*)(void*); + auto sync = reinterpret_cast(dlsym(library, "cs1_stream_sync")); + if (!sync || sync == &cs1_stream_sync) std::abort(); + int result = sync(stream); + if (result == 0 && pending.exchange(false)) { + const char* path = std::getenv("CUA_S1_TEST_COPY_FAILURE_EVIDENCE"); + if (!path) std::abort(); + FILE* evidence = std::fopen(path, "w"); + if (!evidence) std::abort(); + std::fputs("drained\n", evidence); + if (std::fclose(evidence) != 0) std::abort(); + } + dlclose(library); + return result; +} diff --git a/tests/qwen3_5/device_cuda.rs b/tests/qwen3_5/device_cuda.rs new file mode 100644 index 00000000..39553f75 --- /dev/null +++ b/tests/qwen3_5/device_cuda.rs @@ -0,0 +1,33 @@ +//! Device boundary smoke tests need a GPU/library, but no model weights. +use super::*; +#[test] +#[ignore = "requires CUDA and CUA_S1_CUDA_LIB; run serially"] +fn device_finite_scan_rejects_all_nonfinite_encodings_and_recovers() { + let library = std::path::PathBuf::from(std::env::var_os("CUA_S1_CUDA_LIB").unwrap()); + load(&library).unwrap(); + set_device(0).unwrap(); + let stream = new_stream().unwrap(); + let buffer = DeviceBuffer::new(514).unwrap(); + let mut values = vec![half::bf16::from_f32(1.); 257]; + for bad in [0x7f80, 0xff80, 0x7fc0, 0xffc1] { + values[256] = half::bf16::from_bits(bad); + let bytes: Vec<_> = values.iter().flat_map(|v| v.to_le_bytes()).collect(); + unsafe { + upload(buffer.at(0), &bytes, stream).unwrap(); + } + let view = buffer.bf16_matrix(257, 1).unwrap(); + assert!(view.validate_finite(stream).is_err()); + values[256] = half::bf16::from_f32(-0.); + let bytes: Vec<_> = values.iter().flat_map(|v| v.to_le_bytes()).collect(); + unsafe { + upload(buffer.at(0), &bytes, stream).unwrap(); + } + let view = buffer.bf16_matrix(257, 1).unwrap(); + view.validate_finite(stream).unwrap(); + assert_eq!(view.download().unwrap(), values); + } + synchronize(stream).unwrap(); + unsafe { + (api().cs1_stream_destroy)(stream); + } +} diff --git a/tests/qwen3_5/device_features.rs b/tests/qwen3_5/device_features.rs new file mode 100644 index 00000000..caf9d56b --- /dev/null +++ b/tests/qwen3_5/device_features.rs @@ -0,0 +1,26 @@ +use omni_qwen3_5_native::cuda::DeviceFeatureLayout; + +#[test] +fn device_layout_requires_exact_nonempty_bf16_matrix_and_same_device() { + assert!( + DeviceFeatureLayout::new(3, 2560, 15360, 0) + .unwrap() + .validate(3, 2560, 0) + .is_ok() + ); + for (rows, columns, bytes, device) in [ + (0, 2560, 0, 0), + (3, 0, 0, 0), + (3, 2560, 15359, 0), + (3, 2560, 15362, 0), + (3, 2560, 15360, -1), + (usize::MAX, 2, 0, 0), + (usize::MAX / 2 + 1, 1, 0, 0), + ] { + assert!(DeviceFeatureLayout::new(rows, columns, bytes, device).is_err()); + } + let layout = DeviceFeatureLayout::new(3, 2560, 15360, 1).unwrap(); + assert!(layout.validate(3, 2560, 0).is_err()); + assert!(layout.validate(2, 2560, 1).is_err()); + assert!(layout.validate(3, 1024, 1).is_err()); +} diff --git a/tests/qwen3_5/device_inputs.rs b/tests/qwen3_5/device_inputs.rs new file mode 100644 index 00000000..caa767b0 --- /dev/null +++ b/tests/qwen3_5/device_inputs.rs @@ -0,0 +1,57 @@ +//! Weight-free checks of metadata shared by host and resident input validation. +//! No synthetic device pointer or CUDA allocation is needed for these checks. +use super::*; + +fn input() -> MultimodalInput<'static> { + MultimodalInput { + token_ids: &[1, 9, 9, 2], + image_token_indices: &[1, 2], + image_embeddings: &[], + position_ids: [&[0, 1, 1, 2], &[0, 1, 2, 3], &[0, 1, 2, 3]], + } +} + +#[test] +fn resident_metadata_accepts_valid_placeholder_rows_and_positions() { + let input = input(); + input.validate_metadata(2, 10, 9, 8, 4).unwrap(); + input.validate_positions(8).unwrap(); +} + +#[test] +fn resident_metadata_rejects_missing_reordered_and_extra_placeholders() { + for indices in [&[1][..], &[2, 1], &[1, 2, 3], &[1, 1]] { + let mut input = input(); + input.image_token_indices = indices; + let error = input + .validate_metadata(2, 10, 9, 8, indices.len() * 2) + .unwrap_err(); + assert!(error.to_string().contains("ordered placeholders")); + } +} + +#[test] +fn resident_metadata_rejects_bad_prompt_bounds_vocabulary_and_extent() { + let mut invalid = input(); + invalid.token_ids = &[]; + assert!(invalid.validate_metadata(2, 10, 9, 8, 4).is_err()); + assert!(input().validate_metadata(2, 10, 9, 3, 4).is_err()); + assert!(input().validate_metadata(2, 9, 9, 8, 4).is_err()); + for elements in [0, 3, 5] { + assert!(input().validate_metadata(2, 10, 9, 8, elements).is_err()); + } + assert!(input().validate_metadata(usize::MAX, 10, 9, 8, 0).is_err()); + input().validate_metadata(2, 10, 9, 8, 4).unwrap(); +} + +#[test] +fn resident_positions_reject_bad_axis_length_negative_and_limit() { + for axis in [&[0, 1][..], &[0, -1, 2, 3], &[0, 1, 2, 8]] { + for changed_axis in 0..3 { + let mut input = input(); + input.position_ids[changed_axis] = axis; + assert!(input.validate_positions(8).is_err()); + } + } + input().validate_positions(8).unwrap(); +} diff --git a/tests/qwen3_5/vision_graph.rs b/tests/qwen3_5/vision_graph.rs index e7b31da6..89933718 100644 --- a/tests/qwen3_5/vision_graph.rs +++ b/tests/qwen3_5/vision_graph.rs @@ -203,3 +203,68 @@ fn vision_capture_failure_keeps_eager_result_and_disables_replay() { assert!(model.scratch.iter().all(|s| s.graph.is_none())); assert_eq!(bits(&model.forward(&input).unwrap()), expected); } + +#[test] +#[ignore = "requires CUDA and pinned CUA_S1_BASE/CUA_S1_VISION_ADAPTER/CUA_S1_CUDA_LIB"] +fn device_vision_graph_replays_changed_pixels_and_replaces_geometry() { + let mut model = load(); + let inputs = [ + image([1, 16, 16], 1), + image([1, 16, 16], 2), + image([1, 8, 32], 2), + image([1, 16, 32], 3), + image([1, 16, 16], 1), + ]; + model.graph_enabled = false; + let expected: Vec<_> = inputs + .iter() + .map(|i| bits(&model.forward(i).unwrap())) + .collect(); + model.graph_enabled = true; + for (input, expected) in inputs.iter().zip(&expected) { + for _ in 0..2 { + let features = model.forward_device(input).unwrap(); + assert_eq!(bits(&features.download().unwrap()), *expected); + } + let scratch = model + .scratch + .iter() + .find(|s| s.grid == input.image_grid_thw) + .unwrap(); + assert_eq!(scratch.grid, input.image_grid_thw); + assert!( + scratch.graph.is_some(), + "resident path must capture/replay vision" + ); + } +} + +#[test] +#[ignore = "requires the test-only failing CUDA wrapper; see native_vision_graph.md"] +fn device_vision_capture_failure_preserves_resident_eager_result_and_recovery() { + assert_eq!( + std::env::var("CUA_S1_TEST_CAPTURE_FAILURE").as_deref(), + Ok("1") + ); + let mut model = load(); + model.graph_enabled = false; + let input = image([1, 16, 16], 9); + let changed = image([1, 8, 32], 10); + let expected = bits(&model.forward(&input).unwrap()); + let expected_changed = bits(&model.forward(&changed).unwrap()); + model.graph_enabled = true; + assert_eq!( + bits(&model.forward_device(&input).unwrap().download().unwrap()), + expected + ); + assert!(!model.graph_enabled); + assert!(model.scratch.iter().all(|s| s.graph.is_none())); + assert_eq!( + bits(&model.forward_device(&changed).unwrap().download().unwrap()), + expected_changed + ); + assert_eq!( + bits(&model.forward_device(&input).unwrap().download().unwrap()), + expected + ); +}