diff --git a/docs/evidence/windows-alpha-phase-v1.md b/docs/evidence/windows-alpha-phase-v1.md new file mode 100644 index 0000000..381bb58 --- /dev/null +++ b/docs/evidence/windows-alpha-phase-v1.md @@ -0,0 +1,84 @@ +# Portable cutout coverage phase + +Sponza and skinned/alpha motion now pass the existing portable image thresholds +on Radeon 760M/Vulkan using the modern tier and software GI. The correction makes +the approved material-space coverage phase explicit with integer arithmetic. +Approved images, quality thresholds, and sampled alpha textures are unchanged. + +## First divergent stage + +Exact attachment captures at diagnostic source `30e7625` expose the cutout +decision on the nearest opaque card, including rejected fragments. Separate U +and V captures preserve each coordinate's complete f32 bit pattern in the albedo +MRT. Both channels have identical depth, HDR inputs, and decision flags on each +backend. On Windows, all 51,820 accepted nearest-card fragments match the +unmodified canonical scene's depth within `1e-6`. These probes change occlusion +and therefore describe the nearest card, not every layer of the original leaf. + +Across 103,644 corresponding card pixels, most UV differences are tiny: the +median absolute difference is about `1.49e-8` in each coordinate. Despite that, +69,212 Bayer thresholds and 11,131 survival decisions differ. Only 267 survival +disagreements occur where the thresholds agree. Coplanar geometry can select +different triangles at the same depth, so large UV outliers are retained. + +The shader computed its phase extent with: + +```wgsl +floor(vec2(dimensions) * exp2(-max(floor(lod), 1.0))) +``` + +On the tested Radeon, powers of two produced the full mip dimension. On hosted +Metal, the observed phase extent was one smaller for the power-of-two foliage +texture. The smaller extent predicts all 101,403 inspected Metal thresholds +away from integer LOD boundaries; the full extent disagrees at 67,103 of those +pixels. Conversely, the full extent predicts all 101,419 inspected Windows +thresholds. The exclusions account for half-precision LOD in the HDR capture; +UVs retain full precision. + +WGSL permits error in `exp2`, so an approximate result immediately below an +integer boundary can change the subsequent floor. The old expression required +more numerical precision than its operation guarantees. This observation does +not establish a driver defect. See the [WGSL floating-point accuracy contract](https://www.w3.org/TR/2026/CRD-WGSL-20260831/#floating-point-accuracy). + +## Correction + +Scene and shadow shaders now derive the repeating phase extent from the final +source texel index using integer arithmetic: + +```text +level = clamp(floor(lod), 1, 31) +phase_extent = max((texture_dimensions - 1) >> level, 1) +``` + +This preserves the grid observed in the approved Metal captures and makes it +identical across backends. Reducing the final source texel index also handles +odd texture dimensions: subtracting one from the already reduced mip count +would change those cases incorrectly. The phase calculation leaves the alpha +sample's coordinates, filtering, mip selection, and coverage probability intact. +The level bound and minimum extent handle very coarse LODs and one-texel axes. + +## Local image results + +The same frozen example commands, modern capability tier, software GI, 120 +warm-up frames, 240 measured frames, and raw attachment export produced: + +| Scene | Previous SSIM | Corrected SSIM | Corrected luminance RMSE | Image gate | +| --- | ---: | ---: | ---: | --- | +| Sponza | 0.969274342 | 0.986160457 | 0.009024118 | Pass | +| Skinned/alpha motion | 0.931530774 | 0.990073442 | 0.017359400 | Pass | + +These focused captures establish image behavior, not hardware timing budgets. +Their executable hashes, commands, raw attachments, metrics, and logs are retained +under `tools/quality/out/windows-engine-plan/alpha-phase-fix/`. + +The GPU regression executes the production scene and shadow threshold functions +against twelve fixed reference points. It covers power-of-two and odd extents, +UV repetition, fractional LOD, one-texel axes, and coarse LODs. The corrected +Radeon/Vulkan and DX12 paths pass all 24 values. The old Vulkan shader fails the +first reference point (`0.03125` instead of `0.34375`). Local contracts, strict +lint, and the complete shared suite pass, including 91 reported golden-test +passes with four ignored tests; optional external-input fixtures remain separate +acceptance requirements. Full-corpus, repeated-run, and hosted validation are +recorded separately as they complete. This change does not close +the remaining temporal, packaging, platform, or named RTX 4080 requirements in +the [engine completion plan](../windows-engine-plan.md). diff --git a/docs/windows-engine-plan.md b/docs/windows-engine-plan.md index c07897a..0078685 100644 --- a/docs/windows-engine-plan.md +++ b/docs/windows-engine-plan.md @@ -15,7 +15,7 @@ first nine-scene Radeon evidence are in draft PR #154. Follow-up work starts at | Work | Required completion evidence | Current state | | --- | --- | --- | | #127 Vulkan PT correctness | Three deterministic progressive and motion runs, both negative controls, finite intermediates, reset/lighting/rigid-motion checks, retained report | Canonical hardware gate, all four focused temporal tests, and CPU reference sanity check pass on Radeon/Vulkan; [report](evidence/issue-127-windows-vulkan-v1.md) and [raw evidence](https://github.com/Bloom-Engine/engine/releases/tag/quality-evidence-155-windows-vulkan-20260910) published | -| #128 Windows image discrepancies | Identify the first incorrect stage or document a reviewed backend-specific baseline decision; rerun the full strict corpus and reproducibility checks | Sponza and skinned/alpha still fail against portable baselines; original baseline source reproduces both failures | +| #128 Windows image discrepancies | Identify the first incorrect stage or document a reviewed backend-specific baseline decision; rerun the full strict corpus and reproducibility checks | [Cutout phase correction](evidence/windows-alpha-phase-v1.md) makes both focused Windows images pass unchanged portable baselines; full strict corpus and reproducibility checks are pending | | #135 / #149 temporal reconstruction | Enforced motion/producer/quality-preset corpus, representative scenes, fractional/native and frozen A/B timing, memory/resize checks, platform evidence | Device lifetime and resource gates repaired; [stationary SSGI fix](evidence/windows-ssgi-stationary-v1.md) and [profiler correction](evidence/windows-profiler-integrity-v1.md) pass 90 golden tests with 4 ignored and 2 optional external-input skips. Complete-phase/lighting control and corrected frozen A/B recorded. HD TAA-jitter and the full representative corpus remain open | | #140 integration gates | Same required local/hosted lanes pass on exact source; release package startup and all-example evidence | All 23 hosted checks pass at #155 source `64d5eed`, including macOS shared/golden tests, all mobile target builds, native/web builds, and browser startup. Scheduled physical-hardware checks, all-example compilation, and release-install acceptance remain separate requirements | | #138 capability fallback | Actual constrained-adapter startup and relevant forced-tier corpus, truthful capability outputs | Existing implementation/evidence preserved; physical constrained-limit acceptance still needs proof | @@ -88,9 +88,16 @@ audit are saved in `tools/quality/out/windows-engine-plan/plan-requirements.json Skinned/alpha has 9,745 depth coverage disagreements before TAA, while albedo RGB closely agrees on matching surfaces. Disabling foliage shadows and an isolated isotropic alpha-sampling control retain the failure. - The cutout decision is the next diagnostic target; its precise cause remains - unresolved. Raw export leaves both Windows final PNGs byte-identical. + Exact cutout-input probes at `30e7625` identify a different Bayer phase + extent: Metal's observed extent predicts every inspected threshold away from + integer LOD boundaries. The [integer phase correction](evidence/windows-alpha-phase-v1.md) + preserves that approved grid on both backends. Focused Windows images now + pass at SSIM 0.986160457 and 0.990073442; full strict corpus and repeated-run + validation follow. Raw export leaves unmodified Windows final PNGs byte-identical. Shared-runner timing cannot qualify hardware budgets. + Separately, the HD TAA fixture fails at its required 16-frame warm-up but + passes diagnostic controls at 32, 64, and 128 frames. The startup/settling + requirement remains open; its warm-up and thresholds are unchanged. 4. Continue starter/all-example and release-install checks, asset/world streaming, schema-generated APIs, components, and runtime UI against each issue's full acceptance criteria. Hardware-specific acceptance remains open while local diff --git a/native/shared/src/renderer/shaders/core.rs b/native/shared/src/renderer/shaders/core.rs index 5d8131b..2fc3fdd 100644 --- a/native/shared/src/renderer/shaders/core.rs +++ b/native/shared/src/renderer/shaders/core.rs @@ -321,12 +321,16 @@ fn mask_coverage_threshold( // instead. Quantising the phase LOD keeps it fixed between mip boundaries; // TAA only has to absorb the occasional boundary transition, not a new // decision on every subpixel camera step. - let phase_lod = max(floor(lod), 1.0); - let mip_scale = exp2(-phase_lod); - let mip_dimensions = max( - floor(vec2(dimensions) * mip_scale), - vec2(1.0), - ); + // Preserve the phase extent in the approved Metal captures: the final + // source texel index reduced to the owning mip. exp2 followed by floor + // made that extent one texel larger on Vulkan when its power-of-two result + // was exact. Integer arithmetic fixes the phase on every backend, including + // non-power-of-two textures, without changing the sampled alpha mip. + let phase_lod = u32(clamp(floor(lod), 1.0, 31.0)); + let mip_dimensions = vec2(max( + (dimensions - vec2(1u)) >> vec2(phase_lod), + vec2(1u), + )); let texel = vec2(floor(wrapped_uv * mip_dimensions)); let x = texel.x & 3u; let y = texel.y & 3u; diff --git a/native/shared/src/renderer/shaders/core_tests.rs b/native/shared/src/renderer/shaders/core_tests.rs index d3376c1..bd09072 100644 --- a/native/shared/src/renderer/shaders/core_tests.rs +++ b/native/shared/src/renderer/shaders/core_tests.rs @@ -74,7 +74,7 @@ fn selected_shadow_cascade_miss_hands_off_instead_of_punching_a_lit_hole() { fn masked_coverage_phase_follows_authored_texture_coordinates() { wgpu::naga::front::wgsl::parse_str(SCENE_SHADER) .unwrap_or_else(|error| panic!("ordinary scene WGSL failed: {error:?}")); - assert!(SCENE_SHADER.contains("let phase_lod = max(floor(lod), 1.0);")); + assert!(SCENE_SHADER.contains("let phase_lod = u32(clamp(floor(lod), 1.0, 31.0));")); assert!(SCENE_SHADER.contains("wrapped_uv * mip_dimensions")); assert!(!SCENE_SHADER.contains("mask_coverage_threshold(in.clip_position.xy")); } diff --git a/native/shared/src/shadows.rs b/native/shared/src/shadows.rs index af896b8..a728752 100644 --- a/native/shared/src/shadows.rs +++ b/native/shared/src/shadows.rs @@ -136,12 +136,13 @@ fn mask_coverage_threshold( // Match the scene pass's sampled-mip footprint. A level-zero phase at // distance makes one shadow texel cross many binary leaf decisions during // even a sub-texel cascade translation, which appears as bright sparkle. - let phase_lod = max(floor(lod), 1.0); - let mip_scale = exp2(-phase_lod); - let mip_dimensions = max( - floor(vec2(dimensions) * mip_scale), - vec2(1.0), - ); + // Match the approved material-space phase with exact integer extents. + // exp2/floor rounded the phase differently on Metal and Vulkan. + let phase_lod = u32(clamp(floor(lod), 1.0, 31.0)); + let mip_dimensions = vec2(max( + (dimensions - vec2(1u)) >> vec2(phase_lod), + vec2(1u), + )); let texel = vec2(floor(wrapped_uv * mip_dimensions)); let x = texel.x & 3u; let y = texel.y & 3u; @@ -1178,7 +1179,7 @@ mod shader_tests { #[test] fn cutout_shadow_coverage_phase_follows_authored_texture_coordinates() { - assert!(SHADOW_SHADER_CUTOUT.contains("let phase_lod = max(floor(lod), 1.0);")); + assert!(SHADOW_SHADER_CUTOUT.contains("let phase_lod = u32(clamp(floor(lod), 1.0, 31.0));")); assert!(SHADOW_SHADER_CUTOUT.contains("wrapped_uv * mip_dimensions")); assert!(!SHADOW_SHADER_CUTOUT.contains("mask_coverage_threshold(in.pos.xy")); } diff --git a/native/shared/tests/golden_render.rs b/native/shared/tests/golden_render.rs index 4909550..98f5f13 100644 --- a/native/shared/tests/golden_render.rs +++ b/native/shared/tests/golden_render.rs @@ -32,6 +32,8 @@ use golden_device::{requested_backends, try_engine, try_isolated_engine}; #[path = "golden_render/metrics.rs"] mod metrics; use metrics::{calculate_diff_metrics, select_outlier_gate}; +#[path = "golden_render/alpha_coverage_phase.rs"] +mod alpha_coverage_phase; #[path = "golden_render/fractional_dynamic.rs"] mod fractional_dynamic; #[path = "golden_render/layered_pbr_motion.rs"] diff --git a/native/shared/tests/golden_render/alpha_coverage_phase.rs b/native/shared/tests/golden_render/alpha_coverage_phase.rs new file mode 100644 index 0000000..f679309 --- /dev/null +++ b/native/shared/tests/golden_render/alpha_coverage_phase.rs @@ -0,0 +1,135 @@ +use super::*; +use wgpu::util::DeviceExt; + +fn phase_function(source: &str, name: &str) -> String { + let source = source.replace("\r\n", "\n"); + let begin = source.find("fn mask_coverage_threshold(").unwrap(); + let end = begin + source[begin..].find("\n}\n").unwrap() + 3; + source[begin..end].replace("mask_coverage_threshold", name) +} + +#[test] +fn cutout_phase_preserves_the_approved_grid_on_scene_and_shadow_paths() { + let Some(eng) = try_engine() else { + eprintln!("skip: no GPU adapter"); + return; + }; + let device = &eng.renderer.device; + let queue = &eng.renderer.queue; + // UV, LOD, padding, texture dimensions, approved Bayer rank, padding. + // Fixed reference points cover power-of-two and odd extents, UV wrapping, + // fractional LOD, one-texel axes, and shifts beyond the integer bit width. + let cases: [[f32; 8]; 12] = [ + [0.25, 0.25, 1.0, 0.0, 1024.0, 1024.0, 5.0, 0.0], + [0.25, 0.5, 2.0, 0.0, 1024.0, 512.0, 5.0, 0.0], + [0.25, 0.25, 1.0, 0.0, 513.0, 257.0, 0.0, 0.0], + [0.75, 0.75, 1.0, 0.0, 1.0, 1.0, 0.0, 0.0], + [0.25, 0.25, 1000.0, 0.0, 1024.0, 1024.0, 0.0, 0.0], + [-0.75, 1.25, 1.0, 0.0, 1024.0, 1024.0, 5.0, 0.0], + [0.25, 0.25, 1.99, 0.0, 1024.0, 1024.0, 5.0, 0.0], + [0.25, 0.25, 0.0, 0.0, 1024.0, 1024.0, 5.0, 0.0], + [0.5, 0.5, 1.0, 0.0, 8.0, 8.0, 4.0, 0.0], + [0.5, 0.5, 2.0, 0.0, 8.0, 8.0, 0.0, 0.0], + [0.5, 0.5, 3.0, 0.0, 193.0, 65.0, 0.0, 0.0], + [0.25, 0.25, 1.0, 0.0, 1.0, 1024.0, 15.0, 0.0], + ]; + let source = format!( + "{}\n{}\n{}", + phase_function( + include_str!("../../src/renderer/shaders/core.rs"), + "scene_phase" + ), + phase_function(include_str!("../../src/shadows.rs"), "shadow_phase"), + r#" +struct Case { coordinate: vec4, extent: vec4 }; +@group(0) @binding(0) var cases: array; +@group(0) @binding(1) var output: array>; +@compute @workgroup_size(64) +fn check_phase(@builtin(global_invocation_id) id: vec3) { + if (id.x >= arrayLength(&cases)) { return; } + let c = cases[id.x]; + output[id.x] = vec2( + scene_phase(c.coordinate.xy, vec2(c.extent.xy), c.coordinate.z), + shadow_phase(c.coordinate.xy, vec2(c.extent.xy), c.coordinate.z), + ); +} +"#, + ); + let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("cutout_phase_reference"), + source: wgpu::ShaderSource::Wgsl(source.into()), + }); + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("cutout_phase_reference"), + layout: None, + module: &shader, + entry_point: Some("check_phase"), + compilation_options: Default::default(), + cache: None, + }); + let input = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("cutout_phase_cases"), + contents: bytemuck::cast_slice(&cases), + usage: wgpu::BufferUsages::STORAGE, + }); + let size = cases.len() as u64 * 8; + let output = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("cutout_phase_output"), + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let staging = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("cutout_phase_readback"), + size, + usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ, + mapped_at_creation: false, + }); + let bindings = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("cutout_phase_reference"), + layout: &pipeline.get_bind_group_layout(0), + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: input.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: output.as_entire_binding(), + }, + ], + }); + let mut encoder = device.create_command_encoder(&Default::default()); + { + let mut pass = encoder.begin_compute_pass(&Default::default()); + pass.set_pipeline(&pipeline); + pass.set_bind_group(0, &bindings, &[]); + pass.dispatch_workgroups(1, 1, 1); + } + encoder.copy_buffer_to_buffer(&output, 0, &staging, 0, size); + queue.submit([encoder.finish()]); + let slice = staging.slice(..); + let (tx, rx) = std::sync::mpsc::channel(); + slice.map_async(wgpu::MapMode::Read, move |result| tx.send(result).unwrap()); + device + .poll(wgpu::PollType::Wait { + submission_index: None, + timeout: None, + }) + .unwrap(); + rx.recv().unwrap().expect("cutout phase readback"); + let bytes = slice.get_mapped_range(); + for (index, case) in cases.iter().enumerate() { + let expected = (case[6] + 0.5) / 16.0; + for (path, offset) in [("scene", index * 8), ("shadow", index * 8 + 4)] { + let actual = f32::from_le_bytes(bytes[offset..offset + 4].try_into().unwrap()); + assert_eq!(actual, expected, "{path} phase case {index}: {case:?}"); + } + } + drop(bytes); + staging.unmap(); + eprintln!( + "cutout phase: {} approved reference points pass on both paths", + cases.len() + ); +}