From 501c2f2241e36d965f943189eb24dd3bad191ace Mon Sep 17 00:00:00 2001 From: rsasaki0109 Date: Mon, 12 Oct 2026 08:20:41 +0900 Subject: [PATCH] Bound GPU gray workgroups across dispatch rows --- CHANGELOG.md | 4 + .../spatialrust-gpu/src/image/kernels/gray.rs | 124 +++++- .../src/image/kernels/gray_benchmark.rs | 64 ++++ .../src/image/kernels/gray_grid.wgsl | 29 ++ .../src/image/kernels/gray_tests.rs | 361 ++++++++++++++++++ .../tests/fixtures/gray_previous.wgsl | 20 + docs/ARCHITECTURE.md | 9 + docs/FEATURE_MATRIX.md | 5 + docs/ROADMAP.md | 1 + notes/2026-10-12_bounded_gpu_gray_dispatch.md | 163 ++++++++ 10 files changed, 764 insertions(+), 16 deletions(-) create mode 100644 crates/spatialrust-gpu/src/image/kernels/gray_benchmark.rs create mode 100644 crates/spatialrust-gpu/src/image/kernels/gray_grid.wgsl create mode 100644 crates/spatialrust-gpu/src/image/kernels/gray_tests.rs create mode 100644 crates/spatialrust-gpu/tests/fixtures/gray_previous.wgsl create mode 100644 notes/2026-10-12_bounded_gpu_gray_dispatch.md diff --git a/CHANGELOG.md b/CHANGELOG.md index 331f27f..0e10f0e 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 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 + reject before allocation, with unchanged luma bytes and explicit transfers. - FPFH point-cloud descriptors share one KD-tree across bounded scoped CPU workers for both SPFH construction and neighbor weighting. Small/one-worker clouds stay serial; output order and exact descriptor words are preserved diff --git a/crates/spatialrust-gpu/src/image/kernels/gray.rs b/crates/spatialrust-gpu/src/image/kernels/gray.rs index 006cc2e..9b46eb9 100644 --- a/crates/spatialrust-gpu/src/image/kernels/gray.rs +++ b/crates/spatialrust-gpu/src/image/kernels/gray.rs @@ -3,6 +3,7 @@ use bytemuck::{Pod, Zeroable}; use spatialrust_core::{SpatialError, SpatialResult}; use spatialrust_image::{ColorSpace, ImageMetadata}; +use std::sync::OnceLock; use wgpu::util::DeviceExt; use crate::image::gpu_image::{create_texture, GpuImage, GpuImageReceipt}; @@ -15,8 +16,53 @@ const WORKGROUP: u32 = 256; struct GrayParams { width: u32, height: u32, - _pad0: u32, - _pad1: u32, + groups_per_row: u32, + group_count: u32, +} + +#[derive(Clone, Copy, Debug)] +struct GrayPlan { + groups_x: u32, + groups_y: u32, + group_count: u32, +} + +fn gray_plan(width: u32, height: u32, limits: &wgpu::Limits) -> SpatialResult { + let invalid = || { + SpatialError::InvalidArgument( + "GPU gray requires positive dimensions, u32 pixel counts and supported texture/compute/binding limits".to_owned(), + ) + }; + let pixels = u64::from(width) * u64::from(height); + if width == 0 + || height == 0 + || 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 + || limits.max_compute_workgroup_size_x < WORKGROUP + || limits.max_compute_workgroup_size_y == 0 + || limits.max_compute_workgroup_size_z == 0 + || limits.max_compute_workgroups_per_dimension == 0 + || limits.max_sampled_textures_per_shader_stage == 0 + || limits.max_storage_textures_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 group_count = pixels.div_ceil(u64::from(WORKGROUP)) as u32; + let groups_y = group_count.div_ceil(limits.max_compute_workgroups_per_dimension); + if groups_y > limits.max_compute_workgroups_per_dimension { + return Err(invalid()); + } + let groups_x = group_count.div_ceil(groups_y); + Ok(GrayPlan { groups_x, groups_y, group_count }) } const GRAY_WGSL: &str = r#" @@ -43,7 +89,8 @@ fn main(@builtin(global_invocation_id) gid: vec3) { pub(crate) struct GrayPipeline { bind_group_layout: wgpu::BindGroupLayout, - pipeline: wgpu::ComputePipeline, + pipeline: OnceLock, + grid_pipeline: OnceLock, } pub(crate) fn create_gray_pipeline(device: &wgpu::Device) -> GrayPipeline { @@ -82,27 +129,37 @@ pub(crate) fn create_gray_pipeline(device: &wgpu::Device) -> GrayPipeline { }, ], }); + GrayPipeline { bind_group_layout, pipeline: OnceLock::new(), grid_pipeline: OnceLock::new() } +} + +fn create_pipeline( + device: &wgpu::Device, + bind_group_layout: &wgpu::BindGroupLayout, + source: &str, +) -> wgpu::ComputePipeline { let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("gpu-image-gray-shader"), - source: wgpu::ShaderSource::Wgsl(GRAY_WGSL.into()), + source: wgpu::ShaderSource::Wgsl(source.into()), }); let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { label: Some("gpu-image-gray-pl"), - bind_group_layouts: &[&bind_group_layout], + bind_group_layouts: &[bind_group_layout], push_constant_ranges: &[], }); - let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some("gpu-image-gray-pipeline"), layout: Some(&pipeline_layout), module: &shader, entry_point: Some("main"), compilation_options: wgpu::PipelineCompilationOptions::default(), cache: None, - }); - GrayPipeline { bind_group_layout, pipeline } + }) } /// Converts an RGB `GpuImage` to a gray `GpuImage` using BT.601 fixed-point luma. +/// +/// Large images distribute flat workgroups across bounded dispatch rows. Runtime, +/// dimensions and device limits are checked before allocation or pipeline setup. pub fn rgb_to_gray_gpu(runtime: &WgpuRuntime, source: &GpuImage) -> SpatialResult { rgb_to_gray_gpu_with_profiler(runtime, source, None) } @@ -110,7 +167,7 @@ pub fn rgb_to_gray_gpu(runtime: &WgpuRuntime, source: &GpuImage) -> SpatialResul pub(crate) fn rgb_to_gray_gpu_with_profiler( runtime: &WgpuRuntime, source: &GpuImage, - mut profiler: Option<&mut crate::GpuStageProfiler>, + profiler: Option<&mut crate::GpuStageProfiler>, ) -> SpatialResult { source.validate_runtime(runtime)?; if source.channels() != 3 { @@ -118,22 +175,53 @@ pub(crate) fn rgb_to_gray_gpu_with_profiler( "rgb_to_gray_gpu requires a 3-channel RGB GpuImage".to_owned(), )); } - let pixel_count = (source.width() as usize).saturating_mul(source.height() as usize); + let plan = gray_plan(source.width(), source.height(), &runtime.device().limits())?; + if let Some(profiler) = &profiler { + profiler.validate(runtime, 1)?; + } + let cache = runtime.image_gray_pipeline.get_or_init(|| create_gray_pipeline(runtime.device())); + let pipeline = if plan.groups_y == 1 { + cache + .pipeline + .get_or_init(|| create_pipeline(runtime.device(), &cache.bind_group_layout, GRAY_WGSL)) + } else { + cache.grid_pipeline.get_or_init(|| { + create_pipeline( + runtime.device(), + &cache.bind_group_layout, + include_str!("gray_grid.wgsl"), + ) + }) + }; + gray_dispatch(runtime, source, plan, cache, pipeline, profiler) +} + +fn gray_dispatch( + runtime: &WgpuRuntime, + source: &GpuImage, + plan: GrayPlan, + cache: &GrayPipeline, + pipeline: &wgpu::ComputePipeline, + mut profiler: Option<&mut crate::GpuStageProfiler>, +) -> SpatialResult { let output = create_texture(runtime, source.width(), source.height(), "gpu-image-gray-out"); let out_bytes = u64::from(source.width()) * u64::from(source.height()) * 4; - let params = GrayParams { width: source.width(), height: source.height(), _pad0: 0, _pad1: 0 }; + let params = GrayParams { + width: source.width(), + height: source.height(), + groups_per_row: plan.groups_x, + group_count: plan.group_count, + }; let uniform = runtime.device().create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("gpu-image-gray-params"), contents: bytemuck::bytes_of(¶ms), usage: wgpu::BufferUsages::UNIFORM, }); - let pipeline = - runtime.image_gray_pipeline.get_or_init(|| create_gray_pipeline(runtime.device())); let source_view = source.view(); let output_view = output.create_view(&wgpu::TextureViewDescriptor::default()); let bind_group = runtime.device().create_bind_group(&wgpu::BindGroupDescriptor { label: Some("gpu-image-gray-bg"), - layout: &pipeline.bind_group_layout, + layout: &cache.bind_group_layout, entries: &[ wgpu::BindGroupEntry { binding: 0, resource: uniform.as_entire_binding() }, wgpu::BindGroupEntry { @@ -154,9 +242,9 @@ pub(crate) fn rgb_to_gray_gpu_with_profiler( label: Some("gpu-image-gray-pass"), timestamp_writes: crate::profiling::timestamp_writes(&mut profiler, "rgb_to_gray_gpu"), }); - pass.set_pipeline(&pipeline.pipeline); + pass.set_pipeline(pipeline); pass.set_bind_group(0, &bind_group, &[]); - pass.dispatch_workgroups(pixel_count.div_ceil(WORKGROUP as usize) as u32, 1, 1); + pass.dispatch_workgroups(plan.groups_x, plan.groups_y, 1); } runtime.queue().submit(Some(encoder.finish())); let mut receipt = GpuImageReceipt::default(); @@ -165,3 +253,7 @@ pub(crate) fn rgb_to_gray_gpu_with_profiler( let metadata = ImageMetadata { color_space: ColorSpace::Gray, ..source.metadata() }; GpuImage::from_parts(runtime, source.width(), source.height(), 1, output, metadata, receipt) } + +#[cfg(all(test, not(target_arch = "wasm32")))] +#[path = "gray_tests.rs"] +mod tests; diff --git a/crates/spatialrust-gpu/src/image/kernels/gray_benchmark.rs b/crates/spatialrust-gpu/src/image/kernels/gray_benchmark.rs new file mode 100644 index 0000000..6737669 --- /dev/null +++ b/crates/spatialrust-gpu/src/image/kernels/gray_benchmark.rs @@ -0,0 +1,64 @@ +use super::*; + +#[test] +#[ignore = "explicit native resident GPU gray diagnostic"] +fn paired_frozen_current_and_large_gray() { + let runtime = + WgpuRuntime::new_headless_with_timestamps(crate::WgpuPowerPreference::HighPerformance) + .unwrap(); + let cache = runtime.image_gray_pipeline.get_or_init(|| create_gray_pipeline(runtime.device())); + let frozen = create_pipeline( + runtime.device(), + &cache.bind_group_layout, + include_str!("../../../tests/fixtures/gray_previous.wgsl"), + ); + eprintln!("backend={}; device_type={}; 2 warmups + 20 alternating rounds; shared host executor; timings include validation/allocation/encoding/submit/wait, exclude 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), (640, 480), (1920, 1080), (3840, 2160), (4096, 4095), (4096, 4096)] + { + let image = fixture(w, h, 0); + let expected = oracle(image.as_slice()); + let source = GpuImage::upload_u8(&runtime, image.view()).unwrap(); + let plan = gray_plan(w as u32, h as u32, &runtime.device().limits()).unwrap(); + // The frozen linear kernel has no valid dispatch for the large case. + let modes = if plan.groups_y == 1 { 2 } else { 1 }; + 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..modes { + let mode = (round + offset) % modes; + let current = modes == 1 || mode == 1; + let mut profiler = crate::GpuStageProfiler::new(&runtime, 1).unwrap(); + let start = std::time::Instant::now(); + let output = if current { + rgb_to_gray_gpu_with_profiler(&runtime, &source, Some(&mut profiler)).unwrap() + } else { + source.validate_runtime(&runtime).unwrap(); + gray_plan(w as u32, h as u32, &runtime.device().limits()).unwrap(); + profiler.validate(&runtime, 1).unwrap(); + gray_dispatch(&runtime, &source, plan, cache, &frozen, 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.receipt().stages(), ["upload_u8_texture", "rgb_to_gray_gpu"]); + assert_eq!(output.receipt().gpu_to_gpu_bytes(), (w * h * 4) as u64); + assert_eq!(output.receipt().device_to_host_bytes(), 0); + assert_rgba(&runtime, &output, &expected); + assert_eq!(output.metadata().color_space, ColorSpace::Gray); + assert_eq!(output.metadata().color_range, image.metadata().color_range); + if round >= 2 { + let i = usize::from(current); + stage[i].push(report.stages[0].elapsed_ns / 1000.0); + wall[i].push(ms); + } + output.recycle(&runtime); + } + } + eprintln!("GRAY {{\"width\":{w},\"height\":{h},\"groups_x\":{},\"groups_y\":{},\"previous_us\":{:?},\"current_us\":{:?},\"previous_ms\":{:?},\"current_ms\":{:?}}}",plan.groups_x,plan.groups_y,stage[0],stage[1],wall[0],wall[1]); + source.recycle(&runtime); + } +} diff --git a/crates/spatialrust-gpu/src/image/kernels/gray_grid.wgsl b/crates/spatialrust-gpu/src/image/kernels/gray_grid.wgsl new file mode 100644 index 0000000..aaaf9e8 --- /dev/null +++ b/crates/spatialrust-gpu/src/image/kernels/gray_grid.wgsl @@ -0,0 +1,29 @@ +struct Params { + width: u32, + height: u32, + groups_per_row: u32, + group_count: u32, +}; + +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var input_px: texture_2d; +@group(0) @binding(2) var output_px: texture_storage_2d; + +fn flat_pixel(group: vec3, lane: u32) -> u32 { + let flat_group = group.y * params.groups_per_row + group.x; + // A balanced rectangle can contain padding groups. Reject them before + // multiplying by 256 so the u32::MAX pixel boundary cannot wrap. + if (flat_group >= params.group_count) { return 0xffffffffu; } + return flat_group * 256u + lane; +} + +@compute @workgroup_size(256) +fn main(@builtin(workgroup_id) group: vec3, @builtin(local_invocation_index) lane: u32) { + let index = flat_pixel(group, lane); + let pixel_count = params.width * params.height; + if (index >= pixel_count) { return; } + let xy = vec2(i32(index % params.width), i32(index / params.width)); + let rgb = textureLoad(input_px, xy, 0).rgb; + let gray = (77u * rgb.r + 150u * rgb.g + 29u * rgb.b + 128u) >> 8u; + textureStore(output_px, xy, vec4(gray, 0u, 0u, 0u)); +} diff --git a/crates/spatialrust-gpu/src/image/kernels/gray_tests.rs b/crates/spatialrust-gpu/src/image/kernels/gray_tests.rs new file mode 100644 index 0000000..8243525 --- /dev/null +++ b/crates/spatialrust-gpu/src/image/kernels/gray_tests.rs @@ -0,0 +1,361 @@ +use super::*; +use crate::image::gpu_image::read_staging_bytes; +use spatialrust_image::{ColorRange, Image, ImageRegion, ImageView}; + +pub(super) fn fixture(w: usize, h: usize, kind: usize) -> Image { + let data = (0..w * h) + .flat_map(|i| match kind { + 0 => [(i * 37) as u8, (i * 11 / 3) as u8, (i * 71) as u8], + 1 => [0; 3], + 2 => [255; 3], + _ => [(i >> 16) as u8, (i >> 8) as u8, i as u8], + }) + .collect(); + Image::try_new_with_metadata( + w, + h, + data, + ImageMetadata { color_range: ColorRange::Limited, ..Default::default() }, + ) + .unwrap() +} + +pub(super) fn oracle(rgb: &[u8]) -> Vec { + rgb.chunks_exact(3) + .map(|p| { + ((77 * u32::from(p[0]) + 150 * u32::from(p[1]) + 29 * u32::from(p[2]) + 128) / 256) + as u8 + }) + .collect() +} + +fn physical_rgba(runtime: &WgpuRuntime, output: &GpuImage) -> Image { + // A temporary handle to the same texture exposes physical unused channels + // for explicit test readback; it performs no texture allocation or copy. + let mut physical = GpuImage::from_parts( + runtime, + output.width(), + output.height(), + 4, + output.texture().clone(), + ImageMetadata { color_space: ColorSpace::Unknown, ..output.metadata() }, + output.receipt().clone(), + ) + .unwrap(); + physical.readback_u8::<4>(runtime).unwrap() +} + +fn assert_rgba(runtime: &WgpuRuntime, output: &GpuImage, expected: &[u8]) { + let actual = physical_rgba(runtime, output); + assert_eq!(actual.as_slice().len(), expected.len() * 4); + for (pixel, &gray) in actual.as_slice().chunks_exact(4).zip(expected) { + assert_eq!(pixel, [gray, 0, 0, 0]); + } +} + +#[test] +fn planner_bounds_both_axes_and_u32_coordinate_arithmetic() { + let mut limits = wgpu::Limits::default(); + for (w, h) in [(0, 1), (1, 0), (u32::MAX, 1), (1, u32::MAX)] { + assert!(gray_plan(w, h, &limits).is_err()); + } + for field in 0..12 { + let mut changed = limits.clone(); + match field { + 0 => changed.max_compute_invocations_per_workgroup = 255, + 1 => changed.max_compute_workgroup_size_x = 255, + 2 => changed.max_compute_workgroup_size_y = 0, + 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_textures_per_shader_stage = 0, + 7 => changed.max_uniform_buffers_per_shader_stage = 0, + 8 => changed.max_uniform_buffer_binding_size = 15, + 9 => changed.max_buffer_size = 15, + 10 => changed.max_bind_groups = 0, + _ => changed.max_bindings_per_bind_group = 2, + } + assert!(gray_plan(17, 19, &changed).is_err(), "field {field}"); + } + limits.max_texture_dimension_2d = u32::MAX; + limits.max_compute_workgroup_storage_size = 0; + limits.max_storage_buffers_per_shader_stage = 0; + assert!(gray_plan(65536, 65536, &limits).is_err()); + let maximum = gray_plan(65535, 65537, &limits).unwrap(); + assert_eq!(maximum.group_count, 16_777_216); + assert!( + maximum.groups_x <= limits.max_compute_workgroups_per_dimension + && maximum.groups_y <= limits.max_compute_workgroups_per_dimension + ); + assert!(gray_plan(i32::MAX as u32, 2, &limits).is_ok()); + assert!(gray_plan(i32::MAX as u32 + 1, 1, &limits).is_err()); + limits.max_compute_workgroups_per_dimension = 2; + assert_eq!(gray_plan(1, 1024, &limits).unwrap().groups_y, 2); + assert!(gray_plan(1, 1025, &limits).is_err()); + limits.max_texture_dimension_2d = 16; + assert!(gray_plan(17, 1, &limits).is_err()); +} + +#[test] +fn forced_grid_and_frozen_linear_match_independent_full_rgba_oracle() { + let runtime = WgpuRuntime::new_headless().unwrap(); + let cache = runtime.image_gray_pipeline.get_or_init(|| create_gray_pipeline(runtime.device())); + let frozen = create_pipeline( + runtime.device(), + &cache.bind_group_layout, + include_str!("../../../tests/fixtures/gray_previous.wgsl"), + ); + let grid = + create_pipeline(runtime.device(), &cache.bind_group_layout, include_str!("gray_grid.wgsl")); + let mut bytes = 0; + let mut cases = 0; + for (w, h) in + [(1, 1), (1, 255), (1, 256), (1, 257), (257, 1), (17, 19), (31, 33), (65, 67), (257, 259)] + { + for kind in 0..3 { + let image = fixture(w, h, kind); + let source = GpuImage::upload_u8(&runtime, image.view()).unwrap(); + let expected = oracle(image.as_slice()); + let linear = gray_plan(w as u32, h as u32, &runtime.device().limits()).unwrap(); + let old = gray_dispatch(&runtime, &source, linear, cache, &frozen, None).unwrap(); + assert_rgba(&runtime, &old, &expected); + old.recycle(&runtime); + let current = rgb_to_gray_gpu(&runtime, &source).unwrap(); + assert_rgba(&runtime, ¤t, &expected); + current.recycle(&runtime); + bytes += w * h * 8; + for cap in [2, 3, 7, 17, 31] { + let mut limits = runtime.device().limits(); + limits.max_compute_workgroups_per_dimension = cap; + if let Ok(plan) = gray_plan(w as u32, h as u32, &limits) { + let output = + gray_dispatch(&runtime, &source, plan, cache, &grid, None).unwrap(); + assert_rgba(&runtime, &output, &expected); + assert_eq!(output.metadata().color_space, ColorSpace::Gray); + assert_eq!(output.metadata().color_range, image.metadata().color_range); + assert_eq!(output.receipt().stages(), ["upload_u8_texture", "rgb_to_gray_gpu"]); + assert_eq!(output.receipt().gpu_to_gpu_bytes(), (w * h * 4) as u64); + assert_eq!(output.receipt().device_to_host_bytes(), 0); + output.recycle(&runtime); + bytes += w * h * 4; + cases += 1; + } + } + source.recycle(&runtime); + } + } + eprintln!("Verified {cases} forced grid mappings and {bytes} exact RGBA bytes across frozen/public/grid variants."); +} + +#[test] +fn maximum_pixel_index_and_padding_groups_do_not_wrap_on_device() { + let runtime = WgpuRuntime::new_headless().unwrap(); + let shader = format!( + "{}\n{}", + include_str!("gray_grid.wgsl"), + r#" +@group(0) @binding(3) var cases: array>; +@group(0) @binding(4) var answers: array>; +@compute @workgroup_size(64) +fn probe(@builtin(global_invocation_id) id: vec3) { + if (id.x >= arrayLength(&cases)) { return; } + let c = cases[id.x]; + let index = flat_pixel(vec3(c.xy, 0u), c.z); + answers[id.x] = vec2(index, select(0u, 1u, index < params.width * params.height)); +}"# + ); + let module = runtime.device().create_shader_module(wgpu::ShaderModuleDescriptor { + label: None, + source: wgpu::ShaderSource::Wgsl(shader.into()), + }); + let pipeline = runtime.device().create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: None, + layout: None, + module: &module, + entry_point: Some("probe"), + compilation_options: Default::default(), + cache: None, + }); + let mut words = 0; + for (w, h, cap) in + [(1u32, 1u32, 2u32), (17, 33, 2), (65535, 65537, 65535), (i32::MAX as u32, 2, 65535)] + { + let mut limits = runtime.device().limits(); + limits.max_texture_dimension_2d = u32::MAX; + limits.max_compute_workgroups_per_dimension = cap; + let plan = gray_plan(w, h, &limits).unwrap(); + let mut cases = Vec::<[u32; 4]>::new(); + let mut expected = Vec::<[u32; 2]>::new(); + for flat in [0, plan.group_count - 1, plan.group_count] { + for lane in [0, 1, 254, 255] { + cases.push([flat % plan.groups_x, flat / plan.groups_x, lane, 0]); + let index = if flat >= plan.group_count { + u32::MAX + } else { + (u64::from(flat) * 256 + u64::from(lane)) as u32 + }; + expected.push([index, u32::from(u64::from(index) < u64::from(w) * u64::from(h))]); + } + } + let params = GrayParams { + width: w, + height: h, + groups_per_row: plan.groups_x, + group_count: plan.group_count, + }; + let uniform = runtime.device().create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: None, + contents: bytemuck::bytes_of(¶ms), + usage: wgpu::BufferUsages::UNIFORM, + }); + let input = runtime.device().create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: None, + contents: bytemuck::cast_slice(&cases), + usage: wgpu::BufferUsages::STORAGE, + }); + let size = expected.len() as u64 * 8; + let output = runtime.device().create_buffer(&wgpu::BufferDescriptor { + label: None, + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let staging = runtime.device().create_buffer(&wgpu::BufferDescriptor { + label: None, + size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let group = runtime.device().create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &pipeline.get_bind_group_layout(0), + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: uniform.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: input.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 4, resource: output.as_entire_binding() }, + ], + }); + let mut encoder = runtime.device().create_command_encoder(&Default::default()); + { + let mut pass = encoder.begin_compute_pass(&Default::default()); + pass.set_pipeline(&pipeline); + pass.set_bind_group(0, &group, &[]); + pass.dispatch_workgroups(1, 1, 1); + } + encoder.copy_buffer_to_buffer(&output, 0, &staging, 0, size); + runtime.queue().submit(Some(encoder.finish())); + let actual = read_staging_bytes(runtime.device(), &staging, size as usize).unwrap(); + assert_eq!(actual, bytemuck::cast_slice(&expected)); + words += expected.len() * 2; + } + eprintln!( + "Verified {words} exact device index/validity words including u32::MAX and padding groups." + ); +} + +#[test] +fn exhaustive_rgb_large_public_grid_is_lazy_and_exact() { + let runtime = WgpuRuntime::new_headless().unwrap(); + let image = fixture(4096, 4096, 3); // All 2^24 RGB combinations, exactly once. + let plan = gray_plan(4096, 4096, &runtime.device().limits()).unwrap(); + assert!(plan.groups_y > 1); + let source = GpuImage::upload_u8(&runtime, image.view()).unwrap(); + let output = rgb_to_gray_gpu(&runtime, &source).unwrap(); + assert_rgba(&runtime, &output, &oracle(image.as_slice())); + let cache = runtime.image_gray_pipeline.get().unwrap(); + assert!(cache.grid_pipeline.get().is_some() && cache.pipeline.get().is_none()); + assert_eq!(output.receipt().gpu_to_gpu_bytes(), 67_108_864); + assert_eq!(output.receipt().device_to_host_bytes(), 0); + output.recycle(&runtime); + source.recycle(&runtime); + eprintln!("Verified all 16777216 RGB values and 67108864 exact RGBA bytes above the single-axis dispatch limit."); +} + +#[test] +fn invalid_ownership_channels_and_capacity_leave_gray_work_unallocated() { + 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(17, 19, 0).view()).unwrap(); + let gray = + GpuImage::upload_u8(&runtime, Image::::try_new(1, 1, vec![0]).unwrap().view()) + .unwrap(); + let before = runtime.image_texture_pool_stats(); + assert!(rgb_to_gray_gpu(&other, &source).is_err()); + assert!(rgb_to_gray_gpu(&runtime, &gray).is_err()); + let mut foreign = crate::GpuStageProfiler::new(&other, 1).unwrap(); + assert!(rgb_to_gray_gpu_with_profiler(&runtime, &source, Some(&mut foreign)).is_err()); + assert_eq!(runtime.image_texture_pool_stats(), before); + assert!( + runtime.image_gray_pipeline.get().is_none() && other.image_gray_pipeline.get().is_none() + ); + 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.image_texture_pool_stats(); + assert!(rgb_to_gray_gpu_with_profiler(&runtime, &source, Some(&mut profiler)).is_err()); + assert_eq!(runtime.image_texture_pool_stats(), before); + assert!(runtime.image_gray_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 strided_roi_cold_concurrency_live_outputs_and_reuse_are_exact() { + let runtime = WgpuRuntime::new_headless().unwrap(); + let data: Vec<_> = (0..(67 * 3 + 11) * 79).map(|i| (i * 31 + 7) as u8).collect(); + let view = ImageView::::new(67, 79, 67 * 3 + 11, &data) + .unwrap() + .subview(ImageRegion::new(2, 3, 17, 19)) + .unwrap(); + let expected = + oracle(&(0..19).flat_map(|y| view.row(y).unwrap().iter().copied()).collect::>()); + 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 || rgb_to_gray_gpu(runtime, source).unwrap()) + }) + .collect(); + jobs.into_iter().map(|j| j.join().unwrap()).collect::>() + }); + let cache = runtime.image_gray_pipeline.get().unwrap(); + assert!(cache.pipeline.get().is_some() && cache.grid_pipeline.get().is_none()); + source.recycle(&runtime); + let source = GpuImage::upload_u8(&runtime, fixture(17, 19, 2).view()).unwrap(); + for _ in 0..12 { + rgb_to_gray_gpu(&runtime, &source).unwrap().recycle(&runtime); + } + let runtime = Box::new(runtime); + for output in outputs { + assert_rgba(&runtime, &output, &expected); + output.recycle(&runtime); + } + let before = runtime.image_texture_pool_stats(); + for _ in 0..12 { + rgb_to_gray_gpu(&runtime, &source).unwrap().recycle(&runtime); + } + assert_eq!(runtime.image_texture_pool_stats().allocation_count, before.allocation_count); + runtime.set_image_texture_pool_budget(0); + let output = rgb_to_gray_gpu(&runtime, &source).unwrap(); + assert_rgba(&runtime, &output, &vec![255; 17 * 19]); + output.recycle(&runtime); + source.recycle(&runtime); + assert_eq!(runtime.image_texture_pool_stats().cached_bytes, 0); +} + +#[path = "gray_benchmark.rs"] +mod benchmark; diff --git a/crates/spatialrust-gpu/tests/fixtures/gray_previous.wgsl b/crates/spatialrust-gpu/tests/fixtures/gray_previous.wgsl new file mode 100644 index 0000000..88c86f4 --- /dev/null +++ b/crates/spatialrust-gpu/tests/fixtures/gray_previous.wgsl @@ -0,0 +1,20 @@ + +struct Params { width: u32, height: u32, pad0: u32, pad1: u32, }; + +@group(0) @binding(0) var params: Params; +@group(0) @binding(1) var input_px: texture_2d; +@group(0) @binding(2) var output_px: texture_storage_2d; + +@compute @workgroup_size(256) +fn main(@builtin(global_invocation_id) gid: vec3) { + let index = gid.x; + let pixel_count = params.width * params.height; + if (index >= pixel_count) { + return; + } + let xy = vec2(i32(index % params.width), i32(index / params.width)); + let rgb = textureLoad(input_px, xy, 0).rgb; + // Match CPU BT.601 fixed-point: (77*R + 150*G + 29*B + 128) >> 8 + let gray = (77u * rgb.r + 150u * rgb.g + 29u * rgb.b + 128u) >> 8u; + textureStore(output_px, xy, vec4(gray, 0u, 0u, 0u)); +} diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md index c45371d..fab4a69 100644 --- a/docs/ARCHITECTURE.md +++ b/docs/ARCHITECTURE.md @@ -74,6 +74,15 @@ channel count. Kernels return device-resident images and append named receipt stages; only `upload_u8` and `readback_u8` cross the host/device boundary. Runtime adapter identity and synchronization are explicit (`adapter_info`, `wait_idle`), and `recycle` returns textures to the runtime pool. +RGB-to-gray retains the original 256-lane linear shader when its group count +fits one dispatch axis. Larger inputs balance those flat groups across bounded +dispatch rows. Padding groups return before pixel-index multiplication to prevent +u32 wrap; source dimensions, pixel count, compute/binding limits and one free +profiler stage are validated before work allocation or pipeline initialization. +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. 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 3549108..3aa0512 100644 --- a/docs/FEATURE_MATRIX.md +++ b/docs/FEATURE_MATRIX.md @@ -186,6 +186,11 @@ workspace because its build requires a Python toolchain. overflow. Valid shader math, metadata, receipts and transfers remain unchanged. Shared-input Sobel tiling remains test-only after inconsistent paired results. See [spatial dispatch validation and measurements](../notes/2026-10-12_gpu_spatial_dispatch_validation.md). +- GPU RGB-to-gray distributes large flat workloads across bounded dispatch rows, + while single-row workloads retain their original shader. Both pipelines are + 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). - `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 b595e2b..b09b056 100644 --- a/docs/ROADMAP.md +++ b/docs/ROADMAP.md @@ -1096,5 +1096,6 @@ v1.3 admission gates. | Bounded CPU MLS smoothing | Complete | Independent point fits on shared KD-tree with existing bounded staging and disjoint ordered output bands | Frozen exact-coordinate parity with 1/2/6 workers, curved/duplicate/degenerate/support cases, all attribute types and metadata, feature/facade/WASM checks and twelve completed-position-API comparisons; `notes/2026-10-12_parallel_mls_smoothing.md` | | 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 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_bounded_gpu_gray_dispatch.md b/notes/2026-10-12_bounded_gpu_gray_dispatch.md new file mode 100644 index 0000000..eeb945d --- /dev/null +++ b/notes/2026-10-12_bounded_gpu_gray_dispatch.md @@ -0,0 +1,163 @@ +# Bounded GPU RGB-to-gray dispatch + +Date: 2026-10-12. This named slice follows merged +[PR #145](https://github.com/rsasaki0109/SpatialRust/pull/145), main +`95fb8e854333272c866b2439727c5fe35c360f72`. + +## Problem and resulting behavior + +`rgb_to_gray_gpu` previously dispatched `ceil(width * height / 256)` workgroups +on the X axis. With the requested device limit of 65,535 groups per dimension, +the largest legal one-axis workload is 16,776,960 pixels. A valid 4096x4096 +texture has 16,777,216 pixels and requires 65,536 groups, so the previous call +reached wgpu dispatch validation with an unsupported X count. + +The operation now retains the original linear shader when the group count fits +one axis. Larger inputs use a balanced rectangle of flat 256-lane groups: + +```text +pixels = u64(width) * u64(height) +groups = ceil(pixels / 256) +rows = ceil(groups / max_groups_per_dimension) +columns = ceil(groups / rows) +``` + +Both dispatch axes must fit the device limit. For 4096x4096 the plan is exactly +32,768 columns and two rows, with no padding group. Each group's flat index is +`row * columns + column`. The grid shader rejects padding groups before +multiplying their index by 256, then rejects tail lanes against the pixel count. +This order prevents a padding group at the u32 pixel boundary from wrapping +back to pixel zero. Fewer than `rows` groups are padding; no partial-result +array or second pass is introduced. + +The pure planner checks positive dimensions, texture limits, signed texture +coordinates, a pixel count at most `u32::MAX`, the existing 256-lane workgroup, +both dispatch dimensions, sampled/storage textures, uniform/buffer sizes and +bind-group capacity. It requires no workgroup storage or storage buffer. Runtime +ownership, RGB channel count and one free profiler stage are checked before +texture/uniform allocation or cache initialization. Synthetic counter/limit +tests do not imply that such textures can be allocated on the local device. + +Linear and grid compute pipelines are independently lazy in one runtime-owned +gray family. They share the original bind layout and 16-byte uniform: two +previously unused words now carry the group-row stride and total group count. +The original linear WGSL matches the frozen parent exactly. Every valid pixel +retains `(77*R + 150*G + 29*B + 128) >> 8`, gray metadata with the original color +range, `[gray, 0, 0, 0]` physical output and the `rgb_to_gray_gpu` receipt stage. +One pass still writes four GPU bytes per pixel; explicit timestamp collection +returns 16 query bytes. There is no implicit payload readback, host wait, global +scratch, new dependency, feature or public signature. + +This expands the standalone gray operation's supported dispatch range. Other +kernels retain their own limits; full large-image Vision-chain admission is +not established here. Existing valid six-to-eight-pass Vision-chain regressions +continue to pass. The texture pool still bounds retained bytes, not active +source/output textures or peak driver/process memory. + +## Correctness and feature validation + +Six new native tests cover: + +- **Every one of the 16,777,216 possible RGB values**, packed into the public + 4096x4096 grid path, checked against independent widened integer CPU luma. + This checks **67,108,864 exact physical RGBA bytes**, including zero unused + channels. The cold call initializes only the grid pipeline. +- Nine small/odd/singleton shapes crossed with mixed, zero and full-scale + patterns. Frozen linear and current public outputs match the CPU reference; + **120 forced grid mappings** under 2/3/7/17/31-group limits check tail/padding + groups over **3,655,572 exact RGBA bytes** across variants. +- A GPU probe calls the production grid-index function for tiny plans and + synthetic `u32::MAX`/`u32::MAX-1` pixel counts. **96 exact index/validity words** + cover first/last/padding groups and lanes 0/1/254/255 without allocating giant + textures. Returning before multiplication distinguishes padding from pixel zero. +- Pure dimension, texture, pixel-overflow, compute, binding and uniform-limit + rejection. A two-by-two dispatch admits four groups and rejects five. Zero + workgroup-storage/storage-buffer capacities remain acceptable. +- Wrong runtime/channel/profile ownership and exhausted capacity leave texture + counters and cold gray caches unchanged. Rejection preserves an earlier Sobel + stage and its 16-byte explicit query collection. +- A strided RGB ROI, four concurrent cold linear calls, live outputs while the + original source texture is reused, rapid submissions, moved runtime, warm + allocation reuse and zero texture-retention budget preserve exact results. + The cold small path initializes only the linear pipeline. + +The physical-channel test helper temporarily aliases the same texture with +unknown color-space metadata to allow explicit four-channel readback. It +allocates no texture and copies no texture payload; the public output retains +its original gray metadata and logical channel count. Test readback is explicit +and is outside production receipts and benchmark timers. + +Final native GPU regressions pass: **148 tests**, zero failures, 11 ignored +diagnostics. This includes 121 unit tests and 27 CDF/CLAHE/equalization/histogram/ +shared-stage integration tests, covering images and points. Strict all-feature +and default all-target Clippy, formatting and diff checks pass. Minimal +`gpu-image` and combined `gpu-image,gpu-aoso-staging` WASM compilation pass; +browser execution/performance is not established. + +```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 +``` + +## Resident timing diagnostic + +Release native Windows, Rust/Cargo 1.98.1 and wgpu 24.0.5, Vulkan discrete +backend with pass timestamps. The adapter model is omitted. No tests or builds +launched by this task overlap measurement; desktop activity, clocks and external +GPU load are not controlled. This change expands dispatch support; no speedup +is claimed. + +Seven shapes use two warmups and twenty alternating rounds. Six legal linear +shapes compare frozen and current kernels; the new large shape records only +current results, because its frozen X dispatch is unsupported. Each warmup and +measured result checks all physical output bytes against the CPU oracle, +metadata, the receipt, one pass and 16-byte explicit query collection. + +Completed timers include runtime/plan/profiler validation, output/uniform +allocation, encoding, submission and `wait_idle`. Upload, warmed pipeline setup, +profiler creation, query collection, physical readback/checking and recycling +are outside timers. The frozen shader uses the current host executor and +validation, so this **does not measure historical full-API overhead**. +One runtime and its default 64 MiB texture-retention budget persist across all +conditions; the pool is not cleared between shapes. Outputs are recycled after +checking, so shape and retained-pool occupancy can affect allocation reuse. + +Times below are variant medians. Ratios are medians of twenty paired +frozen/current samples, with 10,000 paired median bootstrap resamples +(seed 20261012, sorted indices 249 and 9749). Above one favors current. These +are descriptive intervals for this run, not independent-repeat confidence or +a portable performance guarantee. Every condition and sample range is retained. + +| Size | Dispatch x/y | Frozen/current GPU us | GPU ratio [95% CI] | Frozen/current completed ms | Completed ratio [95% CI] | +| --- | --- | --- | --- | --- | --- | +| 17x19 | 2/1 | 4.224 / 4.272 | 0.981 [0.974, 1.023] | 0.185 / 0.186 | 0.969 [0.897, 1.015] | +| 128x128 | 64/1 | 4.544 / 4.544 | 1.000 [0.976, 1.025] | 0.180 / 0.184 | 1.001 [0.948, 1.039] | +| 640x480 | 1200/1 | 18.752 / 18.832 | 0.995 [0.989, 1.003] | 0.342 / 0.348 | 0.983 [0.944, 1.116] | +| 1920x1080 | 8100/1 | 94.512 / 94.624 | 1.000 [0.994, 1.004] | 0.605 / 0.585 | 1.014 [0.976, 1.057] | +| 3840x2160 | 32400/1 | 5279.440 / 5290.928 | 0.999 [0.996, 1.002] | 6.750 / 6.835 | 0.996 [0.983, 1.014] | +| 4096x4095 | 65520/1 | 12116.400 / 12109.872 | 1.001 [1.000, 1.002] | 23.599 / 23.741 | 0.993 [0.991, 1.001] | +| 4096x4096 | 32768/2 | N/A / 12098.672 | N/A | N/A / 23.633 | N/A | + +Source and reproduction paths are absolute: + +- `C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/src/image/kernels/gray.rs` +- `C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/src/image/kernels/gray_grid.wgsl` +- `C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/src/image/kernels/gray_tests.rs` +- `C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/src/image/kernels/gray_benchmark.rs` +- `C:/Users/rsasa/Workspace/SpatialRust/crates/spatialrust-gpu/tests/fixtures/gray_previous.wgsl` + +```text +cargo test -p spatialrust-gpu --all-features --release --lib paired_frozen_current_and_large_gray -- --ignored --nocapture --test-threads=1 +``` + +Private logs under `C:/Users/rsasa/Workspace/SpatialRust/target/automation/`: +`bounded-gpu-gray-regressions.log`, `bounded-gpu-gray-clippy.log`, +`bounded-gpu-gray-default-clippy.log`, `bounded-gpu-gray-wasm.log`, +`bounded-gpu-gray-wasm-combined.log`, `bounded-gpu-gray-release-build.log`, +`bounded-gpu-gray-paired.log`, `bounded-gpu-gray-estimates.json`, +`bounded-gpu-gray-table.md` and `summarize-bounded-gpu-gray.py`.