oxigdal-gpu 0.1.6

# TODO: oxigdal-gpu

> **Purpose:** wgpu 29 backed GPU acceleration for OxiGDAL raster ops — WGSL compute kernels (element-wise, resampling, convolution, statistics, FFT), pipelines, memory pool, multi-GPU primitives, backend-specific paths (CUDA/Vulkan/Metal stubs).
> **Status (2026-05-16):** 13,242 LoC · 324 tests (hardware-gated GPU tests remain `#[ignore]`) · 5 real stubs (multi-GPU gather, CUDA workgroup memory, three Metal MPS shaders, Vulkan future-support stubs).
> **Roadmap:** v0.1.5 → v0.2.0 → v1.0.0

## High Priority (verified gaps)
- [x] Lanczos resampling with configurable window size (completed 2026-05-16)
  - **Done:** `ResamplingMethod::Lanczos { a: u32 }` (a in [1,8]) with full WGSL sinc-based 2D separable kernel. `sinc(x)·sinc(x/a)` evaluated per-pixel, 2a×2a neighbourhood, clamp-to-edge boundary, weight-normalised output. Validation (`a == 0` or `a > 8`) returns `GpuError::invalid_kernel_params`.
  - **Files:** `src/kernels/resampling.rs`, `tests/lanczos_test.rs` (new, 14 tests — 8 CPU-only pass immediately, 6 GPU `#[ignore]`).

- [x] Real cross-GPU `gather` implementation (completed 2026-05-17)
  - **Done:** `InterGpuTransfer` in `src/multi_gpu.rs` gains `per_device_buffers: Vec<Arc<Mutex<Option<(Arc<wgpu::Buffer>, u64)>>>>` tracking last-submitted result buffer + size per device. `set_device_buffer(idx, buffer, size)` registers a result. `gather(dst_device)` iterates source devices: allocates `MAP_READ | COPY_DST` staging buffer, encodes `copy_buffer_to_buffer`, submits, `device.poll(PollType::wait_indefinitely())`, `map_async` via mpsc channel, `get_mapped_range().to_vec()`, `unmap`. `gather_blocking` wraps via `pollster::block_on`. Devices with no registered buffer return `Vec::new()`. `MultiGpuManager::new_empty_for_testing()` added for CPU-only tests.
  - **Tests added (3):** test_gather_single_device_returns_empty_vec (single device → Ok([])), test_gather_invalid_dst_device_returns_error, test_gather_blocking_matches_async. Total: passes all CPU-only tests (GPU-gated tests remain `#[ignore]`).

- [x] Async readback with callback support (completed 2026-05-17)
  - **Done:** `GpuBuffer::read_async()` added to `src/buffer.rs` — named async method delegating to the existing oneshot-based `read()` internal path. `ComputePipeline::read_async()` added to `src/compute.rs` — mirrors `read_blocking` with async delegation. `GpuContext::spawn_poll_task()` added to `src/context.rs` — spawns a background thread calling `device.poll(PollType::Poll)` in a 1 ms loop; documented as an advanced helper for runtimes without automatic polling. All tests CPU-only; hardware tests remain `#[ignore]`.
  - **Tests added (5):** test_read_async_future_type_parameters_satisfy_send, test_read_async_and_blocking_api_signatures_compatible, test_oneshot_channel_resolves_on_send, test_poll_type_enum_accessible, test_read_async_error_type_matches_blocking. Total: 315 tests pass (47 pre-existing GPU backend failures unchanged).

- [x] GPU device-lost recovery mid-pipeline (completed 2026-05-17).
  - **Done:** `src/context.rs` extended: `Arc<AtomicBool> device_lost` field added to `GpuContext`; `set_device_lost_callback` call in `with_config()` wires the flag via `device.on_uncaptured_error` / lost callback (wgpu 29 API); `is_device_lost() -> bool` and `check_device_lost() -> GpuResult<()>` public inspectors; `with_device_lost_flag(flag)` builder for mock injection; `reinitialize(&self) -> GpuResult<GpuContext>` re-enumerates adapters via `Instance::enumerate_adapters` with same backend preference and recreates device+queue. `src/compute.rs` gains `check_device_lost()?` before each `queue.submit()` call. `GpuError::DeviceLost` variant added to `src/error.rs`. Tests in `src/device_lost_test.rs` (declared as `#[cfg(test)] mod device_lost_test` in `lib.rs`).
  - **Tests added (5):** test_device_lost_flag_starts_false, test_device_lost_flag_can_be_set_externally, test_device_lost_check_returns_err_when_flag_set, test_device_lost_with_flag_builder_works, test_device_lost_is_device_lost_accessor. Total: 320 tests pass (47 pre-existing GPU backend failures unchanged).

- [x] Workgroup-size auto-tuning per adapter limits (completed 2026-05-17)
  - **Done:** New `src/workgroup_tuner.rs` — `WorkgroupTuner { raster_2d: (u32, u32), reduction: u32, fft: u32 }` with `derive_from_limits(limits: &wgpu::Limits) -> Self` selecting tiers: `(16,16)` if `≥256` total invocations, `(8,8)` if `≥64`, else `(4,4)`; `reduction = min(256, max_x, max_total)`; `fft = min(64, max_x)`. `GpuContext.tuner: WorkgroupTuner` field populated from `adapter.limits()` in `GpuContext::new`; `with_tuner(tuner)` builder for test injection. `src/kernels/resampling.rs` WGSL `@workgroup_size(16, 16)` hardcoded strings replaced with `format!("@workgroup_size({}, {})", w, h)` from `tuner.raster_2d`; `new_with_workgroup_size()` added for explicit override. `WorkgroupTuner` re-exported from `src/lib.rs`. All tests CPU-only (no GPU hardware required).
  - **Tests added (14):** test_tuner_picks_16x16_when_limits_unbounded, test_tuner_falls_back_8x8_on_low_end_adapter, test_tuner_falls_back_4x4_on_minimum_adapter, test_tuner_respects_max_invocations_per_workgroup, test_tuner_reduction_capped_at_256, plus 9 additional tier/boundary/builder tests. Total: 310 tests pass (47 pre-existing GPU backend failures unchanged).

- [x] f16 buffer support for bandwidth-sensitive operations (completed 2026-05-17)
  - **Done:** `src/buffer.rs` extended with `BufferElementType { F32, F16, U8, U16, U32, I32 }` enum with `byte_size()` and `wgsl_type()`. Free functions: `f16_to_f32_slice`, `f32_to_f16_slice`, `from_f16_slice_widening` (uploads as f32, transparent fallback path), `from_f16_slice_native` (raw u8 storage when adapter supports SHADER_F16), `read_f16_from_f32_buffer`. New `src/shader_helpers.rs` — `make_element_wise_shader_source(element_type) -> String` emits `enable f16;` directive and `array<f16>` bindings for F16 type. `GpuContextConfig::with_f16_support()` adds `Features::SHADER_F16` to requested set (best-effort). Re-exported `BufferElementType`, f16 helpers, and `half` re-export from `lib.rs`. Workspace `Cargo.toml` gains `half = { version = "2.7.1", features = ["bytemuck"] }`.
  - **Tests added (8, CPU-only):** test_buffer_element_type_byte_sizes, test_buffer_element_type_wgsl_type_strings, test_f16_round_trip_within_1bit_lsb_via_widening, test_f16_kernel_source_contains_enable_directive_when_f16, test_f16_fallback_widens_to_f32_on_unsupported, test_from_f16_slice_empty, test_f16_to_f32_to_f16_idempotent_for_representable_values, test_f16_f32_roundtrip_special_values. Total: 434 pass / 47 pre-existing GPU backend failures.

- [ ] Backend-specific shader paths beyond placeholders
  - **Verified gap:** `src/backends/metal.rs:374` — `"// MPS image filter shader placeholder\n".to_string()`; `:378` `"// MPS reduction shader placeholder\n"`; `:382` `"// MPS neural network shader placeholder\n"`; `src/backends/cuda.rs:402` — `// In WGSL, this is a placeholder that needs workgroup memory`; `src/backends/vulkan.rs:148` — `// these are placeholders for future support`.
  - **Goal:** Either delete these unused backend hooks (if all paths already go through wgpu's WGSL) or wire them to real Metal-shader-source loading via the existing `backends/metal.rs::compile_metal_shader` stub.
  - **Design:** Audit call sites of each `*_placeholder` constant; if no caller, remove the dead module. Otherwise, expose Metal-Performance-Shaders via `metal-rs` (already an optional dep elsewhere, gated by `metal` feature on macOS).
  - **Files:** `crates/oxigdal-gpu/src/backends/{cuda,metal,vulkan}.rs`.
  - **Tests:** (proposed) `test_metal_path_unused_or_real_compile`, `test_cuda_workgroup_kernel_compiles_when_feature_enabled`, `test_vulkan_subgroup_op_when_supported`.
  - **Risk:** Backend stubs may be exercised by integration tests elsewhere — search workspace before deletion.
  - **Prerequisites:** None.

## Medium Priority
- [x] Radix-2 1D FFT GPU compute kernel (forward + inverse) (completed 2026-05-19).
  - **Done:** New `src/fft.rs` (627 LoC). `bit_reverse(x, log2_n)` pure-Rust bit-reversal. `twiddle_factor(k, n, inverse)` pure-Rust twiddle reference. `make_fft_shader_source(size, inverse) -> String` emits monomorphic WGSL with workgroup-shared Cooley-Tukey DIT butterfly stages fully unrolled (required for WGSL shared-memory constant bounds). `Fft1d { pipeline, bind_group_layout, size, inverse }` — validates size: power-of-two, 4–2048; uses `PipelineCache`; `execute(ctx, real, imag)` dispatches one workgroup per transform; `execute_batch(ctx, batch)` dispatches N workgroups. GPU-side inverse includes 1/N normalization. Re-exported from `lib.rs`.
  - **Tests added (12+9 inline):** 7 pure-Rust (bit_reverse, twiddle_factor, shader source content, size validation), 5 GPU-conditional (catch_unwind: impulse, DC, round-trip, cosine bins, batch). Total: 21 targeted pass.
- [x] FFT-based convolution shader (point-spread-function / large kernel convolution) (completed 2026-05-19).
  - **Done:** New `src/convolution_fft.rs` (~300 LoC). `FftConvolution { ctx: Arc<GpuContext> }`. `convolve(signal, kernel) -> GpuResult<Vec<f32>>` — zero-pads to next power-of-two (capped at `MAX_FFT_CONVOLUTION_SIZE = 2048`), runs `Fft1d::execute` on both, complex-multiplies CPU-side, runs IFFT via `Fft1d::execute(inverse=true)`, trims output to `signal.len() + kernel.len() - 1`. `correlate(signal, kernel)` reverses kernel then convolves. `convolve_batch(signals, kernel)` amortizes the kernel FFT across multiple signals. Pure-Rust `convolve_reference(signal, kernel)` for test parity. `complex_multiply(ar, ai, br, bi) -> (f32, f32)` helper. Re-exported from `lib.rs`.
  - **Tests added (12 in `tests/convolution_fft_test.rs`):** test_convolve_reference_impulse_kernel_identity, test_convolve_reference_two_element_known, test_convolve_reference_shift_by_delta, test_complex_multiply_zero_imaginary, test_complex_multiply_pure_imaginary, test_fft_convolution_new_builds_without_gpu_when_ctx_present, test_fft_convolution_impulse_kernel_identity_when_backend_present, test_fft_convolution_box_kernel_smooths_step_when_backend_present, test_fft_convolution_round_trip_matches_reference_when_backend_present, test_fft_convolution_correlate_shifts_opposite_direction_when_backend_present, test_fft_convolution_batch_matches_single_results_when_backend_present, test_fft_convolution_large_kernel_exceeds_max_returns_error. Total: 552 pass, 47 pre-existing headless failures.
- [x] Texture-based resampling using hardware samplers (completed 2026-05-18).
  - **Done:** New `src/texture_resample.rs` (~470 LoC). `TextureFilterMethod { Nearest, Linear }` with `wgpu_filter()` mapping to `wgpu::FilterMode`. `texture_filter_for_resampling(method) -> Option<TextureFilterMethod>` — NearestNeighbor → Nearest, Bilinear/Bicubic → Linear (bicubic falls back to linear, documented), Lanczos { .. } → None (caller uses compute-buffer path). `TextureResampler { pipeline, bind_group_layout, sampler, method, dst_format }` — wraps a wgpu sampler-bound compute pipeline; uses `Features::TEXTURE_BINDING | StorageTexture write`. `new(ctx, method, dst_format)` builds pipeline with shader from `make_texture_resample_shader_source`. `make_texture_resample_shader_source(filter, dst_fmt) -> String` — emits WGSL with `texture_2d<f32>`, `sampler` binding (filterable), `textureSampleLevel(src, samp, uv, 0.0)`, `textureStore` for output. `new_input_texture_r32float(ctx, w, h, &[f32]) -> GpuResult<wgpu::Texture>` uploads via `queue.write_texture` with `TexelCopyTextureInfo`/`TexelCopyBufferLayout`. `dispatch()` method for end-to-end sampling. Re-exported from `lib.rs`.
  - **Tests added (22):** 8 inline unit tests + 14 integration tests (3 GPU-required using Slice 13 `try_gpu_context` + `catch_unwind` pattern) covering shader source contents, filter mapping for all resampling methods, format validation, size-mismatch errors, and end-to-end texture round-trip on live backend. Total: 471 pass.
- [x] Pipeline caching keyed by `(shader_hash, layout, constants)` (completed 2026-05-17).
  - **Done:** New `src/pipeline_cache.rs` (~345 LoC). `fnv1a_64(data) -> u64` — self-contained FNV-1a 64-bit hash. `PipelineCacheKey { shader_hash, entry_point, layout_tag }` — `Hash+Eq` key with `new(source, entry, tag)` (hashes source inline), `from_hash`, and `Display` impl. `PipelineCache` — `HashMap`-backed store of `Arc<wgpu::ComputePipeline>`; `get_or_insert_with(key, factory) -> Result<Arc<Pipeline>, E>` caches on hit, calls factory on miss; `len()`, `is_empty()`, `clear()`, `evict(key)`, `keys()`, `retain(pred)`. `SharedPipelineCache = Arc<Mutex<PipelineCache>>` + `new_shared_pipeline_cache()`. `GpuContext` gains `pipeline_cache: SharedPipelineCache` field (initialized in `with_config`; cleared in `reinitialize()`); `pipeline_cache()` accessor. Re-exported from `lib.rs`.
  - **Tests added (49 total new):** 26 in `tests/pipeline_cache_test.rs` + 18 inline unit tests in `pipeline_cache.rs` + 5 inherited. All CPU-only. Total: 369 pass (47 pre-existing GPU backend failures unchanged).
- [x] Tiled processing for rasters exceeding VRAM budget (completed 2026-05-17).
  - **Done:** New `src/tiled.rs` (~504 LoC). `TiledConfig { tile_width, tile_height, overlap_pixels, vram_safety_margin }` with builder methods. `RasterTile { data, width, height, overlap_*, origin_*, raster_*, tile_index }` with `padded_width/height/len()`. `split_into_tiles(raster, w, h, config) -> Vec<RasterTile>` — row-major tile extraction with edge-replication halo (clamp-to-edge at raster boundary). `stitch_tiles(tiles, raster_w, raster_h) -> Vec<f32>` — writes only core interior of each tile, discards halo. `vram_per_tile(tile) -> usize` — `padded_len×4×2+256`. `auto_tile_size(preferred_w, preferred_h, overlap, vram_budget, safety_margin) -> (usize,usize)` — halves wider dimension iteratively until tile fits, floors at (16,16). `execute_tiled<F>(raster, w, h, config, tile_fn) -> GpuResult<Vec<f32>>` — split→per-tile fn→stitch. Re-exported from `lib.rs`.
  - **Tests added (18):** test_split_into_tiles_exact_fit, test_split_into_tiles_non_exact_fit, test_split_overlap_adds_halo_pixels, test_split_overlap_edge_replication_at_corners, test_stitch_reconstructs_identity, test_stitch_non_overlapping_tiles, test_vram_per_tile_formula, test_auto_tile_size_halves_to_fit_budget, test_auto_tile_size_preferred_fits_in_large_budget, test_execute_tiled_passthrough_matches_original, test_execute_tiled_scale_fn_applies_per_tile, test_split_single_tile_covers_all, plus 6 additional. Total: 426 pass (47 pre-existing GPU backend failures unchanged).
- [x] Storage-texture support for direct image output (skips final readback) (completed 2026-05-17).
  - **Done:** New `src/storage_texture.rs` (~729 LoC). `StorageTextureBinding { texture, view, format, width, height }`. `StorageTextureKernel { pipeline: Arc<ComputePipeline>, bind_group_layout }`. `is_supported_storage_format(format)` — predicate for Rgba32Float/Rgba8Unorm/R32Float. `make_storage_texture_shader_source(format) -> String` — WGSL kernel that copies flat `array<f32>` input into `texture_storage_2d<fmt, write>` output at 16×16 workgroup. `new_storage_texture(ctx, w, h, format) -> GpuResult<StorageTextureBinding>` — validates format, creates texture with `STORAGE_BINDING | COPY_SRC` usage. `build_storage_texture_kernel(ctx, wgsl, entry, format)` — pipeline cached via `SharedPipelineCache`. `StorageTextureKernel::dispatch_to_texture(ctx, input_buf, texture)` — encodes bind group + `dispatch_workgroups(⌈w/16⌉, ⌈h/16⌉, 1)`, submits without CPU readback. `read_texture_to_vec_f32(ctx, tex)` — opt-in parity-test fallback (texture→staging→host). `GpuError::UnsupportedFormat(String)` added to `src/error.rs`. Re-exported from `lib.rs`.
  - **Tests added (17):** 12 unit tests (pure-Rust, no GPU required) + 5 integration tests (4 GPU-skippable via `catch_unwind`). Total: 451 pass (47 pre-existing GPU backend failures unchanged).
- [x] Band-math expression compiler that generates optimized WGSL from algebraic strings (completed 2026-05-18).
  - **Done:** New `src/band_math/mod.rs` + `src/band_math/{parser.rs (~558 LoC), codegen.rs (~343 LoC)}`. `parser.rs` — hand-rolled `Lexer` over byte stream + precedence-climbing `Parser` with grammar `expr → term → factor → power → unary → primary`; `BandMathError` enum (UnexpectedToken, UnknownFunction, InvalidBandIndex, DivByConstantZero, TrailingTokens, EmptyExpression, UnmatchedParen). Band syntax 1-based externally (`B1..B999`), 0-based in AST. Supported functions: `log`/`ln` (natural log), `exp`, `sqrt`, `abs`, `min`, `max`, `pow`, `clamp`. `^` parsed and lowered to `pow(a,b)` (WGSL has no `^`); right-associative. Literal `B/0.0` rejected at parse time. `codegen.rs` — `band_expression_to_wgsl(expr, &input_bindings)` emits bind groups sorted by ascending band index, fixed `@workgroup_size(64)`, bounds-checked write to `output[idx]`. Recursive `constant_fold` runs before emission; `f32` literals always include decimal point. `src/algebra.rs` extended with `Min`, `Max`, `Pow`, `Log`, `Exp`, `Clamp { value, lo, hi }` variants on `BandExpression`; `evaluate` updated. Re-exported `parse_band_expression`, `band_expression_to_wgsl`, `BandMathError` from `lib.rs`.
  - **Tests added (24):** 8 inline unit tests + 16 integration tests covering parse, constant folding, error paths, NDVI formula, function calls, round-trip AST eval. Total: 458 pass (47 pre-existing GPU backend failures unchanged).
- [x] Compute-pipeline profiling with GPU timestamp queries (completed 2026-05-18).
  - **Done:** New `src/profiling.rs` (~376 LoC). `GpuTimestampProfiler { enabled, capacity, next_slot, period_ns, labels, query_set, resolve_buffer, staging_buffer }`. `try_new(ctx, capacity) -> Option<Self>` — returns `None` if `Features::TIMESTAMP_QUERY` and/or `Features::TIMESTAMP_QUERY_INSIDE_ENCODERS` absent; rounds capacity up to even (min 2); creates `QuerySet` with `QueryType::Timestamp`, plus `QUERY_RESOLVE | COPY_SRC` resolve buffer and `MAP_READ | COPY_DST` staging buffer (8 bytes per slot). `dummy(period_ns)` — test-only no-GPU constructor. `begin_pass(encoder, label) -> Option<u32>` writes start timestamp via `encoder.write_timestamp(qs, slot)`, returns slot. `end_pass(encoder, start_slot)` writes end timestamp at `start_slot + 1`. `resolve(ctx) -> GpuResult<Vec<PassTiming>>` — encodes `resolve_query_set` + `copy_buffer_to_buffer`, submits, polls device with `PollType::wait_indefinitely()`, maps staging buffer synchronously via mpsc channel, decodes u64 timestamps, computes `(end - start) * period_ns / 1000.0` µs per pair. `PassTiming { label, start_ns, end_ns, duration_us }`. Accessors `is_enabled`, `period_ns`, `capacity`, `next_slot`, `pass_labels`, `reset`. Re-exported from `lib.rs`.
  - **Tests added (13):** 4 inline unit tests (dummy constructor, period, capacity rounding, multi-pair) + 9 integration tests in `tests/profiling_test.rs` using `try_gpu_context()` catch_unwind pattern for graceful no-backend skip. Total: 471 pass (47 pre-existing GPU backend failures unchanged).
- [x] Automatic CPU fallback on GPU failure/timeout (completed 2026-05-17).
  - **Done:** New `src/cpu_fallback.rs` (~370 LoC). `FallbackConfig { enabled, auto_fallback_on_error, log_fallbacks, max_gpu_retries }`. `ExecutionPath` enum (Gpu/Cpu). `FallbackResult<T> { value, path, gpu_error }`. `FallbackConfig::should_fallback(err)` checks `GpuError::DeviceLost/NoAdapter/BackendNotAvailable`. `execute_with_fallback(config, gpu_op, cpu_op) -> FallbackResult<T>` retries GPU up to `max_gpu_retries`, falls back to CPU on matching errors. `execute_with_fallback_timed` adds microsecond timing. `FallbackContext` wraps `GpuContext` + config + stats tracking. `cpu` submodule exposes pure-Rust implementations: `add`, `sub`, `mul`, `div`, `scalar_mul`, `clamp`, `mean`, `min`, `max`, `sum`, `normalize`, `dot`, `dot_product`, `transpose_2d`. Re-exported from `lib.rs`.
  - **Tests added (36):** test_fallback_config_default, test_fallback_config_disabled, test_should_fallback_device_lost, test_should_fallback_no_adapter, test_should_fallback_backend_not_available, test_no_fallback_other_error, test_execute_gpu_success, test_execute_fallback_on_error, test_cpu_add, test_cpu_mul, test_cpu_clamp, test_cpu_mean, test_cpu_normalize, test_cpu_transpose_2d, test_fallback_context_creation, plus 21 additional. Total: 405 pass (47 pre-existing GPU backend failures unchanged).
- [ ] WebGPU browser compatibility testing on wasm32-unknown-unknown.
  - **Files:** `.github/workflows` (out of scope per CI policy), local `cargo test --target wasm32-unknown-unknown`.
  - **Why deferred:** CI policy restricts new yaml files.

## Low Priority / Future (one-liners)
- [x] Ray-marching shader for volumetric DEM rendering (completed 2026-05-19).
  - **Done:** New `src/ray_march.rs` (~430 LoC). `RayMarchConfig { view_dir, light_dir, step_size, max_steps, vertical_exaggeration, pixel_size_world }` with `validate()`. `RayMarchResult { width, height, shaded, depth }`. `DemRayMarcher` GPU compute struct — builds WGSL pipeline in `new()`, dispatches 16×16 workgroups in `march()`. `make_ray_march_shader_source()` WGSL shader with 4 bindings (`@binding(0)` DEM, `@binding(1)` uniform, `@binding(2)` shaded out, `@binding(3)` depth out). CPU reference: `ray_march_cpu()`, `normalize3()`, `sample_dem_bilinear()`. Re-exported from `lib.rs`.
  - **Tests added (13):** 12 pure-CPU + 1 GPU catch-unwind. Total: passes.
- [x] wgpu push constants when adapter supports them (completed 2026-05-19).
  - **Done:** New `src/push_constants.rs` (~450 LoC). Adapted for wgpu 29 which renamed `PUSH_CONSTANTS` to `IMMEDIATES` and `var<push_constant>` to `var<immediate>`. `PushConstantRange { stages, start, end }` with `compute(size)`, `validate()`, `to_wgpu()`. `PushConstantsLayout { ranges, total_size }` with `compute_only(size)`, `validate()`, `to_wgpu_ranges()`, `immediate_size_for_wgpu()`. `PushConstantsBuffer { data, layout }` with `write<T: NoUninit>`, `write_u32/f32/vec4_f32`. `supports_push_constants(ctx)`, `max_push_constants_size(ctx)`, `make_push_constants_shader_source`, `build_push_constants_pipeline`, `dispatch_with_push_constants`. `GpuContextConfig::with_push_constants()` requests `Features::IMMEDIATES`. Re-exported from `lib.rs`.
  - **Tests added (30 in `tests/push_constants_test.rs`):** 8 unit + 22 integration; pure-Rust tests run unconditionally, GPU tests use `try_gpu_context()` + catch-unwind.
- [x] Cooperative-matrix operations for ML inference on GPU. (completed 2026-05-19)
  - **Done:** New `src/cooperative_matrix.rs` (~430 LoC). `CoopMatrixComponentType { F16, F32, I8, U8 }` with `as_wgsl()`, `byte_size()`. `CoopMatrixUse { A, B, Accumulator }`. `CoopMatrixDim { m, n, k }` (default 16×16×16). `CoopMatrixGemmConfig { dim, a_type, b_type, accum_type, workgroup_size }`. `supports_cooperative_matrix()` checks `Features::EXPERIMENTAL_COOPERATIVE_MATRIX | SUBGROUP`. `make_gemm_wgsl()` — fully functional workgroup-tiled GEMM WGSL kernel; emits cooperative-matrix comment for detectability; `make_gemm_wgsl_fallback()` for simpler path. `build_cooperative_matrix_gemm_pipeline()`, `dispatch_cooperative_gemm()`. `GpuContextConfig::with_cooperative_matrix()`. Re-exported from `lib.rs`.
  - **Tests added (13 in `tests/cooperative_matrix_test.rs` + 5 unit):** 8 pure struct/WGSL-string tests; 5 GPU tests with catch-unwind fallback for headless envs.
- [x] Indirect dispatch for adaptive workload sizing. (completed 2026-05-19)
  - **Done:** New `src/indirect_dispatch.rs` (~220 LoC). `#[repr(C)] DispatchIndirectArgs { x, y, z: u32 }` (12-byte little-endian, stable layout). `IndirectDispatchBuffer` wrapping `Arc<wgpu::Buffer>` with `INDIRECT | COPY_DST | STORAGE` usage; `new()`, `new_with_capacity()`, `update()`, `update_at()`. `dispatch_indirect_on_pass(pass, pipeline, bind_groups, indirect)` convenience one-shot helper. Pure helpers: `workgroup_count_1d/2d/3d`, `args_for_elements`. Re-exported from `lib.rs`.
  - **Tests added (12):** 8 pure math/layout tests (no GPU required) + 4 catch-unwind GPU tests (skip gracefully on headless). Full end-to-end test uses no-op WGSL shader writing 42u via `dispatch_workgroups_indirect`, reads back and verifies. Total: 12 pass.
- [ ] Benchmarking suite (GPU vs CPU for each kernel).
- [x] Shader hot-reload file-watcher (skeleton already in `src/shader_reload.rs`) (completed 2026-05-20).
  - **Done:** New `src/shader_reload_native.rs` (~272 LoC, `#[cfg(feature = "shader-hot-reload")]`). `FilesystemPoller { mtimes, poll_interval, last_poll }` polls `std::fs::metadata().modified()`: `register_path` (seeds mtime, propagates `NotFound`), `track_expected` (UNIX_EPOCH sentinel for not-yet-created files), `deregister_path`, `watched_paths` (sorted), `poll` (throttled to `poll_interval`), `force_poll` (bypasses throttle). Emits `PolledChange { path, kind: PolledChangeKind::{Modified, Deleted, Created} }`; a shared `diff_against_disk` helper iterates sorted keys, distinguishes Created (sentinel→exists) from Modified, and never reports a still-missing sentinel as Deleted. `read_shader_source` strips a leading UTF-8 BOM. `ShaderWatcher::poll_filesystem(&mut self, &mut FilesystemPoller)` (one method added to `shader_reload.rs`) maps Modified/Created paths to source labels, re-reads, calls `update_source`, returns aggregated `ShaderChangeEvent`s; unmatched labels / deleted files skipped gracefully. New `shader-hot-reload = []` feature (std-only, no new deps).
  - **Tests added (15):** 13 in `tests/shader_reload_native_test.rs` (default interval, register seeds / nonexistent-io-error, deregister, sorted watched paths, modify/delete/create detection via `File::set_modified`, throttle / force_poll, watcher round-trip, BOM strip, missing-label) + 2 inline. All 15 pass (the crate's pre-existing headless-GPU hardware tests are environment-dependent and unaffected by this change).
- [x] Multi-format texture compression output (BC1, BC4) (completed 2026-05-19).
  - **Done:** New `src/texture_compress.rs` (~580 LoC). `TextureFormat::{Bc1RgbUnorm, Bc4RUnorm}`. `TextureCompressor { format, pipeline, width, height }`. `new(ctx, format, w, h)` validates `w % 4 == 0 && h % 4 == 0`, compiles monomorphic WGSL shader. `compress(ctx, input) -> GpuResult<Vec<u8>>` — dispatches one thread per 4×4 block (BC1: RGBA8 → 8 bytes/block; BC4: R8 → 8 bytes/block); falls back to CPU loop on headless. `validate_texture_dimensions(w, h)`. Pure-CPU helpers: `compress_bc1_block_cpu(&[u8;64]) -> [u8;8]` (min/max-luminance RGB565 endpoints + 2-bit palette), `compress_bc4_block_cpu(&[u8;16]) -> [u8;8]` (min/max R endpoints + 3-bit 8-interpolation indices), `quantize_rgb565`, `dequantize_rgb565`, `nearest_index_4`. ASTC/ETC2 deferred — require optional adapter features. Re-exported from `lib.rs`.
  - **Tests added (22 total: 11 inline + 11 integration in `tests/texture_compress_test.rs`):** test_quantize_dequantize_round_trip_basics, test_validate_texture_dimensions_rejects_non_multiple, test_compress_bc1_block_cpu_solid_color_emits_equal_endpoints, test_compress_bc1_block_cpu_two_color_block, test_compress_bc4_block_cpu_solid_value, test_compress_bc4_block_cpu_two_value, test_nearest_index_4_black_endpoint, test_texture_compressor_rejects_non_multiple_of_4_width/height, test_texture_compressor_bc1/bc4_8x8_when_backend_present, test_output_length_correct_when_backend_present. Total: 574 passed, 47 pre-existing headless-backend failures skipped.

## Cross-crate dependencies
- **Blocks:** oxigdal-gpu-advanced (extends this crate), oxigdal-services (GPU-accelerated raster ops), oxigdal-ml (CUDA/Vulkan/Metal EP coverage).
- **Blocked by:** wgpu 29 (workspace pin), oxionnx (for ML interop).

## Recently completed (verbatim)
*(No `[x]` entries on previous TODO. Memory notes wgpu 29 API migration fixes for this crate are project-level commits.)*

---
*Last audited: 2026-05-17*