diff --git a/docs/benchmarks/cua-s1-cuda-graphs/README.md b/docs/benchmarks/cua-s1-cuda-graphs/README.md new file mode 100644 index 0000000..7652259 --- /dev/null +++ b/docs/benchmarks/cua-s1-cuda-graphs/README.md @@ -0,0 +1,90 @@ +# Cua-S1 exact-length CUDA Graph experiment + +This compares the native worker from PR #19, baseline +`8367333a5d0e115c4c4ad366b67a91a0553449ce`, with the Graph patch in this PR. +The same modified executable and CUDA library serve both configurations: +`CUA_S1_GRAPH=0` (eager) versus `CUA_S1_GRAPH=1` (Graph replay). There is no +GEMM tuning or change to kernel arithmetic. Since these measurements, the patch +has added capture-failure error cleanup, regression tests, comments and +documentation. The latency benchmark was not rerun for that recovery fix. + +## Protocol and controls + +The hypothesis was that Graph replay reduces CPU launch overhead on repeated +prompt lengths. Success required exact response equality and repeatable latency +reduction beyond run variability. Stop conditions were any failed correctness +check, 120-second readiness timeout, or 20-minute reservation timeout. No runs +were added after inspecting results. + +- One scheduler-reserved GPU 3: NVIDIA L20X, compute capability 8.9; CUDA 13.0; + driver 570.133.20; Rust 1.98.1. CPU affinity was inherited (CPUs 0–223) in both configurations; + NUMA placement was not explicitly bound. Other GPUs had unrelated workloads. +- Pinned Qwen3.5-4B base `851bf6e806efd8d0a36b00ddf55e13ccb7b8cd0a`, Cua-S1 + text adapter `16818868b0cc7813808aae4e87b417657046ab79`, merged BF16 weights. + Both modes reused exactly the same export. Export dependencies were torch + 2.13.0+cu129, transformers 5.14.1, peft 0.20.0, tokenizers 0.22.2, + safetensors 0.8.0, huggingface-hub 1.30.0: an existing environment, differing + from the recipe pins. The native runtime does not use these Python packages. +- One real model-loaded/warmed server per mode, reused for correctness, + feasibility and two measured runs. No cache clearing. HTTP directly to the + worker, concurrency 1; no frontend or Python baseline comparison. +- Each case: three warmups, twenty measured requests per run. A separate + feasibility run with one request per case was excluded from measured results. +- Seventeen correctness requests: fourteen fixtures followed by the positive + fixture with `Submit` replaced by `Cancel` in its state, the one-option + fixture, then the changed fixture again. This exercises changed token ids, + scratch growth and later shorter prompts, eviction and recapture. Scratch + capacity only grows; shorter prompts reuse that allocation. All seventeen response bodies + matched exactly, including probabilities and choices. + +## Reproduction + +Use the build and merged-weight steps in +[the native recipe](../../../recipe/cua_s1/native.md), reserving one exact GPU +through the host's scheduler for all device work. Retrieve the author's frozen +benchmark and earlier fixtures: + +```sh +curl -fL https://raw.githubusercontent.com/twu3202/system1-omni/f9ab3fc808e27d92d26c45c9179ee2a0faf57963/recipe/cua_s1/bench_text.py -o /tmp/bench_text.py +curl -fL https://raw.githubusercontent.com/twu3202/system1-omni/a086a316babe796d41ba96cc4b4be0cc5123f15d/tests/cua_s1/data/text_inputs.json -o /tmp/text_inputs.json +``` + +For each mode, start one worker with `CUA_S1_GRAPH=0` or `1`, wait for its real +`/health`, validate the first inference, and run: + +```sh +python /tmp/bench_text.py --direct http://127.0.0.1:8000 --inputs /tmp/text_inputs.json --warmup 3 --repeat 1 --out feasibility.json +python /tmp/bench_text.py --direct http://127.0.0.1:8000 --inputs /tmp/text_inputs.json --warmup 3 --repeat 20 --out run1.json +python /tmp/bench_text.py --direct http://127.0.0.1:8000 --inputs /tmp/text_inputs.json --warmup 3 --repeat 20 --out run2.json +``` + +The runner uses only Python's standard library. Preserve both measured runs; +compare the same case across modes. [results.json](results.json) contains all +measured samples, warmups, correctness summary and startup metadata. Full +response bodies are omitted; the correctness summary records 17/17 identical +responses. + +## Results + +P50 HTTP milliseconds, showing both measured runs: + +| Prompt tokens | Eager | Graph replay | +| ---: | ---: | ---: | +| 139 | 6.31 / 6.24 | 5.73 / 5.74 | +| 154 | 6.60 / 6.60 | 5.82 / 5.87 | +| 218 | 7.25 / 7.27 | 6.63 / 6.93 | +| 292 | 8.58 / 8.69 | 7.99 / 8.01 | +| 712 | 16.17 / 16.11 | 16.04 / 16.07 | +| 15,446 | 355.10 / 355.05 | 355.50 / 355.73 | + +For these short cases (139–292 tokens), the reduction in the mean of run +medians is 6.6–11.4%. The long case has no meaningful gain. Graph capture adds +first-use work: the first post-readiness fixture took 25.23 ms eager and +33.66 ms with Graphs. Process-to-readiness was 5.58 / 4.93 seconds respectively; +these single observations are not controlled startup comparisons. Export time +is excluded and the files were already prepared for both modes. + +This is one GPU and a fixed fixture workload, not a general speedup claim. +Mixed lengths can evict captures; first-use and capture costs are excluded from +warm request measurements. Full fp32-reference accuracy, concurrent serving, +other CUDA versions and other hardware were not evaluated for this change. diff --git a/docs/benchmarks/cua-s1-cuda-graphs/results.json b/docs/benchmarks/cua-s1-cuda-graphs/results.json new file mode 100644 index 0000000..642afe3 --- /dev/null +++ b/docs/benchmarks/cua-s1-cuda-graphs/results.json @@ -0,0 +1,86 @@ +{ + "metadata": {"sha": "8367333a5d0e115c4c4ad366b67a91a0553449ce", "variable": "CUA_S1_GRAPH=0 versus 1; identical executable/library", "gpu": "3", "affinity": "0-223", "warmup": 3, "repeat": 20, "runs": 2, "cache": "no cache clearing; server reused within each configuration; same merged checkpoint", "preparation": "reuse previously exported merged weights", "eager": {"ready_seconds": 5.5847480073571205, "first_inference_ms": 25.23132972419262}, "graph": {"ready_seconds": 4.929868537001312, "first_inference_ms": 33.65746792405844}}, + "correctness": {"compared_requests": 17, "identical_response_bodies": 17}, + "cases": { + "fixture_positive": {"input_tokens": 218, "status": 200}, + "fixture_negative": {"input_tokens": 292, "status": 200}, + "one_option": {"input_tokens": 139, "status": 200}, + "two_options": {"input_tokens": 154, "status": 200}, + "max_26_options": {"input_tokens": 712, "status": 200}, + "long_state": {"input_tokens": 15446, "status": 200}, + "structured": {"input_tokens": 287, "status": 200}, + "null_criteria": {"input_tokens": 163, "status": 200}, + "non_ascii": {"input_tokens": 188, "status": 200}, + "escaping": {"input_tokens": 188, "status": 200}, + "special_token_text": {"input_tokens": 165, "status": 200}, + "multi_question": {"input_tokens": 346, "status": 200}, + "no_goal": {"input_tokens": 296, "status": 200}, + "array_state": {"input_tokens": 189, "status": 200} + }, + "runs": { + "eager-run1": { + "fixture_positive": {"warmup_ms": [7.63, 7.32, 7.38], "p50_ms": 7.25, "p95_ms": 7.39, "min_ms": 7.04, "raw_ms": [7.3, 7.37, 7.24, 7.42, 7.37, 7.3, 7.26, 7.18, 7.31, 7.31, 7.16, 7.23, 7.24, 7.19, 7.17, 7.04, 7.2, 7.16, 7.27, 7.39]}, + "fixture_negative": {"warmup_ms": [8.66, 8.55, 8.51], "p50_ms": 8.58, "p95_ms": 8.85, "min_ms": 8.47, "raw_ms": [8.63, 8.55, 8.53, 8.75, 8.82, 8.6, 8.49, 8.57, 8.48, 8.53, 8.51, 8.58, 8.47, 8.47, 8.59, 8.67, 8.83, 8.88, 8.85, 8.69]}, + "one_option": {"warmup_ms": [6.35, 6.38, 6.35], "p50_ms": 6.31, "p95_ms": 6.61, "min_ms": 6.12, "raw_ms": [6.61, 6.51, 6.51, 6.35, 6.17, 6.2, 6.21, 6.25, 6.5, 7.37, 6.32, 6.28, 6.33, 6.17, 6.12, 6.28, 6.31, 6.31, 6.37, 6.23]}, + "two_options": {"warmup_ms": [6.44, 6.38, 6.4], "p50_ms": 6.6, "p95_ms": 6.77, "min_ms": 6.35, "raw_ms": [6.4, 6.35, 6.4, 6.44, 6.44, 6.65, 6.61, 6.77, 6.69, 6.66, 6.59, 6.65, 6.73, 6.8, 6.53, 6.6, 6.6, 6.72, 6.59, 6.6]}, + "max_26_options": {"warmup_ms": [16.34, 16.15, 16.2], "p50_ms": 16.17, "p95_ms": 16.36, "min_ms": 16.01, "raw_ms": [16.14, 16.12, 16.13, 16.18, 16.01, 16.16, 16.09, 16.04, 16.07, 16.09, 16.21, 16.36, 16.26, 16.17, 16.26, 16.39, 16.36, 16.17, 16.31, 16.34]}, + "long_state": {"warmup_ms": [352.39, 355.07, 353.5], "p50_ms": 355.1, "p95_ms": 365.56, "min_ms": 352.25, "raw_ms": [352.25, 354.6, 353.88, 353.29, 355.67, 353.29, 354.41, 360.28, 354.4, 364.22, 355.6, 353.95, 365.56, 365.75, 353.16, 356.22, 355.84, 359.18, 354.21, 358.47]}, + "structured": {"warmup_ms": [8.63, 8.51, 8.49], "p50_ms": 8.41, "p95_ms": 8.79, "min_ms": 8.21, "raw_ms": [8.6, 8.25, 8.68, 8.44, 8.25, 8.38, 8.37, 8.29, 8.21, 8.91, 8.45, 8.27, 8.65, 8.57, 8.35, 8.79, 8.49, 8.37, 8.33, 8.48]}, + "null_criteria": {"warmup_ms": [6.58, 7.48, 7.81], "p50_ms": 6.55, "p95_ms": 6.67, "min_ms": 6.41, "raw_ms": [6.67, 6.53, 6.59, 6.58, 6.66, 6.48, 6.41, 6.47, 6.5, 6.49, 6.68, 6.65, 6.66, 6.57, 6.42, 6.58, 6.58, 6.51, 6.44, 6.46]}, + "non_ascii": {"warmup_ms": [6.87, 6.72, 6.72], "p50_ms": 6.99, "p95_ms": 7.22, "min_ms": 6.82, "raw_ms": [6.88, 6.82, 6.82, 6.83, 6.87, 7.04, 7.03, 7.22, 7.12, 6.82, 6.88, 6.99, 7.04, 6.9, 7.01, 7.23, 6.98, 6.86, 7.0, 7.2]}, + "escaping": {"warmup_ms": [6.92, 6.86, 7.18], "p50_ms": 6.89, "p95_ms": 7.39, "min_ms": 6.71, "raw_ms": [7.03, 6.82, 6.8, 6.96, 7.0, 6.8, 6.72, 6.94, 6.85, 6.89, 6.89, 6.88, 7.5, 7.07, 6.92, 6.74, 6.71, 7.39, 7.2, 6.89]}, + "special_token_text": {"warmup_ms": [6.78, 6.51, 6.68], "p50_ms": 6.48, "p95_ms": 6.58, "min_ms": 6.37, "raw_ms": [6.58, 6.37, 6.46, 6.52, 6.47, 6.43, 6.48, 6.47, 6.48, 6.48, 6.57, 6.48, 6.64, 6.57, 6.47, 6.55, 6.56, 6.49, 6.46, 6.51]}, + "multi_question": {"warmup_ms": [13.03, 13.11, 13.28], "p50_ms": 13.35, "p95_ms": 13.51, "min_ms": 13.16, "raw_ms": [13.34, 13.37, 13.34, 13.51, 13.45, 13.43, 13.39, 13.51, 13.37, 13.35, 13.27, 13.32, 13.27, 13.35, 13.16, 13.24, 13.42, 13.67, 13.18, 13.3]}, + "no_goal": {"warmup_ms": [12.8, 12.69, 12.65], "p50_ms": 12.59, "p95_ms": 12.76, "min_ms": 12.49, "raw_ms": [12.63, 12.77, 12.59, 12.59, 12.61, 12.7, 12.76, 12.6, 12.55, 12.61, 12.68, 12.55, 12.49, 12.56, 12.56, 12.55, 12.59, 12.53, 12.53, 12.51]}, + "array_state": {"warmup_ms": [6.88, 6.93, 6.88], "p50_ms": 6.86, "p95_ms": 7.0, "min_ms": 6.68, "raw_ms": [6.74, 6.68, 6.83, 6.89, 7.37, 7.0, 6.97, 6.86, 6.89, 6.98, 6.78, 6.74, 6.87, 6.8, 6.79, 6.81, 6.87, 6.77, 6.98, 6.91]} + }, + "eager-run2": { + "fixture_positive": {"warmup_ms": [7.27, 7.44, 7.27], "p50_ms": 7.27, "p95_ms": 7.53, "min_ms": 7.15, "raw_ms": [7.15, 7.21, 7.32, 7.73, 7.41, 7.23, 7.26, 7.28, 7.26, 7.32, 7.21, 7.5, 7.27, 7.36, 7.2, 7.16, 7.24, 7.53, 7.28, 7.27]}, + "fixture_negative": {"warmup_ms": [8.77, 8.92, 8.71], "p50_ms": 8.69, "p95_ms": 9.0, "min_ms": 8.5, "raw_ms": [8.62, 8.6, 8.83, 8.63, 8.88, 8.72, 8.7, 8.59, 8.94, 8.85, 8.54, 9.0, 9.02, 8.69, 8.6, 8.86, 8.67, 8.5, 8.94, 8.67]}, + "one_option": {"warmup_ms": [6.19, 6.31, 6.51], "p50_ms": 6.24, "p95_ms": 6.42, "min_ms": 6.13, "raw_ms": [6.37, 6.32, 6.42, 6.3, 6.22, 6.24, 6.21, 6.22, 6.28, 6.5, 6.19, 6.13, 6.25, 6.25, 6.27, 6.3, 6.2, 6.17, 6.18, 6.13]}, + "two_options": {"warmup_ms": [6.47, 6.54, 6.41], "p50_ms": 6.6, "p95_ms": 6.84, "min_ms": 6.23, "raw_ms": [6.54, 6.58, 6.55, 6.41, 6.23, 6.37, 6.47, 6.35, 6.6, 6.84, 6.99, 6.81, 6.67, 6.38, 6.63, 6.83, 6.64, 6.61, 6.62, 6.78]}, + "max_26_options": {"warmup_ms": [16.2, 16.31, 17.2], "p50_ms": 16.11, "p95_ms": 16.36, "min_ms": 16.0, "raw_ms": [16.42, 16.36, 16.26, 16.13, 16.0, 16.11, 16.07, 16.06, 16.11, 16.04, 16.05, 16.01, 16.14, 16.14, 16.13, 16.04, 16.07, 16.01, 16.2, 16.22]}, + "long_state": {"warmup_ms": [357.71, 354.68, 355.53], "p50_ms": 355.05, "p95_ms": 356.88, "min_ms": 353.2, "raw_ms": [356.11, 353.2, 356.88, 355.64, 356.32, 354.13, 355.61, 354.62, 356.17, 363.43, 354.71, 354.73, 355.81, 355.0, 354.23, 354.23, 355.77, 353.63, 354.97, 355.1]}, + "structured": {"warmup_ms": [8.69, 8.69, 8.7], "p50_ms": 8.47, "p95_ms": 8.73, "min_ms": 8.33, "raw_ms": [8.73, 8.44, 8.5, 9.13, 8.58, 8.41, 8.36, 8.61, 8.41, 8.4, 8.68, 8.55, 8.44, 8.44, 8.52, 8.45, 8.35, 8.5, 8.49, 8.33]}, + "null_criteria": {"warmup_ms": [6.43, 6.59, 6.51], "p50_ms": 6.39, "p95_ms": 6.52, "min_ms": 6.33, "raw_ms": [6.35, 6.37, 6.41, 6.4, 6.36, 6.39, 6.39, 6.42, 6.45, 6.52, 6.33, 6.36, 6.39, 6.4, 6.35, 6.47, 6.6, 6.47, 6.48, 6.39]}, + "non_ascii": {"warmup_ms": [6.73, 6.78, 6.68], "p50_ms": 6.73, "p95_ms": 6.92, "min_ms": 6.61, "raw_ms": [6.68, 6.72, 6.74, 6.73, 6.66, 6.61, 6.68, 6.71, 6.7, 6.65, 6.72, 6.77, 6.73, 6.81, 6.67, 6.97, 6.92, 6.78, 6.8, 6.77]}, + "escaping": {"warmup_ms": [6.73, 6.84, 6.72], "p50_ms": 7.07, "p95_ms": 7.22, "min_ms": 6.72, "raw_ms": [6.97, 7.09, 7.16, 7.13, 7.31, 7.07, 7.1, 7.16, 7.02, 6.88, 7.12, 7.06, 6.89, 6.74, 6.72, 7.21, 6.91, 6.74, 7.08, 7.22]}, + "special_token_text": {"warmup_ms": [6.61, 6.63, 6.79], "p50_ms": 6.55, "p95_ms": 6.83, "min_ms": 6.31, "raw_ms": [6.86, 6.59, 6.74, 6.5, 6.57, 6.68, 6.31, 6.42, 6.7, 6.48, 6.51, 6.7, 6.71, 6.52, 6.65, 6.83, 6.39, 6.45, 6.53, 6.53]}, + "multi_question": {"warmup_ms": [13.04, 13.08, 13.04], "p50_ms": 13.21, "p95_ms": 13.72, "min_ms": 12.93, "raw_ms": [12.93, 13.09, 13.6, 13.24, 13.07, 13.19, 13.14, 13.64, 13.72, 13.11, 13.35, 13.16, 13.03, 12.96, 13.07, 13.32, 13.97, 13.32, 13.4, 13.27]}, + "no_goal": {"warmup_ms": [12.71, 12.68, 12.63], "p50_ms": 12.59, "p95_ms": 12.67, "min_ms": 12.41, "raw_ms": [12.77, 12.53, 12.53, 12.61, 12.65, 12.54, 12.5, 12.59, 12.64, 12.64, 12.53, 12.64, 12.63, 12.67, 12.6, 12.61, 12.58, 12.46, 12.41, 12.44]}, + "array_state": {"warmup_ms": [6.86, 6.81, 6.88], "p50_ms": 6.72, "p95_ms": 6.81, "min_ms": 6.58, "raw_ms": [6.71, 6.77, 6.81, 6.74, 6.67, 6.72, 6.73, 6.72, 6.58, 6.64, 6.83, 6.72, 6.78, 6.67, 6.78, 6.73, 6.64, 6.66, 6.77, 6.66]} + }, + "graph-run1": { + "fixture_positive": {"warmup_ms": [6.91, 6.72, 6.69], "p50_ms": 6.63, "p95_ms": 6.92, "min_ms": 6.55, "raw_ms": [6.74, 6.66, 6.67, 6.61, 6.64, 6.57, 6.62, 6.64, 6.67, 6.67, 6.58, 6.63, 6.63, 6.78, 6.61, 7.33, 6.61, 6.92, 6.55, 6.58]}, + "fixture_negative": {"warmup_ms": [8.05, 7.92, 7.95], "p50_ms": 7.99, "p95_ms": 8.87, "min_ms": 7.86, "raw_ms": [7.89, 7.87, 7.88, 7.98, 8.02, 8.13, 8.2, 7.99, 7.86, 7.95, 8.11, 8.03, 8.24, 7.9, 8.97, 8.87, 7.92, 7.99, 7.88, 7.88]}, + "one_option": {"warmup_ms": [6.22, 5.85, 5.69], "p50_ms": 5.73, "p95_ms": 5.93, "min_ms": 5.63, "raw_ms": [5.72, 5.81, 5.77, 5.7, 5.67, 5.98, 5.82, 5.75, 5.66, 5.7, 5.72, 5.73, 5.74, 5.93, 5.73, 5.84, 5.71, 5.66, 5.63, 5.74]}, + "two_options": {"warmup_ms": [6.59, 5.88, 5.87], "p50_ms": 5.82, "p95_ms": 5.98, "min_ms": 5.71, "raw_ms": [5.88, 5.82, 5.8, 5.84, 5.77, 5.83, 5.78, 5.82, 6.04, 5.8, 5.84, 5.71, 5.84, 5.72, 5.75, 5.91, 5.98, 5.82, 5.81, 5.81]}, + "max_26_options": {"warmup_ms": [15.85, 16.41, 16.35], "p50_ms": 16.04, "p95_ms": 16.84, "min_ms": 15.68, "raw_ms": [15.68, 15.75, 16.02, 16.84, 16.73, 17.07, 16.42, 16.05, 16.4, 15.97, 16.4, 16.82, 16.24, 15.82, 16.42, 15.8, 15.72, 15.83, 15.75, 15.77]}, + "long_state": {"warmup_ms": [355.68, 355.8, 354.17], "p50_ms": 355.5, "p95_ms": 362.27, "min_ms": 352.76, "raw_ms": [355.59, 363.25, 354.52, 356.0, 356.21, 353.35, 357.57, 352.76, 353.39, 362.27, 356.25, 354.4, 355.42, 354.21, 354.64, 356.54, 355.38, 355.24, 357.17, 357.89]}, + "structured": {"warmup_ms": [8.36, 8.17, 8.02], "p50_ms": 7.86, "p95_ms": 8.41, "min_ms": 7.7, "raw_ms": [7.87, 7.79, 7.76, 7.84, 7.89, 7.86, 7.93, 8.07, 8.41, 7.73, 7.86, 7.85, 7.7, 7.84, 8.04, 7.85, 7.86, 7.82, 8.02, 8.5]}, + "null_criteria": {"warmup_ms": [6.56, 6.49, 6.18], "p50_ms": 5.98, "p95_ms": 6.51, "min_ms": 5.88, "raw_ms": [5.98, 6.27, 5.99, 5.97, 6.04, 6.08, 6.51, 5.97, 6.03, 6.62, 6.0, 5.98, 5.97, 5.97, 5.88, 5.9, 5.92, 5.98, 5.88, 5.95]}, + "non_ascii": {"warmup_ms": [6.32, 6.2, 6.39], "p50_ms": 6.34, "p95_ms": 7.03, "min_ms": 6.09, "raw_ms": [6.3, 6.28, 6.09, 6.09, 6.22, 6.09, 6.23, 6.15, 6.32, 6.3, 7.03, 6.83, 6.62, 6.59, 6.8, 7.12, 6.55, 6.37, 6.41, 6.64]}, + "escaping": {"warmup_ms": [6.65, 6.66, 6.63], "p50_ms": 6.62, "p95_ms": 6.9, "min_ms": 6.19, "raw_ms": [7.22, 6.8, 6.73, 6.61, 6.56, 6.9, 6.68, 6.43, 6.46, 6.78, 6.76, 6.63, 6.64, 6.53, 6.49, 6.38, 6.19, 6.4, 6.61, 6.66]}, + "special_token_text": {"warmup_ms": [6.65, 6.59, 7.18], "p50_ms": 6.3, "p95_ms": 6.52, "min_ms": 6.07, "raw_ms": [6.51, 6.27, 6.15, 6.16, 6.35, 6.39, 6.44, 6.49, 6.4, 6.26, 6.52, 6.62, 6.4, 6.07, 6.14, 6.07, 6.13, 6.33, 6.15, 6.21]}, + "multi_question": {"warmup_ms": [12.58, 12.67, 12.6], "p50_ms": 12.49, "p95_ms": 13.33, "min_ms": 12.26, "raw_ms": [12.49, 12.83, 12.64, 12.42, 12.26, 12.38, 12.49, 12.4, 12.34, 13.19, 13.33, 12.54, 12.5, 12.54, 12.6, 13.93, 12.49, 12.42, 12.32, 12.49]}, + "no_goal": {"warmup_ms": [12.01, 12.0, 11.87], "p50_ms": 11.9, "p95_ms": 12.81, "min_ms": 11.76, "raw_ms": [11.86, 12.0, 11.87, 11.87, 11.96, 11.89, 11.83, 13.15, 11.84, 12.02, 11.89, 11.83, 12.11, 12.02, 12.81, 12.09, 11.91, 11.86, 11.95, 11.76]}, + "array_state": {"warmup_ms": [6.84, 6.34, 6.39], "p50_ms": 6.37, "p95_ms": 6.85, "min_ms": 6.22, "raw_ms": [6.38, 6.37, 7.01, 6.23, 6.22, 6.33, 6.46, 6.35, 6.36, 6.8, 6.47, 6.85, 6.37, 6.37, 6.3, 6.43, 6.58, 6.34, 6.47, 6.53]} + }, + "graph-run2": { + "fixture_positive": {"warmup_ms": [7.53, 6.8, 6.68], "p50_ms": 6.93, "p95_ms": 7.52, "min_ms": 6.61, "raw_ms": [6.73, 6.99, 7.51, 6.83, 6.94, 6.78, 6.61, 6.97, 7.52, 6.84, 7.75, 6.98, 7.09, 6.89, 6.73, 6.93, 6.93, 6.95, 6.85, 6.95]}, + "fixture_negative": {"warmup_ms": [8.56, 8.12, 8.11], "p50_ms": 8.01, "p95_ms": 8.55, "min_ms": 7.79, "raw_ms": [8.04, 8.01, 7.89, 7.93, 7.79, 7.85, 7.79, 7.99, 8.11, 8.05, 7.94, 8.1, 8.0, 8.06, 8.02, 7.88, 8.65, 8.55, 8.11, 8.05]}, + "one_option": {"warmup_ms": [6.03, 5.72, 5.8], "p50_ms": 5.74, "p95_ms": 5.84, "min_ms": 5.61, "raw_ms": [5.86, 5.83, 5.83, 5.75, 5.78, 5.73, 5.81, 5.84, 5.84, 5.8, 5.83, 5.66, 5.65, 5.72, 5.64, 5.61, 5.64, 5.62, 5.62, 5.61]}, + "two_options": {"warmup_ms": [5.93, 5.86, 5.79], "p50_ms": 5.87, "p95_ms": 6.36, "min_ms": 5.74, "raw_ms": [5.75, 5.77, 5.83, 5.8, 5.77, 5.74, 5.84, 6.14, 5.99, 5.86, 6.36, 5.8, 5.99, 6.41, 5.91, 5.84, 5.92, 5.88, 5.88, 5.94]}, + "max_26_options": {"warmup_ms": [15.98, 15.94, 15.93], "p50_ms": 16.07, "p95_ms": 17.22, "min_ms": 15.79, "raw_ms": [15.99, 16.09, 17.27, 15.94, 16.54, 16.37, 16.16, 16.15, 17.22, 16.85, 15.89, 15.79, 16.05, 16.07, 16.2, 16.02, 16.02, 15.94, 16.08, 15.92]}, + "long_state": {"warmup_ms": [355.49, 364.25, 356.84], "p50_ms": 355.73, "p95_ms": 358.11, "min_ms": 354.54, "raw_ms": [356.4, 366.18, 354.86, 354.54, 357.05, 356.94, 355.55, 355.38, 355.84, 357.43, 355.72, 355.81, 355.5, 355.16, 354.97, 354.7, 355.74, 356.9, 355.55, 358.11]}, + "structured": {"warmup_ms": [8.2, 8.1, 8.1], "p50_ms": 7.84, "p95_ms": 8.15, "min_ms": 7.7, "raw_ms": [7.83, 7.84, 7.79, 7.91, 7.75, 7.88, 7.74, 7.7, 7.79, 7.82, 7.92, 7.87, 7.91, 8.45, 8.02, 7.81, 7.79, 7.85, 8.01, 8.15]}, + "null_criteria": {"warmup_ms": [7.01, 6.84, 6.18], "p50_ms": 5.99, "p95_ms": 6.11, "min_ms": 5.86, "raw_ms": [5.94, 6.11, 5.99, 5.94, 5.94, 5.98, 6.01, 6.11, 6.14, 6.04, 6.0, 6.03, 5.99, 5.98, 5.98, 6.02, 5.93, 5.92, 5.86, 5.95]}, + "non_ascii": {"warmup_ms": [6.42, 6.29, 6.39], "p50_ms": 6.39, "p95_ms": 6.73, "min_ms": 6.17, "raw_ms": [6.31, 6.24, 6.39, 6.33, 6.39, 6.17, 6.19, 6.26, 6.39, 6.5, 7.06, 6.4, 6.4, 6.24, 6.43, 6.73, 6.2, 6.37, 6.38, 6.7]}, + "escaping": {"warmup_ms": [6.31, 6.13, 6.29], "p50_ms": 6.46, "p95_ms": 6.61, "min_ms": 6.26, "raw_ms": [6.49, 6.47, 6.45, 6.41, 6.26, 6.44, 6.6, 6.7, 6.61, 6.47, 6.39, 6.32, 6.47, 6.37, 6.38, 6.38, 6.36, 6.47, 6.46, 6.51]}, + "special_token_text": {"warmup_ms": [6.4, 6.29, 6.46], "p50_ms": 6.26, "p95_ms": 6.37, "min_ms": 6.08, "raw_ms": [6.32, 6.11, 6.35, 6.27, 6.37, 6.32, 6.24, 6.24, 6.16, 6.22, 6.32, 6.27, 6.19, 6.12, 6.11, 6.08, 6.09, 6.34, 6.33, 6.68]}, + "multi_question": {"warmup_ms": [13.25, 12.42, 12.41], "p50_ms": 12.54, "p95_ms": 13.22, "min_ms": 12.38, "raw_ms": [12.56, 12.78, 12.46, 12.61, 12.45, 12.52, 13.25, 12.44, 13.05, 12.38, 12.51, 12.48, 13.22, 12.93, 12.41, 12.95, 12.45, 13.07, 12.54, 12.53]}, + "no_goal": {"warmup_ms": [12.17, 12.06, 12.65], "p50_ms": 11.38, "p95_ms": 11.99, "min_ms": 11.24, "raw_ms": [11.99, 11.45, 11.46, 11.57, 11.46, 11.42, 11.38, 11.34, 11.4, 11.3, 11.35, 11.38, 11.24, 12.46, 11.34, 11.32, 11.3, 11.24, 11.63, 11.25]}, + "array_state": {"warmup_ms": [6.14, 6.15, 6.08], "p50_ms": 6.04, "p95_ms": 6.11, "min_ms": 6.03, "raw_ms": [6.05, 6.05, 6.03, 6.04, 6.04, 6.05, 6.03, 6.07, 6.04, 6.04, 6.03, 6.08, 6.03, 6.1, 6.11, 6.2, 6.08, 6.03, 6.04, 6.03]} + } + } +} diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index cf9e996..018feff 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -23,7 +23,16 @@ Start the worker (`CUA_S1_HOST` and `CUA_S1_PORT` default to `127.0.0.1` and `80 CUA_S1_MODEL=weights/cua-s1-4b-0.2-text-merged target/release/omni-cua-s1-native ``` -Each question is one eager forward pass over its prompt; 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. +For the local CUDA Graph experiment, also set `CUA_S1_GRAPH=1`. The first use of +each exact prompt length warms the GEMM plans and captures the forward pass; +later requests replay it with freshly uploaded token ids. At most eight lengths +are cached. 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 3). +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. The request tests need no GPU; the kernel tests compare attention and the chunked Gated DeltaNet prefill with float64 references: diff --git a/src/backends/cuda/qwen3_5/ops.h b/src/backends/cuda/qwen3_5/ops.h index 6f63488..ec0dcd6 100644 --- a/src/backends/cuda/qwen3_5/ops.h +++ b/src/backends/cuda/qwen3_5/ops.h @@ -13,7 +13,7 @@ #include // Bumped whenever a signature below changes. -#define CS1_ABI_VERSION 2 +#define CS1_ABI_VERSION 3 #ifdef __cplusplus extern "C" { @@ -28,6 +28,10 @@ int cs1_malloc(void** ptr, size_t bytes); int cs1_free(void* ptr); int cs1_stream_create(void** stream); int cs1_stream_sync(void* stream); +int cs1_graph_begin(void* stream); +int cs1_graph_end(void* stream, void** exec); +int cs1_graph_launch(void* exec, void* stream); +int cs1_graph_destroy(void* exec); // Copy and wait for the copy. int cs1_upload(void* dst, const void* src, size_t bytes, void* stream); int cs1_download(void* dst, const void* src, size_t bytes, void* stream); diff --git a/src/backends/cuda/qwen3_5/runtime.cu b/src/backends/cuda/qwen3_5/runtime.cu index 8361b83..02a7db6 100644 --- a/src/backends/cuda/qwen3_5/runtime.cu +++ b/src/backends/cuda/qwen3_5/runtime.cu @@ -34,4 +34,28 @@ int cs1_download(void* dst, const void* src, size_t bytes, void* stream) { return e != cudaSuccess ? e : cudaStreamSynchronize(st); } +int cs1_graph_begin(void* stream) { + const cudaError_t e = cudaStreamBeginCapture(static_cast(stream), cudaStreamCaptureModeThreadLocal); + if (e != cudaSuccess) (void)cudaGetLastError(); + return e; +} + +int cs1_graph_end(void* stream, void** exec) { + cudaGraph_t graph = nullptr; + cudaError_t e = cudaStreamEndCapture(static_cast(stream), &graph); + if (e == cudaSuccess) e = cudaGraphInstantiate(reinterpret_cast(exec), graph, 0); + if (graph) cudaGraphDestroy(graph); + // The returned error is already reported; do not poison the next capture. + if (e != cudaSuccess) (void)cudaGetLastError(); + return e; +} + +int cs1_graph_launch(void* exec, void* stream) { + const cudaError_t e = cudaGraphLaunch(static_cast(exec), static_cast(stream)); + if (e != cudaSuccess) (void)cudaGetLastError(); + return e; +} + +int cs1_graph_destroy(void* exec) { return cudaGraphExecDestroy(static_cast(exec)); } + } // extern "C" diff --git a/src/models/cua_s1/native/src/cuda.rs b/src/models/cua_s1/native/src/cuda.rs index a15fef2..719dc45 100644 --- a/src/models/cua_s1/native/src/cuda.rs +++ b/src/models/cua_s1/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 = 2; +const ABI_VERSION: u32 = 3; pub const LIBRARY: &str = "libqwen3_5_cuda.so"; /// A `cudaStream_t`. @@ -55,6 +55,10 @@ api! { cs1_free(ptr: *mut c_void) -> c_int; cs1_stream_create(stream: *mut Stream) -> c_int; cs1_stream_sync(stream: Stream) -> c_int; + cs1_graph_begin(stream: Stream) -> c_int; + cs1_graph_end(stream: Stream, exec: *mut *mut c_void) -> c_int; + cs1_graph_launch(exec: *mut c_void, stream: Stream) -> c_int; + cs1_graph_destroy(exec: *mut c_void) -> c_int; cs1_upload(dst: *mut c_void, src: *const c_void, bytes: usize, stream: Stream) -> c_int; cs1_download(dst: *mut c_void, src: *const c_void, bytes: usize, stream: Stream) -> c_int; cs1_embed(ids: *const i32, table: *const c_void, out: *mut c_void, t: c_int, d: c_int, stream: Stream) -> c_int; @@ -237,3 +241,51 @@ pub unsafe fn download(dst: &mut [u8], src: *const c_void, stream: Stream) -> Re "copy to host", ) } + +pub struct Graph { + exec: *mut c_void, +} + +// SAFETY: the executable graph is only launched by its owner, one launch at a time. +unsafe impl Send for Graph {} + +impl Graph { + /// Capture the work `record` queues on `stream` (nothing runs) and instantiate it. + pub fn capture(stream: Stream, record: impl FnOnce() -> Result<()>) -> Result { + // SAFETY: plain runtime calls on a stream from new_stream; the capture is + // always ended, also when `record` fails. + unsafe { + check((api().cs1_graph_begin)(stream), "cudaStreamBeginCapture")?; + let recorded = record(); + let mut exec = std::ptr::null_mut(); + let ended = check( + (api().cs1_graph_end)(stream, &mut exec), + "capturing a CUDA graph", + ); + match recorded.and(ended) { + Ok(()) => Ok(Graph { exec }), + Err(e) => { + if !exec.is_null() { + (api().cs1_graph_destroy)(exec); + } + Err(e) + } + } + } + } + + pub fn launch(&self, stream: Stream) -> Result<()> { + // SAFETY: an instantiated graph whose buffers outlive it (see Model). + check( + unsafe { (api().cs1_graph_launch)(self.exec, stream) }, + "cudaGraphLaunch", + ) + } +} + +impl Drop for Graph { + fn drop(&mut self) { + // SAFETY: instantiated by capture and not destroyed before. + unsafe { (api().cs1_graph_destroy)(self.exec) }; + } +} diff --git a/src/models/cua_s1/native/src/model.rs b/src/models/cua_s1/native/src/model.rs index 2b04077..ae14d55 100644 --- a/src/models/cua_s1/native/src/model.rs +++ b/src/models/cua_s1/native/src/model.rs @@ -9,7 +9,7 @@ //! token, so the multimodal rotary sections all get the same position and the //! rotary embedding is the plain one. -use std::collections::HashMap; +use std::collections::{HashMap, VecDeque}; use std::ffi::c_void; use std::path::Path; @@ -514,6 +514,9 @@ pub struct Model { gemm: *mut c_void, /// Buffers for the longest prompt so far; grows as needed. scratch: Option, + /// Opt-in replay with at most eight exact-length captures. + graph_enabled: bool, + graphs: VecDeque<(usize, cuda::Graph)>, } // SAFETY: the raw pointers are device addresses and a cuBLASLt handle owned by the @@ -522,6 +525,7 @@ unsafe impl Send for Model {} impl Drop for Model { fn drop(&mut self) { + self.graphs.clear(); // SAFETY: created by cs1_gemm_create and not destroyed before. unsafe { (cuda::api().cs1_gemm_destroy)(self.gemm) }; } @@ -601,6 +605,8 @@ impl Model { stream, gemm, scratch: None, + graph_enabled: std::env::var("CUA_S1_GRAPH").as_deref() == Ok("1"), + graphs: VecDeque::new(), }; Ok(model) } @@ -637,6 +643,7 @@ impl Model { ); cuda::set_device(0)?; if self.scratch.as_ref().is_none_or(|s| t > s.cap) { + self.graphs.clear(); self.scratch = None; self.scratch = Some(Scratch::new( &self.cfg, @@ -648,7 +655,38 @@ impl Model { let ids32: Vec = ids.iter().flat_map(|&i| (i as i32).to_le_bytes()).collect(); // SAFETY: the ids buffer holds at least t int32 values. unsafe { cuda::upload(s.at(s.ids), &ids32, self.stream)? }; - self.run(s, t)?; + if self.graph_enabled { + if !self.graphs.iter().any(|(length, _)| *length == t) { + // Initialize every cuBLASLt plan before stream capture. + self.run(s, t)?; + cuda::synchronize(self.stream)?; + match cuda::Graph::capture(self.stream, || self.run(s, t)) { + Ok(graph) => { + if self.graphs.len() == 8 { + self.graphs.pop_front(); + } + self.graphs.push_back((t, graph)); + } + Err(error) => { + // Capture records without executing: the eager result is valid. + // Disable graphs for this worker rather than retrying failures. + eprintln!("CUDA Graph capture failed; using eager execution: {error:#}"); + self.graph_enabled = false; + self.graphs.clear(); + } + } + } + if self.graph_enabled { + self.graphs + .iter() + .find(|(length, _)| *length == t) + .unwrap() + .1 + .launch(self.stream)?; + } + } else { + self.run(s, t)?; + } let mut last = vec![0u8; h * BF16]; // SAFETY: x holds at least t rows of the hidden size. unsafe { cuda::download(&mut last, s.at(s.x + (t - 1) * h * BF16), self.stream)? }; diff --git a/src/models/cua_s1/native/tests/kernels.rs b/src/models/cua_s1/native/tests/kernels.rs index cd044c8..692f1b5 100644 --- a/src/models/cua_s1/native/tests/kernels.rs +++ b/src/models/cua_s1/native/tests/kernels.rs @@ -59,6 +59,78 @@ fn from_device(buf: &DeviceBuffer, n: usize, st: Stream) -> Vec { .collect() } +#[test] +#[ignore = "needs a GPU and CUA_S1_CUDA_LIB"] +fn graph_replay_reads_updated_inputs_after_failed_capture() { + let st = setup(); + // A failed recording must end capture so this stream can be captured again. + assert!(cuda::Graph::capture(st, || anyhow::bail!("recording failed")).is_err()); + let table: Vec = (0..24).map(|i| bf16::from_f32(i as f32)).collect(); + let weights = to_device(&table, st); + let ids = DeviceBuffer::new(8).unwrap(); + let output = DeviceBuffer::new(32).unwrap(); + let embed = || { + // SAFETY: two int32 ids, three embedding rows of width eight, two output rows. + check( + unsafe { (api().cs1_embed)(ids.at(0).cast(), weights.at(0), output.at(0), 2, 8, st) }, + "capture embed", + ) + }; + let graph = cuda::Graph::capture(st, embed).unwrap(); + let assert_replay = |graph: &cuda::Graph, rows: [i32; 2]| { + let bytes: Vec = rows.iter().flat_map(|id| id.to_le_bytes()).collect(); + // SAFETY: ids holds two int32 values; every id is a valid embedding row. + unsafe { cuda::upload(ids.at(0), &bytes, st).unwrap() }; + graph.launch(st).unwrap(); + let expected: Vec = rows + .iter() + .flat_map(|&row| (row * 8..row * 8 + 8).map(|i| i as f32)) + .collect(); + assert_eq!(from_device(&output, 16, st), expected); + }; + for rows in [[0i32, 1], [2, 0], [1, 2]] { + assert_replay(&graph, rows); + } + // SAFETY: zero is CUDA's valid legacy default stream handle. Capturing it + // is unsupported and must report an error without poisoning this thread. + let default_stream: Stream = unsafe { std::mem::zeroed() }; + assert_ne!(unsafe { (api().cs1_graph_begin)(default_stream) }, 0); + let recovered = cuda::Graph::capture(st, embed).unwrap(); + assert_replay(&recovered, [0, 1]); + // SAFETY: a null graph handle deliberately exercises CUDA's argument error. + assert_ne!( + unsafe { (api().cs1_graph_launch)(std::ptr::null_mut(), st) }, + 0 + ); + let recovered = cuda::Graph::capture(st, embed).unwrap(); + assert_replay(&recovered, [2, 0]); + for propagate in [true, false] { + let error = cuda::Graph::capture(st, || { + embed()?; + // Synchronizing a capturing stream invalidates the capture (900). + // EndCapture then reports 901, even if the closure returns Ok. + // SAFETY: st is a live stream created by setup. + let code = unsafe { (api().cs1_stream_sync)(st) }; + assert_eq!(code, 900); + if propagate { + check(code, "invalidate capture") + } else { + Ok(()) + } + }) + .err() + .expect("synchronization must invalidate capture"); + let expected_code = if propagate { "(900)" } else { "(901)" }; + assert!(error.to_string().contains(expected_code), "{error}"); + + // Retained-graph replay and download do not consume CUDA's last error. + // Recapture must work on this same thread without clearing it here. + assert_replay(&graph, [2, 1]); + let recovered = cuda::Graph::capture(st, embed).unwrap(); + assert_replay(&recovered, [0, 2]); + } +} + #[test] #[ignore = "needs a GPU and CUA_S1_CUDA_LIB"] fn flash_attention_matches_float64_reference() {