From 98d0465de57d4d67c5a49f00cc5fb95526df5360 Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Sun, 11 Oct 2026 11:00:25 +0800 Subject: [PATCH 1/2] feat(cua-s1): keep vision features resident for language execution --- src/backends/cuda/qwen3_5/elementwise.cu | 16 +++ src/backends/cuda/qwen3_5/ops.h | 5 +- src/backends/cuda/qwen3_5/runtime.cu | 2 + src/models/cua_s1/native/src/vision/model.rs | 66 ++++++++-- src/models/cua_s1/native/src/vision_engine.rs | 109 +++++++++++++-- src/models/qwen3_5/native/src/cuda.rs | 124 +++++++++++++++++- src/models/qwen3_5/native/src/inputs.rs | 72 +++++++++- src/models/qwen3_5/native/src/model.rs | 56 +++++++- 8 files changed, 415 insertions(+), 35 deletions(-) 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 50d3618a..e29704dc 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/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.run(image, Some(&mut callback)) } + /// Synchronized resident merger output, borrowed from geometry-specific scratch. + /// Holding the view prevents another forward or scratch retirement on this model. + pub fn forward_device(&mut self, image: &ProcessedImage) -> Result> { + let result: Result = (|| { + let (n, capture) = self.enqueue_image(image, &mut None)?; + self.synchronize()?; + if capture { + self.capture_current(); + } + Ok(n) + })(); + // Drain errors while every allocation/Graph referenced by queued work lives. + if result.is_err() { + self.synchronize()?; + } + let n = result?; + self.scratch.as_ref().unwrap().out.bf16_matrix(n / 4, 2560) + } + fn run(&mut self, image: &ProcessedImage, mut callback: Trace<'_>) -> Result> { + let result = (|| { + let (n, capture) = self.enqueue_image(image, &mut callback)?; + // This read is the host path's completion boundary, as before. + let output = self.read(&self.scratch.as_ref().unwrap().out, n / 4 * 2560)?; + if let Some(callback) = callback.as_mut() { + callback("merger.output", &output)?; + } else if capture { + self.capture_current(); + } + Ok(output) + })(); + if result.is_err() { + self.synchronize()?; + } + result + } + + fn capture_current(&mut self) { + let scratch = self.scratch.as_ref().unwrap(); + // Warmup is completed before capture. Recording never reruns the math. + let captured = cuda::Graph::capture(self.stream.0, || self.enqueue(scratch, &mut None)); + self.finish_capture(captured); + } + + // Enqueue only; each caller supplies its completion boundary before capture. + fn enqueue_image( + &mut self, + image: &ProcessedImage, + callback: &mut Trace<'_>, + ) -> Result<(usize, bool)> { cuda::set_device(0)?; let geo = if self .scratch @@ -282,19 +331,10 @@ impl VisionModel { { graph.launch(self.stream.0)?; graph_trace("replayed", scratch.grid); - return self.read(&scratch.out, n / 4 * 2560); - } - self.enqueue(scratch, &mut callback)?; - let result = self.read(&scratch.out, n / 4 * 2560)?; - if let Some(callback) = callback.as_mut() { - callback("merger.output", &result)?; - } else if self.graph_enabled { - // The eager read completed warmup. Capture does not execute the operators; - // return the already completed result even if recording fails. - let captured = cuda::Graph::capture(self.stream.0, || self.enqueue(scratch, &mut None)); - self.finish_capture(captured); + return Ok((n, false)); } - Ok(result) + self.enqueue(scratch, callback)?; + Ok((n, callback.is_none() && self.graph_enabled)) } fn finish_capture(&mut self, captured: Result) { match captured { diff --git a/src/models/cua_s1/native/src/vision_engine.rs b/src/models/cua_s1/native/src/vision_engine.rs index 9d621767..0365de90 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", diff --git a/src/models/qwen3_5/native/src/cuda.rs b/src/models/qwen3_5/native/src/cuda.rs index 5e8141de..39b12d94 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`. @@ -62,6 +62,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; @@ -195,6 +197,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. @@ -203,12 +206,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. @@ -218,6 +232,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() { diff --git a/src/models/qwen3_5/native/src/inputs.rs b/src/models/qwen3_5/native/src/inputs.rs index bf719eca..f67d45f0 100644 --- a/src/models/qwen3_5/native/src/inputs.rs +++ b/src/models/qwen3_5/native/src/inputs.rs @@ -1,5 +1,6 @@ //! Batch-one, unpadded inputs at the adapted-vision / language-model boundary. +use crate::cuda::DeviceBf16Matrix; use anyhow::{Result, ensure}; use half::bf16; @@ -21,6 +22,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 +62,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 +85,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 304c3083..d434ba7c 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; @@ -849,6 +849,60 @@ impl Model { self.last_hidden(s, t) } + /// 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, From f548d01c94698a70d82d49056f8ee9730dc231bd Mon Sep 17 00:00:00 2001 From: levius <2114377220@qq.com> Date: Sun, 11 Oct 2026 13:00:00 +0800 Subject: [PATCH 2/2] test(cua-s1): cover resident feature validation and recovery --- recipe/cua_s1/native_vision_graph.md | 89 +++++++- src/backends/cuda/qwen3_5/README.md | 9 +- src/models/cua_s1/native/Cargo.toml | 4 + src/models/cua_s1/native/src/vision_engine.rs | 8 + src/models/qwen3_5/native/Cargo.toml | 4 + src/models/qwen3_5/native/src/cuda.rs | 4 + src/models/qwen3_5/native/src/inputs.rs | 4 + tests/cua_s1/device_controls.rs | 46 ++++ tests/cua_s1/device_copy_failure_gpu.rs | 42 ++++ tests/cua_s1/device_features_gpu.rs | 202 ++++++++++++++++++ tests/cua_s1/fail_device_copy.cpp | 49 +++++ tests/cua_s1/vision_graph.rs | 61 ++++++ tests/qwen3_5/device_cuda.rs | 33 +++ tests/qwen3_5/device_features.rs | 26 +++ tests/qwen3_5/device_inputs.rs | 57 +++++ 15 files changed, 635 insertions(+), 3 deletions(-) create mode 100644 tests/cua_s1/device_controls.rs create mode 100644 tests/cua_s1/device_copy_failure_gpu.rs create mode 100644 tests/cua_s1/device_features_gpu.rs create mode 100644 tests/cua_s1/fail_device_copy.cpp create mode 100644 tests/qwen3_5/device_cuda.rs create mode 100644 tests/qwen3_5/device_features.rs create mode 100644 tests/qwen3_5/device_inputs.rs diff --git a/recipe/cua_s1/native_vision_graph.md b/recipe/cua_s1/native_vision_graph.md index e6fcc71c..ac50c38b 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,6 +51,91 @@ 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, while multimodal language execution remains eager in both modes. +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-cua-s1-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-cua-s1-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 diff --git a/src/backends/cuda/qwen3_5/README.md b/src/backends/cuda/qwen3_5/README.md index 63825533..b9002ee1 100644 --- a/src/backends/cuda/qwen3_5/README.md +++ b/src/backends/cuda/qwen3_5/README.md @@ -11,8 +11,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/models/cua_s1/native/Cargo.toml b/src/models/cua_s1/native/Cargo.toml index b06b6dc4..b489947d 100644 --- a/src/models/cua_s1/native/Cargo.toml +++ b/src/models/cua_s1/native/Cargo.toml @@ -34,3 +34,7 @@ path = "src/vision_main.rs" [[example]] name = "vision_graph_bench" path = "../../../../tests/cua_s1/vision_graph_bench.rs" + +[[test]] +name = "device_controls" +path = "../../../../tests/cua_s1/device_controls.rs" diff --git a/src/models/cua_s1/native/src/vision_engine.rs b/src/models/cua_s1/native/src/vision_engine.rs index 0365de90..69ff501e 100644 --- a/src/models/cua_s1/native/src/vision_engine.rs +++ b/src/models/cua_s1/native/src/vision_engine.rs @@ -312,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 57c2428f..e7e95993 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 39b12d94..994d0039 100644 --- a/src/models/qwen3_5/native/src/cuda.rs +++ b/src/models/qwen3_5/native/src/cuda.rs @@ -471,3 +471,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 f67d45f0..f43b8380 100644 --- a/src/models/qwen3_5/native/src/inputs.rs +++ b/src/models/qwen3_5/native/src/inputs.rs @@ -4,6 +4,10 @@ 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. 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/cua_s1/vision_graph.rs b/tests/cua_s1/vision_graph.rs index ae3e737a..4a345904 100644 --- a/tests/cua_s1/vision_graph.rs +++ b/tests/cua_s1/vision_graph.rs @@ -131,3 +131,64 @@ fn vision_capture_failure_keeps_eager_result_and_disables_replay() { assert!(model.scratch.as_ref().unwrap().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.as_ref().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.as_ref().unwrap().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 + ); +} 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(); +}