From d717e497fac143cae37f9aa8c18bab00e3497c8b Mon Sep 17 00:00:00 2001 From: Tianyao Wu Date: Tue, 6 Oct 2026 13:48:45 +0800 Subject: [PATCH 1/2] qwen3_5: add a GEMM handle with one algorithm per weight shape cuBLASLt's heuristic picks an algorithm per M, so a row's result can change with the number of rows in the call, and a prefix and its branch run separately round differently from one pass over the same tokens. Add cs1_gemm_create_fixed: one algorithm per weight shape (N, K, ldy), the heuristic's first choice at a reference M among algorithms without split-K, for every M. The default handle is unchanged. Bump the ABI to 7 and add a GPU test that each projection's rows match across M. --- docs/supported-models.md | 2 +- recipe/cua_s1/native.md | 2 +- recipe/jev_vl/README.md | 4 +- recipe/open_jev/native.md | 2 +- src/backends/cuda/contract.md | 4 +- src/backends/cuda/qwen3_5/README.md | 19 ++++- src/backends/cuda/qwen3_5/gemm.cu | 67 +++++++++++++-- src/backends/cuda/qwen3_5/ops.h | 8 +- src/models/qwen3_5/native/src/cuda.rs | 3 +- tests/qwen3_5/kernels.rs | 117 ++++++++++++++++++++++++++ 10 files changed, 210 insertions(+), 18 deletions(-) diff --git a/docs/supported-models.md b/docs/supported-models.md index 34d7b465..55249138 100644 --- a/docs/supported-models.md +++ b/docs/supported-models.md @@ -17,7 +17,7 @@ Models that are being added are also tracked in issues labeled [new model](https | 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 | | 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 7 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 | - **Validated:** covered by the recipe on `main` or by the checks in the linked merged pull request. diff --git a/recipe/cua_s1/native.md b/recipe/cua_s1/native.md index 83322b25..55e43929 100644 --- a/recipe/cua_s1/native.md +++ b/recipe/cua_s1/native.md @@ -28,7 +28,7 @@ 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 6). +the eager control. Rebuild both the worker and CUDA library together (ABI 7). If capture fails, the worker returns the completed eager result and disables Graph capture/replay for its remaining lifetime, logging the failure to stderr. diff --git a/recipe/jev_vl/README.md b/recipe/jev_vl/README.md index 6304c364..03ddfe1e 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 7** 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 7 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..331f0d1e 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 7, 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/contract.md b/src/backends/cuda/contract.md index 41301af1..e8ccade6 100644 --- a/src/backends/cuda/contract.md +++ b/src/backends/cuda/contract.md @@ -49,7 +49,7 @@ change. Schema: ```json { "name": "qwen3_5", - "abi_version": 4, + "abi_version": 7, "status": "validated", "sources": ["common.cuh", "mma.cuh", "ops.h", "norm.cu", "elementwise.cu", "attention.cu", "gdn_prefill.cu", "gemm.cu", "runtime.cu"], @@ -136,7 +136,7 @@ reference for the shape. The contract fixes four points: whose value it does not know. The manifest's `abi_version` is **that number**, not a version of this document, and the checker reads the macro out of the declared sources and requires the two to agree. `#19`'s `ops.h` says - `CS1_ABI_VERSION 4` today, so a `qwen3_5` manifest declares 4. + `CS1_ABI_VERSION 7` today, so a `qwen3_5` manifest declares 7. Two backends may declare different values. They are independent libraries, and the repository layout says CUDA and Metal implementations need not share diff --git a/src/backends/cuda/qwen3_5/README.md b/src/backends/cuda/qwen3_5/README.md index 63825533..e1745051 100644 --- a/src/backends/cuda/qwen3_5/README.md +++ b/src/backends/cuda/qwen3_5/README.md @@ -11,8 +11,20 @@ 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 workers together for ABI version 7, which +includes vision, CUDA Graph, gated attention, prefix-continuation entry points and +the fixed-algorithm GEMM handle below. + +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. @@ -24,6 +36,7 @@ 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 7 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. diff --git a/src/backends/cuda/qwen3_5/gemm.cu b/src/backends/cuda/qwen3_5/gemm.cu index 154ad698..f0febca4 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 @@ -28,7 +35,9 @@ struct Gemm { cublasLtHandle_t handle = nullptr; void* workspace = nullptr; size_t workspace_bytes = 0; + 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; } @@ -58,15 +67,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 = 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{}; @@ -79,6 +86,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}; @@ -89,7 +130,14 @@ 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) { + // Vision's FP32 and biased GEMMs exclude all split-K to avoid BF16 intermediate + // reductions; the language GEMMs exclude only in-place reductions. + rc = first_choice(g, p, fp32 || bias ? CUBLASLT_REDUCTION_SCHEME_NONE + : CUBLASLT_REDUCTION_SCHEME_MASK & ~CUBLASLT_REDUCTION_SCHEME_INPLACE); + } if (rc != 0) { destroy(p); return rc; @@ -100,6 +148,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 50d3618a..5a4d03cd 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 7 #ifdef __cplusplus extern "C" { @@ -118,6 +118,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/models/qwen3_5/native/src/cuda.rs b/src/models/qwen3_5/native/src/cuda.rs index 5e8141de..b9194e8d 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 = 7; pub const LIBRARY: &str = "libqwen3_5_cuda.so"; /// A `cudaStream_t`. @@ -125,6 +125,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 7a217c81..cce51549 100644 --- a/tests/qwen3_5/kernels.rs +++ b/tests/qwen3_5/kernels.rs @@ -1086,3 +1086,120 @@ 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) }; +} From d8cd04fe77bf4ab282de3ac3ee66555e13d6c9c6 Mon Sep 17 00:00:00 2001 From: Tianyao Wu Date: Sun, 11 Oct 2026 22:16:39 +0800 Subject: [PATCH 2/2] qwen3_5: declare ABI 7 in the backend manifest --- src/backends/cuda/qwen3_5/qwen3_5.backend.json | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/backends/cuda/qwen3_5/qwen3_5.backend.json b/src/backends/cuda/qwen3_5/qwen3_5.backend.json index 37484cef..cf06c19d 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": 7, "status": "experimental", "sources": [ "ops.h",