From 6eedf8e2c5347165348e73979a65ea4b6177e138 Mon Sep 17 00:00:00 2001 From: Brandon Thomas Date: Mon, 5 Oct 2026 10:08:21 -0400 Subject: [PATCH 1/2] GPU on wgpu's OpenGL backend: renderable depth target, portable shaders, serialised tests FreeBSD's CI VM (Mesa llvmpipe over GL) failed 22 effectcraft-gpu tests: - Advanced 3D's camera-depth target was R32Float, which GLES can't render to without EXT_color_buffer_float: create_render_pipeline failed validation (release builds only log that and draw nothing). The target is now R32Uint holding the f32 bits (renderable on every backend, lossless), and GpuContext checks the adapter's format features (render attachment, blendable colour) and independent blending up front; when they're missing the render pipelines are never built and Advanced 3D renders on the CPU. - Signed `%` in the shaders: naga emits GLSL `%`, undefined for negative operands, which broke the VR effects' horizontal wrap and Cell Pattern's tiling. A shared `imod` never takes `%` of a negative number. - asin/acos: Mesa's are coarse approximations (~1e-4 rad), visible in every lens/sphere/cylinder mapping. `asin_p`/`acos_p` (Cephes' asinf polynomial) are f32-accurate on every backend. - "Could not lock adapter context": the GL backend has one context lock per adapter with a timeout, so tests driving devices concurrently panic. Every gpu test helper that hands out a device holds one process-wide lock for the calling test; ui_web_job_gpu's two device tests share a lock too. - On GL, up to 1 % of pixels may flip across a decision boundary (division through a reciprocal, unfused fma); the GPU-vs-CPU comparisons allow that there only. Also adds the FreeBSD workflow (from the FreeBSD CI PR) so this branch is checked there. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/freebsd.yml | 48 ++++++++++++++++++ crates/gpu/src/adv3d.rs | 61 +++++++++++++++-------- crates/gpu/src/context.rs | 8 +++ crates/gpu/src/shaders/adv3d.wgsl | 6 ++- crates/gpu/src/shaders/advanced3d.wgsl | 25 ++++++---- crates/gpu/src/shaders/classic3d.wgsl | 2 +- crates/gpu/src/shaders/common.wgsl | 34 +++++++++++++ crates/gpu/src/shaders/fx_distort.wgsl | 2 +- crates/gpu/src/shaders/fx_gen2.wgsl | 2 +- crates/gpu/src/shaders/fx_generate.wgsl | 6 +-- crates/gpu/src/shaders/fx_pixel2.wgsl | 2 +- crates/gpu/src/shaders/fx_stylize.wgsl | 4 +- crates/gpu/src/shaders/fx_transition.wgsl | 6 +-- crates/gpu/src/shaders/fx_vr.wgsl | 18 +++---- crates/gpu/src/shaders/fx_warp.wgsl | 2 +- crates/gpu/src/shaders/sky.wgsl | 2 +- crates/gpu/src/tests.rs | 48 +++++++++++++++++- crates/gpu/src/tests_adv3d.rs | 23 +++++++++ crates/gpu/src/tests_deferred.rs | 2 + crates/ui-egui/tests/ui_web_job_gpu.rs | 6 +++ 20 files changed, 250 insertions(+), 57 deletions(-) create mode 100644 .github/workflows/freebsd.yml diff --git a/.github/workflows/freebsd.yml b/.github/workflows/freebsd.yml new file mode 100644 index 000000000..9dc8fe0b9 --- /dev/null +++ b/.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 diff --git a/crates/gpu/src/adv3d.rs b/crates/gpu/src/adv3d.rs index ebcfcca7c..f64f5e2d4 100644 --- a/crates/gpu/src/adv3d.rs +++ b/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 { + 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 = 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), diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs index b494e680a..95478949c 100644 --- a/crates/gpu/src/context.rs +++ b/crates/gpu/src/context.rs @@ -259,6 +259,9 @@ pub struct GpuContext { zeros: Mutex>, /// Advanced 3D render pipelines (built on first use). pub(crate) adv3d: std::sync::OnceLock, + /// 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, @@ -303,6 +306,10 @@ impl GpuContext { } } 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 +473,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, diff --git a/crates/gpu/src/shaders/adv3d.wgsl b/crates/gpu/src/shaders/adv3d.wgsl index 1a2d3c81a..6d63d88d6 100644 --- a/crates/gpu/src/shaders/adv3d.wgsl +++ b/crates/gpu/src/shaders/adv3d.wgsl @@ -25,7 +25,9 @@ struct Post { @group(0) @binding(0) var P: Post; @group(0) @binding(1) var color_in: texture_2d; -@group(0) @binding(2) var z_in: texture_2d; +// 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; @group(0) @binding(3) var acc: array>; @group(0) @binding(4) var depth: array; @group(0) @binding(5) var src: array>; @@ -55,7 +57,7 @@ fn resolve(@builtin(global_invocation_id) id: vec3) { for (var i = 0u; i < k; i = i + 1u) { let p = vec2(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(textureLoad(z_in, p, 0).x); if (z >= 0.0) { dz = dz + z; dn = dn + 1.0; diff --git a/crates/gpu/src/shaders/advanced3d.wgsl b/crates/gpu/src/shaders/advanced3d.wgsl index f7c94ff54..e1dbc8f58 100644 --- a/crates/gpu/src/shaders/advanced3d.wgsl +++ b/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, x: i32, y: i32) -> vec4 { @@ -282,7 +284,8 @@ fn attenuation(kind: u32, falloff: u32, radius: f32, falloff_distance: f32, cos_ struct FOut { @location(0) color: vec4, - @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(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(c * alpha, alpha); diff --git a/crates/gpu/src/shaders/classic3d.wgsl b/crates/gpu/src/shaders/classic3d.wgsl index 4e8d21397..d2cad8324 100644 --- a/crates/gpu/src/shaders/classic3d.wgsl +++ b/crates/gpu/src/shaders/classic3d.wgsl @@ -186,7 +186,7 @@ fn attenuation(l: u32, p: vec3) -> 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) { diff --git a/crates/gpu/src/shaders/common.wgsl b/crates/gpu/src/shaders/common.wgsl index 2949c0dc0..00ac946d5 100644 --- a/crates/gpu/src/shaders/common.wgsl +++ b/crates/gpu/src/shaders/common.wgsl @@ -16,6 +16,40 @@ struct Params { @group(0) @binding(2) var aux: texture_2d; @group(0) @binding(4) var data: array; +// `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, x: i32, y: i32) -> vec4 { diff --git a/crates/gpu/src/shaders/fx_distort.wgsl b/crates/gpu/src/shaders/fx_distort.wgsl index a5493621b..5fbc76b4b 100644 --- a/crates/gpu/src/shaders/fx_distort.wgsl +++ b/crates/gpu/src/shaders/fx_distort.wgsl @@ -215,7 +215,7 @@ fn dst_cclens(@builtin(global_invocation_id) gid: vec3) { } 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(c.x + dx * k, c.y + dy * k); } diff --git a/crates/gpu/src/shaders/fx_gen2.wgsl b/crates/gpu/src/shaders/fx_gen2.wgsl index fe3ed11b0..b7458f87d 100644 --- a/crates/gpu/src/shaders/fx_gen2.wgsl +++ b/crates/gpu/src/shaders/fx_gen2.wgsl @@ -653,7 +653,7 @@ fn fg2_threads(@builtin(global_invocation_id) gid: vec3) { 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; diff --git a/crates/gpu/src/shaders/fx_generate.wgsl b/crates/gpu/src/shaders/fx_generate.wgsl index e5e2f4111..60fc2f9d1 100644 --- a/crates/gpu/src/shaders/fx_generate.wgsl +++ b/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 { 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(tx); let hy = bitcast(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); } diff --git a/crates/gpu/src/shaders/fx_pixel2.wgsl b/crates/gpu/src/shaders/fx_pixel2.wgsl index 5916ed865..c2ff0c870 100644 --- a/crates/gpu/src/shaders/fx_pixel2.wgsl +++ b/crates/gpu/src/shaders/fx_pixel2.wgsl @@ -661,7 +661,7 @@ fn fp2_hextile(@builtin(global_invocation_id) gid: vec3) { 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; } diff --git a/crates/gpu/src/shaders/fx_stylize.wgsl b/crates/gpu/src/shaders/fx_stylize.wgsl index d39f0f7a4..68b693cf0 100644 --- a/crates/gpu/src/shaders/fx_stylize.wgsl +++ b/crates/gpu/src/shaders/fx_stylize.wgsl @@ -44,7 +44,7 @@ fn fxs_qdiv(x: vec3, a: f32) -> vec3 { } 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) { 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(c.x + dx * k, c.y + dy * k); } } diff --git a/crates/gpu/src/shaders/fx_transition.wgsl b/crates/gpu/src/shaders/fx_transition.wgsl index 2b9a0b9f8..66278a681 100644 --- a/crates/gpu/src/shaders/fx_transition.wgsl +++ b/crates/gpu/src/shaders/fx_transition.wgsl @@ -681,7 +681,7 @@ fn ftr_cylinder(@builtin(global_invocation_id) gid: vec3) { 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(s, 0.0, cz)); @@ -695,7 +695,7 @@ fn ftr_sphere_tex(q: vec3) -> vec4 { 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) -> vec2 { } return vec2((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(ftr_lon_u(r.x, r.z), 0.5 + lat / 3.141592653589793); } diff --git a/crates/gpu/src/shaders/fx_vr.wgsl b/crates/gpu/src/shaders/fx_vr.wgsl index 3a1d06d06..a2bb364c8 100644 --- a/crates/gpu/src/shaders/fx_vr.wgsl +++ b/crates/gpu/src/shaders/fx_vr.wgsl @@ -22,7 +22,7 @@ fn fxv_equi_dir(x: f32, y: f32, w: f32, h: f32) -> vec3 { // vr::equi_px. fn fxv_equi_px(d: vec3, w: f32, h: f32) -> vec2 { 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((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 { yi = 2 * h - 1 - yi; xi += w / 2; } - let xm = ((xi % w) + w) % w; + let xm = imod(xi, w); return textureLoad(src, vec2(xm, clamp(yi, 0, h - 1)), 0); } @@ -233,7 +233,7 @@ fn fxv_proj_sample(kind: u32, fov: f32, d: vec3) -> vec4 { 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(0.0); } @@ -250,7 +250,7 @@ fn fxv_proj_sample(kind: u32, fov: f32, d: vec3) -> vec4 { } 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(0.0); } @@ -346,12 +346,12 @@ fn fxv_box_h(@builtin(global_invocation_id) gid: vec3) { let norm = 1.0 / f32(2 * r + 1); var acc = vec4(0.0); for (var i = b0 - r; i <= b0 + r; i++) { - acc += textureLoad(src, vec2(((i % w) + w) % w, y), 0); + acc += textureLoad(src, vec2(imod(i, w), y), 0); } for (var x = b0; x < b1; x++) { textureStore(out, vec2(x, y), acc * norm); let add = textureLoad(src, vec2((x + r + 1) % w, y), 0); - let sub = textureLoad(src, vec2(((x - r) % w + w) % w, y), 0); + let sub = textureLoad(src, vec2(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 { 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(x, clamp(y, 0, h - 1)), 0); } @@ -413,7 +413,7 @@ fn fxv_wrap_pad(@builtin(global_invocation_id) gid: vec3) { sy = 2 * s.y - 1 - sy; sx += s.x / 2; } - textureStore(out, p, textureLoad(src, vec2(((sx % s.x) + s.x) % s.x, clamp(sy, 0, s.y - 1)), 0)); + textureStore(out, p, textureLoad(src, vec2(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, diff --git a/crates/gpu/src/shaders/fx_warp.wgsl b/crates/gpu/src/shaders/fx_warp.wgsl index 7cfd86489..9d7ccefa1 100644 --- a/crates/gpu/src/shaders/fx_warp.wgsl +++ b/crates/gpu/src/shaders/fx_warp.wgsl @@ -237,7 +237,7 @@ fn fxw_pageturn(@builtin(global_invocation_id) gid: vec3) { cand[0] = vec4(FXW_PI * r - d, 1.0, 0.9, 1.0); cand[3] = vec4(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(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(r * a, 0.0, 1.0 - 0.4 * psi * (0.5 + 0.5 * ldot), 1.0); diff --git a/crates/gpu/src/shaders/sky.wgsl b/crates/gpu/src/shaders/sky.wgsl index 4ee6b4603..f37565cc9 100644 --- a/crates/gpu/src/shaders/sky.wgsl +++ b/crates/gpu/src/shaders/sky.wgsl @@ -21,7 +21,7 @@ fn adv_sky(@builtin(global_invocation_id) gid: vec3) { // 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; diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs index 4b848da4a..106d66827 100644 --- a/crates/gpu/src/tests.rs +++ b/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>> = 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> = 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, 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 + ); } } diff --git a/crates/gpu/src/tests_adv3d.rs b/crates/gpu/src/tests_adv3d.rs index 030a2ec67..ee99c5c08 100644 --- a/crates/gpu/src/tests_adv3d.rs +++ b/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 { + 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}"); +} diff --git a/crates/gpu/src/tests_deferred.rs b/crates/gpu/src/tests_deferred.rs index 3c50b9c9f..5f80affe4 100644 --- a/crates/gpu/src/tests_deferred.rs +++ b/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> = 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()))) } diff --git a/crates/ui-egui/tests/ui_web_job_gpu.rs b/crates/ui-egui/tests/ui_web_job_gpu.rs index e24829a86..05630dbe5 100644 --- a/crates/ui-egui/tests/ui_web_job_gpu.rs +++ b/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)>>); @@ -59,6 +63,7 @@ fn page() -> (Session, Arc) { #[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; From 3767332efd349b561d8a2d7b50e6dafe91fdd19e Mon Sep 17 00:00:00 2001 From: Brandon Thomas Date: Mon, 5 Oct 2026 10:52:04 -0400 Subject: [PATCH 2/2] GPU compositor declines devices whose limits its pipelines exceed egui-wgpu creates its device on GL with WebGL2's limits (no storage buffers), so the desktop app on FreeBSD built the compositor's pipelines on it: validation errors (only logged in release builds) and wrong frames. GpuContext::new now checks the limits its pipelines bind (storage buffers and textures per stage, bind groups, workgroup sizes) and returns Err, which keeps the app on the CPU. ui_gpu skips its GPU-vs-CPU comparison when that happens on GL. Co-Authored-By: Claude Opus 5.5 --- crates/gpu/src/context.rs | 17 ++++++++++++++++- crates/gpu/src/tests.rs | 14 ++++++++++++++ crates/ui-egui/tests/ui_gpu.rs | 11 +++++++++-- 3 files changed, 39 insertions(+), 3 deletions(-) diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs index 95478949c..d252dcf68 100644 --- a/crates/gpu/src/context.rs +++ b/crates/gpu/src/context.rs @@ -294,12 +294,27 @@ 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 { 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)); diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs index 106d66827..3c6fdfa35 100644 --- a/crates/gpu/src/tests.rs +++ b/crates/gpu/src/tests.rs @@ -650,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}"); +} diff --git a/crates/ui-egui/tests/ui_gpu.rs b/crates/ui-egui/tests/ui_gpu.rs index d97829e9c..a6966f83d 100644 --- a/crates/ui-egui/tests/ui_gpu.rs +++ b/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");