Merge pull request #69 from storytold/pr/gpu-gl-fallback

GPU on wgpu's OpenGL backend (FreeBSD / Linux without Vulkan)

.github/workflows/freebsd.yml

@@ -0,0 +1,48 @@
+# Builds and tests the workspace on FreeBSD, inside a FreeBSD VM on a Linux runner
+# (GitHub has no FreeBSD runners). GPU tests skip themselves when the VM has no adapter.
+name: FreeBSD
+
+on:
+  push:
+    branches: [main]
+  pull_request:
+    paths:
+      - Cargo.toml
+      - Cargo.lock
+      - "**/Cargo.toml"
+      - .github/workflows/freebsd.yml
+  workflow_dispatch:
+
+concurrency:
+  group: freebsd-${{ github.ref }}
+  cancel-in-progress: true
+
+permissions:
+  contents: read
+
+jobs:
+  freebsd:
+    name: FreeBSD 14 (build + test)
+    runs-on: ubuntu-latest
+    timeout-minutes: 300
+    steps:
+      - uses: actions/checkout@v4
+      - name: Build and test in a FreeBSD VM
+        uses: vmactions/freebsd-vm@c46abacb49f09938ca4e1702d15d836285d694cc # v1
+        with:
+          release: "14.3"
+          usesh: true
+          cpu: 4
+          mem: 12288
+          # Nothing to bring back; copying the target dir out of the VM fails on large builds.
+          copyback: false
+          prepare: |
+            pkg install -y git curl pkgconf cmake python3 libxkbcommon wayland libX11 libXcursor libXrandr libXi libxcb mesa-libs vulkan-loader gtk3 fontconfig freetype2 alsa-lib
+            curl -fsSL https://sh.rustup.rs | sh -s -- -y --profile minimal
+          run: |
+            . "$HOME/.cargo/env"
+            # The checkout is synced in under another owner; let tests that call git read it.
+            git config --global --add safe.directory "*"
+            rustc --version
+            CARGO_PROFILE_DEV_DEBUG=0 cargo build --workspace --all-targets --locked
+            CARGO_PROFILE_DEV_DEBUG=0 cargo test --workspace --locked

crates/gpu/src/adv3d.rs

@@ -2,7 +2,7 @@
 //! (`Depth32Float`, Greater), the scene's materials, textures, lights and shadow maps in
 //! storage buffers, and `advanced3d.wgsl`'s physically based fragment shader. Opaque triangles
 //! are drawn with depth writes, then the sorted transparent ones blended over (premultiplied).
-//! Colour (`Rgba16Float`) and camera depth (`R32Float`) stay on the GPU for `adv3d.wgsl`'s
+//! Colour (`Rgba16Float`) and camera depth (`R32Uint` holding f32 bits) stay on the GPU for `adv3d.wgsl`'s
 //! compute kernels, which run the rest of `effectcraft_render::three_d::adv::render_prepared`:
 //! the 2×2 supersampling resolve, the motion-blur sub-sample average (nearest depth), the
 //! depth-based iris depth of field (the Classic 3D bokeh kernel's row spans and prefix-sum
@@ -22,20 +22,44 @@ use crate::context::{Enc, GpuContext, GpuImage};
 /// Pipelines (built on first use).
 pub(crate) struct Pipes {
     bgl: wgpu::BindGroupLayout,
-    opaque: wgpu::RenderPipeline,
-    transparent: wgpu::RenderPipeline,
+    /// (opaque, transparent); `None` when the adapter can't rasterise ([`raster_unsupported`]):
+    /// scenes then render on the CPU.
+    raster: Option<(wgpu::RenderPipeline, wgpu::RenderPipeline)>,
     post: Post,
 }
 
 fn pipes(g: &GpuContext) -> &Pipes {
-    g.adv3d.get_or_init(|| Pipes::new(&g.device))
+    g.adv3d.get_or_init(|| Pipes::new(&g.device, g.adv3d_raster))
 }
 
 const COLOR: wgpu::TextureFormat = wgpu::TextureFormat::Rgba16Float;
-const DEPTH_OUT: wgpu::TextureFormat = wgpu::TextureFormat::R32Float;
+/// Camera depth as the bits of an f32: integer targets render on every backend, while `R32Float`
+/// needs `EXT_color_buffer_float` on GLES (absent on Mesa's llvmpipe, for one).
+const DEPTH_OUT: wgpu::TextureFormat = wgpu::TextureFormat::R32Uint;
+/// [`DEPTH_OUT`]'s clear value: the bits of −1.0 ("nothing drawn").
+const NO_DEPTH: f64 = 0xBF80_0000u32 as f64;
 const DEPTH: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float;
 const VERTEX_SIZE: u64 = 52;
 
+/// Why `adapter` can't run the rasteriser, or `None` when it can. Building the pipelines anyway
+/// would be a validation error, which release builds only log (and then draw nothing), so the
+/// caller keeps the scenes on the CPU instead.
+pub(crate) fn raster_unsupported(adapter: &wgpu::Adapter) -> Option<String> {
+    use wgpu::TextureFormatFeatureFlags as F;
+    let needs = [(COLOR, F::BLENDABLE), (DEPTH_OUT, F::empty()), (DEPTH, F::empty())];
+    for (format, flags) in needs {
+        let f = adapter.get_texture_format_features(format);
+        if !f.allowed_usages.contains(wgpu::TextureUsages::RENDER_ATTACHMENT) || !f.flags.contains(flags) {
+            return Some(format!("{format:?} is not a {}render target", if flags.is_empty() { "" } else { "blendable " }));
+        }
+    }
+    // The transparent pipeline blends colour and masks the depth target off.
+    if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::INDEPENDENT_BLEND) {
+        return Some("no independent blending".into());
+    }
+    None
+}
+
 fn storage(binding: u32) -> wgpu::BindGroupLayoutEntry {
     wgpu::BindGroupLayoutEntry {
         binding,
@@ -46,7 +70,7 @@ fn storage(binding: u32) -> wgpu::BindGroupLayoutEntry {
 }
 
 impl Pipes {
-    fn new(device: &wgpu::Device) -> Pipes {
+    fn new(device: &wgpu::Device, raster: bool) -> Pipes {
         let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
             label: Some("effectcraft advanced 3d"),
             source: wgpu::ShaderSource::Wgsl(include_str!("shaders/advanced3d.wgsl").into()),
@@ -110,7 +134,7 @@ impl Pipes {
                 cache: None,
             })
         };
-        Pipes { opaque: make(false), transparent: make(true), bgl, post: Post::new(device) }
+        Pipes { raster: raster.then(|| (make(false), make(true))), bgl, post: Post::new(device) }
     }
 }
 
@@ -222,9 +246,9 @@ fn half_to_f32(h: u16) -> f32 {
     }
 }
 
-/// Record the render pass of one scene into `e`: (colour `Rgba16Float`, camera depth `R32Float`
+/// Record the render pass of one scene into `e`: (colour `Rgba16Float`, camera depth `R32Uint`
 /// with −1 where nothing was drawn) at the scene's raster size. `None` when the device can't
-/// (size or buffer limits).
+/// (render targets, size or buffer limits).
 fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)> {
     let g = e.g;
     if !g.fits(s.width, s.height) || s.indices.is_empty() {
@@ -236,6 +260,7 @@ fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)>
         return None;
     }
     let p = pipes(g);
+    let Some((opaque, transparent)) = &p.raster else { return None };
     let dev = &g.device;
     let vb: Vec<u8> = s
         .vertices
@@ -294,7 +319,7 @@ fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)>
                     view: &zv,
                     depth_slice: None,
                     resolve_target: None,
-                    ops: wgpu::Operations { load: wgpu::LoadOp::Clear(wgpu::Color { r: -1.0, g: 0.0, b: 0.0, a: 0.0 }), store: wgpu::StoreOp::Store },
+                    ops: wgpu::Operations { load: wgpu::LoadOp::Clear(wgpu::Color { r: NO_DEPTH, g: 0.0, b: 0.0, a: 0.0 }), store: wgpu::StoreOp::Store },
                 }),
             ],
             depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
@@ -311,11 +336,11 @@ fn raster_into(e: &mut Enc, s: &Scene) -> Option<(wgpu::Texture, wgpu::Texture)>
         pass.set_index_buffer(ibuf.slice(..), wgpu::IndexFormat::Uint32);
         let n = s.indices.len() as u32;
         if s.opaque_count > 0 {
-            pass.set_pipeline(&p.opaque);
+            pass.set_pipeline(opaque);
             pass.draw_indexed(0..s.opaque_count, 0, 0..1);
         }
         if n > s.opaque_count {
-            pass.set_pipeline(&p.transparent);
+            pass.set_pipeline(transparent);
             pass.draw_indexed(s.opaque_count..n, 0, 0..1);
         }
     }
@@ -424,14 +449,10 @@ impl Post {
             label: Some("effectcraft advanced 3d post"),
             source: wgpu::ShaderSource::Wgsl(include_str!("shaders/adv3d.wgsl").into()),
         });
-        let tex = |binding| wgpu::BindGroupLayoutEntry {
+        let tex = |binding, sample_type| wgpu::BindGroupLayoutEntry {
             binding,
             visibility: wgpu::ShaderStages::COMPUTE,
-            ty: wgpu::BindingType::Texture {
-                sample_type: wgpu::TextureSampleType::Float { filterable: false },
-                view_dimension: wgpu::TextureViewDimension::D2,
-                multisampled: false,
-            },
+            ty: wgpu::BindingType::Texture { sample_type, view_dimension: wgpu::TextureViewDimension::D2, multisampled: false },
             count: None,
         };
         let buf = |binding, read_only| wgpu::BindGroupLayoutEntry {
@@ -447,8 +468,8 @@ impl Post {
                 ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, min_binding_size: None },
                 count: None,
             },
-            tex(1),
-            tex(2),
+            tex(1, wgpu::TextureSampleType::Float { filterable: false }),
+            tex(2, wgpu::TextureSampleType::Uint),
             buf(3, false),
             buf(4, false),
             buf(5, false),

crates/gpu/src/context.rs

@@ -259,6 +259,9 @@ pub struct GpuContext {
     zeros: Mutex<HashMap<(u32, u32), GpuImage>>,
     /// Advanced 3D render pipelines (built on first use).
     pub(crate) adv3d: std::sync::OnceLock<crate::adv3d::Pipes>,
+    /// The adapter can run the Advanced 3D rasteriser (`adv3d::raster_unsupported`); without
+    /// it Advanced 3D scenes render on the CPU.
+    pub(crate) adv3d_raster: bool,
     /// GPU particle pipeline (built on first use) and simulation checkpoints.
     pub(crate) particles: crate::particles::PipesCell,
     pub(crate) particle_states: crate::particles::StatesCell,
@@ -291,18 +294,37 @@ fn storage_tex_entry(binding: u32, format: wgpu::TextureFormat) -> wgpu::BindGro
 impl GpuContext {
     /// Build on an existing device (the desktop app shares egui-wgpu's). `Err` when the
     /// adapter cannot run the compositor (no compute shaders, e.g. WebGL2; no float storage
-    /// textures).
+    /// textures) or the device was created with limits below what its pipelines bind.
     pub fn new(adapter: &wgpu::Adapter, device: wgpu::Device, queue: wgpu::Queue) -> Result<GpuContext, String> {
         let info = adapter.get_info();
         if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
             return Err(format!("{} ({:?}): no compute shaders", info.name, info.backend));
         }
+        // egui-wgpu asks for WebGL2's limits on GL (no storage buffers at all): building the
+        // pipelines there would be a validation error, which release builds only log.
+        let l = device.limits();
+        let needs = [
+            // `adv3d`'s post kernels bind 7 (its rasteriser's fragment stage 6).
+            ("max_storage_buffers_per_shader_stage", l.max_storage_buffers_per_shader_stage, 7),
+            ("max_storage_textures_per_shader_stage", l.max_storage_textures_per_shader_stage, 1),
+            ("max_bind_groups", l.max_bind_groups, 2),
+            ("max_compute_workgroup_size_x", l.max_compute_workgroup_size_x, 64),
+            ("max_compute_workgroup_size_y", l.max_compute_workgroup_size_y, 16),
+            ("max_compute_invocations_per_workgroup", l.max_compute_invocations_per_workgroup, 256),
+        ];
+        if let Some((limit, have, need)) = needs.iter().find(|(_, have, need)| have < need) {
+            return Err(format!("{} ({:?}): device limit {limit} is {have}, needs {need}", info.name, info.backend));
+        }
         for f in [FORMAT, wgpu::TextureFormat::Rgba8Unorm] {
             if !adapter.get_texture_format_features(f).allowed_usages.contains(wgpu::TextureUsages::STORAGE_BINDING) {
                 return Err(format!("{} ({:?}): {f:?} storage textures unsupported", info.name, info.backend));
             }
         }
         let name = format!("{} ({:?})", info.name, info.backend);
+        let adv3d_unsupported = crate::adv3d::raster_unsupported(adapter);
+        if let Some(why) = &adv3d_unsupported {
+            log::info!("gpu {name}: Advanced 3D renders on the CPU ({why})");
+        }
         let src = [
             include_str!("shaders/common.wgsl"),
             include_str!("shaders/kernels.wgsl"),
@@ -466,6 +488,7 @@ impl GpuContext {
             transfers: Default::default(),
             zeros: Default::default(),
             adv3d: std::sync::OnceLock::new(),
+            adv3d_raster: adv3d_unsupported.is_none(),
             particles: Default::default(),
             particle_states: Default::default(),
             deferred: None,

crates/gpu/src/shaders/adv3d.wgsl

@@ -25,7 +25,9 @@ struct Post {
 
 @group(0) @binding(0) var<uniform> P: Post;
 @group(0) @binding(1) var color_in: texture_2d<f32>;
-@group(0) @binding(2) var z_in: texture_2d<f32>;
+// Camera depth as f32 bits (an `R32Uint` target: integer targets render on every backend,
+// `R32Float` doesn't on GLES without EXT_color_buffer_float).
+@group(0) @binding(2) var z_in: texture_2d<u32>;
 @group(0) @binding(3) var<storage, read_write> acc: array<vec4<f32>>;
 @group(0) @binding(4) var<storage, read_write> depth: array<f32>;
 @group(0) @binding(5) var<storage, read_write> src: array<vec4<f32>>;
@@ -55,7 +57,7 @@ fn resolve(@builtin(global_invocation_id) id: vec3<u32>) {
         for (var i = 0u; i < k; i = i + 1u) {
             let p = vec2<i32>(i32(id.x * k + i), i32(id.y * k + j));
             a = a + textureLoad(color_in, p, 0);
-            let z = textureLoad(z_in, p, 0).x;
+            let z = bitcast<f32>(textureLoad(z_in, p, 0).x);
             if (z >= 0.0) {
                 dz = dz + z;
                 dn = dn + 1.0;

crates/gpu/src/shaders/advanced3d.wgsl

@@ -75,26 +75,28 @@ fn srgb_to_linear(c: f32) -> f32 {
     return pow((c + 0.055) / 1.055, 2.4);
 }
 
+// `a` mod `n` in [0, n) for n > 0 (floored). Never `%` on a negative operand: WGSL's `%` truncates,
+// but naga's GLSL output is plain `%`, undefined for negative operands (wrong on Mesa's GL).
+fn imod(a: i32, n: i32) -> i32 {
+    if (a >= 0) {
+        return a % n;
+    }
+    return n - 1 - (-(a + 1)) % n;
+}
+
 fn wrap_i(i: i32, n: i32, mode: u32) -> i32 {
     if (mode == 1u) {
         return clamp(i, 0, n - 1);
     }
     if (mode == 2u) {
         let p = 2 * n;
-        var m = i % p;
-        if (m < 0) {
-            m = m + p;
-        }
+        let m = imod(i, p);
         if (m >= n) {
             return p - 1 - m;
         }
         return m;
     }
-    var m = i % n;
-    if (m < 0) {
-        m = m + n;
-    }
-    return m;
+    return imod(i, n);
 }
 
 fn texel(t: vec4<u32>, x: i32, y: i32) -> vec4<f32> {
@@ -282,7 +284,8 @@ fn attenuation(kind: u32, falloff: u32, radius: f32, falloff_distance: f32, cos_
 
 struct FOut {
     @location(0) color: vec4<f32>,
-    @location(1) depth: f32,
+    // Camera depth's f32 bits (an `R32Uint` target, see `adv3d.rs`).
+    @location(1) depth: u32,
 };
 
 @fragment
@@ -326,7 +329,7 @@ fn fs(i: VOut, @builtin(front_facing) front: bool) -> FOut {
         emissive = emissive * sample_tex(tex_em, i.uv).rgb;
     }
     var out: FOut;
-    out.depth = dot(g.view_z.xyz, i.pos) + g.view_z.w;
+    out.depth = bitcast<u32>(dot(g.view_z.xyz, i.pos) + g.view_z.w);
     if (unlit || !accepts_lights || g.env2.y < 0.5) {
         let c = rgb + emissive;
         out.color = vec4<f32>(c * alpha, alpha);

crates/gpu/src/shaders/classic3d.wgsl

@@ -186,7 +186,7 @@ fn attenuation(l: u32, p: vec3<f32>) -> f32 {
     }
     if (li.ints.x == 3u) {
         let cos_a = clamp(dot(-tl.xyz, li.dir.xyz), -1.0, 1.0);
-        let a = acos(cos_a);
+        let a = acos_p(cos_a);
         let outer = li.cone.x;
         let inner = outer * (1.0 - li.cone.y);
         if (a <= inner) {

crates/gpu/src/shaders/common.wgsl

@@ -16,6 +16,40 @@ struct Params {
 @group(0) @binding(2) var aux: texture_2d<f32>;
 @group(0) @binding(4) var<storage, read> data: array<f32>;
 
+// `a` mod `n` in [0, n) for n > 0 (floored). Never `%` on a negative operand: WGSL's `%` truncates,
+// but naga's GLSL output is plain `%`, undefined for negative operands (wrong on Mesa's GL).
+fn imod(a: i32, n: i32) -> i32 {
+    if (a >= 0) {
+        return a % n;
+    }
+    return n - 1 - (-(a + 1)) % n;
+}
+
+// asin / acos to f32 precision on every backend (Cephes' asinf polynomial). WGSL leaves their
+// precision to the implementation, and Mesa's GL drivers use a coarse approximation (errors near
+// 1e-4 rad), which shows in every effect that maps through a sphere or a lens.
+fn asin_p(x: f32) -> f32 {
+    let a = abs(x);
+    let big = a > 0.5;
+    let z = select(a * a, 0.5 * (1.0 - a), big);
+    let s = select(a, sqrt(z), big);
+    var p = ((((4.2163199048e-2 * z + 2.4181311049e-2) * z + 4.5470025998e-2) * z + 7.4953002686e-2) * z + 1.6666752422e-1) * z * s + s;
+    if (big) {
+        p = 1.5707963267948966 - 2.0 * p;
+    }
+    return select(p, -p, x < 0.0);
+}
+
+fn acos_p(x: f32) -> f32 {
+    if (x > 0.5) {
+        return 2.0 * asin_p(sqrt(0.5 * (1.0 - x)));
+    }
+    if (x < -0.5) {
+        return 3.141592653589793 - 2.0 * asin_p(sqrt(0.5 * (1.0 + x)));
+    }
+    return 1.5707963267948966 - asin_p(x);
+}
+
 // ---------------------------------------------------------------- sampling (raster::Image)
 
 fn tex_get(t: texture_2d<f32>, x: i32, y: i32) -> vec4<f32> {

crates/gpu/src/shaders/fx_distort.wgsl

@@ -215,7 +215,7 @@ fn dst_cclens(@builtin(global_invocation_id) gid: vec3<u32>) {
     }
     var s = q;
     if (d >= 1e-9) {
-        let g = powz(asin(min(d / r, 1.0)) / DST_HALF_PI, P.f[0].w);
+        let g = powz(asin_p(min(d / r, 1.0)) / DST_HALF_PI, P.f[0].w);
         let k = r * g / d;
         s = vec2<f32>(c.x + dx * k, c.y + dy * k);
     }

crates/gpu/src/shaders/fx_gen2.wgsl

@@ -653,7 +653,7 @@ fn fg2_threads(@builtin(global_invocation_id) gid: vec3<u32>) {
     let fv = v / th - f32(cj);
     let in_h = abs(fv - 0.5) < coverage * 0.5;
     let in_v = abs(fu - 0.5) < coverage * 0.5;
-    let h_top = ((ci + cj) % (2 * overlaps) + 2 * overlaps) % (2 * overlaps) < overlaps;
+    let h_top = imod(ci + cj, 2 * overlaps) < overlaps;
     var horizontal: bool;
     if (in_h && in_v) {
         horizontal = h_top;

crates/gpu/src/shaders/fx_generate.wgsl

@@ -349,8 +349,8 @@ fn gen_worley(u: f32, v: f32, disperse: f32, evo: f32, seed: u32) -> vec3<f32> {
             var tx = gx;
             var ty = gy;
             if (tile) {
-                tx = ((gx % nx) + nx) % nx;
-                ty = ((gy % ny) + ny) % ny;
+                tx = imod(gx, nx);
+                ty = imod(gy, ny);
             }
             let hx = bitcast<u32>(tx);
             let hy = bitcast<u32>(ty);
@@ -428,7 +428,7 @@ fn gen_axis(q: f32, size: f32, f: f32) -> f32 {
     let fr = u - fl;
     let d = min(fr, 1.0 - fr) * size;
     let n = i32(fl);
-    let s = select(-1.0, 1.0, ((n % 2) + 2) % 2 == 0);
+    let s = select(-1.0, 1.0, imod(n, 2) == 0);
     return s * min(d / (f * 0.5), 1.0);
 }
 

crates/gpu/src/shaders/fx_pixel2.wgsl

@@ -661,7 +661,7 @@ fn fp2_hextile(@builtin(global_invocation_id) gid: vec3<u32>) {
     let hy = r * 1.5 * rz;
     var lu = u - hx;
     var lv = v - hy;
-    let odd = ((i32(rx) + i32(rz)) % 2 + 2) % 2 == 1;
+    let odd = imod(i32(rx) + i32(rz), 2) == 1;
     if ((mode == 1u || mode == 2u) && odd) {
         lu = -lu;
     }

crates/gpu/src/shaders/fx_stylize.wgsl

@@ -44,7 +44,7 @@ fn fxs_qdiv(x: vec3<f32>, a: f32) -> vec3<f32> {
 }
 
 fn fxs_rem2(i: i32) -> i32 {
-    return ((i % 2) + 2) % 2;
+    return imod(i, 2);
 }
 
 // ---------------------------------------------------------------- per-pixel colour effects
@@ -257,7 +257,7 @@ fn fxs_warp(@builtin(global_invocation_id) gid: vec3<u32>) {
             let d = sqrt(dx * dx + dy * dy);
             if (d < r && d != 0.0) {
                 let nd = d / r;
-                let k = (asin(nd) / FXS_HALF_PI) / nd;
+                let k = (asin_p(nd) / FXS_HALF_PI) / nd;
                 s = vec2<f32>(c.x + dx * k, c.y + dy * k);
             }
         }

crates/gpu/src/shaders/fx_transition.wgsl

@@ -681,7 +681,7 @@ fn ftr_cylinder(@builtin(global_invocation_id) gid: vec3<u32>) {
         return;
     }
     let tau = 6.283185307179586;
-    let th = asin(s);
+    let th = asin_p(s);
     let cz = cos(th);
     let spin = P.f[2].x;
     let front = ftr_shade(ftr_texel((th + spin) / tau + 0.5, v), vec3<f32>(s, 0.0, cz));
@@ -695,7 +695,7 @@ fn ftr_sphere_tex(q: vec3<f32>) -> vec4<f32> {
     let m1 = P.f[8].xyz;
     let m2 = P.f[9].xyz;
     let t = m0 * q.x + m1 * q.y + m2 * q.z;
-    let lat = asin(clamp(t.y, -1.0, 1.0));
+    let lat = asin_p(clamp(t.y, -1.0, 1.0));
     return ftr_texel(ftr_lon_u(t.x, t.z), lat / 3.141592653589793 + 0.5);
 }
 
@@ -775,7 +775,7 @@ fn ftr_env_uv(mapping: u32, r: vec3<f32>) -> vec2<f32> {
         }
         return vec2<f32>((cr.x + (cr.z + 1.0) * 0.5) / 3.0, (cr.y + (cr.w + 1.0) * 0.5) / 4.0);
     }
-    let lat = asin(clamp(r.y, -1.0, 1.0));
+    let lat = asin_p(clamp(r.y, -1.0, 1.0));
     return vec2<f32>(ftr_lon_u(r.x, r.z), 0.5 + lat / 3.141592653589793);
 }
 

crates/gpu/src/shaders/fx_vr.wgsl

@@ -22,7 +22,7 @@ fn fxv_equi_dir(x: f32, y: f32, w: f32, h: f32) -> vec3<f32> {
 // vr::equi_px.
 fn fxv_equi_px(d: vec3<f32>, w: f32, h: f32) -> vec2<f32> {
     let lon = atan2(d.x, d.z);
-    let lat = asin(clamp(d.y, -1.0, 1.0));
+    let lat = asin_p(clamp(d.y, -1.0, 1.0));
     return vec2<f32>((lon / (2.0 * FXV_PI) + 0.5) * w, (0.5 - lat / FXV_PI) * h);
 }
 
@@ -42,7 +42,7 @@ fn fxv_equi_at(xi0: i32, yi0: i32, w: i32, h: i32) -> vec4<f32> {
         yi = 2 * h - 1 - yi;
         xi += w / 2;
     }
-    let xm = ((xi % w) + w) % w;
+    let xm = imod(xi, w);
     return textureLoad(src, vec2<i32>(xm, clamp(yi, 0, h - 1)), 0);
 }
 
@@ -233,7 +233,7 @@ fn fxv_proj_sample(kind: u32, fov: f32, d: vec3<f32>) -> vec4<f32> {
         return sample_bilinear_clamped(src, x, y);
     }
     if (kind == 8u) {
-        let th = acos(clamp(d.y, -1.0, 1.0));
+        let th = acos_p(clamp(d.y, -1.0, 1.0));
         if (th > FXV_PI / 2.0 + 1e-6) {
             return vec4<f32>(0.0);
         }
@@ -250,7 +250,7 @@ fn fxv_proj_sample(kind: u32, fov: f32, d: vec3<f32>) -> vec4<f32> {
     }
     if (kind == 4u) {
         let half = radians(clamp(fov, 1.0, 360.0) / 2.0);
-        let th = acos(clamp(d.z, -1.0, 1.0));
+        let th = acos_p(clamp(d.z, -1.0, 1.0));
         if (th > half + 1e-6) {
             return vec4<f32>(0.0);
         }
@@ -346,12 +346,12 @@ fn fxv_box_h(@builtin(global_invocation_id) gid: vec3<u32>) {
     let norm = 1.0 / f32(2 * r + 1);
     var acc = vec4<f32>(0.0);
     for (var i = b0 - r; i <= b0 + r; i++) {
-        acc += textureLoad(src, vec2<i32>(((i % w) + w) % w, y), 0);
+        acc += textureLoad(src, vec2<i32>(imod(i, w), y), 0);
     }
     for (var x = b0; x < b1; x++) {
         textureStore(out, vec2<i32>(x, y), acc * norm);
         let add = textureLoad(src, vec2<i32>((x + r + 1) % w, y), 0);
-        let sub = textureLoad(src, vec2<i32>(((x - r) % w + w) % w, y), 0);
+        let sub = textureLoad(src, vec2<i32>(imod(x - r, w), y), 0);
         acc += add - sub;
     }
 }
@@ -362,10 +362,10 @@ fn fxv_pole_at(x0: i32, y0: i32, w: i32, h: i32) -> vec4<f32> {
     var y = y0;
     if (y < 0) {
         y = -1 - y;
-        x = (x + w / 2) % w;
+        x = imod(x + w / 2, w);
     } else if (y >= h) {
         y = 2 * h - 1 - y;
-        x = (x + w / 2) % w;
+        x = imod(x + w / 2, w);
     }
     return textureLoad(src, vec2<i32>(x, clamp(y, 0, h - 1)), 0);
 }
@@ -413,7 +413,7 @@ fn fxv_wrap_pad(@builtin(global_invocation_id) gid: vec3<u32>) {
         sy = 2 * s.y - 1 - sy;
         sx += s.x / 2;
     }
-    textureStore(out, p, textureLoad(src, vec2<i32>(((sx % s.x) + s.x) % s.x, clamp(sy, 0, s.y - 1)), 0));
+    textureStore(out, p, textureLoad(src, vec2<i32>(imod(sx, s.x), clamp(sy, 0, s.y - 1)), 0));
 }
 
 // Per-pixel steps. u[0].x = 0: VR Glow's highlights (u[0].y tint; f[0] = (threshold,

crates/gpu/src/shaders/fx_warp.wgsl

@@ -237,7 +237,7 @@ fn fxw_pageturn(@builtin(global_invocation_id) gid: vec3<u32>) {
         cand[0] = vec4<f32>(FXW_PI * r - d, 1.0, 0.9, 1.0);
         cand[3] = vec4<f32>(d, 0.0, 1.0, 1.0);
     } else {
-        let a = asin(clamp(d / r, -1.0, 1.0));
+        let a = asin_p(clamp(d / r, -1.0, 1.0));
         cand[1] = vec4<f32>(FXW_PI * r - r * a, 1.0, 0.6 + 0.4 * cos(a), 1.0);
         let psi = a / (0.5 * FXW_PI);
         cand[2] = vec4<f32>(r * a, 0.0, 1.0 - 0.4 * psi * (0.5 + 0.5 * ldot), 1.0);

crates/gpu/src/shaders/sky.wgsl

@@ -21,7 +21,7 @@ fn adv_sky(@builtin(global_invocation_id) gid: vec3<u32>) {
     // shade::equirect_uv.
     let pi = 3.141592653589793;
     var u = 0.5 + (atan2(dir.x, dir.z) + P.f[2].w) / (2.0 * pi);
-    let v = acos(clamp(-dir.y, -1.0, 1.0)) / pi;
+    let v = acos_p(clamp(-dir.y, -1.0, 1.0)) / pi;
     u = u - floor(u);
     let iw = P.f[4].x;
     let ih = P.f[4].y;

crates/gpu/src/tests.rs

@@ -26,8 +26,31 @@ use effectcraft_time::{FrameRate, Tick};
 
 use crate::Gpu;
 
+/// 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
+/// thread waits on it for more than a few seconds ("Could not lock adapter context"), and
+/// concurrent devices on llvmpipe can crash. Every test helper that hands out a device calls this
+/// first; calling it again on the same thread is a no-op.
+pub(crate) fn hold_gpu_lock() {
+    use std::cell::RefCell;
+    use std::sync::{Mutex, MutexGuard, PoisonError};
+    static GPU_LOCK: Mutex<()> = Mutex::new(());
+    thread_local! {
+        static HELD: RefCell<Option<MutexGuard<'static, ()>>> = const { RefCell::new(None) };
+    }
+    HELD.with(|h| {
+        let mut h = h.borrow_mut();
+        if h.is_none() {
+            *h = Some(GPU_LOCK.lock().unwrap_or_else(PoisonError::into_inner));
+        }
+    });
+}
+
+/// The shared test device (the calling test holds [`hold_gpu_lock`]).
 pub(crate) fn gpu() -> Option<&'static Gpu> {
     static G: OnceLock<Option<Gpu>> = OnceLock::new();
+    hold_gpu_lock();
     G.get_or_init(|| {
         let g = Gpu::headless();
         if g.is_none() {
@@ -38,6 +61,20 @@ pub(crate) fn gpu() -> Option<&'static Gpu> {
     .as_ref()
 }
 
+/// The test device is on wgpu's OpenGL backend (Mesa's llvmpipe on the FreeBSD CI job).
+pub(crate) fn on_gl() -> bool {
+    gpu().is_some_and(|g| g.ctx.name.ends_with("(Gl)"))
+}
+
+/// On GL, the fraction of pixels allowed past the tolerance: pixels on a decision boundary (a
+/// hex cell's edge, a threshold, a corner pin's coverage edge, a median's rank tie) may land on
+/// the other side, since GL drivers may divide through a reciprocal and GLSL's `fma` need not be
+/// fused (so `fxs_qdiv`-style corrections don't apply). Measured on llvmpipe: at most 0.53 % (CC
+/// HexTile at half resolution, many cell edges in a small frame). A wrong wrap (signed `%`) showed
+/// in 3–70 % of pixels and still fails; a precision loss may show in fewer (Mesa's `asin`: 0.3–3 %),
+/// so the shaders avoid imprecise builtins themselves (`asin_p`) rather than lean on this.
+pub(crate) const GL_BOUNDARY_FLIPS: f64 = 0.01;
+
 /// Procedural footage: smooth colour gradients, fine noise and an alpha with soft, opaque and
 /// fully transparent regions (different per item).
 pub(crate) struct Pattern;
@@ -205,6 +242,7 @@ pub(crate) fn check(label: &str, d: Option<Diff>, allow: f64) {
     let frac = d.over as f64 / d.total as f64;
     // Quantised depths: rounding-boundary flips (see the module docs).
     let (allow, cap) = if d.quantized && allow == 0.0 { (0.003, 4.0 / 255.0 + 1e-6) } else { (allow, f32::INFINITY) };
+    let (allow, cap) = if on_gl() { (allow.max(GL_BOUNDARY_FLIPS), f32::INFINITY) } else { (allow, cap) };
     assert!(frac <= allow && d.max <= cap, "{label}: {} of {} pixels over tolerance (max {:.6}); worst {:?}", d.over, d.total, d.max, d.worst);
     if d.over > 0 {
         eprintln!("{label}: {} of {} pixels over tolerance (max {:.6})", d.over, d.total, d.max);
@@ -403,7 +441,15 @@ pub(crate) fn effect_direct(id: &str, vals: &[(&str, Value)], adjustment: bool)
         let out = effectcraft_render::Accelerator::effects(g, &[effectcraft_render::FxStep { spec, ctx: ctx() }], &buf, None).expect("the GPU runs the effect");
         assert_eq!((out.offset, out.scale), (cpu.offset, cpu.scale), "{id}: geometry");
         let d = diff(&cpu.img, &out.img, 1e-3);
-        assert!(d.over == 0, "{id} {vals:?} adj {adjustment} scale {scale}: {} of {} pixels over 1e-3 (max {}); worst {:?}", d.over, d.total, d.max, d.worst);
+        let allow = if on_gl() { (d.total as f64 * GL_BOUNDARY_FLIPS) as usize } else { 0 };
+        assert!(
+            d.over <= allow,
+            "{id} {vals:?} adj {adjustment} scale {scale}: {} of {} pixels over 1e-3 (max {}); worst {:?}",
+            d.over,
+            d.total,
+            d.max,
+            d.worst
+        );
     }
 }
 
@@ -604,3 +650,17 @@ fn wgsl_float_literals_fit_f32() {
     }
     assert!(checked > 0);
 }
+
+/// A device created with WebGL2's limits (what egui-wgpu asks for on GL, as the desktop app on
+/// FreeBSD gets): the compositor declines it instead of building pipelines that fail validation
+/// (wgpu's default error handler panics here; release builds would only log it).
+#[test]
+fn devices_with_webgl2_limits_are_declined() {
+    hold_gpu_lock();
+    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle());
+    let Ok(adapter) = pollster::block_on(instance.request_adapter(&Default::default())) else { return };
+    let desc = wgpu::DeviceDescriptor { required_limits: wgpu::Limits::downlevel_webgl2_defaults(), ..Default::default() };
+    let Ok((device, queue)) = pollster::block_on(adapter.request_device(&desc)) else { return };
+    let err = crate::context::GpuContext::new(&adapter, device, queue).err().expect("declined");
+    assert!(err.contains("max_storage_buffers_per_shader_stage"), "{err}");
+}

crates/gpu/src/tests_adv3d.rs

@@ -14,7 +14,9 @@ use effectcraft_time::{FrameRate, Tick};
 
 use crate::Gpu;
 
+/// A device of the test's own (the test holds [`crate::tests::hold_gpu_lock`]).
 fn gpu() -> Option<Gpu> {
+    crate::tests::hold_gpu_lock();
     let g = Gpu::headless();
     if g.is_none() {
         eprintln!("effectcraft-gpu adv3d tests: no GPU adapter, skipping");
@@ -527,3 +529,24 @@ fn gpu_compositor_draws_split_advanced_3d_runs_and_skies() {
         assert!(share <= 0.015, "case {case}: {:.3} % of pixels differ from the CPU", share * 100.0);
     }
 }
+
+/// An adapter whose render targets don't qualify (`R32Float` was not renderable on GL's llvmpipe;
+/// the pipelines failed validation, which release builds only log): Advanced 3D renders on the
+/// CPU, the GPU compositor's frame matches the CPU's, and no render pipeline is built (a
+/// validation error panics under test).
+#[test]
+fn adapters_that_cannot_rasterise_render_advanced_3d_on_the_cpu() {
+    crate::tests::hold_gpu_lock();
+    let Some(mut ctx) = crate::context::GpuContext::headless() else { return };
+    ctx.adv3d_raster = false;
+    let g = Gpu::from_context(ctx);
+    let (p, cid) = scene_project();
+    let s = scene(&p, cid);
+    assert!(g.raster_3d(&s).is_none(), "no GPU raster");
+    let cpu = Renderer::new(&p, &NoFootage, RenderOpts::default()).comp_frame(cid, Tick::ZERO);
+    let mut r = Renderer::new(&p, &NoFootage, RenderOpts { backend: effectcraft_render::Backend::Gpu, ..Default::default() });
+    r.accel = Some(&g);
+    let gimg = r.comp_frame(cid, Tick::ZERO);
+    let worst = cpu.data.iter().zip(&gimg.data).flat_map(|(a, b)| (0..4).map(move |k| (a[k] - b[k]).abs())).fold(0.0f32, f32::max);
+    assert!(worst <= 2.0 / 255.0, "GPU compositor with CPU Advanced 3D vs CPU: worst {worst}");
+}

crates/gpu/src/tests_deferred.rs

@@ -19,6 +19,8 @@ use crate::tests::{Pattern, Scene, diff, n, set, tolerance, v3};
 pub(crate) fn deferred_gpu() -> Option<(&'static Gpu, std::sync::MutexGuard<'static, ()>)> {
     static G: OnceLock<Option<Gpu>> = OnceLock::new();
     static LOCK: std::sync::Mutex<()> = std::sync::Mutex::new(());
+    // Also no other GPU test while it runs (see `hold_gpu_lock`).
+    crate::tests::hold_gpu_lock();
     let g = G.get_or_init(|| pollster::block_on(Gpu::request_deferred()).map_err(|e| eprintln!("deferred gpu tests skipped: {e}")).ok()).as_ref()?;
     Some((g, LOCK.lock().unwrap_or_else(|e| e.into_inner())))
 }

crates/ui-egui/tests/ui_gpu.rs

@@ -21,15 +21,22 @@ fn settle(h: &mut Harness<'_, EffectcraftApp>) {
 
 #[test]
 fn viewer_draws_gpu_frames_that_match_the_cpu() {
-    if effectcraft_gpu::Gpu::headless().is_none() {
+    let Some(probe) = effectcraft_gpu::Gpu::headless() else {
         eprintln!("no GPU adapter: skipping");
         return;
-    }
+    };
+    // egui-wgpu creates its GL device with WebGL2's limits, which the compositor declines.
+    let gl = effectcraft_engine::render::Accelerator::name(&probe).ends_with("(Gl)");
+    drop(probe);
     let mut s = Session::default();
     s.execute("file.openDemoProject", json!({})).unwrap();
     s.execute("time.set", json!({"time": 2.0})).unwrap();
     let mut h = Harness::builder().with_size(egui::vec2(1400.0, 900.0)).with_pixels_per_point(1.0).wgpu().build_eframe(|_| EffectcraftApp::new(s));
     settle(&mut h);
+    if gl && h.state().gpu_adapter().is_none() {
+        eprintln!("egui-wgpu's GL device has WebGL2 limits: the viewer stays on the CPU, skipping");
+        return;
+    }
     assert!(h.state().gpu_adapter().is_some(), "GPU compositor on egui-wgpu's device");
     let gpu_px = h.state_mut().viewer_pixels().expect("read back the GPU frame");
     let gpu_shot = h.render().expect("render");

crates/ui-egui/tests/ui_web_job_gpu.rs

@@ -10,6 +10,10 @@ use effectcraft_engine::offload::{Inbox, Offload, Post, WorkerReply, WorkerReque
 use effectcraft_gpu::Gpu;
 use serde_json::json;
 
+/// One device at a time in this test binary: wgpu's OpenGL backend (Mesa llvmpipe on FreeBSD,
+/// Linux without Vulkan) panics or crashes when tests create and drive devices concurrently.
+static GPU_LOCK: Mutex<()> = Mutex::new(());
+
 /// Holds offloaded jobs until the test runs them.
 #[derive(Default)]
 struct Manual(Mutex<Vec<(WorkerRequest, Arc<Inbox>)>>);
@@ -59,6 +63,7 @@ fn page() -> (Session, Arc<Manual>) {
 
 #[test]
 fn render_queue_jobs_render_on_the_worker_gpu() {
+    let _gpu = GPU_LOCK.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
     let Some(g) = Gpu::headless_deferred() else {
         eprintln!("no GPU adapter: skipped");
         return;
@@ -97,6 +102,7 @@ fn render_queue_jobs_render_on_the_worker_gpu() {
 
 #[test]
 fn analyses_render_their_input_on_the_worker_gpu() {
+    let _gpu = GPU_LOCK.lock().unwrap_or_else(std::sync::PoisonError::into_inner);
     let Some(g) = Gpu::headless_deferred() else {
         eprintln!("no GPU adapter: skipped");
         return;