Make cutout coverage phase deterministic across backends

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<f32>(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).

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

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<f32>(dimensions) * mip_scale),
-        vec2<f32>(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<f32>(max(
+        (dimensions - vec2<u32>(1u)) >> vec2<u32>(phase_lod),
+        vec2<u32>(1u),
+    ));
     let texel = vec2<u32>(floor(wrapped_uv * mip_dimensions));
     let x = texel.x & 3u;
     let y = texel.y & 3u;

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"));
 }

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<f32>(dimensions) * mip_scale),
-        vec2<f32>(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<f32>(max(
+        (dimensions - vec2<u32>(1u)) >> vec2<u32>(phase_lod),
+        vec2<u32>(1u),
+    ));
     let texel = vec2<u32>(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"));
     }

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"]

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<f32>, extent: vec4<f32> };
+@group(0) @binding(0) var<storage, read> cases: array<Case>;
+@group(0) @binding(1) var<storage, read_write> output: array<vec2<f32>>;
+@compute @workgroup_size(64)
+fn check_phase(@builtin(global_invocation_id) id: vec3<u32>) {
+    if (id.x >= arrayLength(&cases)) { return; }
+    let c = cases[id.x];
+    output[id.x] = vec2<f32>(
+        scene_phase(c.coordinate.xy, vec2<u32>(c.extent.xy), c.coordinate.z),
+        shadow_phase(c.coordinate.xy, vec2<u32>(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()
+    );
+}