facett-core 0.1.18

//! **GFX_V2 item 7 — TAA really CONVERGES, and stays sharp** (needs `--features wgpu`).
//!
//! The claim is temporal: accumulating N *jittered* frames antialiases an edge that a
//! single frame renders as a staircase. So it is measured, not eyeballed, and the metric
//! is chosen so that the two ways of faking it both fail.
//!
//! **THE METRIC — per-row coverage residual.** A slightly tilted hard edge is rendered
//! and, for each row `y`, the luma is summed across a fixed horizontal band that the
//! edge crosses. That sum is proportional to how much of the band the shape covers, and
//! for a *straight* edge the true coverage is **linear in `y`**. A single aliased frame
//! can only step by whole pixels, so its row sums form a staircase — fit a line and the
//! RMS residual is the aliasing, in luma units. Properly antialiased, the residual
//! collapses.
//!
//! **WHY THAT ALONE IS HOLLOW, AND WHAT CATCHES IT.** A plain blur also flattens the
//! staircase. So the same frame is measured for **transition width** — pixels per row
//! strictly between background and foreground. Real antialiasing of a hard edge gives
//! about one; a blur gives three or more; a dead pass gives zero. The pair
//! (residual collapsed AND width ≈ 1) is satisfied only by genuine sub-pixel
//! convergence.
//!
//! **THE NON-VACUITY THE WHOLE FILE RESTS ON.** The single-frame residual is asserted
//! *large* first. Every "TAA reduced it" assertion is a ratio against a control measured
//! in the same run on the same geometry — never against a constant. And the mutation is
//! **zeroing the jitter**: N identical frames average to the same aliased frame, so the
//! residual does not move. That mutation is what proves the metric measures the jitter
//! rather than the mere presence of a blend.
//!
//! The frames come from the OIT lane rendering an opaque polygon — reusing a pass that
//! already exists rather than adding a test-only rasteriser (LAW #5). Alpha is 1, so the
//! OIT resolve emits a hard-edged shape, which is exactly the aliased input TAA is for.
//!
//! Run: `cargo test -p facett-core --features wgpu --test gpu_taa -- --nocapture`
#![cfg(feature = "wgpu")]

use facett_core::render::gpu::oit::{OitBatch, OitPass};
use facett_core::render::gpu::taa::{jitter_px, TaaPass, TAA_PHASES};
use facett_core::render::gpu::{preferred_backends, read_texture_region, request_best_adapter};
use egui::pos2;

/// The component every row and every SKIP in this file is filed under.
const GATE_COMPONENT: &str = "facett-core";
/// What a skipped arm of this file would have proven, so a SKIPPED matrix row is
/// readable without opening the test.
const GATE_DETAIL: &str = "TAA jitter accumulation on a real device — convergence, sharpness and the no-jitter control";

/// `Rgba8Unorm` throughout, so the readback is the same bytes the resolve wrote and the
/// measurement is not mediated by a float conversion.
const FMT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba8Unorm;
const W: u32 = 128;
const H: u32 = 128;

/// The edge sweeps `EDGE_RUN` px horizontally over the full height — a slope of one
/// pixel every 16 rows, so an aliased render has 16-row-long flat treads and a 255-luma
/// riser between them. Shallow on purpose: a steep edge aliases less and would make the
/// control weak.
const EDGE_X: f32 = 52.0;
const EDGE_RUN: f32 = 8.0;
/// The band the row sums are taken over. Wide enough to contain the edge at every row
/// with slack either side.
const BAND_X0: u32 = 44;
const BAND_X1: u32 = 68;

fn device() -> Option<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> {
    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
        backends: preferred_backends(),
        flags: wgpu::InstanceFlags::from_build_config().with_env(),
        backend_options: wgpu::BackendOptions::from_env_or_default(),
        memory_budget_thresholds: wgpu::MemoryBudgetThresholds::default(),
        display: None,
    });
    let adapter = request_best_adapter(&instance, preferred_backends())?;
    let info = adapter.get_info();
    // ATTRIBUTION, not decoration: a graceful skip is how a GPU test file goes green
    // having run no shader.
    eprintln!(
        "[gpu_taa] device: {} ({:?}, {:?}) fragment_writable_storage={}",
        info.name,
        info.backend,
        info.device_type,
        OitPass::supported(&adapter)
    );
    let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
        label: Some("gpu_taa_test"),
        required_features: wgpu::Features::empty(),
        required_limits: wgpu::Limits::downlevel_defaults(),
        ..Default::default()
    }))
    .ok()?;
    Some((adapter, device, queue))
}

fn color_target(device: &wgpu::Device) -> (wgpu::Texture, wgpu::TextureView) {
    let tex = device.create_texture(&wgpu::TextureDescriptor {
        label: Some("gpu_taa_scene"),
        size: wgpu::Extent3d { width: W, height: H, depth_or_array_layers: 1 },
        mip_level_count: 1,
        sample_count: facett_core::render::gpu::NO_MSAA_SAMPLES,
        dimension: wgpu::TextureDimension::D2,
        format: FMT,
        usage: wgpu::TextureUsages::RENDER_ATTACHMENT
            | wgpu::TextureUsages::TEXTURE_BINDING
            | wgpu::TextureUsages::COPY_SRC,
        view_formats: &[],
    });
    let view = tex.create_view(&Default::default());
    (tex, view)
}

/// A white opaque wedge whose right edge is tilted, offset by `(jx, jy)` sub-pixels.
fn wedge(jx: f32, jy: f32) -> OitBatch {
    let mut b = OitBatch::new();
    let h = H as f32;
    let p = |x: f32, y: f32| pos2(x + jx, y + jy);
    let white = [1.0, 1.0, 1.0, 1.0];
    // (0,0) → (EDGE_X,0) → (EDGE_X+EDGE_RUN,H) → (0,H)
    b.push_tri(p(-4.0, -4.0), p(EDGE_X, -4.0), p(EDGE_X + EDGE_RUN, h + 4.0), white, 1.0);
    b.push_tri(p(-4.0, -4.0), p(EDGE_X + EDGE_RUN, h + 4.0), p(-4.0, h + 4.0), white, 1.0);
    b
}

fn luma(p: &[u8]) -> f64 {
    0.2126 * f64::from(p[0]) + 0.7152 * f64::from(p[1]) + 0.0722 * f64::from(p[2])
}

/// Sum of luma across the band at each row — proportional to the covered fraction.
fn row_sums(bytes: &[u8]) -> Vec<f64> {
    (0..H)
        .map(|y| {
            (BAND_X0..BAND_X1)
                .map(|x| {
                    let o = ((y * W + x) * 4) as usize;
                    luma(&bytes[o..o + 4])
                })
                .sum()
        })
        .collect()
}

/// **The aliasing, in luma units.** RMS residual of the row sums about their
/// least-squares line. A straight edge's true coverage is linear in `y`; only
/// quantisation makes it deviate.
fn edge_residual_rms(bytes: &[u8]) -> f64 {
    let s = row_sums(bytes);
    let n = s.len() as f64;
    let mean_x = (s.len() as f64 - 1.0) / 2.0;
    let mean_y = s.iter().sum::<f64>() / n;
    let (mut sxy, mut sxx) = (0.0, 0.0);
    for (i, v) in s.iter().enumerate() {
        let dx = i as f64 - mean_x;
        sxy += dx * (v - mean_y);
        sxx += dx * dx;
    }
    let slope = if sxx == 0.0 { 0.0 } else { sxy / sxx };
    let intercept = mean_y - slope * mean_x;
    let sq: f64 = s
        .iter()
        .enumerate()
        .map(|(i, v)| {
            let r = v - (slope * i as f64 + intercept);
            r * r
        })
        .sum();
    (sq / n).sqrt()
}

/// Mean pixels per row strictly between background and foreground — the edge's softness.
/// 0 = a hard staircase, ~1 = antialiased, >2 = smeared.
fn transition_width(bytes: &[u8]) -> f64 {
    let mut total = 0usize;
    for y in 0..H {
        for x in BAND_X0..BAND_X1 {
            let o = ((y * W + x) * 4) as usize;
            let l = luma(&bytes[o..o + 4]);
            if l > 16.0 && l < 239.0 {
                total += 1;
            }
        }
    }
    total as f64 / f64::from(H)
}

/// Render one jittered frame of the wedge through the OIT lane into `scene`.
fn render_frame(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
    oit: &mut OitPass,
    scene: &wgpu::TextureView,
    jx: f32,
    jy: f32,
) {
    let batch = wedge(jx, jy);
    oit.set_frame(queue, [0.0, 0.0, 0.0, 1.0]);
    oit.upload(device, queue, &batch);
    let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
        label: Some("gpu_taa_scene_enc"),
    });
    oit.record(&mut enc, scene);
    queue.submit(Some(enc.finish()));
}

/// IEEE-754 binary16 bits of a normal `f32` — enough to fill an `Rg16Float` velocity
/// texture with a couple of exact constants.
fn f16_bits(v: f32) -> u16 {
    if v == 0.0 {
        return if v.is_sign_negative() { 0x8000 } else { 0 };
    }
    let b = v.to_bits();
    let sign = ((b >> 31) & 1) as u16;
    let exp = ((b >> 23) & 0xFF) as i32 - 127 + 15;
    assert!((1..=30).contains(&exp), "f16_bits only handles normals ({v})");
    let mant = ((b & 0x007F_FFFF) >> 13) as u16;
    (sign << 15) | ((exp as u16) << 10) | (mant & 0x03FF)
}

/// A full-screen constant motion-vector texture, in **UV per frame**.
fn velocity_texture(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
    vx: f32,
    vy: f32,
) -> (wgpu::Texture, wgpu::TextureView) {
    let tex = device.create_texture(&wgpu::TextureDescriptor {
        label: Some("gpu_taa_velocity"),
        size: wgpu::Extent3d { width: W, height: H, depth_or_array_layers: 1 },
        mip_level_count: 1,
        sample_count: facett_core::render::gpu::NO_MSAA_SAMPLES,
        dimension: wgpu::TextureDimension::D2,
        format: wgpu::TextureFormat::Rg16Float,
        usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST,
        view_formats: &[],
    });
    let (bx, by) = (f16_bits(vx), f16_bits(vy));
    let mut data = Vec::with_capacity((W * H * 4) as usize);
    for _ in 0..(W * H) {
        data.extend_from_slice(&bx.to_le_bytes());
        data.extend_from_slice(&by.to_le_bytes());
    }
    queue.write_texture(
        wgpu::TexelCopyTextureInfo {
            texture: &tex,
            mip_level: 0,
            origin: wgpu::Origin3d::ZERO,
            aspect: wgpu::TextureAspect::All,
        },
        &data,
        wgpu::TexelCopyBufferLayout {
            offset: 0,
            bytes_per_row: Some(W * 4),
            rows_per_image: Some(H),
        },
        wgpu::Extent3d { width: W, height: H, depth_or_array_layers: 1 },
    );
    let view = tex.create_view(&Default::default());
    (tex, view)
}

/// Accumulate `frames` frames through TAA and read back the resolved image.
///
/// `jitter` selects the sample pattern: `true` = the real Halton phases, `false` = the
/// MUTATION-shaped control that renders every frame at the same position.
fn accumulate(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
    frames: u64,
    jitter: bool,
) -> Vec<u8> {
    accumulate_with_velocity(device, queue, frames, jitter, None)
}

fn accumulate_with_velocity(
    device: &wgpu::Device,
    queue: &wgpu::Queue,
    frames: u64,
    jitter: bool,
    velocity: Option<&wgpu::TextureView>,
) -> Vec<u8> {
    let (_scene_tex, scene) = color_target(device);
    let mut oit = OitPass::new(device, FMT);
    oit.ensure(device, W, H, 4);
    let mut taa = TaaPass::new(device, FMT);
    taa.ensure(device, W, H);

    for _ in 0..frames {
        let [jx, jy] = if jitter { taa.jitter_px() } else { [0.0, 0.0] };
        render_frame(device, queue, &mut oit, &scene, jx, jy);
        let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
            label: Some("gpu_taa_resolve_enc"),
        });
        let ok = taa.record(device, queue, &mut enc, &scene, velocity);
        assert!(ok, "the TAA pass must actually resolve, not silently decline");
        queue.submit(Some(enc.finish()));
    }
    assert_eq!(taa.frame(), frames, "the pass counted every frame it resolved");
    let tex = taa.resolved().expect("a resolved target after accumulating").clone();
    read_texture_region(device, queue, &tex, 4, 0, 0, W, H)
}

/// A single unjittered frame straight out of the rasteriser — the aliased control.
fn single_frame(device: &wgpu::Device, queue: &wgpu::Queue) -> Vec<u8> {
    let (tex, view) = color_target(device);
    let mut oit = OitPass::new(device, FMT);
    oit.ensure(device, W, H, 4);
    render_frame(device, queue, &mut oit, &view, 0.0, 0.0);
    read_texture_region(device, queue, &tex, 4, 0, 0, W, H)
}

const FRAMES: u64 = 64;

#[test]
fn jittered_accumulation_converges_the_edge_and_keeps_it_sharp() {
    let Some((adapter, device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "jittered_accumulation_converges_the_edge_and_keeps_it_sharp", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    if !OitPass::supported(&adapter) {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "jittered_accumulation_converges_the_edge_and_keeps_it_sharp", "adapter lacks FRAGMENT_WRITABLE_STORAGE", GATE_DETAIL);
        return;
    }
    assert_eq!(FRAMES % u64::from(TAA_PHASES), 0, "accumulate whole jitter cycles");

    let one = single_frame(&device, &queue);
    let taa = accumulate(&device, &queue, FRAMES, true);

    let (r1, rt) = (edge_residual_rms(&one), edge_residual_rms(&taa));
    let (w1, wt) = (transition_width(&one), transition_width(&taa));
    eprintln!(
        "[gpu_taa] single frame: residual={r1:.2} width={w1:.2}px | {FRAMES} jittered frames: \
         residual={rt:.2} width={wt:.2}px | reduction={:.2}x",
        r1 / rt.max(1e-9)
    );

    // ── NON-VACUITY: the control is genuinely aliased. ──
    // 16-row treads with a 255-luma riser give a staircase residual near 255/sqrt(12).
    assert!(
        r1 > 40.0,
        "the single-frame control must actually be aliased (residual {r1:.2}); if it is not, \
         every 'TAA reduced it' assertion below is measuring nothing"
    );
    assert!(w1 < 0.2, "and hard-edged — a rasterised opaque edge has no partial pixels ({w1:.2}px)");
    // The frame is a real picture, not blank: the band must be part covered, part not.
    let sums = row_sums(&taa);
    let (lo, hi) = sums
        .iter()
        .fold((f64::MAX, f64::MIN), |(l, h), &v| (l.min(v), h.max(v)));
    assert!(
        hi - lo > 255.0 * 4.0,
        "the wedge must sweep across the band ({lo:.0}..{hi:.0}) — a uniform frame proves nothing"
    );

    // ── THE HEADLINE: jitter across N frames reduces the edge residual. ──
    assert!(
        rt < r1 * 0.25,
        "{FRAMES} jittered frames must cut the edge residual to under a quarter of one frame's \
         ({rt:.2} vs {r1:.2})"
    );

    // ── THE SHARPNESS CONTROL: a blur would also flatten the staircase. ──
    assert!(
        wt > 0.5,
        "the converged edge must actually be soft ({wt:.2}px) — zero partial pixels means the \
         resolve produced no sub-pixel detail at all"
    );
    assert!(
        wt < 2.0,
        "…but no wider than ~1 px ({wt:.2}px). A box blur reduces the residual too; this is the \
         assertion it fails"
    );
}

#[test]
fn without_jitter_the_same_accumulation_does_not_converge() {
    let Some((adapter, device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "without_jitter_the_same_accumulation_does_not_converge", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    if !OitPass::supported(&adapter) {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "without_jitter_the_same_accumulation_does_not_converge", "adapter lacks FRAGMENT_WRITABLE_STORAGE", GATE_DETAIL);
        return;
    }
    // **The mutation, kept as a permanent test.** Same TAA pass, same frame count, same
    // geometry — only the sub-pixel offsets are zeroed. Averaging N identical frames is
    // still that frame, so if this ever converged, the test above would be measuring the
    // blend rather than the jitter.
    let one = single_frame(&device, &queue);
    let flat = accumulate(&device, &queue, FRAMES, false);
    let jittered = accumulate(&device, &queue, FRAMES, true);

    let (r1, rf, rj) = (
        edge_residual_rms(&one),
        edge_residual_rms(&flat),
        edge_residual_rms(&jittered),
    );
    eprintln!(
        "[gpu_taa] residual — single {r1:.2} | {FRAMES} UNJITTERED {rf:.2} | {FRAMES} jittered {rj:.2}"
    );
    assert!(
        rf > r1 * 0.9,
        "accumulating {FRAMES} IDENTICAL frames must leave the aliasing alone ({rf:.2} vs {r1:.2}) \
         — a temporal blend that antialiases without jitter is blurring, not resolving"
    );
    assert!(
        rj < rf * 0.25,
        "and the only difference between these two runs is the jitter, so the jitter is what \
         converges the edge ({rj:.2} jittered vs {rf:.2} not)"
    );
    assert!(
        transition_width(&flat) < 0.2,
        "the unjittered accumulation stays hard-edged, i.e. it really did not antialias"
    );
}

#[test]
fn the_first_frame_is_taken_outright_and_a_reset_restarts_the_sequence() {
    let Some((adapter, device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "the_first_frame_is_taken_outright_and_a_reset_restarts_the_sequence", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    if !OitPass::supported(&adapter) {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "the_first_frame_is_taken_outright_and_a_reset_restarts_the_sequence", "adapter lacks FRAGMENT_WRITABLE_STORAGE", GATE_DETAIL);
        return;
    }
    // Frame 0 has no history, so the resolve must pass the current frame through
    // UNCHANGED. If it blended against an uninitialised history, the first frame after
    // every teleport would be a dark or garbage flash.
    let (scene_tex, scene) = color_target(&device);
    let mut oit = OitPass::new(&device, FMT);
    oit.ensure(&device, W, H, 4);
    let mut taa = TaaPass::new(&device, FMT);
    taa.ensure(&device, W, H);
    assert_eq!(taa.weight(), 0.0, "no history ⇒ no history weight");

    render_frame(&device, &queue, &mut oit, &scene, 0.0, 0.0);
    let raw = read_texture_region(&device, &queue, &scene_tex, 4, 0, 0, W, H);

    let mut enc = device.create_command_encoder(&Default::default());
    assert!(taa.record(&device, &queue, &mut enc, &scene, None));
    queue.submit(Some(enc.finish()));
    let out = read_texture_region(
        &device,
        &queue,
        taa.resolved().expect("resolved"),
        4,
        0,
        0,
        W,
        H,
    );
    assert_eq!(out, raw, "frame 0 must pass through byte-identically");

    // After accumulating, a reset must return the pass to that state — same phase, no
    // history — otherwise a preset jump reprojects across a discontinuity.
    for _ in 0..5 {
        let [jx, jy] = taa.jitter_px();
        render_frame(&device, &queue, &mut oit, &scene, jx, jy);
        let mut enc = device.create_command_encoder(&Default::default());
        taa.record(&device, &queue, &mut enc, &scene, None);
        queue.submit(Some(enc.finish()));
    }
    assert_eq!(taa.frame(), 6);
    assert!(taa.weight() > 0.5, "the ramp is well under way");
    assert_ne!(taa.jitter_px(), jitter_px(0), "…at a later phase");

    taa.reset();
    assert_eq!(taa.frame(), 0);
    assert_eq!(taa.weight(), 0.0, "a reset drops the history weight");
    assert_eq!(taa.jitter_px(), jitter_px(0), "and restarts the jitter sequence");

    render_frame(&device, &queue, &mut oit, &scene, 0.0, 0.0);
    let mut enc = device.create_command_encoder(&Default::default());
    taa.record(&device, &queue, &mut enc, &scene, None);
    queue.submit(Some(enc.finish()));
    let after = read_texture_region(
        &device,
        &queue,
        taa.resolved().expect("resolved"),
        4,
        0,
        0,
        W,
        H,
    );
    assert_eq!(
        after, raw,
        "and the first frame after a reset is again taken outright — no stale history bleeds in"
    );
}

#[test]
fn a_history_whose_reprojection_leaves_the_frame_cannot_converge_the_edge() {
    let Some((adapter, device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "a_history_whose_reprojection_leaves_the_frame_cannot_converge_the_edge", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    if !OitPass::supported(&adapter) {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "a_history_whose_reprojection_leaves_the_frame_cannot_converge_the_edge", "adapter lacks FRAGMENT_WRITABLE_STORAGE", GATE_DETAIL);
        return;
    }
    // **The velocity path, exercised.** With zero velocity `huv` never leaves `[0,1]` and
    // the reprojection is the identity, so the tests above say nothing about motion at
    // all. A velocity of 2 UV per frame puts EVERY pixel's reprojection outside the
    // frame: no usable history anywhere, so the resolve must fall back to the current
    // frame and therefore must NOT converge, however many frames it accumulates.
    //
    // WHAT THIS DOES AND DOES NOT PROVE, measured. It proves the rejection happens. It
    // does NOT isolate the explicit bounds test in the shader: deleting that test moves
    // the residual below from 71.70 to 71.28 and nothing here notices, because the 3x3
    // box clamp already bounds an off-screen sample into the current frame's own range.
    // So this asserts the REJECTION, and the bounds test is documented as belt on braces
    // in `taa.wgsl` rather than claimed to be under test.
    let (_vt, vel) = velocity_texture(&device, &queue, 2.0, 2.0);

    let still = accumulate(&device, &queue, FRAMES, true);
    let flung = accumulate_with_velocity(&device, &queue, FRAMES, true, Some(&vel));

    let (rs, rf) = (edge_residual_rms(&still), edge_residual_rms(&flung));
    let one = edge_residual_rms(&single_frame(&device, &queue));
    eprintln!(
        "[gpu_taa] residual — single {one:.2} | {FRAMES} frames, no motion {rs:.2} | \
         {FRAMES} frames, all reprojections OFF SCREEN {rf:.2}"
    );
    assert!(
        rs < one * 0.25,
        "the zero-velocity control must still converge ({rs:.2} vs {one:.2}) — without it, a \
         resolve that was simply broken would satisfy the assertion below"
    );
    assert!(
        rf > one * 0.9,
        "every reprojection lands off screen, so no history is usable and the aliasing must \
         survive ({rf:.2} vs a single frame's {one:.2})"
    );
    assert!(
        transition_width(&flung) < 0.2,
        "…and the edge must stay hard: a rejected history contributes nothing, it does not \
         contribute a clamped border texel"
    );
}