diff --git a/docs/supported-models.md b/docs/supported-models.md index e6c646f5..aecd6670 100644 --- a/docs/supported-models.md +++ b/docs/supported-models.md @@ -15,12 +15,12 @@ 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` | +| 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 9 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 | | JEMM / Qwen3.8-27B | [Native Rust/CUDA worker](../recipe/jemm/README.md) | Host checks only; no CPU inference | A800 (`sm_80`) [text and 1–4-image frozen corpus](../src/models/jemm/README.md#validation) | Not supported | Pinned `jemm-native/2` export with BF16 base and unmerged FP32 LoRA, selected untied LM head, cuDNN 9 and additive reference CUDA symbols | -| 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 | +| 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 9 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/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index 0737bca7..3a89e9d9 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -5,7 +5,7 @@ The unmerged precision profile has source-pinned A800 80 GB (sm_80), CUDA 13.0 evidence at ABI-8 source `7293c71cfb738504a30ff033ae3567f3368333c1`, plus earlier cumulative ABI-8 and independent ABI-6 campaigns linked in [PR #143](https://github.com/ThinkFlowLab/system1-omni/pull/143). -The current branch uses ABI 8. The legacy-readout compatibility fix was separately +The current branch uses ABI 9. The legacy-readout compatibility fix was separately GPU-tested at `f212c56f23ac750ebfa38f4c72e9fcdc914bf1cd` on the same pinned A800. See the [source-pinned repair evidence](https://github.com/Levius-Fubuki/system1-omni/releases/tag/cua143-merged-readout-f212c56f-20261015). The earlier merged path was checked on an RTX 6000 Ada (sm_89) with CUDA 13.2. @@ -35,7 +35,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 8). +the eager control. Rebuild both the worker and CUDA library together (ABI 9). If capture fails, the worker returns the completed eager result and disables Graph capture/replay for its remaining lifetime, logging the failure to stderr. @@ -62,7 +62,7 @@ against the pinned reference. Results remain scoped to that source and artifact. Merging the adapter, TF32/BF16 GDN intermediates, and different GEMM reduction policies can change decisions on close or tied options. Re-export with the command above to use the precision profile. Rebuild worker and library -together: this branch now uses ABI 8, and the precision profile also requires the new optional +together: this branch now uses ABI 9, and the precision profile also requires the new optional FP32 GDN symbols. No cuDNN or Python is needed at serving time. The validation scope and tolerances are in diff --git a/recipe/cua_s1/native_screenshot.md b/recipe/cua_s1/native_screenshot.md index 82709a6e..d591a6df 100644 --- a/recipe/cua_s1/native_screenshot.md +++ b/recipe/cua_s1/native_screenshot.md @@ -28,7 +28,7 @@ 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. +ABI 9, including the default host-feature path. ## Start the worker and frontend diff --git a/recipe/cua_s1/native_vision_graph.md b/recipe/cua_s1/native_vision_graph.md index feb7a179..ac999256 100644 --- a/recipe/cua_s1/native_vision_graph.md +++ b/recipe/cua_s1/native_vision_graph.md @@ -33,7 +33,7 @@ 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 for **ABI 8**. Resident feature validation adds +CUDA library together for **ABI 9**. Resident feature validation adds `cs1_get_device` and `cs1_bf16_finite`; older libraries are rejected even when resident features are disabled. diff --git a/recipe/jev_vl/README.md b/recipe/jev_vl/README.md index 940dc7ab..9887101e 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 8** and older libraries must be rebuilt. +the integrated backend uses **ABI 9** 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 8 library: +CUDA kernel tests require an allocated GPU and rebuilt ABI 9 library: ```sh JEV_VL_EXPORT=$PWD/weights/jev-vl-merged \ diff --git a/recipe/open_jev/native.md b/recipe/open_jev/native.md index 761c4f28..463e36c0 100644 --- a/recipe/open_jev/native.md +++ b/recipe/open_jev/native.md @@ -67,7 +67,7 @@ or an incomplete export. The saved limit defaults to 4096 tokens per candidate; The CUDA kernels require compute capability 8.0 or newer. The current build target below is Ada (`89`); pass your GPU's compute capability explicitly. The CUDA shared library and Rust workers must be rebuilt together for ABI -version 6, which includes vision, CUDA Graph and prefix-continuation entry points. +version 9, which includes vision, CUDA Graph, prefix-continuation and fixed-GEMM entry points. ```sh src/backends/cuda/qwen3_5/build.sh target/release 89 diff --git a/src/backends/cuda/qwen3_5/README.md b/src/backends/cuda/qwen3_5/README.md index 13d48921..56df461f 100644 --- a/src/backends/cuda/qwen3_5/README.md +++ b/src/backends/cuda/qwen3_5/README.md @@ -17,13 +17,24 @@ The norm, elementwise and q/k preparation kernels round to bfloat16 where Transf 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 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 +**ABI version 9**, even with resident features disabled. ABI 9 retains the vision, +CUDA Graph, gated attention and prefix-continuation entry points and ABI 8's +`cs1_get_device` and `cs1_bf16_finite` for device feature validation, and adds the +fixed-algorithm GEMM handle below. 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. +cuBLASLt's heuristic picks a GEMM algorithm per M, so a row's result can change with +the number of rows in the call: a prefix and its branch, run separately, round +differently from the same tokens in one pass. `cs1_gemm_create_fixed` returns a handle +that keeps one algorithm per weight shape for every M, the heuristic's first choice at a +reference M among algorithms without split-K, so each row's result is the same whatever +M is and wherever the row sits. cuBLASLt doesn't document this property; +`fixed_gemm_rows_do_not_depend_on_m` checks it, so run it on a new GPU before relying +on it (it passed on sm_89). It is slower for some shapes and faster for others: on an +RTX 6000 Ada, the down and output projections take up to about 4 times as long at small +M without split-K. `cs1_gemm_create` remains the default. + 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. The shared Rust model can pack independent sequences for input and gate/up GEMMs. @@ -34,7 +45,8 @@ entry points. Open-Jev uses this path within requests; Cua-S1 keeps single-promp Prefix continuation retains full-attention KV, three convolution input rows, and FP32 Gated DeltaNet state at 64-token chunk boundaries. It is used by the experimental JEV-VL worker; existing full-prefill callers keep their own paths. -ABI 6 adds mandatory copy and prefix entry points after ABI 5 native vision. +ABI 6 adds mandatory copy and prefix entry points after ABI 5 native vision, +and ABI 9 adds the fixed-algorithm GEMM handle. Rebuild the library and **all** Qwen worker binaries together; older libraries are rejected. Current-branch GPU regression is required before release. @@ -66,7 +78,7 @@ Torch 2.14 cumulative-decay association, and a final BF16 output cast. Each call starts with zero state; prefix continuation is rejected by the model adapter. The context drains its last stream before destruction, and is retired before that stream. Rebuild the library and all consuming workers together for the -current ABI 8, including this optional group. The independent A800 qualification +current ABI 9, including this optional group. The independent A800 qualification at `27f77e49868d84b091aa804f02ccb4d3474a9a00` used ABI 6; the earlier cumulative campaign used ABI 8. Neither campaign qualifies the current integrated source. Builds now link both cuBLAS and cuBLASLt. Every consumer of the rebuilt shared diff --git a/src/backends/cuda/qwen3_5/gemm.cu b/src/backends/cuda/qwen3_5/gemm.cu index b7f8d63d..d761f862 100644 --- a/src/backends/cuda/qwen3_5/gemm.cu +++ b/src/backends/cuda/qwen3_5/gemm.cu @@ -6,7 +6,14 @@ // // Each shape uses cuBLASLt's first heuristic choice, excluding split-K reductions that // accumulate into the output in place, since their order, and so the rounding, is not -// fixed. +// fixed (vision's FP32 and biased GEMMs exclude all split-K). That choice depends on M, +// so a row's result can change with the number of rows in the call. A handle from +// cs1_gemm_create_fixed instead keeps one algorithm per weight shape (N, K, ldy, and for +// vision's GEMMs the data type and bias) for every M: the heuristic's first choice at a +// reference M among algorithms without split-K. Each output row then takes the same +// path whatever the other rows, so its result does not depend on M or on its row index. +// cuBLASLt does not document this; tests/qwen3_5/kernels.rs checks it on the GPU it +// runs on. #include #include @@ -29,7 +36,9 @@ struct Gemm { void* workspace = nullptr; size_t workspace_bytes = 0; bool reference = false; + int reference_m = 0; // > 0: one algorithm per weight shape, chosen at this M std::map plans; + std::map, cublasLtMatmulAlgo_t> fixed; // N, K, ldy, FP32, bias }; int status(cublasStatus_t s) { return s == CUBLAS_STATUS_SUCCESS ? 0 : 1000 + (int)s; } @@ -59,15 +68,13 @@ int describe(int M, int N, int K, int ldy, Plan& p, bool fp32, bool bias) { return 0; } -// The heuristic's first choice. Vision disables all split-K to avoid BF16 -// intermediate reductions; existing language GEMMs exclude only in-place reductions. -int first_choice(Gemm& g, Plan& p, bool vision) { +// The heuristic's first choice among the given reduction schemes. +int first_choice(Gemm& g, Plan& p, uint32_t schemes) { cublasLtMatmulPreference_t pref; cublasStatus_t s = cublasLtMatmulPreferenceCreate(&pref); if (s != CUBLAS_STATUS_SUCCESS) return status(s); cublasLtMatmulPreferenceSetAttribute(pref, CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &g.workspace_bytes, sizeof(g.workspace_bytes)); - const uint32_t schemes = g.reference ? CUBLASLT_REDUCTION_SCHEME_MASK : vision ? CUBLASLT_REDUCTION_SCHEME_NONE : (CUBLASLT_REDUCTION_SCHEME_MASK & ~CUBLASLT_REDUCTION_SCHEME_INPLACE); cublasLtMatmulPreferenceSetAttribute(pref, CUBLASLT_MATMUL_PREF_REDUCTION_SCHEME_MASK, &schemes, sizeof(schemes)); cublasLtMatmulHeuristicResult_t r{}; @@ -80,6 +87,40 @@ int first_choice(Gemm& g, Plan& p, bool vision) { return 0; } +// The weight shape's algorithm, chosen at the reference M on first use. An M it cannot +// serve is an error, not a reason to switch algorithms. +int fixed_choice(Gemm& g, int N, int K, int ldy, bool fp32, bool bias, Plan& p) { + const std::tuple key{N, K, ldy, fp32, bias}; + auto it = g.fixed.find(key); + if (it == g.fixed.end()) { + Plan r; + int rc = describe(g.reference_m, N, K, ldy, r, fp32, bias); + if (rc == 0) rc = first_choice(g, r, CUBLASLT_REDUCTION_SCHEME_NONE); + if (rc == 0) { + // Check that the choice really has no split-K reduction. + int32_t splits = 0; + uint32_t scheme = 0; + size_t written = 0; + const bool read = + cublasLtMatmulAlgoConfigGetAttribute(&r.algo, CUBLASLT_ALGO_CONFIG_SPLITK_NUM, &splits, + sizeof(splits), &written) == CUBLAS_STATUS_SUCCESS && + cublasLtMatmulAlgoConfigGetAttribute(&r.algo, CUBLASLT_ALGO_CONFIG_REDUCTION_SCHEME, &scheme, + sizeof(scheme), &written) == CUBLAS_STATUS_SUCCESS; + if (!read || splits != 1 || scheme != CUBLASLT_REDUCTION_SCHEME_NONE) + rc = status(CUBLAS_STATUS_NOT_SUPPORTED); + } + if (rc == 0) it = g.fixed.emplace(key, r.algo).first; + destroy(r); + if (rc != 0) return rc; + } + p.algo = it->second; + cublasLtMatmulHeuristicResult_t check{}; + const cublasStatus_t s = cublasLtMatmulAlgoCheck(g.handle, p.op, p.a, p.b, p.c, p.c, &p.algo, &check); + if (s != CUBLAS_STATUS_SUCCESS) return status(s); + if (check.workspaceSize > g.workspace_bytes) return status(CUBLAS_STATUS_NOT_SUPPORTED); + return 0; +} + // The plan for a shape, created on first use. int plan_for(Gemm& g, int M, int N, int K, int ldy, Plan*& out, bool fp32 = false, bool bias = false) { const Key key{M, N, K, ldy, fp32, bias}; @@ -90,7 +131,17 @@ int plan_for(Gemm& g, int M, int N, int K, int ldy, Plan*& out, bool fp32 = fals } Plan p; int rc = describe(M, N, K, ldy, p, fp32, bias); - if (rc == 0) rc = first_choice(g, p, fp32 || bias); + if (rc == 0 && g.reference_m > 0) { + rc = fixed_choice(g, N, K, ldy, fp32, bias, p); + } else if (rc == 0) { + // The reference handle allows every reduction scheme. Otherwise vision's FP32 and + // biased GEMMs exclude all split-K to avoid BF16 intermediate reductions, and the + // language GEMMs exclude only in-place reductions. + const uint32_t schemes = g.reference ? CUBLASLT_REDUCTION_SCHEME_MASK + : fp32 || bias ? CUBLASLT_REDUCTION_SCHEME_NONE + : CUBLASLT_REDUCTION_SCHEME_MASK & ~CUBLASLT_REDUCTION_SCHEME_INPLACE; + rc = first_choice(g, p, schemes); + } if (rc != 0) { destroy(p); return rc; @@ -101,6 +152,13 @@ int plan_for(Gemm& g, int M, int N, int K, int ldy, Plan*& out, bool fp32 = fals } // namespace +extern "C" void* cs1_gemm_create_fixed(size_t workspace_bytes, int reference_m) { + if (reference_m <= 0) return nullptr; + Gemm* g = static_cast(cs1_gemm_create(workspace_bytes)); + if (g) g->reference_m = reference_m; + return g; +} + extern "C" void* cs1_gemm_create(size_t workspace_bytes) { Gemm* g = new Gemm(); if (cublasLtCreate(&g->handle) != CUBLAS_STATUS_SUCCESS || diff --git a/src/backends/cuda/qwen3_5/ops.h b/src/backends/cuda/qwen3_5/ops.h index 568eabfe..1a4e6695 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 8 +#define CS1_ABI_VERSION 9 #ifdef __cplusplus extern "C" { @@ -129,6 +129,12 @@ int cs1_silu_mul(const void* gate_up, int ld, void* out, int T, int I, void* str // y [M, N] (rows of ldy) = x [M, K] * w [N, K]^T through cuBLASLt, float32 accumulation, // with cuBLASLt's first heuristic choice for each shape (see gemm.cu). void* cs1_gemm_create(size_t workspace_bytes); +// A handle that keeps one algorithm per weight shape (N, K, ldy) for every M, chosen at +// reference_m rows among algorithms without split-K, so that a row's result depends +// neither on M nor on its row index (checked by tests/qwen3_5/kernels.rs on the GPU it +// runs on). Null if reference_m <= 0 or setup fails. cs1_gemm with this handle returns a +// cuBLAS status for an M the algorithm cannot serve, rather than switching algorithms. +void* cs1_gemm_create_fixed(size_t workspace_bytes, int reference_m); void cs1_gemm_destroy(void* gemm); int cs1_gemm(void* gemm, const void* x, const void* w, void* y, int M, int N, int K, int ldy, 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 c80c2dac..e49e1457 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": 8, + "abi_version": 9, "status": "experimental", "sources": [ "ops.h", diff --git a/src/models/qwen3_5/native/src/cuda.rs b/src/models/qwen3_5/native/src/cuda.rs index cb29f545..3ac14eac 100644 --- a/src/models/qwen3_5/native/src/cuda.rs +++ b/src/models/qwen3_5/native/src/cuda.rs @@ -12,7 +12,7 @@ mod reference; pub use reference::Reference; /// `CS1_ABI_VERSION` in ops.h. -const ABI_VERSION: u32 = 8; +const ABI_VERSION: u32 = 9; pub const LIBRARY: &str = "libqwen3_5_cuda.so"; /// A `cudaStream_t`. @@ -193,6 +193,7 @@ api! { cs1_sigmoid_gate(x: *mut c_void, gate: *const c_void, n: usize, stream: Stream) -> c_int; cs1_silu_mul(gate_up: *const c_void, ld: c_int, out: *mut c_void, t: c_int, i: c_int, stream: Stream) -> c_int; cs1_gemm_create(workspace_bytes: usize) -> *mut c_void; + cs1_gemm_create_fixed(workspace_bytes: usize, reference_m: c_int) -> *mut c_void; cs1_gemm_destroy(gemm: *mut c_void); cs1_gemm( gemm: *mut c_void, x: *const c_void, w: *const c_void, y: *mut c_void, m: c_int, n: c_int, k: c_int, diff --git a/tests/qwen3_5/kernels.rs b/tests/qwen3_5/kernels.rs index 58a0a1ac..0ae5b406 100644 --- a/tests/qwen3_5/kernels.rs +++ b/tests/qwen3_5/kernels.rs @@ -1087,5 +1087,121 @@ fn gdn_prefill_two_phase_matches_one_shot_bit_for_bit() { } } +#[test] +#[ignore = "needs a GPU and CUA_S1_CUDA_LIB"] +fn fixed_gemm_rows_do_not_depend_on_m() { + let st = setup(); + // Every projection of the 4B, 9B and 27B backbones as (N, K): the GDN input and + // output, the attention input and output, and the MLP's gate|up and down. + let shapes = [ + (12352usize, 2560usize), + (2560, 4096), + (10240, 2560), + (18432, 2560), + (2560, 9216), + (12352, 4096), + (4096, 4096), + (10240, 4096), + (24576, 4096), + (4096, 12288), + (16480, 5120), + (5120, 6144), + (14336, 5120), + (34816, 5120), + (5120, 17408), + ]; + let max_m = 4096usize; + // SAFETY: plain handle creation; checked for null below. + let fixed = unsafe { (api().cs1_gemm_create_fixed)(32 << 20, 64) }; + assert!(!fixed.is_null()); + // SAFETY: a non-positive reference M is refused before anything is allocated. + assert!(unsafe { (api().cs1_gemm_create_fixed)(32 << 20, 0) }.is_null()); + let nan = vec![0xffu8; max_m * 34816 * 2]; + for (i, (n, k)) in shapes.into_iter().enumerate() { + let x = random(max_m * k, 70 + i as u64, 2.0); + let w = random(n * k, 90 + i as u64, 0.05); + let (xd, wd) = (to_device(&x, st), to_device(&w, st)); + let y = DeviceBuffer::new(max_m * n * 2).unwrap(); + // the bfloat16 output of m rows starting at input row `at`, over a NaN-filled buffer + let rows = |at: usize, m: usize| { + let mut out = vec![0u8; m * n * 2]; + // SAFETY: x holds max_m rows of k, w is [n, k], y has room for max_m rows of n, + // and at + m <= max_m. + unsafe { + cuda::upload(y.at(0), &nan[..max_m * n * 2], st).unwrap(); + check( + (api().cs1_gemm)( + fixed, + xd.at(at * k * 2), + wd.at(0), + y.at(0), + m as i32, + n as i32, + k as i32, + n as i32, + st, + ), + "fixed gemm", + ) + .unwrap(); + cuda::download(&mut out, y.at(0), st).unwrap(); + } + out + }; + let all = rows(0, max_m); + // every output was written (the comparisons below are byte for byte against this) + let (values, _) = all.as_chunks::<2>(); + assert!(values.iter().all(|b| !bf16::from_le_bytes(*b).is_nan())); + let row = |r: usize| &all[r * n * 2..(r + 1) * n * 2]; + // fewer rows, and rows that sit lower in the full call, as a branch's tokens do + for (at, m) in [ + (0usize, 1usize), + (0, 2), + (0, 3), + (0, 8), + (0, 63), + (0, 64), + (0, 65), + (0, 256), + (0, 1000), + (0, 1024), + (0, 3109), + (1, 2), + (37, 64), + (64, 65), + (1000, 200), + (3000, 1096), + ] { + assert!( + rows(at, m) == all[at * n * 2..(at + m) * n * 2], + "rows {at}..{} changed at {n} x {k}", + at + m + ); + } + // the first and last rows against float64, so the product itself is right + for r in [0, max_m - 1] { + let want: Vec = (0..n) + .map(|j| { + (0..k) + .map(|c| x[r * k + c].to_f64() * w[j * k + c].to_f64()) + .sum() + }) + .collect(); + let (got, _) = row(r).as_chunks::<2>(); + let scale = want.iter().fold(0f64, |m, v| m.max(v.abs())); + let worst = got + .iter() + .zip(&want) + .map(|(b, v)| (bf16::from_le_bytes(*b).to_f64() - v).abs()) + .fold(0f64, f64::max); + assert!( + worst <= 1e-2 * scale, + "{n} x {k}, row {r}: {worst} vs {scale}" + ); + } + } + // SAFETY: created above and not destroyed before. + unsafe { (api().cs1_gemm_destroy)(fixed) }; +} #[path = "vision_v2_kernels.rs"] mod vision_v2; diff --git a/tests/qwen3_5/vision_batch_fake.c b/tests/qwen3_5/vision_batch_fake.c index b4e60213..cf5d0aff 100644 --- a/tests/qwen3_5/vision_batch_fake.c +++ b/tests/qwen3_5/vision_batch_fake.c @@ -5,7 +5,7 @@ static int live, linear_count, attention_count, captures, launches, destroyed, f static int rows[16384], outs[16384], lens[16384]; static uintptr_t qkv_base; static uintptr_t qs[16384], ks[16384], vs[16384], os[16384]; -unsigned cs1_abi_version(void){return 8;} +unsigned cs1_abi_version(void){return 9;} const char *cs1_error_string(int c){(void)c;return "fake CUDA error";} int cs1_get_device(int *d){if(!d)return 1;*d=0;return 0;} int cs1_set_device(int d){(void)d;return 0;}