GPU.1: Preserve CPU parity in numerical effect kernels

Cargo.lock

@@ -1331,6 +1331,7 @@ dependencies = [
  "js-sys",
  "log",
  "pollster 0.4.0",
+ "rayon",
  "wasm-bindgen",
  "web-time",
  "wgpu",

crates/effects/src/lib.rs

@@ -99,12 +99,12 @@ pub use generate3::{
 };
 pub use noise::GrainLook;
 pub use textfx::{BoltPlan, TextLook, TextPass, bolt_plan, text_passes};
-pub use transition::place_layer;
+pub use transition::{GridWipe, place_layer};
 // effectcraft-gpu fx_noise.
 pub use keying2::{InnerOuterPlan, inner_outer_plan};
 pub use noise::remove_grain_level;
 pub use noise3::{FractalGpu, fractal_gpu};
-pub use stylize3::{BumpLight, bump_layer_id};
+pub use stylize3::{BumpLight, HexTile, bump_layer_id};
 pub use textfx::{TextLayer, text_layer};
 pub use time_fx::{PixelMotion, TimewarpPlan, TwStep, pmb_plan, posterized_time, time_frames, timewarp_crop, timewarp_plan};
 // effectcraft-gpu fx_tone.

crates/effects/src/stylize3.rs

@@ -364,32 +364,54 @@ fn hex_cell(x: f64, y: f64, r: f64) -> ((f64, f64), (i64, i64)) {
     ((px, py), (rx as i64, rz as i64))
 }
 
+/// Stable f64 hex-cell selection and source geometry for CPU and GPU renders.
+pub struct HexTile {
+    r: f64,
+    c: (f64, f64),
+    sn: f64,
+    cs: f64,
+    smear: f64,
+    mode: u32,
+}
+impl HexTile {
+    /// No image pixels are read; only buffer geometry and parameters are used.
+    pub fn new(ctx: &EffectCtx, b: &Buf) -> Self {
+        let r = (ctx.params.f("radius") * b.scale).max(0.5);
+        let c = b.to_px(ctx.params.v2("center"));
+        let rot = ctx.params.f("rotate").to_radians();
+        let (sn, cs) = rot.sin_cos();
+        let smear = (ctx.params.f("smearing") / 100.0).clamp(0.0, 1.0);
+        let mode = ctx.params.e("render");
+        Self { r, c, sn, cs, smear, mode }
+    }
+    /// Bilinear source position at a buffer pixel centre.
+    pub fn source_at(&self, x: u32, y: u32) -> (f64, f64) {
+        let Self { r, c: (cx, cy), sn, cs, smear, mode } = *self;
+        let (dx, dy) = (x as f64 + 0.5 - cx, y as f64 + 0.5 - cy);
+        // Into the rotated tiling frame.
+        let (u, v) = (dx * cs + dy * sn, -dx * sn + dy * cs);
+        let ((hx, hy), (qi, ri)) = hex_cell(u, v, r);
+        let (mut lu, lv) = (u - hx, v - hy);
+        let odd = (qi + ri).rem_euclid(2) == 1;
+        match mode {
+            1 if odd => lu = -lu,
+            2 if odd => lu = -lu,
+            _ => {}
+        }
+        let lv = if mode == 2 && odd { -lv } else { lv };
+        // Every tile shows the hexagon around the centre point.
+        let (ox, oy) = (lu * cs - lv * sn, lu * sn + lv * cs);
+        let (tx, ty) = (cx + ox, cy + oy);
+        let (sx, sy) = (tx + (x as f64 + 0.5 - tx) * smear, ty + (y as f64 + 0.5 - ty) * smear);
+        (sx, sy)
+    }
+}
 fn hextile(ctx: &EffectCtx, mut b: Buf) -> Buf {
-    let r = (ctx.params.f("radius") * b.scale).max(0.5);
-    let (cx, cy) = b.to_px(ctx.params.v2("center"));
-    let rot = ctx.params.f("rotate").to_radians();
-    let (sn, cs) = rot.sin_cos();
-    let smear = (ctx.params.f("smearing") / 100.0).clamp(0.0, 1.0);
-    let mode = ctx.params.e("render");
+    let geometry = HexTile::new(ctx, &b);
     let src = b.img.clone();
     b.img.rows_mut().for_each(|(y, row)| {
         for (x, px) in row.iter_mut().enumerate() {
-            let (dx, dy) = (x as f64 + 0.5 - cx, y as f64 + 0.5 - cy);
-            // Into the rotated tiling frame.
-            let (u, v) = (dx * cs + dy * sn, -dx * sn + dy * cs);
-            let ((hx, hy), (qi, ri)) = hex_cell(u, v, r);
-            let (mut lu, lv) = (u - hx, v - hy);
-            let odd = (qi + ri).rem_euclid(2) == 1;
-            match mode {
-                1 if odd => lu = -lu,
-                2 if odd => lu = -lu,
-                _ => {}
-            }
-            let lv = if mode == 2 && odd { -lv } else { lv };
-            // Every tile shows the hexagon around the centre point.
-            let (ox, oy) = (lu * cs - lv * sn, lu * sn + lv * cs);
-            let (tx, ty) = (cx + ox, cy + oy);
-            let (sx, sy) = (tx + (x as f64 + 0.5 - tx) * smear, ty + (y as f64 + 0.5 - ty) * smear);
+            let (sx, sy) = geometry.source_at(x as u32, y as u32);
             *px = src.sample_bilinear_clamped(sx, sy);
         }
     });

crates/effects/src/transition.rs

@@ -314,53 +314,86 @@ pub(crate) fn card_wipe_plan(ctx: &EffectCtx, b: &Buf) -> Option<PiecePlan> {
     Some(PiecePlan { pieces, front: PieceTex::Layer, back })
 }
 
+/// Double-precision grid geometry, shared with the GPU mask upload. A tiny
+/// rounding error at a tile boundary can otherwise flip a fully opaque pixel.
+pub struct GridWipe {
+    c: (f64, f64),
+    s: f64,
+    co: f64,
+    cs: f64,
+    maxd: f64,
+    shape: u32,
+    reverse: bool,
+    done: f64,
+    border: f64,
+}
+
+impl GridWipe {
+    /// Construct from buffer geometry; no pixel data is accessed.
+    pub fn new(ctx: &EffectCtx, b: &Buf) -> Self {
+        let done = completion(ctx);
+        let c = b.to_px(ctx.params.v2("center"));
+        let rot = ctx.params.f("rotation").to_radians();
+        let border = ctx.params.f("border").max(0.0) * b.scale;
+        let tiles = ctx.params.f("tiles").round().clamp(1.0, 500.0);
+        let shape = ctx.params.e("shape");
+        let reverse = ctx.params.b("reverse");
+        let (_, _, w, _) = layer_rect(ctx, b);
+        let cs = (w / tiles).max(1.0);
+        let (s, co) = rot.sin_cos();
+        let metric = |dx: f64, dy: f64| match shape {
+            0 => dx.abs(),
+            2 => dx.abs().max(dy.abs()),
+            _ => dx.hypot(dy),
+        };
+        let maxd = max_corner_dist(b, c, |dx, dy| {
+            let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
+            metric(qx, qy)
+        });
+        GridWipe { c, s, co, cs, maxd, shape, reverse, done, border }
+    }
+
+    /// Alpha multiplier at a buffer pixel centre.
+    pub fn coverage(&self, x: u32, y: u32) -> f32 {
+        let Self { c, s, co, cs, maxd, shape, reverse, done, border } = *self;
+        let metric = |dx: f64, dy: f64| match shape {
+            0 => dx.abs(),
+            2 => dx.abs().max(dy.abs()),
+            _ => dx.hypot(dy),
+        };
+        let spread = 0.25;
+        let (dx, dy) = (x as f64 + 0.5 - c.0, y as f64 + 0.5 - c.1);
+        let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
+        let (cx, cy) = (((qx / cs).floor() + 0.5) * cs, ((qy / cs).floor() + 0.5) * cs);
+        let mut o = (metric(cx, cy) / maxd).clamp(0.0, 1.0);
+        if reverse {
+            o = 1.0 - o;
+        }
+        let p = ((done * (1.0 + spread) - o) / spread).clamp(0.0, 1.0);
+        if p <= 0.0 {
+            return 1.0;
+        }
+        let half = cs * 0.5 * (1.0 - p);
+        let m = (qx - cx).abs().max((qy - cy).abs());
+        let k = if border > 0.0 {
+            ((half - m) / border + 0.5).clamp(0.0, 1.0)
+        } else if m < half {
+            1.0
+        } else {
+            0.0
+        };
+        k as f32
+    }
+}
+
 fn grid_wipe(ctx: &EffectCtx, mut b: Buf) -> Buf {
-    let done = completion(ctx);
-    if done <= 0.0 {
+    if completion(ctx) <= 0.0 {
         return b;
     }
-    let c = b.to_px(ctx.params.v2("center"));
-    let rot = ctx.params.f("rotation").to_radians();
-    let border = ctx.params.f("border").max(0.0) * b.scale;
-    let tiles = ctx.params.f("tiles").round().clamp(1.0, 500.0);
-    let shape = ctx.params.e("shape");
-    let reverse = ctx.params.b("reverse");
-    let (_, _, w, _) = layer_rect(ctx, &b);
-    let cs = (w / tiles).max(1.0);
-    let (s, co) = rot.sin_cos();
-    let metric = |dx: f64, dy: f64| match shape {
-        0 => dx.abs(),
-        2 => dx.abs().max(dy.abs()),
-        _ => dx.hypot(dy),
-    };
-    let maxd = max_corner_dist(&b, c, |dx, dy| {
-        let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
-        metric(qx, qy)
-    });
-    let spread = 0.25;
+    let grid = GridWipe::new(ctx, &b);
     b.img.rows_mut().for_each(|(y, row)| {
         for (x, px) in row.iter_mut().enumerate() {
-            let (dx, dy) = (x as f64 + 0.5 - c.0, y as f64 + 0.5 - c.1);
-            let (qx, qy) = (dx * co + dy * s, -dx * s + dy * co);
-            let (cx, cy) = (((qx / cs).floor() + 0.5) * cs, ((qy / cs).floor() + 0.5) * cs);
-            let mut o = (metric(cx, cy) / maxd).clamp(0.0, 1.0);
-            if reverse {
-                o = 1.0 - o;
-            }
-            let p = ((done * (1.0 + spread) - o) / spread).clamp(0.0, 1.0);
-            if p <= 0.0 {
-                continue;
-            }
-            let half = cs * 0.5 * (1.0 - p);
-            let m = (qx - cx).abs().max((qy - cy).abs());
-            let k = if border > 0.0 {
-                ((half - m) / border + 0.5).clamp(0.0, 1.0)
-            } else if m < half {
-                1.0
-            } else {
-                0.0
-            };
-            scale_px(px, k as f32);
+            scale_px(px, grid.coverage(x as u32, y as u32));
         }
     });
     b

crates/gpu/Cargo.toml

@@ -21,6 +21,7 @@ effectcraft-keyframe = { workspace = true }
 wgpu = { workspace = true }
 bytemuck = { workspace = true }
 log = { workspace = true }
+rayon = { workspace = true }
 web-time = { workspace = true }
 
 [target.'cfg(not(target_arch = "wasm32"))'.dependencies]

crates/gpu/src/context.rs

@@ -468,6 +468,8 @@ impl GpuContext {
             include_str!("shaders/sky.wgsl"),
         ]
         .concat();
+        #[cfg(test)]
+        let src = src + crate::tests::MATH_PROBE_SHADER;
         let module = init_resource(&device, "kernels shader module", || {
             device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("effectcraft kernels"), source: wgpu::ShaderSource::Wgsl(src.into()) })
         })?;
@@ -545,6 +547,8 @@ impl GpuContext {
             .map(|(e, layout)| Ok((*e, build(e, layout)?)))
             .collect::<Result<HashMap<_, _>, String>>()?;
         let mut unavailable = vec![];
+        #[cfg(test)]
+        pipelines.insert("math_probe", build("math_probe", &layout)?);
         for e in EFFECT_ENTRIES
             .iter()
             .chain(crate::fx_color::KERNELS)

crates/gpu/src/fx_pixel2.rs

@@ -12,9 +12,11 @@
 //! Layer is chosen; the colouring runs here.
 
 use effectcraft_color::BlendMode;
-use effectcraft_effects::EffectCtx;
+use effectcraft_effects::{Buf, EffectCtx};
 use effectcraft_raster::Image;
 
+use rayon::prelude::*;
+
 use crate::context::{Enc, GpuImage, Params};
 use crate::effects::{GBuf, box_passes, gaussian_blur};
 use crate::ops::mode_id;
@@ -380,14 +382,24 @@ fn clamp_plane(e: &mut Enc, pl: &GpuImage, lo: f32, hi: f32) -> GpuImage {
 // ---------------------------------------------------------------- CC HexTile / CC Mr. Smoothie (stylize3)
 
 fn hextile(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
-    let r = (ctx.params.f("radius") * b.scale).max(0.5);
-    let c = b.to_px(ctx.params.v2("center"));
-    let (sn, cs) = ctx.params.f("rotate").to_radians().sin_cos();
-    let mut p = Params::default();
-    p.u[0][0] = ctx.params.e("render");
-    p.f[0] = [c.0 as f32, c.1 as f32, r as f32, (ctx.params.f("smearing") / 100.0).clamp(0.0, 1.0) as f32];
-    p.f[1] = [sn as f32, cs as f32, 0.0, 0.0];
-    run(e, "fp2_hextile", &p, b, None, None)
+    let count = (b.img.width as usize).checked_mul(b.img.height as usize)?;
+    if count.checked_mul(8)? > e.g.device.limits().max_storage_buffer_binding_size as usize {
+        return None;
+    }
+    let geometry = effectcraft_effects::HexTile::new(
+        ctx,
+        &Buf { img: Image { width: b.img.width, height: b.img.height, data: vec![] }, offset: b.offset, scale: b.scale },
+    );
+    let stride = (b.img.width as usize).checked_mul(2)?.max(2);
+    let mut coords = vec![0.0; count.checked_mul(2)?];
+    coords.par_chunks_mut(stride).enumerate().for_each(|(y, row)| {
+        for (x, xy) in row.as_chunks_mut::<2>().0.iter_mut().enumerate() {
+            let (sx, sy) = geometry.source_at(x as u32, y as u32);
+            *xy = [sx as f32, sy as f32];
+        }
+    });
+    let data = e.data(&coords);
+    run(e, "fp2_hextile", &Params::default(), b, None, Some(&data))
 }
 
 fn smoothie(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {

crates/gpu/src/fx_tone.rs

@@ -174,6 +174,62 @@ fn parse_map(s: &str) -> Vec<f32> {
     vals.iter().map(|v| (v * k).clamp(0.0, 1.0)).collect()
 }
 
+/// Piecewise-linear lookup in the CPU effect's f64 phase domain. Store the
+/// first representable f32 at each boundary, its value and the interval slope.
+/// Computing floor/wrap in shader f32 can select the end of the map instead of
+/// its beginning at a wrap (a full-scale colour error).
+fn arbitrary_map_table(map: &[f32], phase: f64) -> Option<Vec<f32>> {
+    let n = map.len();
+    if !(2..=4096).contains(&n) || !phase.is_finite() {
+        return None;
+    }
+    let scale = (n - 1) as f64;
+    let offset = phase * scale / 255.0;
+    if !offset.is_finite() || offset.abs() >= (1u64 << 52) as f64 {
+        return None;
+    }
+    let at = |v: f32| {
+        let x = (v as f64 * scale + offset).rem_euclid(n as f64);
+        let i = x.floor() as usize;
+        let a = *map.get(i).unwrap_or(&0.0);
+        let b = *map.get((i + 1).min(n - 1)).unwrap_or(&a);
+        (a + (b - a) * (x - x.floor()) as f32, (b - a) * scale as f32)
+    };
+    let mut points = vec![0.0f32];
+    // Search ordered positive f32 encodings using the CPU's actual f64 sum.
+    // Merely rounding a mathematical boundary upward misses the half-ulp
+    // region where f64 addition itself rounds across a wrap.
+    let first = offset.floor() + 1.0;
+    for i in 0..=n {
+        let boundary = first + i as f64;
+        if boundary > scale + offset {
+            break;
+        }
+        let (mut lo, mut hi) = (0u32, 1.0f32.to_bits());
+        while lo < hi {
+            let mid = lo + (hi - lo) / 2;
+            if f32::from_bits(mid) as f64 * scale + offset < boundary {
+                lo = mid + 1;
+            } else {
+                hi = mid;
+            }
+        }
+        let start = f32::from_bits(lo);
+        if points.last().is_none_or(|last| start > *last) {
+            points.push(start);
+        }
+    }
+    if points.last() != Some(&1.0) {
+        points.push(1.0);
+    }
+    let mut table = Vec::with_capacity(points.len() * 3);
+    for v in points {
+        let (value, slope) = at(v);
+        table.extend([v, value, slope]);
+    }
+    Some(table)
+}
+
 fn point(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
     let pr = ctx.params;
     let f = |k: &str| pr.f(k) as f32;
@@ -308,9 +364,9 @@ fn point(e: &mut Enc, id: &str, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
         }
         "ec.color.psarbitrarymap" => {
             let map = parse_map(pr.s("map"));
-            p.u[0] = [12, pr.b("applyPhaseMapToAlpha") as u32, map.len() as u32, 0];
-            p.f[0][0] = (pr.f("phase") * (map.len() - 1) as f64 / 255.0) as f32;
-            data = Some(map);
+            let table = arbitrary_map_table(&map, pr.f("phase"))?;
+            p.u[0] = [12, pr.b("applyPhaseMapToAlpha") as u32, (table.len() / 3) as u32, 0];
+            data = Some(table);
         }
         "ec.color.videolimiter" => {
             const CLIP_LEVELS: [f32; 5] = [0.9, 0.95, 1.0, 1.05, 1.09];

crates/gpu/src/fx_transition.rs

@@ -15,6 +15,8 @@
 use effectcraft_effects::{Buf, EffectCtx};
 use effectcraft_raster::Image;
 
+use rayon::prelude::*;
+
 use crate::context::{Enc, GpuImage, Params};
 use crate::effects::{GBuf, gaussian_blur};
 
@@ -203,25 +205,22 @@ fn grid_wipe(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {
     if done <= 0.0 {
         return Some(b);
     }
-    let pr = ctx.params;
-    let c = b.to_px(pr.v2("center"));
-    let rot = pr.f("rotation").to_radians();
-    let border = pr.f("border").max(0.0) * b.scale;
-    let tiles = pr.f("tiles").round().clamp(1.0, 500.0);
-    let shape = pr.e("shape");
-    let cs = (ctx.layer_size[0] * b.scale / tiles).max(1.0);
-    let (s, co) = rot.sin_cos();
-    let metric = |dx: f64, dy: f64| match shape {
-        0 => dx.abs(),
-        2 => dx.abs().max(dy.abs()),
-        _ => dx.hypot(dy),
-    };
-    let maxd = max_corner_dist(&b, c, |dx, dy| metric(dx * co + dy * s, -dx * s + dy * co));
-    let mut p = Params::default();
-    p.u[0] = [shape, pr.b("reverse") as u32, 0, 0];
-    p.f[0] = f4([c.0, c.1, s, co]);
-    p.f[1] = f4([cs, maxd, done, border]);
-    run(e, "ftr_grid", &p, b, None, None)
+    // Decide discontinuous tile coverage in the CPU effect's f64 geometry.
+    // Only a scalar mask is uploaded; source pixels and multiplication stay resident.
+    let count = (b.img.width as usize).checked_mul(b.img.height as usize)?;
+    if count.checked_mul(4)? > e.g.device.limits().max_storage_buffer_binding_size as usize {
+        return None;
+    }
+    let geometry = shape(&b);
+    let grid = effectcraft_effects::GridWipe::new(ctx, &geometry);
+    let mut mask = vec![0.0; count];
+    mask.par_chunks_mut((b.img.width as usize).max(1)).enumerate().for_each(|(y, row)| {
+        for (x, value) in row.iter_mut().enumerate() {
+            *value = grid.coverage(x as u32, y as u32);
+        }
+    });
+    let data = e.data(&mask);
+    run(e, "ftr_grid", &Params::default(), b, None, Some(&data))
 }
 
 fn radial_scale_wipe(e: &mut Enc, ctx: &EffectCtx, b: GBuf) -> Option<GBuf> {

crates/gpu/src/shaders/common.wgsl

@@ -50,6 +50,91 @@ fn acos_p(x: f32) -> f32 {
     return 1.5707963267948966 - asin_p(x);
 }
 
+// Some hardware's native trig instructions saturate cos near a full turn and
+// have errors large enough to move particles and sphere samples visibly. Reduce
+// to [-pi/4, pi/4] with a split pi/2, then evaluate the Taylor polynomials with
+// explicit fused operations. The effect angles use this bounded range; retain
+// the builtins for non-finite or very large arguments rather than cast them.
+fn sincos_p(x: f32) -> vec2<f32> {
+    if (!(abs(x) <= 8192.0)) {
+        return vec2<f32>(sin(x), cos(x));
+    }
+    let k = round(x * 0.6366197723675814);
+    // Keep the coarse reduction in integers. Some drivers reassociate even
+    // explicit fma chains into x - k*pi/2, losing low bits for large angles.
+    // 3216/2048 is exact; the bounded scaled input and product fit in i32.
+    let scaled = x * 2048.0;
+    let whole = floor(scaled);
+    var r = f32(i32(whole) - i32(k) * 3216) / 2048.0;
+    r += (scaled - whole) / 2048.0;
+    r = fma(-k, 0.0004838267893064022, r);
+    r = fma(-k, 5.564431554416771e-12, r);
+    let z = r * r;
+    var sp = fma(z, -1.0 / 39916800.0, 1.0 / 362880.0);
+    sp = fma(z, sp, -1.0 / 5040.0);
+    sp = fma(z, sp, 1.0 / 120.0);
+    sp = fma(z, sp, -1.0 / 6.0);
+    let sn = fma(r * z, sp, r);
+    var cp = fma(z, -1.0 / 3628800.0, 1.0 / 40320.0);
+    cp = fma(z, cp, -1.0 / 720.0);
+    cp = fma(z, cp, 1.0 / 24.0);
+    cp = fma(z, cp, -0.5);
+    let cs = fma(z, cp, 1.0);
+    switch imod(i32(k), 4) {
+        case 1: { return vec2<f32>(cs, -sn); }
+        case 2: { return vec2<f32>(-sn, -cs); }
+        case 3: { return vec2<f32>(-cs, sn); }
+        default: { return vec2<f32>(sn, cs); }
+    }
+}
+
+fn sin_p(x: f32) -> f32 { return sincos_p(x).x; }
+fn cos_p(x: f32) -> f32 { return sincos_p(x).y; }
+
+// Correctly-rounded positive x/a, clamped to [0,1]. Native division and fma
+// need not have IEEE precision on every shader backend. Discontinuous colour
+// effects need the exact CPU f32 value, not just an approximation within an ulp.
+// Long division uses only bounded 24-bit significands (at most 24 iterations).
+fn unit_div(x: f32, a: f32) -> f32 {
+    if (x <= 0.0 || a <= 0.0) { return 0.0; }
+    if (x >= a) { return 1.0; }
+    if (a == 1.0) { return x; }
+    let xb = bitcast<u32>(x);
+    let ab = bitcast<u32>(a);
+    var nx = xb & 0x7fffffu;
+    var na = ab & 0x7fffffu;
+    var ex = i32((xb >> 23u) & 255u) - 127;
+    var ea = i32((ab >> 23u) & 255u) - 127;
+    if (ex == -127) {
+        let shift = 23u - firstLeadingBit(nx);
+        nx <<= shift;
+        ex = -126 - i32(shift);
+    } else { nx |= 0x800000u; }
+    if (ea == -127) {
+        let shift = 23u - firstLeadingBit(na);
+        na <<= shift;
+        ea = -126 - i32(shift);
+    } else { na |= 0x800000u; }
+    var exponent = ex - ea;
+    if (nx < na) { nx <<= 1u; exponent -= 1; }
+    var rem = nx - na;
+    if (exponent < -150) { return 0.0; }
+    if (exponent == -150) { return bitcast<f32>(select(0u, 1u, rem != 0u)); }
+    let fraction = u32(min(23, exponent + 149));
+    var q = 1u;
+    for (var i = 0u; i <= fraction; i++) {
+        rem <<= 1u;
+        q <<= 1u;
+        if (rem >= na) { rem -= na; q |= 1u; }
+    }
+    let guard = q & 1u;
+    q >>= 1u;
+    if (guard != 0u && (rem != 0u || (q & 1u) != 0u)) { q += 1u; }
+    if (exponent < -126) { return bitcast<f32>(q); }
+    if (q == 0x1000000u) { q >>= 1u; exponent += 1; }
+    return bitcast<f32>((u32(exponent + 127) << 23u) | (q & 0x7fffffu));
+}
+
 // ---------------------------------------------------------------- sampling (raster::Image)
 
 fn tex_get(t: texture_2d<f32>, x: i32, y: i32) -> vec4<f32> {

crates/gpu/src/shaders/fx_pixel2.wgsl

@@ -633,50 +633,9 @@ fn fp2_hextile(@builtin(global_invocation_id) gid: vec3<u32>) {
     if (!fp2_inside(p)) {
         return;
     }
-    let cx = P.f[0].x;
-    let cy = P.f[0].y;
-    let r = P.f[0].z;
-    let smear = P.f[0].w;
-    let sn = P.f[1].x;
-    let cs = P.f[1].y;
-    let mode = P.u[0].x;
-    let dx = f32(p.x) + 0.5 - cx;
-    let dy = f32(p.y) + 0.5 - cy;
-    let u = dx * cs + dy * sn;
-    let v = -dx * sn + dy * cs;
-    // stylize3::hex_cell (cube rounding).
-    let q = (0.5773502691896257 * u - v / 3.0) / r;
-    let rr = (2.0 / 3.0 * v) / r;
-    let ccy = -q - rr;
-    var rx = round_away(q);
-    let ry = round_away(ccy);
-    var rz = round_away(rr);
-    let ddx = abs(rx - q);
-    let ddy = abs(ry - ccy);
-    let ddz = abs(rz - rr);
-    if (ddx > ddy && ddx > ddz) {
-        rx = -ry - rz;
-    } else if (ddy <= ddz) {
-        rz = -rx - ry;
-    }
-    let hx = r * 1.7320508075688772 * (rx + rz / 2.0);
-    let hy = r * 1.5 * rz;
-    var lu = u - hx;
-    var lv = v - hy;
-    let odd = imod(i32(rx) + i32(rz), 2) == 1;
-    if ((mode == 1u || mode == 2u) && odd) {
-        lu = -lu;
-    }
-    if (mode == 2u && odd) {
-        lv = -lv;
-    }
-    let ox = lu * cs - lv * sn;
-    let oy = lu * sn + lv * cs;
-    let tx = cx + ox;
-    let ty = cy + oy;
-    let sx = tx + (f32(p.x) + 0.5 - tx) * smear;
-    let sy = ty + (f32(p.y) + 0.5 - ty) * smear;
-    textureStore(out, p, sample_bilinear_clamped(src, sx, sy));
+    // Cell membership and reflection were decided in the CPU effect's f64 geometry.
+    let i = (gid.y * u32(out_dims().x) + gid.x) * 2u;
+    textureStore(out, p, sample_bilinear_clamped(src, data[i], data[i + 1u]));
 }
 
 // ---------------------------------------------------------------- CC Mr. Smoothie (stylize3::smoothie)

crates/gpu/src/shaders/fx_sim.wgsl

@@ -114,8 +114,8 @@ fn fxm_sprite(at: u32, x: f32, y: f32) -> FxmShade {
             cov = min(0.12 * select(0.0, 1.0, d <= 1.0) + ring * 0.8 + hl, 1.0);
         }
         case 5u: {
-            let sn = sin(data[at + 8u]);
-            let cs = cos(data[at + 8u]);
+            let sn = sin_p(data[at + 8u]);
+            let cs = cos_p(data[at + 8u]);
             let ux = (dx * cs + dy * sn) / r;
             let uy = (-dx * sn + dy * cs) / r;
             let d = sqrt(ux * ux + uy * uy);
@@ -125,23 +125,23 @@ fn fxm_sprite(at: u32, x: f32, y: f32) -> FxmShade {
             cov = min(core + stx + sty, 1.0);
         }
         case 6u: {
-            let sn = sin(data[at + 8u]);
-            let cs = cos(data[at + 8u]);
+            let sn = sin_p(data[at + 8u]);
+            let cs = cos_p(data[at + 8u]);
             let ux = dx * cs + dy * sn;
             let uy = -dx * sn + dy * cs;
             let e = r * 0.75;
             cov = fxm_aa(e - abs(ux)) * fxm_aa(e - abs(uy));
         }
         case 7u: {
-            let sn = sin(data[at + 8u]);
-            let cs = cos(data[at + 8u]);
+            let sn = sin_p(data[at + 8u]);
+            let cs = cos_p(data[at + 8u]);
             let ux = dx * cs + dy * sn;
             let uy = -dx * sn + dy * cs;
             let e = r * 0.5;
             var m = 3.4e38;
             for (var k = 0; k < 3; k++) {
                 let a = 1.5707964 + f32(k) * 6.2831855 / 3.0;
-                m = min(m, e - (ux * cos(a) + uy * sin(a)));
+                m = min(m, e - (ux * cos_p(a) + uy * sin_p(a)));
             }
             cov = fxm_aa(m);
         }
@@ -428,8 +428,8 @@ fn fxm_ripples(x: f32, y: f32) -> f32 {
         let r = data[i * 4u + 2u];
         let u = (d - r) / P.f[0].y;
         if (abs(u) < 1.0) {
-            let env = 0.5 + 0.5 * cos(u * 3.1415927);
-            v += data[i * 4u + 3u] * env * cos((d - r) / P.f[0].x * 6.2831855);
+            let env = 0.5 + 0.5 * cos_p(u * 3.1415927);
+            v += data[i * 4u + 3u] * env * cos_p((d - r) / P.f[0].x * 6.2831855);
         }
     }
     return v;

crates/gpu/src/shaders/fx_stylize.wgsl

@@ -42,7 +42,11 @@ fn fxs_qdiv(x: vec3<f32>, a: f32) -> vec3<f32> {
     let r = fma(-q, vec3<f32>(a), x);
     // `q + r / a`, written as an fma so drivers (NVIDIA Vulkan) can't contract it another way
     // and round differently from the CPU.
-    return fma(r / a, vec3<f32>(1.0), q);
+    let corrected = fma(r / a, vec3<f32>(1.0), q);
+    return vec3<f32>(
+        select(corrected.x, unit_div(x.x, a), x.x >= 0.0 && x.x <= a),
+        select(corrected.y, unit_div(x.y, a), x.y >= 0.0 && x.y <= a),
+        select(corrected.z, unit_div(x.z, a), x.z >= 0.0 && x.z <= a));
 }
 
 fn fxs_rem2(i: i32) -> i32 {
@@ -63,6 +67,32 @@ fn fxs_bin(v: f32, t: f32, inv: bool) -> f32 {
     return select(0.0, 1.0, (v >= t) != inv);
 }
 
+// Preserve the CPU's separately-rounded f32 multiply before rounding to a bin.
+// A shader compiler may reassociate division and multiplication across a half
+// threshold. The 24-bit mantissa times at most 254 fits in u32 exactly.
+fn fxs_posterize_bin(v: f32, levels: u32) -> f32 {
+    let bits = bitcast<u32>(clamp(v, 0.0, 1.0));
+    let exponent = (bits >> 23u) & 255u;
+    if (exponent == 0u) { return 0.0; }
+    var product = ((bits & 0x7fffffu) | 0x800000u) * levels;
+    let high = firstLeadingBit(product);
+    if (high > 23u) {
+        let lost = high - 23u;
+        let mask = (1u << lost) - 1u;
+        let remainder = product & mask;
+        let half = 1u << (lost - 1u);
+        var kept = product >> lost;
+        if (remainder > half || (remainder == half && (kept & 1u) != 0u)) { kept += 1u; }
+        product = kept << lost;
+    }
+    let shift = 150u - exponent;
+    if (shift > 32u) { return 0.0; }
+    if (shift == 32u) { return select(0.0, 1.0, product >= 0x80000000u); }
+    let mask = (1u << shift) - 1u;
+    let integer = (product >> shift) + select(0u, 1u, (product & mask) >= (1u << (shift - 1u)));
+    return f32(integer);
+}
+
 fn fxs_strobe_op(op: u32, v: f32, k: f32) -> f32 {
     switch op {
         case 1u: { return v + k; }
@@ -96,8 +126,8 @@ fn fxs_point(@builtin(global_invocation_id) gid: vec3<u32>) {
     switch mode {
         case 1u: {
             let lv = P.f[0].x;
-            let c = clamp(fxs_qdiv(px.xyz, a), vec3<f32>(0.0), vec3<f32>(1.0)) * lv;
-            let o = vec3<f32>(round_away(c.x), round_away(c.y), round_away(c.z)) / lv;
+            let c = vec3<f32>(unit_div(px.x, a), unit_div(px.y, a), unit_div(px.z, a));
+            let o = vec3<f32>(fxs_posterize_bin(c.x, u32(lv)), fxs_posterize_bin(c.y, u32(lv)), fxs_posterize_bin(c.z, u32(lv))) / lv;
             textureStore(out, p, vec4<f32>(o * a, a));
         }
         case 2u: {

crates/gpu/src/shaders/fx_tone.wgsl

@@ -122,16 +122,17 @@ fn fxt_knee(v: f32, start: f32, limit: f32) -> f32 {
     return start + r * (1.0 - exp(-(v - start) / r));
 }
 
-// PS Arbitrary Map lookup (map in data, n entries; f[0].x = phase offset in entries).
+// PS Arbitrary Map: CPU-rounded interval starts, values and slopes in data.
 fn fxt_lookup(v: f32, n: u32) -> f32 {
-    let nf = f32(n);
-    var x = clamp(v, 0.0, 1.0) * (nf - 1.0) + P.f[0].x;
-    x = x - nf * floor(x / nf);
-    let fl = floor(x);
-    let i0 = u32(fl) % n;
-    let i1 = select(i0, i0 + 1u, i0 + 1u < n);
-    let t = x - fl;
-    return data[i0] + (data[i1] - data[i0]) * t;
+    let x = clamp(v, 0.0, 1.0);
+    var lo = 0u;
+    var hi = n;
+    while (lo + 1u < hi) {
+        let mid = lo + (hi - lo) / 2u;
+        if (x < data[mid * 3u]) { hi = mid; } else { lo = mid; }
+    }
+    let i = lo * 3u;
+    return data[i + 1u] + (x - data[i]) * data[i + 2u];
 }
 
 fn fxt_eq_look(table: u32, n: u32, v: f32) -> f32 {
@@ -440,7 +441,11 @@ fn fxt_point(@builtin(global_invocation_id) gid: vec3<u32>) {
         case 12u: {
             let n = u0.z;
             let aa = max(a, 0.0);
-            let r = vec3<f32>(fxt_lookup(c.x, n), fxt_lookup(c.y, n), fxt_lookup(c.z, n));
+            var straight = vec3<f32>(0.0);
+            if (a > 1e-6) {
+                straight = vec3<f32>(unit_div(px.x, a), unit_div(px.y, a), unit_div(px.z, a));
+            }
+            let r = vec3<f32>(fxt_lookup(straight.x, n), fxt_lookup(straight.y, n), fxt_lookup(straight.z, n));
             var na = aa;
             if (u0.y != 0u) {
                 na = fxt_lookup(aa, n);

crates/gpu/src/shaders/fx_transition.wgsl

@@ -224,28 +224,8 @@ fn ftr_grid(@builtin(global_invocation_id) gid: vec3<u32>) {
         return;
     }
     let px = textureLoad(src, p, 0);
-    let d = ftr_centre(p) - P.f[0].xy;
-    let s = P.f[0].z;
-    let co = P.f[0].w;
-    let q = vec2<f32>(d.x * co + d.y * s, -d.x * s + d.y * co);
-    let cs = P.f[1].x;
-    let cc = (floor(q / cs) + 0.5) * cs;
-    var o = clamp(ftr_metric(P.u[0].x, cc) / P.f[1].y, 0.0, 1.0);
-    if (P.u[0].y != 0u) {
-        o = 1.0 - o;
-    }
-    let spread = 0.25;
-    let pr = clamp((P.f[1].z * (1.0 + spread) - o) / spread, 0.0, 1.0);
-    if (pr <= 0.0) {
-        textureStore(out, p, px);
-        return;
-    }
-    let half = cs * 0.5 * (1.0 - pr);
-    let m = max(abs(q.x - cc.x), abs(q.y - cc.y));
-    var k = select(0.0, 1.0, m < half);
-    if (P.f[1].w > 0.0) {
-        k = clamp((half - m) / P.f[1].w + 0.5, 0.0, 1.0);
-    }
+    // CPU-generated coverage preserves exact f64 tile and hard-edge decisions.
+    let k = data[gid.y * u32(out_dims().x) + gid.x];
     textureStore(out, p, px * k);
 }
 

crates/gpu/src/shaders/fx_vr.wgsl

@@ -16,7 +16,7 @@ fn fxv_unpremul(p: vec4<f32>) -> vec4<f32> {
 fn fxv_equi_dir(x: f32, y: f32, w: f32, h: f32) -> vec3<f32> {
     let lon = (x / w - 0.5) * 2.0 * FXV_PI;
     let lat = (0.5 - y / h) * FXV_PI;
-    return vec3<f32>(cos(lat) * sin(lon), sin(lat), cos(lat) * cos(lon));
+    return vec3<f32>(cos_p(lat) * sin_p(lon), sin_p(lat), cos_p(lat) * cos_p(lon));
 }
 
 // vr::equi_px.
@@ -168,7 +168,7 @@ fn fxv_dir(kind: u32, fov: f32, x: f32, y: f32, w: f32, h: f32) -> vec4<f32> {
             cu = u / r;
             cv = v / r;
         }
-        return vec4<f32>(sin(th) * cu, cos(th), sin(th) * cv, 1.0);
+        return vec4<f32>(sin_p(th) * cu, cos_p(th), sin_p(th) * cv, 1.0);
     }
     if (kind == 4u) {
         let half = radians(clamp(fov, 1.0, 360.0) / 2.0);
@@ -186,7 +186,7 @@ fn fxv_dir(kind: u32, fov: f32, x: f32, y: f32, w: f32, h: f32) -> vec4<f32> {
             cu = u / r;
             cv = v / r;
         }
-        return vec4<f32>(sin(th) * cu, sin(th) * cv, cos(th), 1.0);
+        return vec4<f32>(sin_p(th) * cu, sin_p(th) * cv, cos_p(th), 1.0);
     }
     let t = tan(radians(clamp(fov, 1.0, 179.0) / 2.0));
     let u = (x / w * 2.0 - 1.0) * t;
@@ -519,8 +519,8 @@ fn fxv_gen(@builtin(global_invocation_id) gid: vec3<u32>) {
                 var s = px;
                 if (al >= 1e-9) {
                     let new_th = th * (1.0 - P.f[0][c] * weight);
-                    let perp = normalize(d - center * cos(th));
-                    let nd = center * cos(new_th) + perp * sin(new_th);
+                    let perp = normalize(d - center * cos_p(th));
+                    let nd = center * cos_p(new_th) + perp * sin_p(new_th);
                     let q = fxv_equi_px(nd, w, h);
                     s = fxv_sample_equi(q.x, q.y);
                 }

crates/gpu/src/shaders/particles.wgsl

@@ -53,12 +53,12 @@ fn ps_norm3(v: vec3<f32>) -> vec3<f32> {
 fn ps_rand_dir(i: u32, k: u32, flat: bool) -> vec3<f32> {
     if (flat) {
         let a = ps_h(i, k) * PS_TAU;
-        return vec3<f32>(cos(a), sin(a), 0.0);
+        return vec3<f32>(cos_p(a), sin_p(a), 0.0);
     }
     let z = ps_hs(i, k);
     let a = ps_h(i, k + 1u) * PS_TAU;
     let r = sqrt(max(1.0 - z * z, 0.0));
-    return vec3<f32>(r * cos(a), r * sin(a), z);
+    return vec3<f32>(r * cos_p(a), r * sin_p(a), z);
 }
 
 struct Part {
@@ -189,7 +189,7 @@ fn ps_spawn_cannon(id: u32) -> Part {
     let vel = S.f[1].y;
     let vel_spread = S.f[1].z;
     let a = radians(angle + ps_hs(id, 1u) * spread * 0.5);
-    let d = vec2<f32>(sin(a), -cos(a));
+    let d = vec2<f32>(sin_p(a), -cos_p(a));
     let speed = max(vel + ps_hs(id, 3u) * vel_spread, 0.0);
     var off: vec2<f32>;
     if (r >= 0.0) {
@@ -197,7 +197,7 @@ fn ps_spawn_cannon(id: u32) -> Part {
     } else {
         let aa = ps_h(id, 5u) * PS_TAU;
         let rr = sqrt(ps_h(id, 6u)) * -r;
-        off = vec2<f32>(cos(aa) * rr, sin(aa) * rr);
+        off = vec2<f32>(cos_p(aa) * rr, sin_p(aa) * rr);
     }
     let v = d * speed;
     let dt = S.f[2].z;

crates/gpu/src/tests.rs

@@ -26,6 +26,103 @@ use effectcraft_time::{FrameRate, Tick};
 
 use crate::Gpu;
 
+pub(crate) const MATH_PROBE_SHADER: &str = r#"
+@compute @workgroup_size(16, 16)
+fn math_probe(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let p = vec2<i32>(gid.xy);
+    if (p.x >= out_dims().x || p.y >= out_dims().y) { return; }
+    let v = textureLoad(src, p, 0).x;
+    if (P.u[0].x == 1u) {
+        let a = textureLoad(src, p, 0).w;
+        let q = unit_div(v, a);
+        textureStore(out, p, vec4<f32>(q, fxs_posterize_bin(q, P.u[0].y), 0.0, 1.0));
+        return;
+    }
+    textureStore(out, p, vec4<f32>(sin_p(v), cos_p(v), 0.0, 1.0));
+}
+"#;
+
+/// Direction errors accumulate into visible particle offsets and spherical
+/// remapping errors. Check the numerical primitive before rendering a scene.
+#[test]
+fn angular_math_preserves_particle_and_vr_directions() {
+    let Some(g) = gpu() else { return };
+    for extent in [std::f32::consts::TAU, 100.0, 8192.0] {
+        let mut img = Image::new(4097, 1);
+        for (i, px) in img.data.iter_mut().enumerate() {
+            px[0] = (i as f32 - 2048.0) * extent / 2048.0;
+        }
+        let input = g.ctx.upload_image(&img).unwrap();
+        let mut e = crate::context::Enc::new(&g.ctx);
+        let out = e.scratch(img.width, img.height);
+        e.pixels("math_probe", &Default::default(), &input, None, &out, None);
+        let result = e.download(&out).unwrap();
+        for (i, (v, got)) in img.data.iter().zip(result.data).enumerate() {
+            let want = [(v[0] as f64).sin() as f32, (v[0] as f64).cos() as f32];
+            for c in 0..2 {
+                assert!((want[c] - got[c]).abs() <= 2e-7, "sample {i}, angle {}: CPU {want:?}, GPU {got:?}", v[0]);
+            }
+        }
+    }
+}
+
+#[test]
+fn alpha_division_and_posterization_match_f32_rounding() {
+    let Some(g) = gpu() else { return };
+    let mut values = vec![];
+    for levels in [1u32, 3, 8, 254] {
+        values.clear();
+        for a in [1.0f32, 0.1, 0.3, 1.0 / 255.0, 6.0 / 255.0] {
+            for i in 0..levels {
+                let half = (i as f32 + 0.5) / levels as f32;
+                for v in [0.0, 1e-6, half.next_down(), half, half.next_up(), 1.0] {
+                    values.push([v * a, 0.0, 0.0, a]);
+                }
+            }
+        }
+        let img = Image { width: values.len() as u32, height: 1, data: values.clone() };
+        let input = g.ctx.upload_image(&img).unwrap();
+        let mut e = crate::context::Enc::new(&g.ctx);
+        let out = e.scratch(img.width, 1);
+        let mut params = crate::context::Params::default();
+        params.u[0] = [1, levels, 0, 0];
+        e.pixels("math_probe", &params, &input, None, &out, None);
+        let result = e.download(&out).unwrap();
+        for (i, (v, got)) in values.iter().zip(result.data).enumerate() {
+            let q = (v[0] / v[3]).clamp(0.0, 1.0);
+            assert_eq!(q.to_bits(), got[0].to_bits(), "levels {levels}, pixel {i}, input {v:?}");
+            assert_eq!((q * levels as f32).round(), got[1], "levels {levels}, pixel {i}, input {v:?}");
+        }
+    }
+}
+
+/// Prove that the numerical fixes run on the device, without per-effect image
+/// downloads or CPU fallback. Compiler-limited DX12 devices retain their tested
+/// fallback contract; these kernels must compile on the Vulkan test adapters.
+#[test]
+fn numerical_effects_remain_resident_on_vulkan() {
+    let Some(g) = gpu() else { return };
+    if !g.ctx.name.ends_with("(Vulkan)") {
+        return;
+    }
+    for entry in ["fxs_point", "fxt_point", "ftr_grid", "fp2_hextile"] {
+        assert!(!g.ctx.unavailable_kernels().iter().any(|(name, _)| *name == entry), "{entry}: {:?}", g.ctx.unavailable_kernels());
+    }
+    for (id, vals) in [
+        ("ec.stylize.posterize", vec![("level", n(4.0))]),
+        ("ec.color.psarbitrarymap", vec![("phase", n(40.0))]),
+        ("ec.transition.ccgridwipe", vec![("completion", n(40.0))]),
+        ("ec.stylize.cchextile", vec![("radius", n(9.3))]),
+    ] {
+        let before = g.ctx.transfer_stats();
+        effect_direct(id, &vals, false);
+        let after = g.ctx.transfer_stats();
+        // Full- and half-resolution runs each upload once and download once.
+        assert_eq!(after.uploads - before.uploads, 2, "{id}: uploads");
+        assert_eq!(after.readbacks - before.readbacks, 2, "{id}: readbacks");
+    }
+}
+
 /// One GPU test at a time in this process: the calling test thread holds a process-wide lock
 /// until it exits (libtest runs every test on a thread of its own). wgpu's OpenGL backend (Mesa
 /// llvmpipe on FreeBSD, Linux without Vulkan) has one context lock per adapter and panics when a

crates/gpu/src/tests_fx_stylize.rs

@@ -45,6 +45,66 @@ fn posterize_thresholds_strobe_vignette() {
     effect_case("ec.stylize.ccvignette", &[("amount", n(-60.0)), ("pinHighlights", n(50.0)), ("angleOfView", n(120.0)), ("center", pt(20.0, 15.0))]);
 }
 
+/// Premultiplied, quantized footage often lands exactly on a posterization
+/// threshold after division by alpha. Exercise that input before any placement
+/// or compositing can hide a wrong bin behind the normal image tolerance.
+#[test]
+fn posterize_quantized_premultiplied_boundaries() {
+    let Some(g) = crate::tests::gpu() else { return };
+    let spec = effectcraft_effects::find("ec.stylize.posterize").unwrap();
+    for level in [4.0, 9.0] {
+        let params = effectcraft_effects::Params { values: [("level".into(), n(level))].into() };
+        let ctx =
+            || effectcraft_effects::EffectCtx { params: &params, time: 0.0, layer_size: [70.0, 44.0], seed: 11, adjustment: false, env: Default::default() };
+        for seed in 1..=7 {
+            let mut img = crate::tests::pattern(seed, 70, 44);
+            effectcraft_render::color::quantize(&mut img, 255.0);
+            let input = effectcraft_effects::Buf { img, offset: [0.0; 2], scale: 1.0 };
+            let cpu = (spec.render)(&ctx(), input.clone());
+            let gpu = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &input, None).unwrap();
+            for (i, (want, got)) in cpu.img.data.iter().zip(&gpu.img.data).enumerate() {
+                for c in 0..4 {
+                    assert!((want[c] - got[c]).abs() <= 1e-6, "level {level} seed {seed} pixel {i}: input {:?}, CPU {want:?}, GPU {got:?}", input.img.data[i]);
+                }
+            }
+        }
+    }
+}
+
+/// Wrapped ramps have a large discontinuity at their end. Check signed and
+/// fractional phases, including a wrap a few double-precision ulps above zero.
+#[test]
+fn arbitrary_map_signed_fractional_and_near_zero_wraps() {
+    let Some(g) = crate::tests::gpu() else { return };
+    let spec = effectcraft_effects::find("ec.color.psarbitrarymap").unwrap();
+    for phase in [-1000.0, -20.5, 0.0, 40.0, 1000.0, 255.99999999999997] {
+        let params = effectcraft_effects::Params {
+            values: [
+                ("map".into(), Value::Str((0..256).map(|n| n.to_string()).collect::<Vec<_>>().join(","))),
+                ("phase".into(), n(phase)),
+                ("applyPhaseMapToAlpha".into(), on()),
+            ]
+            .into(),
+        };
+        let ctx =
+            || effectcraft_effects::EffectCtx { params: &params, time: 0.0, layer_size: [1024.0, 1.0], seed: 7, adjustment: false, env: Default::default() };
+        let mut img = effectcraft_raster::Image::new(1024, 1);
+        for (i, px) in img.data.iter_mut().enumerate() {
+            let v = if i < 512 { i as f32 * 1e-18 } else { (i - 512) as f32 / 511.0 };
+            let a = [1.0f32, 0.1, 0.3, 6.0 / 255.0][i % 4];
+            *px = [v * a, v * a, v * a, a];
+        }
+        let input = effectcraft_effects::Buf { img, offset: [0.0; 2], scale: 1.0 };
+        let cpu = (spec.render)(&ctx(), input.clone());
+        let gpu = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &input, None).unwrap();
+        for (i, (want, got)) in cpu.img.data.iter().zip(&gpu.img.data).enumerate() {
+            for c in 0..4 {
+                assert!((want[c] - got[c]).abs() <= 1e-6, "phase {phase}, pixel {i}: CPU {want:?}, GPU {got:?}");
+            }
+        }
+    }
+}
+
 #[test]
 fn scatter_brush_strokes() {
     effect_case("ec.stylize.scatter", &[("amount", n(5.0))]);

crates/host/tests/plugins.rs

@@ -69,14 +69,18 @@ fn rust_plugins_apply_by_id() {
 /// Build the example plug-in (skipped when the wasm32 target isn't installed).
 fn build_example() -> Option<std::path::PathBuf> {
     let root = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("../../examples/plugins/posterize-bands");
-    let target = std::env::var_os("CARGO_TARGET_DIR").map(std::path::PathBuf::from).unwrap_or_else(|| root.join("../../../target")).join("plugin-example");
-    let target = if target.is_absolute() { target } else { root.join("../../..").join(target) };
-    let st = std::process::Command::new(env!("CARGO"))
-        .current_dir(&root)
-        .args(["build", "--release", "--target", "wasm32-unknown-unknown", "--target-dir"])
-        .arg(&target)
-        .status()
-        .ok()?;
+    let workspace = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("../..");
+    let assigned = std::env::var_os("CARGO_TARGET_DIR");
+    let target = assigned.as_ref().map(std::path::PathBuf::from).unwrap_or_else(|| workspace.join("target"));
+    let target = if target.is_absolute() { target } else { workspace.join(target) };
+    // The plug-in is part of this test's build. Inherit the assigned target and
+    // compiler guards rather than creating a second target at its root.
+    let mut command = std::process::Command::new(env!("CARGO"));
+    command.current_dir(&workspace).args(["build", "--release", "--target", "wasm32-unknown-unknown", "--manifest-path"]).arg(root.join("Cargo.toml"));
+    if assigned.is_none() {
+        command.env("CARGO_TARGET_DIR", &target);
+    }
+    let st = command.status().ok()?;
     if !st.success() {
         eprintln!("skipping: the example plug-in didn't build (rustup target add wasm32-unknown-unknown)");
         return None;

crates/ui-egui/src/theme.rs

@@ -332,14 +332,20 @@ static JETBRAINS_MONO: &[u8] = include_bytes!("../../../assets/fonts/JetBrainsMo
 /// frame loop), so the fonts do not need a restart.
 pub fn install(ctx: &egui::Context, t: &Tokens, language: &str) {
     let mut fonts = FontDefinitions::default();
+    let proportional_fallbacks = fonts.families.get(&FontFamily::Proportional).cloned().unwrap_or_default();
     fonts.font_data.insert("inter".into(), Arc::new(FontData::from_static(INTER_REGULAR)));
     fonts.font_data.insert("inter-medium".into(), Arc::new(FontData::from_static(INTER_MEDIUM)));
     fonts.font_data.insert("inter-semibold".into(), Arc::new(FontData::from_static(INTER_SEMIBOLD)));
     fonts.font_data.insert("jbmono".into(), Arc::new(FontData::from_static(JETBRAINS_MONO)));
     fonts.families.entry(FontFamily::Proportional).or_default().insert(0, "inter".into());
     fonts.families.entry(FontFamily::Monospace).or_default().insert(0, "jbmono".into());
-    fonts.families.insert(FontFamily::Name("semibold".into()), vec!["inter-semibold".into(), "inter".into()]);
-    fonts.families.insert(FontFamily::Name("medium".into()), vec!["inter-medium".into(), "inter".into()]);
+    // Named weights need the same symbol fallback chain as ordinary labels.
+    // The bundled Inter subset and JetBrains Mono do not cover every UI symbol.
+    for (family, face) in [("semibold", "inter-semibold"), ("medium", "inter-medium")] {
+        let mut stack = vec![face.into(), "inter".into()];
+        stack.extend(proportional_fallbacks.iter().cloned());
+        fonts.families.insert(FontFamily::Name(family.into()), stack);
+    }
     // Inter and egui's own fallbacks have no small triangles (the `▸` of "File ▸ Import"), which
     // would draw as missing-glyph boxes: JetBrains Mono fills in the symbols the others lack.
     for (family, stack) in &mut fonts.families {