From 149e9ed028ce0de9ac8aaf8e276022f6227ff747 Mon Sep 17 00:00:00 2001 From: rsasaki0109 Date: Mon, 12 Oct 2026 10:11:26 +0900 Subject: [PATCH] Validate GPU AI packing before work allocation --- CHANGELOG.md | 4 + .../src/image/ai_pack_benchmark.rs | 104 +++++ .../src/image/ai_pack_tests.rs | 414 ++++++++++++++++++ crates/spatialrust-gpu/src/image/ai_tensor.rs | 105 ++++- .../fixtures/ai_pack_channels_experiment.wgsl | 22 + .../tests/fixtures/ai_pack_previous.wgsl | 36 ++ .../ai_pack_specialized_experiment.wgsl | 37 ++ docs/ARCHITECTURE.md | 8 + docs/FEATURE_MATRIX.md | 5 + docs/ROADMAP.md | 2 + notes/2026-10-12_validated_gpu_ai_packing.md | 197 +++++++++ 11 files changed, 921 insertions(+), 13 deletions(-) create mode 100644 crates/spatialrust-gpu/src/image/ai_pack_benchmark.rs create mode 100644 crates/spatialrust-gpu/src/image/ai_pack_tests.rs create mode 100644 crates/spatialrust-gpu/tests/fixtures/ai_pack_channels_experiment.wgsl create mode 100644 crates/spatialrust-gpu/tests/fixtures/ai_pack_previous.wgsl create mode 100644 crates/spatialrust-gpu/tests/fixtures/ai_pack_specialized_experiment.wgsl create mode 100644 notes/2026-10-12_validated_gpu_ai_packing.md diff --git a/CHANGELOG.md b/CHANGELOG.md index 0e10f0e4..d93eb1b3 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -21,6 +21,10 @@ removed no sooner than the next major (see `docs/API_STABILITY.md`). ### Changed +- GPU CHW image packing validates output indices, buffer/binding/dispatch limits + and profiler ownership/capacity before work allocation or pipeline setup. + The original f32 normalization shader, CHW order, one-pass receipts and + explicit transfers are preserved. - GPU RGB-to-gray supports images beyond one dispatch axis by distributing flat 256-lane workgroups across bounded rows. Single-row work retains its original shader; dimensions, counters, device limits and profiler capacity diff --git a/crates/spatialrust-gpu/src/image/ai_pack_benchmark.rs b/crates/spatialrust-gpu/src/image/ai_pack_benchmark.rs new file mode 100644 index 00000000..08655658 --- /dev/null +++ b/crates/spatialrust-gpu/src/image/ai_pack_benchmark.rs @@ -0,0 +1,104 @@ +use super::*; + +fn measure( + runtime: &WgpuRuntime, + cache: &AiPackPipeline, + previous: &wgpu::ComputePipeline, + w: usize, + h: usize, +) { + let image = fixture::(w, h); + let source = GpuImage::upload_u8(runtime, image.view()).unwrap(); + let plan = pack_plan(w as u32, h as u32, C as u32, &runtime.device().limits()).unwrap(); + let legacy = force_legacy(plan, w as u32, h as u32); + let scale = 1.0 / 255.0; + let mean = [0.485, 0.456, 0.406, 0.5]; + let std = [0.229, 0.224, 0.225, 0.75]; + let mut reference = + pack_dispatch(runtime, &source, scale, mean, std, legacy, cache, previous, None).unwrap(); + let expected = reference.readback_f32(runtime).unwrap(); + reference.recycle(runtime); + let mut stage: [Vec; 2] = std::array::from_fn(|_| Vec::new()); + let mut wall: [Vec; 2] = std::array::from_fn(|_| Vec::new()); + for round in 0..22 { + for offset in 0..2 { + let i = (round + offset) % 2; + let mut profiler = crate::GpuStageProfiler::new(runtime, 1).unwrap(); + let start = std::time::Instant::now(); + let mut output = if i == 0 { + source.validate_runtime(runtime).unwrap(); + pack_plan(w as u32, h as u32, C as u32, &runtime.device().limits()).unwrap(); + profiler.validate(runtime, 1).unwrap(); + pack_dispatch( + runtime, + &source, + scale, + mean, + std, + legacy, + cache, + previous, + Some(&mut profiler), + ) + .unwrap() + } else { + pack_ai_chw_gpu_with_profiler( + runtime, + &source, + scale, + mean, + std, + Some(&mut profiler), + ) + .unwrap() + }; + runtime.wait_idle(); + let ms = start.elapsed().as_secs_f64() * 1000.0; + let report = profiler.collect(runtime).unwrap(); + assert_eq!(report.stages.len(), 1); + assert_eq!(report.device_to_host_bytes, 16); + assert_eq!(output.shape(), [C as u32, h as u32, w as u32]); + assert_eq!(output.receipt().device_to_host_bytes(), 0); + assert_eq!(output.receipt().gpu_to_gpu_bytes(), (w * h * C * 4) as u64); + assert_words(&output.readback_f32(runtime).unwrap(), &expected); + output.recycle(runtime); + if round >= 2 { + stage[i].push(report.stages[0].elapsed_ns / 1000.0); + wall[i].push(ms); + } + } + } + eprintln!("AI {{\"width\":{w},\"height\":{h},\"channels\":{C},\"previous_us\":{:?},\"current_us\":{:?},\"previous_ms\":{:?},\"current_ms\":{:?}}}",stage[0],stage[1],wall[0],wall[1]); + source.recycle(runtime); +} + +#[test] +#[ignore = "explicit native paired resident GPU AI packing diagnostic"] +fn paired_frozen_and_validated_ai_pack() { + let runtime = + WgpuRuntime::new_headless_with_timestamps(crate::WgpuPowerPreference::HighPerformance) + .unwrap(); + let cache = + runtime.image_ai_pack_pipeline.get_or_init(|| create_ai_pack_pipeline(runtime.device())); + let previous = create_probe_pipeline( + runtime.device(), + &cache.layout, + include_str!("../../tests/fixtures/ai_pack_previous.wgsl"), + None, + ); + eprintln!("backend={}; device_type={}; 2 warmups + 20 alternating rounds; shared host executor; includes validation/output/uniform allocation/encoding/submit/wait; excludes upload/pipeline setup/profiler creation/query collection/readback/checks/recycle",runtime.adapter_info().backend,runtime.adapter_info().device_type); + for (w, h) in [ + (17, 19), + (128, 128), + (256, 256), + (640, 480), + (1920, 1080), + (3840, 2160), + (17, 4097), + (4097, 17), + ] { + measure::<1>(&runtime, cache, &previous, w, h); + measure::<3>(&runtime, cache, &previous, w, h); + measure::<4>(&runtime, cache, &previous, w, h); + } +} diff --git a/crates/spatialrust-gpu/src/image/ai_pack_tests.rs b/crates/spatialrust-gpu/src/image/ai_pack_tests.rs new file mode 100644 index 00000000..808dd0e6 --- /dev/null +++ b/crates/spatialrust-gpu/src/image/ai_pack_tests.rs @@ -0,0 +1,414 @@ +use super::*; +use spatialrust_image::{Image, ImageRegion, ImageView}; + +pub(super) fn fixture(w: usize, h: usize) -> Image { + Image::try_new(w, h, (0..w * h * C).map(|i| (i * 37 + i / C * 11 + i % C * 53) as u8).collect()) + .unwrap() +} + +pub(super) fn assert_words(actual: &[f32], expected: &[f32]) { + assert_eq!(actual.len(), expected.len()); + for (i, (&a, &b)) in actual.iter().zip(expected).enumerate() { + assert_eq!(a.to_bits(), b.to_bits(), "word {i}"); + } +} + +fn cpu(data: &[u8], scale: f32, mean: [f32; 4], std: [f32; 4]) -> Vec { + (0..C) + .flat_map(|c| { + data.chunks_exact(C).map(move |pixel| (f32::from(pixel[c]) * scale - mean[c]) / std[c]) + }) + .collect() +} + +pub(super) fn force_legacy(mut plan: AiPackPlan, w: u32, h: u32) -> AiPackPlan { + plan.groups = [w.div_ceil(16), h.div_ceil(16), 1]; + plan +} + +fn check_channel_cases( + runtime: &WgpuRuntime, + cache: &AiPackPipeline, + previous: &wgpu::ComputePipeline, + parallel: &wgpu::ComputePipeline, +) -> usize { + let specialized_pipeline = create_probe_pipeline( + runtime.device(), + &cache.layout, + include_str!("../../tests/fixtures/ai_pack_specialized_experiment.wgsl"), + Some(C as u32), + ); + let mut words = 0; + for (w, h) in [ + (1, 1), + (1, 17), + (17, 1), + (7, 9), + (8, 8), + (9, 17), + (15, 19), + (16, 16), + (17, 19), + (255, 257), + (256, 256), + (257, 259), + ] { + let image = fixture::(w, h); + let source = GpuImage::upload_u8(runtime, image.view()).unwrap(); + let presets = [ + (1.0, [0.0; 4], [1.0; 4]), + (0.25, [-0.0, 4.0, -8.0, 12.0], [0.5, -2.0, 4.0, 0.25]), + (1.0 / 255.0, [0.485, 0.456, 0.406, 0.5], [0.229, 0.224, 0.225, 0.75]), + ( + f32::MAX / 4.0, + [f32::MAX / 2.0, -f32::MAX / 2.0, 0.0, 1.0], + [f32::MIN_POSITIVE, -f32::MIN_POSITIVE, 1.0, -1.0], + ), + ]; + for (p, (scale, mean, std)) in presets.into_iter().enumerate() { + let plan = pack_plan(w as u32, h as u32, C as u32, &runtime.device().limits()).unwrap(); + let mut old = pack_dispatch( + runtime, + &source, + scale, + mean, + std, + force_legacy(plan, w as u32, h as u32), + cache, + previous, + None, + ) + .unwrap(); + let expected = old.readback_f32(runtime).unwrap(); + old.recycle(runtime); + if p < 2 { + assert_words(&expected, &cpu::(image.as_slice(), scale, mean, std)); + } + let forced = + AiPackPlan { groups: [(w as u32).div_ceil(8), (h as u32).div_ceil(8), 1], ..plan }; + let tensors = [ + pack_dispatch( + runtime, + &source, + scale, + mean, + std, + force_legacy(plan, w as u32, h as u32), + cache, + &specialized_pipeline, + None, + ) + .unwrap(), + pack_dispatch(runtime, &source, scale, mean, std, forced, cache, parallel, None) + .unwrap(), + pack_ai_chw_gpu(runtime, &source, scale, mean, std).unwrap(), + ]; + for mut output in tensors { + assert_eq!(output.shape(), [C as u32, h as u32, w as u32]); + assert_eq!(output.receipt().stages(), ["upload_u8_texture", "pack_ai_chw_gpu"]); + assert_eq!(output.receipt().gpu_to_gpu_bytes(), (w * h * C * 4) as u64); + assert_eq!(output.receipt().device_to_host_bytes(), 0); + assert_words(&output.readback_f32(runtime).unwrap(), &expected); + assert_eq!(output.receipt().device_to_host_bytes(), (w * h * C * 4) as u64); + output.recycle(runtime); + words += w * h * C; + } + } + source.recycle(runtime); + } + words +} + +#[test] +fn all_channels_and_normalizations_match_frozen_words_and_independent_exact_math() { + let runtime = WgpuRuntime::new_headless().unwrap(); + let cache = + runtime.image_ai_pack_pipeline.get_or_init(|| create_ai_pack_pipeline(runtime.device())); + let previous = create_probe_pipeline( + runtime.device(), + &cache.layout, + include_str!("../../tests/fixtures/ai_pack_previous.wgsl"), + None, + ); + let parallel = create_probe_pipeline( + runtime.device(), + &cache.layout, + include_str!("../../tests/fixtures/ai_pack_channels_experiment.wgsl"), + None, + ); + let words = check_channel_cases::<1>(&runtime, cache, &previous, ¶llel) + + check_channel_cases::<2>(&runtime, cache, &previous, ¶llel) + + check_channel_cases::<3>(&runtime, cache, &previous, ¶llel) + + check_channel_cases::<4>(&runtime, cache, &previous, ¶llel); + eprintln!("Verified {words} exact output words across 192 channel/shape/normalization cases and forced/current paths."); +} + +#[test] +fn buffer_index_dispatch_and_optional_channel_limits_are_checked() { + let limits = wgpu::Limits::default(); + for (w, h, c) in + [(0, 1, 1), (1, 0, 1), (1, 1, 0), (1, 1, 5), (u32::MAX, 1, 1), (1, u32::MAX, 1)] + { + assert!(pack_plan(w, h, c, &limits).is_err()); + } + for field in 0..14 { + let mut changed = limits.clone(); + match field { + 0 => changed.max_compute_invocations_per_workgroup = 255, + 1 => changed.max_compute_workgroup_size_x = 15, + 2 => changed.max_compute_workgroup_size_y = 15, + 3 => changed.max_compute_workgroup_size_z = 0, + 4 => changed.max_compute_workgroups_per_dimension = 0, + 5 => changed.max_sampled_textures_per_shader_stage = 0, + 6 => changed.max_storage_buffers_per_shader_stage = 0, + 7 => changed.max_uniform_buffers_per_shader_stage = 0, + 8 => changed.max_uniform_buffer_binding_size = 63, + 9 => changed.max_buffer_size = 63, + 10 => changed.max_bind_groups = 0, + 11 => changed.max_bindings_per_bind_group = 2, + 12 => changed.max_storage_buffer_binding_size = 63, + _ => changed.max_texture_dimension_2d = 16, + } + assert!(pack_plan(17, 19, 3, &changed).is_err(), "field {field}"); + } + let mut changed = limits.clone(); + changed.max_storage_buffer_binding_size = 1024; + assert_eq!(pack_plan(16, 16, 1, &changed).unwrap().byte_len, 1024); + assert!(pack_plan(17, 16, 1, &changed).is_err()); + changed.max_buffer_size = 1023; + assert!(pack_plan(16, 16, 1, &changed).is_err()); + let mut optional = limits.clone(); + optional.max_compute_workgroup_storage_size = 0; + optional.max_compute_workgroup_size_z = 1; + assert_eq!(pack_plan(256, 256, 4, &optional).unwrap().groups, [16, 16, 1]); + optional.max_compute_workgroup_size_z = 4; + optional.max_compute_workgroups_per_dimension = 64; + assert_eq!(pack_plan(1024, 64, 3, &optional).unwrap().groups, [64, 4, 1]); + assert!(pack_plan(255, 257, 3, &limits).is_ok()); + assert!(pack_plan(256, 256, 3, &limits).is_ok()); + assert!(pack_plan(256, 256, 2, &limits).is_ok()); + let mut huge = limits; + huge.max_texture_dimension_2d = u32::MAX; + huge.max_compute_workgroups_per_dimension = u32::MAX; + huge.max_buffer_size = u64::MAX; + huge.max_storage_buffer_binding_size = u32::MAX; + assert!(pack_plan(65536, 65536, 1, &huge).is_err()); + assert!(pack_plan(32768, 32768, 4, &huge).is_err()); + assert!(pack_plan(i32::MAX as u32 + 1, 1, 1, &huge).is_err()); + let boundary = pack_plan(16384, 16384, 3, &huge).unwrap(); + assert_eq!(boundary.byte_len, 3_221_225_472); + assert!(pack_plan(16384, 16384, 4, &huge).is_err()); +} + +#[test] +fn invalid_parameters_ownership_and_capacity_do_not_allocate_or_initialize() { + let runtime = + WgpuRuntime::new_headless_with_timestamps(crate::WgpuPowerPreference::HighPerformance) + .unwrap(); + let other = + WgpuRuntime::new_headless_with_timestamps(crate::WgpuPowerPreference::HighPerformance) + .unwrap(); + let source = GpuImage::upload_u8(&runtime, fixture::<3>(256, 256).view()).unwrap(); + let before = runtime.scratch_buffer_pool_stats(); + assert!(pack_ai_chw_gpu(&other, &source, 1.0, [0.0; 4], [1.0; 4]).is_err()); + for (scale, mean, std) in [ + (f32::NAN, [0.0; 4], [1.0; 4]), + (1.0, [f32::INFINITY; 4], [1.0; 4]), + (1.0, [0.0; 4], [0.0; 4]), + (1.0, [0.0; 4], [-0.0; 4]), + (1.0, [0.0; 4], [f32::NAN; 4]), + ] { + assert!(pack_ai_chw_gpu(&runtime, &source, scale, mean, std).is_err()); + } + let mut foreign = crate::GpuStageProfiler::new(&other, 1).unwrap(); + assert!(pack_ai_chw_gpu_with_profiler( + &runtime, + &source, + 1.0, + [0.0; 4], + [1.0; 4], + Some(&mut foreign) + ) + .is_err()); + assert_eq!(runtime.scratch_buffer_pool_stats(), before); + assert!( + runtime.image_ai_pack_pipeline.get().is_none() + && other.image_ai_pack_pipeline.get().is_none() + ); + let gray = GpuImage::upload_u8(&runtime, fixture::<1>(17, 19).view()).unwrap(); + let mut profiler = crate::GpuStageProfiler::new(&runtime, 1).unwrap(); + let first = crate::image::kernels::spatial::sobel_gpu_with_profiler( + &runtime, + &gray, + Some(&mut profiler), + ) + .unwrap(); + let before = runtime.scratch_buffer_pool_stats(); + assert!(pack_ai_chw_gpu_with_profiler( + &runtime, + &source, + 1.0, + [0.0; 4], + [1.0; 4], + Some(&mut profiler) + ) + .is_err()); + assert_eq!(runtime.scratch_buffer_pool_stats(), before); + assert!(runtime.image_ai_pack_pipeline.get().is_none()); + let report = profiler.collect(&runtime).unwrap(); + assert_eq!(report.stages.len(), 1); + assert_eq!(report.device_to_host_bytes, 16); + first.recycle(&runtime); + gray.recycle(&runtime); + source.recycle(&runtime); +} + +#[test] +fn cold_channel_concurrency_strided_roi_and_live_tensor_reuse_preserve_words() { + let runtime = WgpuRuntime::new_headless().unwrap(); + let stride = 273 * 3 + 13; + let data: Vec<_> = (0..stride * 269).map(|i| (i * 31 + 7) as u8).collect(); + let view = ImageView::::new(273, 269, stride, &data) + .unwrap() + .subview(ImageRegion::new(7, 3, 257, 259)) + .unwrap(); + let packed: Vec<_> = (0..259).flat_map(|y| view.row(y).unwrap().iter().copied()).collect(); + let expected = cpu::<3>(&packed, 0.25, [0.0; 4], [1.0; 4]); + let source = GpuImage::upload_u8(&runtime, view).unwrap(); + let outputs = std::thread::scope(|scope| { + let jobs: Vec<_> = (0..4) + .map(|_| { + let runtime = &runtime; + let source = &source; + scope.spawn(move || { + pack_ai_chw_gpu(runtime, source, 0.25, [0.0; 4], [1.0; 4]).unwrap() + }) + }) + .collect(); + jobs.into_iter().map(|j| j.join().unwrap()).collect::>() + }); + assert!(runtime.image_ai_pack_pipeline.get().is_some()); + source.recycle(&runtime); + let source = GpuImage::upload_u8( + &runtime, + Image::::try_new(257, 259, vec![255; 257 * 259 * 3]).unwrap().view(), + ) + .unwrap(); + for _ in 0..12 { + pack_ai_chw_gpu(&runtime, &source, 1.0, [0.0; 4], [1.0; 4]).unwrap().recycle(&runtime); + } + let runtime = Box::new(runtime); + for mut tensor in outputs { + assert_words(&tensor.readback_f32(&runtime).unwrap(), &expected); + tensor.recycle(&runtime); + } + let before = runtime.scratch_buffer_pool_stats(); + for _ in 0..12 { + pack_ai_chw_gpu(&runtime, &source, 1.0, [0.0; 4], [1.0; 4]).unwrap().recycle(&runtime); + } + assert_eq!(runtime.scratch_buffer_pool_stats().allocation_count, before.allocation_count); + runtime.set_scratch_buffer_pool_budget(0); + let mut tensor = pack_ai_chw_gpu(&runtime, &source, 1.0, [0.0; 4], [1.0; 4]).unwrap(); + assert!(tensor.readback_f32(&runtime).unwrap().iter().all(|&v| v == 255.0)); + assert!(tensor.readback_f32(&WgpuRuntime::new_headless().unwrap()).is_err()); + tensor.recycle(&runtime); + source.recycle(&runtime); + assert_eq!(runtime.scratch_buffer_pool_stats().cached_bytes, 0); +} + +#[test] +fn declared_output_limit_failure_happens_before_tensor_work() { + let runtime = WgpuRuntime::new_headless().unwrap(); + let tiny = GpuImage::upload_u8(&runtime, fixture::<4>(1, 1).view()).unwrap(); + let limits = runtime.device().limits(); + let width = limits.max_texture_dimension_2d.min(16_384); + let limit = limits.max_buffer_size.min(u64::from(limits.max_storage_buffer_binding_size)); + let height = (limit / (u64::from(width) * 16) + 1) as u32; + // A metadata-only negative probe avoids allocating a huge source image. + // The pack call must reject these dimensions before touching the texture. + let source = GpuImage::from_parts( + &runtime, + width, + height, + 4, + tiny.texture().clone(), + tiny.metadata(), + tiny.receipt().clone(), + ) + .unwrap(); + let before = runtime.scratch_buffer_pool_stats(); + assert!(pack_ai_chw_gpu(&runtime, &source, 1.0, [0.0; 4], [1.0; 4]).is_err()); + assert_eq!(runtime.scratch_buffer_pool_stats(), before); + assert!(runtime.image_ai_pack_pipeline.get().is_none()); + // The declared dimensions differ from its backing texture: never pool it. + drop(source); + tiny.recycle(&runtime); +} + +#[test] +fn full_large_resident_vision_chain_keeps_six_stages_and_explicit_transfers() { + let runtime = + WgpuRuntime::new_headless_with_timestamps(crate::WgpuPowerPreference::HighPerformance) + .unwrap(); + let image = Image::::try_new(4096, 4096, vec![128; 4096 * 4096 * 3]).unwrap(); + let source = GpuImage::upload_u8(&runtime, image.view()).unwrap(); + let options = crate::GpuVisionChainOptions { + width: 4096, + height: 4096, + blur_kernel: 1, + morphology_kernel: 1, + scale: 1.0, + mean: [0.0; 4], + std: [1.0; 4], + ..Default::default() + }; + let mut profiler = crate::GpuStageProfiler::new(&runtime, 6).unwrap(); + let mut tensor = + crate::run_gpu_vision_chain_profiled(&runtime, &source, options, &mut profiler).unwrap(); + assert_eq!(tensor.shape(), [1, 4096, 4096]); + assert_eq!(tensor.receipt().host_to_device_bytes(), 67_108_864); + assert_eq!(tensor.receipt().gpu_to_gpu_bytes(), 402_653_184); + assert_eq!(tensor.receipt().device_to_host_bytes(), 0); + let report = profiler.collect(&runtime).unwrap(); + assert_eq!(report.stages.len(), 6); + assert_eq!(report.device_to_host_bytes, 96); + assert!(tensor.readback_f32(&runtime).unwrap().iter().all(|v| v.to_bits() == 0)); + assert_eq!(tensor.receipt().device_to_host_bytes(), 67_108_864); + tensor.recycle(&runtime); + source.recycle(&runtime); + eprintln!("Verified 16777216 exact zero tensor words in a 4096x4096 six-stage resident chain, 96 explicit query bytes."); +} + +#[path = "ai_pack_benchmark.rs"] +mod benchmark; + +fn create_probe_pipeline( + device: &wgpu::Device, + layout: &wgpu::BindGroupLayout, + source: &str, + channels: Option, +) -> wgpu::ComputePipeline { + let constants = channels + .map(|c| std::collections::HashMap::from([("PACK_CHANNELS".to_owned(), f64::from(c))])) + .unwrap_or_default(); + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("ai-pack-probe-shader"), + source: wgpu::ShaderSource::Wgsl(source.into()), + }); + let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("gpu-ai-pack-layout"), + bind_group_layouts: &[layout], + push_constant_ranges: &[], + }); + device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("gpu-ai-pack-pipeline"), + layout: Some(&pipeline_layout), + module: &shader, + entry_point: Some("pack_ai_chw"), + compilation_options: wgpu::PipelineCompilationOptions { + constants: &constants, + ..Default::default() + }, + cache: None, + }) +} diff --git a/crates/spatialrust-gpu/src/image/ai_tensor.rs b/crates/spatialrust-gpu/src/image/ai_tensor.rs index 991cdbec..b4c87c53 100644 --- a/crates/spatialrust-gpu/src/image/ai_tensor.rs +++ b/crates/spatialrust-gpu/src/image/ai_tensor.rs @@ -9,6 +9,64 @@ use crate::WgpuRuntime; const WORKGROUP_X: u32 = 16; const WORKGROUP_Y: u32 = 16; +#[derive(Clone, Copy, Debug)] +struct AiPackPlan { + byte_len: u64, + groups: [u32; 3], +} + +fn pack_plan( + width: u32, + height: u32, + channels: u32, + limits: &wgpu::Limits, +) -> SpatialResult { + let invalid = || { + SpatialError::InvalidArgument( + "GPU AI packing requires positive dimensions, supported channels and buffer/compute/index limits".to_owned(), + ) + }; + let pixels = u64::from(width) * u64::from(height); + if width == 0 + || height == 0 + || !(1..=4).contains(&channels) + || width > limits.max_texture_dimension_2d + || height > limits.max_texture_dimension_2d + || width > i32::MAX as u32 + || height > i32::MAX as u32 + || pixels > u64::from(u32::MAX) + || limits.max_compute_invocations_per_workgroup < WORKGROUP_X * WORKGROUP_Y + || limits.max_compute_workgroup_size_x < WORKGROUP_X + || limits.max_compute_workgroup_size_y < WORKGROUP_Y + || limits.max_compute_workgroup_size_z == 0 + || width.div_ceil(WORKGROUP_X) > limits.max_compute_workgroups_per_dimension + || height.div_ceil(WORKGROUP_Y) > limits.max_compute_workgroups_per_dimension + || limits.max_sampled_textures_per_shader_stage == 0 + || limits.max_storage_buffers_per_shader_stage == 0 + || limits.max_uniform_buffers_per_shader_stage == 0 + || limits.max_uniform_buffer_binding_size < std::mem::size_of::() as u32 + || limits.max_buffer_size < std::mem::size_of::() as u64 + || limits.max_bind_groups == 0 + || limits.max_bindings_per_bind_group < 3 + { + return Err(invalid()); + } + let elements = pixels + .checked_mul(u64::from(channels)) + .filter(|&n| n <= u64::from(u32::MAX)) + .ok_or_else(invalid)?; + let byte_len = elements.checked_mul(4).ok_or_else(invalid)?; + if byte_len > limits.max_buffer_size + || byte_len > u64::from(limits.max_storage_buffer_binding_size) + { + return Err(invalid()); + } + Ok(AiPackPlan { + byte_len, + groups: [width.div_ceil(WORKGROUP_X), height.div_ceil(WORKGROUP_Y), 1], + }) +} + const SHADER: &str = r#" struct Params { width: u32, @@ -181,6 +239,9 @@ impl GpuAiTensor { } /// Packs a resident image into normalized planar CHW `f32` storage. +/// +/// Runtime, index, dispatch, binding and output-buffer limits are checked +/// before allocation or pipeline setup; valid shader math remains unchanged. pub fn pack_ai_chw_gpu( runtime: &WgpuRuntime, source: &GpuImage, @@ -197,7 +258,7 @@ pub(crate) fn pack_ai_chw_gpu_with_profiler( scale: f32, mean: [f32; 4], std: [f32; 4], - mut profiler: Option<&mut crate::GpuStageProfiler>, + profiler: Option<&mut crate::GpuStageProfiler>, ) -> SpatialResult { source.validate_runtime(runtime)?; if !scale.is_finite() @@ -208,9 +269,29 @@ pub(crate) fn pack_ai_chw_gpu_with_profiler( "AI packing scale/mean/std must be finite and std non-zero".to_owned(), )); } - let elements = - u64::from(source.width()) * u64::from(source.height()) * u64::from(source.channels()); - let byte_len = elements * std::mem::size_of::() as u64; + let plan = + pack_plan(source.width(), source.height(), source.channels(), &runtime.device().limits())?; + if let Some(profiler) = &profiler { + profiler.validate(runtime, 1)?; + } + let cache = + runtime.image_ai_pack_pipeline.get_or_init(|| create_ai_pack_pipeline(runtime.device())); + pack_dispatch(runtime, source, scale, mean, std, plan, cache, &cache.pipeline, profiler) +} + +#[allow(clippy::too_many_arguments)] +fn pack_dispatch( + runtime: &WgpuRuntime, + source: &GpuImage, + scale: f32, + mean: [f32; 4], + std: [f32; 4], + plan: AiPackPlan, + cache: &AiPackPipeline, + pipeline: &wgpu::ComputePipeline, + mut profiler: Option<&mut crate::GpuStageProfiler>, +) -> SpatialResult { + let byte_len = plan.byte_len; let output = runtime.buffer_pool().acquire_storage(runtime, "gpu-ai-chw-output", byte_len); let params = AiPackParams { width: source.width(), @@ -229,12 +310,10 @@ pub(crate) fn pack_ai_chw_gpu_with_profiler( mapped_at_creation: false, }); runtime.queue().write_buffer(&uniform, 0, bytemuck::bytes_of(¶ms)); - let pipeline = - runtime.image_ai_pack_pipeline.get_or_init(|| create_ai_pack_pipeline(runtime.device())); let source_view = source.view(); let bind_group = runtime.device().create_bind_group(&wgpu::BindGroupDescriptor { label: Some("gpu-ai-pack-bg"), - layout: &pipeline.layout, + layout: &cache.layout, entries: &[ wgpu::BindGroupEntry { binding: 0, resource: uniform.as_entire_binding() }, wgpu::BindGroupEntry { @@ -252,13 +331,9 @@ pub(crate) fn pack_ai_chw_gpu_with_profiler( label: Some("pack_ai_chw_gpu"), timestamp_writes: crate::profiling::timestamp_writes(&mut profiler, "pack_ai_chw_gpu"), }); - pass.set_pipeline(&pipeline.pipeline); + pass.set_pipeline(pipeline); pass.set_bind_group(0, &bind_group, &[]); - pass.dispatch_workgroups( - source.width().div_ceil(WORKGROUP_X), - source.height().div_ceil(WORKGROUP_Y), - 1, - ); + pass.dispatch_workgroups(plan.groups[0], plan.groups[1], plan.groups[2]); } runtime.queue().submit(Some(encoder.finish())); let mut receipt = source.receipt().clone(); @@ -273,3 +348,7 @@ pub(crate) fn pack_ai_chw_gpu_with_profiler( receipt, }) } + +#[cfg(all(test, not(target_arch = "wasm32")))] +#[path = "ai_pack_tests.rs"] +mod tests; diff --git a/crates/spatialrust-gpu/tests/fixtures/ai_pack_channels_experiment.wgsl b/crates/spatialrust-gpu/tests/fixtures/ai_pack_channels_experiment.wgsl new file mode 100644 index 00000000..64540d8a --- /dev/null +++ b/crates/spatialrust-gpu/tests/fixtures/ai_pack_channels_experiment.wgsl @@ -0,0 +1,22 @@ +struct Params { + width: u32, height: u32, channels: u32, _pad: u32, + scale: f32, mean0: f32, mean1: f32, mean2: f32, mean3: f32, + std0: f32, std1: f32, std2: f32, std3: f32, + _tail0: f32, _tail1: f32, _tail2: f32, +} +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var source_px: texture_2d; +@group(0) @binding(2) var output: array; + +@compute @workgroup_size(8, 8, 4) +fn pack_ai_chw(@builtin(global_invocation_id) gid: vec3) { + let channel = gid.z; + if (gid.x >= params.width || gid.y >= params.height || channel >= params.channels) { return; } + let pixel = textureLoad(source_px, vec2(gid.xy), 0); + let values = vec4(pixel); + let means = vec4(params.mean0, params.mean1, params.mean2, params.mean3); + let stds = vec4(params.std0, params.std1, params.std2, params.std3); + let index = gid.y * params.width + gid.x; + let plane = params.width * params.height; + output[channel * plane + index] = (values[channel] * params.scale - means[channel]) / stds[channel]; +} diff --git a/crates/spatialrust-gpu/tests/fixtures/ai_pack_previous.wgsl b/crates/spatialrust-gpu/tests/fixtures/ai_pack_previous.wgsl new file mode 100644 index 00000000..889f08b9 --- /dev/null +++ b/crates/spatialrust-gpu/tests/fixtures/ai_pack_previous.wgsl @@ -0,0 +1,36 @@ + +struct Params { + width: u32, + height: u32, + channels: u32, + _pad: u32, + scale: f32, + mean0: f32, + mean1: f32, + mean2: f32, + mean3: f32, + std0: f32, + std1: f32, + std2: f32, + std3: f32, + _tail0: f32, + _tail1: f32, + _tail2: f32, +} +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var source_px: texture_2d; +@group(0) @binding(2) var output: array; + +@compute @workgroup_size(16, 16) +fn pack_ai_chw(@builtin(global_invocation_id) gid: vec3) { + if (gid.x >= params.width || gid.y >= params.height) { return; } + let pixel = textureLoad(source_px, vec2(gid.xy), 0); + let values = vec4(pixel); + let means = vec4(params.mean0, params.mean1, params.mean2, params.mean3); + let stds = vec4(params.std0, params.std1, params.std2, params.std3); + let index = gid.y * params.width + gid.x; + let plane = params.width * params.height; + for (var channel = 0u; channel < params.channels; channel++) { + output[channel * plane + index] = (values[channel] * params.scale - means[channel]) / stds[channel]; + } +} diff --git a/crates/spatialrust-gpu/tests/fixtures/ai_pack_specialized_experiment.wgsl b/crates/spatialrust-gpu/tests/fixtures/ai_pack_specialized_experiment.wgsl new file mode 100644 index 00000000..cb517ea3 --- /dev/null +++ b/crates/spatialrust-gpu/tests/fixtures/ai_pack_specialized_experiment.wgsl @@ -0,0 +1,37 @@ +override PACK_CHANNELS: u32 = 4u; + +struct Params { + width: u32, + height: u32, + channels: u32, + _pad: u32, + scale: f32, + mean0: f32, + mean1: f32, + mean2: f32, + mean3: f32, + std0: f32, + std1: f32, + std2: f32, + std3: f32, + _tail0: f32, + _tail1: f32, + _tail2: f32, +} +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var source_px: texture_2d; +@group(0) @binding(2) var output: array; + +@compute @workgroup_size(16, 16) +fn pack_ai_chw(@builtin(global_invocation_id) gid: vec3) { + if (gid.x >= params.width || gid.y >= params.height) { return; } + let pixel = textureLoad(source_px, vec2(gid.xy), 0); + let values = vec4(pixel); + let means = vec4(params.mean0, params.mean1, params.mean2, params.mean3); + let stds = vec4(params.std0, params.std1, params.std2, params.std3); + let index = gid.y * params.width + gid.x; + let plane = params.width * params.height; + for (var channel = 0u; channel < PACK_CHANNELS; channel++) { + output[channel * plane + index] = (values[channel] * params.scale - means[channel]) / stds[channel]; + } +} diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md index fab4a69b..650aac16 100644 --- a/docs/ARCHITECTURE.md +++ b/docs/ARCHITECTURE.md @@ -83,6 +83,14 @@ Linear and multi-row pipelines are independently lazy within one runtime-owned family and share the existing 16-byte uniform/layout. Integer BT.601 output, gray metadata, one pass, four physical write bytes per pixel and 16 query bytes on explicit collection remain unchanged. There is no new scratch or payload copy. +GPU AI CHW packing validates positive dimensions, logical channels, signed texture +coordinates, u32 pixel/element indices, storage-binding/output-buffer sizes, +16x16 dispatch limits and the 64-byte uniform before allocation or cache setup. +Profile ownership and one free stage are checked before work begins. Finite +scale/mean and finite nonzero standard deviations retain their existing contract; +the original f32 normalization and CHW shader, single pass, storage retention +and explicit payload/query collection remain unchanged. Experimental channel-lane +and fixed-channel kernels remain native test fixtures after completed-call measurements. Gray GPU box blur retains the integer full-box mean with one final half-up rounding. Boxes with area at least 81 and at least 65,536 total pixel/sample visits use horizontal unrounded u32 sums and a vertical reduction in one diff --git a/docs/FEATURE_MATRIX.md b/docs/FEATURE_MATRIX.md index 3aa05125..2677e921 100644 --- a/docs/FEATURE_MATRIX.md +++ b/docs/FEATURE_MATRIX.md @@ -191,6 +191,11 @@ workspace because its build requires a Python toolchain. independently lazy; checked u32 pixel/index math, pre-allocation limits and profiler checks preserve exact BT.601 bytes, receipts and explicit transfers. See [bounded gray dispatch validation](../notes/2026-10-12_bounded_gpu_gray_dispatch.md). +- GPU CHW packing checks dimension/index/buffer/dispatch limits and profiler + ownership/capacity before allocation. Normalization math, output order, + one-pass receipts and explicit transfers are unchanged. Channel-lane and + fixed-channel optimizations remain test-only after completed-call comparisons. + See [AI packing validation and measurements](../notes/2026-10-12_validated_gpu_ai_packing.md). - `filter-mls` uses bounded scoped CPU workers for independent point fits, sharing one KD-tree and preserving coordinate words, field order and metadata. Existing small/one-worker point staging remains serial; each concurrent call diff --git a/docs/ROADMAP.md b/docs/ROADMAP.md index b09b0560..1d27fd69 100644 --- a/docs/ROADMAP.md +++ b/docs/ROADMAP.md @@ -1097,5 +1097,7 @@ v1.3 admission gates. | Bounded CPU FPFH descriptors | Complete | Shared KD-tree, scoped SPFH/weighting phases and direct disjoint descriptor writes using existing point staging | 18,402,648 frozen exact descriptor words with 1/2/6 workers and public staging, independent planar/degenerate cases, validation/concurrency/seeded alignment, feature/facade/WASM checks and twelve completed-descriptor-API comparisons; `notes/2026-10-12_parallel_fpfh_descriptors.md` | | GPU spatial execution guards | Complete | Pre-allocation runtime/dimension/device/profiler checks for nearest resize and Sobel, including checked resize coordinate products | CPU/frozen full-RGBA parity, invalid-limit/capacity and ROI/concurrent/live-output/pool checks, image/point regressions and feature/WASM builds; `notes/2026-10-12_gpu_spatial_dispatch_validation.md` | | Bounded GPU gray dispatch | Complete | Balanced multi-row flat workgroups beyond the single-axis limit, independently lazy kernels and pre-allocation arithmetic/device/profiler validation | All 16,777,216 RGB values, 120 forced grid mappings, maximum-counter device probes, ownership/capacity/ROI/concurrency/reuse checks, native regressions and resident timing; `notes/2026-10-12_bounded_gpu_gray_dispatch.md` | +| GPU AI execution guards | Complete | Pre-allocation image/index/buffer/dispatch and profiler validation while preserving the original CHW shader | 23,857,560 frozen output words across 192 cases, ROI/concurrency/live tensor reuse, negative-limit/capacity checks, 4096x4096 six-stage chain and feature/native/WASM validation; `notes/2026-10-12_validated_gpu_ai_packing.md` | +| GPU AI channel kernels | Deferred after measurement | Retain channel-lane and fixed-channel prototypes only in native fixtures | Lane kernels regressed across all 12 selected GPU conditions; fixed-channel micro gains did not establish completed-call improvement, with adverse cases retained; `notes/2026-10-12_validated_gpu_ai_packing.md` | | GPU shared Sobel tiles | Deferred after repeat measurement | Preserve the direct production shader while retaining an exact shared-halo prototype in native tests | 54 CPU/frozen/direct/experimental fixtures; initial gains did not survive repeat measurement across 36 conditions, with adverse wide/odd strips; `notes/2026-10-12_gpu_spatial_dispatch_validation.md` | | GPU contrast geometry reuse | Deferred after measurement | Keep per-call device geometry and existing bounded scratch while shortening exact ratios | A 24-condition survey found geometry below 1% of GPU pass time for 3840x2160 / 8x8 mixed and hot fixtures; revisit retention if repeated geometry becomes a material completed-call cost; `notes/2026-10-12_exact_gpu_clahe_geometry.md` | diff --git a/notes/2026-10-12_validated_gpu_ai_packing.md b/notes/2026-10-12_validated_gpu_ai_packing.md new file mode 100644 index 00000000..ad33065e --- /dev/null +++ b/notes/2026-10-12_validated_gpu_ai_packing.md @@ -0,0 +1,197 @@ +# Validated GPU AI packing + +Date: 2026-10-12. This slice follows merged +[PR #146](https://github.com/rsasaki0109/SpatialRust/pull/146), main +`a9ba90e132e1282aa15e8d111af6fa7e5fbd5120`. + +## Delivered behavior + +`pack_ai_chw_gpu` previously computed output sizes and acquired storage before +checking output-binding/device limits or profiler capacity. Unsupported storage +sizes could reach wgpu validation rather than return a spatial error. Packing +now validates its complete work plan before output/uniform allocation or +AI-pipeline initialization. Valid operations retain the original shader, +pipeline factory, normalization, CHW shape/order and transfer receipts. + +The plan requires positive source dimensions, channels in 1..=4, device texture +limits and signed texture coordinates. Widened host arithmetic checks pixel and +element counts against u32 shader indexing and multiplies element count by four +with checked arithmetic. Output bytes must fit both `max_buffer_size` and +`max_storage_buffer_binding_size`. It also checks the existing 16x16/256-lane +workgroup, per-axis dispatch limits, one sampled texture, storage buffer and +uniform buffer, a 64-byte uniform/buffer, one bind group and three bindings. +It imposes no new workgroup-storage or backend-feature requirement. + +Runtime ownership and existing finite scale/mean, finite nonzero std validation +remain ahead of execution; all four parameter slots are still validated, even +for fewer logical channels. Negative std and permitted intermediate overflow +retain their previous behavior. Profile ownership and one free stage are +checked before work. Rejection preserves previously recorded stages. + +Production still computes `(byte * scale - mean[channel]) / std[channel]` in +the original f32 operation order. There is one runtime-owned cached pipeline, +one GPU pass and no extra scratch or payload transfer. Receipts record +`4 * channels * width * height` GPU write bytes, zero implicit payload readback, +and 16 query bytes on explicit one-stage collection. Public signatures, +dependencies, storage retention and tensor ownership/readback remain unchanged. +The buffer pool bounds retained storage, not active outputs or peak driver memory. + +## Correctness and feature validation + +Six new tests in +`C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/src/image/ai_pack_tests.rs` +cover: + +- **23,857,560 exact f32 output words** compared with a frozen parent GPU + shader across 192 channel/shape/normalization cases. Twelve shapes include + singleton axes, 7/8/9 and 15/16/17 boundaries, odd tails and 255x257/256x256/ + 257x259 images. Channels 1/2/3/4 and four parameter sets cover identity, + exact binary scaling/negative std/signed zero, ordinary AI normalization + and extreme finite parameters producing overflow. Forced lane, forced fixed- + channel and public-original outputs retain frozen words. Identity and binary + scaling additionally match an independent CPU CHW lookup. +- Dimension/channel, texture/signed-coordinate, pixel/element overflow, exact + buffer-binding boundaries, output/uniform buffer, compute/dispatch and binding + limits. Zero workgroup storage remains acceptable. Large synthetic limits are + pure planning tests, not claims that giant buffers can be allocated. +- Invalid normalization, foreign image/profile ownership and exhausted capacity + leave storage-pool counters and cold AI caches unchanged. Capacity rejection + preserves an earlier Sobel stage and its 16-byte query collection. +- A metadata-only negative probe derives declared output dimensions from the + real device's buffer limits while backing the probe with one tiny texture. + It confirms rejection before any work or cache creation without allocating a + huge source. The mismatched probe is dropped rather than returned to the pool. +- Four concurrent cold calls on a strided 257x259 RGB ROI, independent live + tensors while the old source texture is reused, rapid submissions, runtime + movement, allocation reuse, zero retention budget and foreign tensor readback. +- A **4096x4096 complete resident Vision chain** on constant RGB input with + 1x1 blur/morphology and identity packing. All **16,777,216 final f32 words** are + zero, matching the independent constant-image/Sobel expectation. It records + six stages, 67,108,864 upload bytes, 402,653,184 GPU write bytes and zero payload + readback before explicit collection. Query collection separately returns + 96 bytes; explicit tensor readback adds 67,108,864 payload bytes. + +The frozen parent shader is +`C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/tests/fixtures/ai_pack_previous.wgsl`. +The production shader, pipeline factory and tensor ownership/readback block are +checked against parent main. Both experimental shaders live only in native +fixtures; production does not compile, cache or select them. + +Final validation: **154 native image/point GPU tests pass**, zero failures and +12 ignored diagnostics (127 unit plus 27 CDF/CLAHE/equalization/histogram/shared- +stage tests). Strict default/all-feature all-target Clippy, two WASM feature +compilation profiles, formatting and diff checks pass. WASM compilation does +not establish browser execution. The original six-to-eight-pass chain regressions +also pass; the new 4096x4096 proof covers the six-pass constant-image configuration. + +```text +cargo test -p spatialrust-gpu --all-features --tests -- --test-threads=1 --nocapture +cargo clippy -p spatialrust-gpu --all-features --all-targets -- -D warnings +cargo clippy -p spatialrust-gpu --no-default-features --all-targets -- -D warnings +cargo check -p spatialrust-gpu --target wasm32-unknown-unknown --no-default-features --features gpu-image +cargo check -p spatialrust-gpu --target wasm32-unknown-unknown --no-default-features --features gpu-image,gpu-aoso-staging +cargo fmt --all --check +git diff --check +``` + +## Why the experimental kernels were deferred + +Two native release pilots tested the same 24 conditions: eight shapes crossed +with channels 1/3/4. Each used two warmups and twenty alternating frozen/current +rounds, checking every output against frozen words. Experiments selected RGB/ +RGBA at 65,536 pixels or more; twelve cases were selected and twelve retained +the original kernel. Raw samples and all controls are retained privately. + +The first prototype assigns channels to separate lanes in an 8x8x4 workgroup. +All twelve selected GPU intervals fell below one: frozen/experimental paired +ratios range **0.551–0.713**. Completed ratios range 0.873–1.329, including noisy +positive samples. The GPU regression was consistent enough to reject this route. +The exact prototype remains +`C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/tests/fixtures/ai_pack_channels_experiment.wgsl`. + +The second prototype retains 16x16 workgroups and one texture read per pixel, +fixing the loop bound with a pipeline override for RGB/RGBA. At 256x256, +GPU ratios were 1.085 [1.056, 1.103] for RGB and 1.058 [1.039, 1.077] for RGBA, +but completed intervals crossed one (1.010 [0.948, 1.067] and +0.948 [0.886, 1.020]). Full-HD/UHD GPU ratios were about 0.997–1.001; UHD RGBA +had a small adverse GPU interval, 0.997 [0.996, 0.998]. The 17x4097 RGBA +completed ratio was 0.968 [0.938, 0.983]. These results do not establish a +completed-call gain worth additional production pipelines. The prototype remains +`C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/tests/fixtures/ai_pack_specialized_experiment.wgsl`. + +No speedup is claimed for the shipped validation change. Both original-kernel +controls and adverse experimental samples are retained; no experimental +selection policy or new pipeline is present in production. + +## Final resident diagnostic + +Release native Windows, Rust/Cargo 1.98.1, wgpu 24.0.5, Vulkan discrete backend +with pass timestamps. No local adapter model is published. No GPU tests or builds +launched by this task overlap any measurement run. Desktop activity, clocks +and external GPU load are not controlled. + +The final run compares frozen and current original kernels for all 24 conditions. +Each has two warmups and twenty alternating rounds. Every warmup/measured output +matches frozen words, CHW shape, write/readback accounting and a one-stage, +16-byte query report. Pipeline setup, upload, profiler creation, query collection, +payload readback/checking and recycling are outside timers. Completed timers +include source/plan/profile validation, output/uniform allocation, encoding, +submission and `wait_idle`. The frozen shader shares the current host executor; +this **does not measure historical API overhead or validation cost**. + +One runtime and its default 64 MiB retained-storage budget persist across shapes. +The pool is not cleared between conditions; live/output allocation and retained +occupancy affect reuse. Times are variant medians. Ratios are medians of twenty +paired frozen/current times, with 10,000 paired median bootstrap resamples +(seed 20261012; sorted indices 249 and 9749). Above one favors current. These +describe one run's samples rather than independent-repeat confidence. + +| Size | Channels | Frozen/current GPU us | GPU ratio [95% CI] | Frozen/current completed ms | Completed ratio [95% CI] | +| --- | ---: | --- | --- | --- | --- | +| 17x19 | 1 | 3.936 / 3.504 | 1.128 [0.944, 1.290] | 0.257 / 0.257 | 1.002 [0.890, 1.059] | +| 17x19 | 3 | 4.160 / 4.128 | 1.000 [0.923, 1.036] | 0.242 / 0.241 | 0.989 [0.948, 1.065] | +| 17x19 | 4 | 4.368 / 4.288 | 1.004 [0.993, 1.041] | 0.221 / 0.222 | 0.991 [0.932, 1.050] | +| 128x128 | 1 | 4.128 / 4.128 | 1.016 [0.953, 1.069] | 0.248 / 0.268 | 0.945 [0.883, 1.086] | +| 128x128 | 3 | 4.416 / 4.448 | 0.971 [0.891, 1.029] | 0.279 / 0.292 | 0.950 [0.898, 1.030] | +| 128x128 | 4 | 4.704 / 4.704 | 1.003 [0.916, 1.065] | 0.310 / 0.301 | 1.030 [0.919, 1.202] | +| 256x256 | 1 | 5.376 / 5.376 | 1.000 [0.953, 1.030] | 0.250 / 0.277 | 0.919 [0.871, 0.964] | +| 256x256 | 3 | 7.296 / 7.328 | 0.996 [0.981, 1.044] | 0.318 / 0.311 | 0.992 [0.928, 1.106] | +| 256x256 | 4 | 8.880 / 8.928 | 0.991 [0.963, 1.000] | 0.329 / 0.342 | 0.961 [0.882, 1.074] | +| 640x480 | 1 | 15.504 / 15.600 | 0.982 [0.964, 1.002] | 0.363 / 0.358 | 0.880 [0.771, 1.097] | +| 640x480 | 3 | 21.536 / 21.760 | 0.993 [0.973, 1.006] | 0.459 / 0.471 | 0.974 [0.889, 1.115] | +| 640x480 | 4 | 28.272 / 28.272 | 1.006 [0.991, 1.018] | 0.474 / 0.467 | 1.002 [0.865, 1.150] | +| 1920x1080 | 1 | 379.904 / 380.512 | 1.001 [0.990, 1.012] | 1.062 / 1.059 | 1.002 [0.972, 1.102] | +| 1920x1080 | 3 | 980.896 / 982.640 | 0.999 [0.998, 1.002] | 2.040 / 1.987 | 1.038 [1.019, 1.056] | +| 1920x1080 | 4 | 1448.000 / 1447.120 | 1.001 [1.000, 1.003] | 3.784 / 3.759 | 0.997 [0.995, 1.013] | +| 3840x2160 | 1 | 4002.096 / 4006.176 | 0.997 [0.995, 1.002] | 13.030 / 13.061 | 0.995 [0.990, 1.005] | +| 3840x2160 | 3 | 9934.272 / 9935.088 | 0.999 [0.999, 1.001] | 26.141 / 25.987 | 1.001 [0.994, 1.010] | +| 3840x2160 | 4 | 12572.480 / 12562.352 | 1.000 [0.999, 1.001] | 33.714 / 33.804 | 0.996 [0.993, 0.999] | +| 17x4097 | 1 | 41.280 / 41.952 | 0.995 [0.907, 1.025] | 1.022 / 1.058 | 0.998 [0.779, 1.119] | +| 17x4097 | 3 | 100.720 / 100.096 | 1.008 [0.986, 1.021] | 1.193 / 1.250 | 0.991 [0.927, 1.013] | +| 17x4097 | 4 | 86.416 / 81.872 | 1.008 [0.987, 1.157] | 1.272 / 1.227 | 0.999 [0.930, 1.052] | +| 4097x17 | 1 | 35.680 / 65.216 | 0.973 [0.543, 1.008] | 0.674 / 1.209 | 0.914 [0.537, 1.053] | +| 4097x17 | 3 | 104.128 / 105.344 | 0.984 [0.952, 1.006] | 1.219 / 1.236 | 0.969 [0.914, 1.048] | +| 4097x17 | 4 | 89.408 / 94.880 | 0.999 [0.986, 1.018] | 1.203 / 1.199 | 0.997 [0.907, 1.042] | + +Reproduction and private log paths are absolute: + +- `C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/src/image/ai_tensor.rs` +- `C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/src/image/ai_pack_benchmark.rs` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-regressions.log` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-clippy.log` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-default-clippy.log` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-wasm.log` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-wasm-combined.log` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-release-build.log` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-paired.log` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-estimates.json` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/validated-gpu-ai-pack-table.md` +- `C:/Users/rsasa/Workspace/SpatialRust/target/automation/summarize-validated-gpu-ai-pack.py` + +Earlier raw/estimate/table sets use `ai-pack-channel-lanes-*` and +`ai-pack-fixed-channel-*` in that private directory. Estimates include all +sample values, minimum/maximum times and paired intervals. + +```text +cargo test -p spatialrust-gpu --all-features --release --lib paired_frozen_and_validated_ai_pack -- --ignored --nocapture --test-threads=1 +```