From 480fd8d517de40620ea576b41b5132e822413eaa Mon Sep 17 00:00:00 2001 From: bherbruck Date: Wed, 14 Oct 2026 10:40:24 -0400 Subject: [PATCH] Fix the startup crash on D3D12 with FXC: CPU fallback for the tessellation pipeline, FXC-friendly shader (#1098) (#1099) * GPU tessellation: a pipeline the device cannot build falls back to CPU meshing instead of crashing (#1098) Build the tessellation compute pipeline inside error scopes; TessGpu::new returns an error, and the viewport logs it and meshes bodies on the CPU. A Windows CI job builds the pipeline with D3D12's FXC on the WARP adapter. Co-Authored-By: Claude Opus 5.5 * Tessellation shader: basis functions in one-dimensional local arrays, so D3D12's FXC compiles it (#1098) FXC has to unroll loops that write an array inside a struct (Basis.n0/n1) or a 2-D array (ndu) by a counter, and cannot with run-time trip counts (X3511). A2.2 basis functions in a local 1-D array returned whole; derivatives from the degree p-1 functions; the pole step offsets computed rather than indexed. Co-Authored-By: Claude Opus 5.5 --------- Co-authored-by: Claude Opus 5.5 --- .github/workflows/ci.yml | 26 ++++++++- Cargo.lock | 1 + Cargo.toml | 1 + crates/gpu/Cargo.toml | 2 + crates/gpu/src/lib.rs | 101 ++++++++++++++++++++++++--------- crates/gpu/src/tess.wgsl | 96 ++++++++++++++++++------------- crates/gpu/src/tests.rs | 79 ++++++++++++++++++++++++++ crates/ui-egui/src/gpu.rs | 63 ++++++++++++++------ crates/ui-egui/src/viewport.rs | 46 ++++++++++++++- 9 files changed, 329 insertions(+), 86 deletions(-) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 3bcaed40..e1730e26 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -286,9 +286,31 @@ jobs: name: oracle-report path: docs/oracle.md + # D3D12's older shader compiler (FXC, used where dxcompiler.dll is absent) is stricter than + # DXC, Vulkan and Metal: a shader it cannot compile once crashed the app at startup. Build the + # GPU tessellation pipeline with FXC on the runner's WARP software adapter and compare its + # results with the CPU. + windows-fxc: + name: Windows D3D12 FXC shaders + needs: changes + if: needs.changes.outputs.native == 'true' + runs-on: windows-latest + timeout-minutes: 45 + env: + WGPU_BACKEND: dx12 + WGPU_DX12_COMPILER: fxc + SOLIDCRAFT_REQUIRE_GPU: "1" + steps: + - uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 + - uses: dtolnay/rust-toolchain@stable + - uses: Swatinem/rust-cache@6323deb102c322ba6fcbdcafc7e3dddab59af2b6 # v2.9.2 + with: + key: windows-fxc + - run: cargo test -p solidcraft-gpu --profile ci --locked + # Preserve the existing check name for branch protection, including docs-only PRs. ci: - needs: [changes, checks, clippy, build-tests, tests, doc-tests, oracle] + needs: [changes, checks, clippy, build-tests, tests, doc-tests, oracle, windows-fxc] if: always() runs-on: ubuntu-24.04 steps: @@ -302,6 +324,7 @@ jobs: TESTS: ${{ needs.tests.result }} DOCS: ${{ needs.doc-tests.result }} ORACLE: ${{ needs.oracle.result }} + WINDOWS_FXC: ${{ needs.windows-fxc.result }} run: | test "$CHANGES" = success test "$CHECKS" = success @@ -311,6 +334,7 @@ jobs: test "$TESTS" = success test "$DOCS" = success test "$ORACLE" = success + test "$WINDOWS_FXC" = success fi wasm: diff --git a/Cargo.lock b/Cargo.lock index 619a70ab..d7c4f0b3 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -3351,6 +3351,7 @@ name = "solidcraft-gpu" version = "0.10.0" dependencies = [ "log", + "naga", "pollster 0.4.0", "solidcraft-geom", "solidcraft-kernel", diff --git a/Cargo.toml b/Cargo.toml index 8faed3d7..cf51a67f 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -61,6 +61,7 @@ web-time = "1" egui-wgpu = "0.36" # the wgpu egui-wgpu uses, for compute (crates/gpu); backends come from the apps' features wgpu = { version = "30", default-features = false, features = ["std", "wgsl"] } +naga = { version = "30", default-features = false } pollster = "0.4" rfd = "0.17" diff --git a/crates/gpu/Cargo.toml b/crates/gpu/Cargo.toml index 1571d746..077909c4 100644 --- a/crates/gpu/Cargo.toml +++ b/crates/gpu/Cargo.toml @@ -21,3 +21,5 @@ wgpu = { workspace = true, features = ["metal", "vulkan", "dx12", "gles"] } [dev-dependencies] solidcraft-kernel = { workspace = true } +# The HLSL the D3D12 backend compiles (tests check it suits FXC). +naga = { workspace = true, features = ["wgsl-in", "hlsl-out"] } diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs index 130c2b35..e069db06 100644 --- a/crates/gpu/src/lib.rs +++ b/crates/gpu/src/lib.rs @@ -17,6 +17,37 @@ use solidcraft_geom::Vec3; use solidcraft_geom::tess::{EvalCurve, EvalSurface, TessPlan}; use wgpu::util::DeviceExt; +/// The tessellation compute shader (WGSL). +pub const SHADER: &str = include_str!("tess.wgsl"); + +/// The error a popped error scope caught. Native devices answer at once; a browser's answer +/// comes later, and is not waited for (no error is assumed). +fn popped(f: impl std::future::Future>) -> Option { + let mut f = std::pin::pin!(f); + match f.as_mut().poll(&mut std::task::Context::from_waker(std::task::Waker::noop())) { + std::task::Poll::Ready(e) => e, + std::task::Poll::Pending => None, + } +} + +/// An adapter and device without a window (native only). The environment may pick the backend +/// and the D3D12 shader compiler (`WGPU_BACKEND`, `WGPU_DX12_COMPILER`). +#[cfg(not(target_arch = "wasm32"))] +fn open_headless() -> Result<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle_from_env()); + let adapter = pollster::block_on( + instance.request_adapter(&wgpu::RequestAdapterOptions { power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default() }), + ) + .map_err(|_| GpuError::NoAdapter)?; + let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + label: Some("solidcraft compute"), + required_limits: wgpu::Limits::downlevel_defaults().using_resolution(adapter.limits()), + ..Default::default() + })) + .map_err(|e| GpuError::Device(e.to_string()))?; + Ok((adapter, device, queue)) +} + /// Most vertices evaluated in one dispatch (bounds the buffers; larger plans go in batches). pub const MAX_BATCH: usize = 1 << 20; const WORKGROUP: u32 = 64; @@ -29,6 +60,8 @@ pub enum GpuError { Device(String), #[error("the GPU evaluation failed: {0}")] Failed(String), + #[error("the GPU tessellation pipeline could not be created: {0}")] + Pipeline(String), #[error("not supported yet: {0}")] Unsupported(String), } @@ -78,11 +111,47 @@ pub struct TessGpu { } impl TessGpu { - /// The pipeline on an existing device (`adapter`: its name, for reports). - pub fn new(adapter: impl Into, device: wgpu::Device, queue: wgpu::Queue) -> Self { + /// The pipeline on an existing device (`adapter`: its name, for reports). An error (never + /// a panic) when the device cannot build it, for example where the D3D12 driver's older shader + /// compiler (FXC) rejects the shader: the caller then builds bodies on the CPU. + pub fn new(adapter: impl Into, device: wgpu::Device, queue: wgpu::Queue) -> Result { + Self::with_shader(adapter, device, queue, SHADER) + } + + /// [`TessGpu::new`] with another shader source: tests use it to make pipeline creation fail + /// as it does on a device whose compiler rejects the real shader. + #[doc(hidden)] + pub fn with_shader(adapter: impl Into, device: wgpu::Device, queue: wgpu::Queue, shader: &str) -> Result { + // Errors here would go to the device's uncaptured-error handler, which panics by default: + // capture them, so a device that cannot run the shader only loses GPU tessellation. + let scopes = [wgpu::ErrorFilter::OutOfMemory, wgpu::ErrorFilter::Internal, wgpu::ErrorFilter::Validation].map(|f| device.push_error_scope(f)); + let (pipeline, layout) = Self::build(&device, shader); + let mut failed = None; + for scope in scopes.into_iter().rev() { + if let Some(e) = popped(scope.pop()) { + failed.get_or_insert(e.to_string()); + } + } + if let Some(e) = failed { + return Err(GpuError::Pipeline(e)); + } + // A viewport records the evaluation into its own frame (see `TessGpu::record`). + #[cfg(target_arch = "wasm32")] + let _ = queue; + Ok(TessGpu { + adapter: adapter.into(), + device, + #[cfg(not(target_arch = "wasm32"))] + queue, + pipeline, + layout, + }) + } + + fn build(device: &wgpu::Device, shader: &str) -> (wgpu::ComputePipeline, wgpu::BindGroupLayout) { let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("solidcraft tessellation"), - source: wgpu::ShaderSource::Wgsl(include_str!("tess.wgsl").into()), + source: wgpu::ShaderSource::Wgsl(shader.into()), }); let entry = |binding: u32, ty: wgpu::BufferBindingType| wgpu::BindGroupLayoutEntry { binding, @@ -114,34 +183,14 @@ impl TessGpu { compilation_options: Default::default(), cache: None, }); - // A viewport records the evaluation into its own frame (see `TessGpu::record`). - #[cfg(target_arch = "wasm32")] - let _ = queue; - TessGpu { - adapter: adapter.into(), - device, - #[cfg(not(target_arch = "wasm32"))] - queue, - pipeline, - layout, - } + (pipeline, layout) } /// A device of its own, without a window (native only). #[cfg(not(target_arch = "wasm32"))] pub fn headless() -> Result { - let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle()); - let adapter = pollster::block_on( - instance.request_adapter(&wgpu::RequestAdapterOptions { power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default() }), - ) - .map_err(|_| GpuError::NoAdapter)?; - let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { - label: Some("solidcraft compute"), - required_limits: wgpu::Limits::downlevel_defaults().using_resolution(adapter.limits()), - ..Default::default() - })) - .map_err(|e| GpuError::Device(e.to_string()))?; - Ok(Self::new(adapter.get_info().name, device, queue)) + let (adapter, device, queue) = open_headless()?; + Self::new(adapter.get_info().name, device, queue) } /// The adapter's name, for reports. diff --git a/crates/gpu/src/tess.wgsl b/crates/gpu/src/tess.wgsl index 66266300..1bb3051d 100644 --- a/crates/gpu/src/tess.wgsl +++ b/crates/gpu/src/tess.wgsl @@ -1,7 +1,7 @@ // Display tessellation: surface points and normals at each vertex's (u, v). // // The GPU twin of solidcraft_geom::tess::EvalSurface::{derivs, eval}: the same surface forms, -// the same algorithms (The NURBS Book A2.1 span search, A2.3 basis functions and first +// the same algorithms (The NURBS Book A2.1 span search, A2.2 basis functions, A2.3 first // derivatives, the rational quotient rule), in f32. const PLANE: u32 = 0u; @@ -60,11 +60,6 @@ struct Derivs { ok: bool, }; -struct Basis { - n0: array, - n1: array, -}; - fn vec3_at(o: u32) -> vec3 { return vec3(data[o], data[o + 1u], data[o + 2u]); } @@ -101,42 +96,58 @@ fn span_of(k0: u32, p: u32, n: u32, t: f32) -> u32 { return lo; } -// Non-zero basis functions of degree p on span `span` and their first derivatives (A2.3, n = 1). -fn basis(k0: u32, span: u32, p: u32, t: f32) -> Basis { - var ndu: array, 8>; - var left: array; - var right: array; - ndu[0][0] = 1.0; +// The arrays below are one-dimensional locals, written by loop counters and returned whole: +// D3D12's older shader compiler (FXC) must unroll a loop that writes an array inside a struct +// (or a two-dimensional one) by a counter, and it cannot unroll loops whose trip count is only +// known at run time, so such a loop fails to compile there ("unable to unroll loop"). + +// The non-zero basis functions of degree p on span `span` (The NURBS Book A2.2). +fn basis_funs(k0: u32, span: u32, p: u32, t: f32) -> array { + var n: array; + n[0] = 1.0; for (var j = 1u; j <= p; j = j + 1u) { - left[j] = t - data[k0 + span + 1u - j]; - right[j] = data[k0 + span + j] - t; var saved = 0.0; for (var r = 0u; r < j; r = r + 1u) { - ndu[j][r] = right[r + 1u] + left[j - r]; + // right[r + 1] and left[j - r] of A2.2 (span ≥ p ≥ j, so no underflow). + let right = data[k0 + span + r + 1u] - t; + let left = t - data[k0 + span + 1u + r - j]; + let den = right + left; var temp = 0.0; - if (ndu[j][r] != 0.0) { - temp = ndu[r][j - 1u] / ndu[j][r]; + if (den != 0.0) { + temp = n[r] / den; } - ndu[r][j] = saved + right[r + 1u] * temp; - saved = left[j - r] * temp; + n[r] = saved + right * temp; + saved = left * temp; } - ndu[j][j] = saved; + n[j] = saved; } - var b: Basis; - for (var j = 0u; j <= p; j = j + 1u) { - b.n0[j] = ndu[j][p]; + return n; +} + +// Their first derivatives, from the degree p - 1 functions on the same span (A2.3, n = 1). +fn basis_ders(k0: u32, span: u32, p: u32, t: f32) -> array { + var d: array; + if (p == 0u) { + return d; } + let lower = basis_funs(k0, span, p - 1u, t); for (var r = 0u; r <= p; r = r + 1u) { - var d = 0.0; - if (r >= 1u && ndu[p][r - 1u] != 0.0) { - d = d + ndu[r - 1u][p - 1u] / ndu[p][r - 1u]; + var x = 0.0; + if (r >= 1u) { + let den = data[k0 + span + r] - data[k0 + span + r - p]; + if (den != 0.0) { + x = x + lower[r - 1u] / den; + } } - if (r < p && ndu[p][r] != 0.0) { - d = d - ndu[r][p - 1u] / ndu[p][r]; + if (r < p) { + let den = data[k0 + span + r + 1u] - data[k0 + span + r + 1u - p]; + if (den != 0.0) { + x = x - lower[r] / den; + } } - b.n1[r] = d * f32(p); + d[r] = x * f32(p); } - return b; + return d; } // Point and derivative of a rational function from its homogeneous value and derivative. @@ -152,17 +163,19 @@ fn nurbs(s: Surface, u: f32, v: f32) -> Derivs { var r: Derivs; let su = span_of(s.knots_u, s.pu, s.nu, u); let sv = span_of(s.knots_v, s.pv, s.nv, v); - let bu = basis(s.knots_u, su, s.pu, u); - let bv = basis(s.knots_v, sv, s.pv, v); + let bu0 = basis_funs(s.knots_u, su, s.pu, u); + let bu1 = basis_ders(s.knots_u, su, s.pu, u); + let bv0 = basis_funs(s.knots_v, sv, s.pv, v); + let bv1 = basis_ders(s.knots_v, sv, s.pv, v); var a = vec4(0.0); var au = vec4(0.0); var av = vec4(0.0); for (var j = 0u; j <= s.pv; j = j + 1u) { for (var i = 0u; i <= s.pu; i = i + 1u) { let c = ctrl_at(s.ctrl + 4u * ((sv + j - s.pv) * s.nu + su + i - s.pu)); - a = a + bu.n0[i] * bv.n0[j] * c; - au = au + bu.n1[i] * bv.n0[j] * c; - av = av + bu.n0[i] * bv.n1[j] * c; + a = a + bu0[i] * bv0[j] * c; + au = au + bu1[i] * bv0[j] * c; + av = av + bu0[i] * bv1[j] * c; } } r.ok = abs(a.w) > 0.0; @@ -189,13 +202,14 @@ fn revolved(s: Surface, u: f32, v: f32) -> Derivs { } // The profile curve: a one-row B-spline. let span = span_of(s.knots_u, s.pu, s.nu, t); - let b = basis(s.knots_u, span, s.pu, t); + let b0 = basis_funs(s.knots_u, span, s.pu, t); + let b1 = basis_ders(s.knots_u, span, s.pu, t); var c = vec4(0.0); var dc = vec4(0.0); for (var i = 0u; i <= s.pu; i = i + 1u) { let h = ctrl_at(s.ctrl + 4u * (span + i - s.pu)); - c = c + b.n0[i] * h; - dc = dc + b.n1[i] * h; + c = c + b0[i] * h; + dc = dc + b1[i] * h; } r.ok = abs(c.w) > 0.0; if (!r.ok) { @@ -263,9 +277,11 @@ fn main(@builtin(global_invocation_id) gid: vec3) { if (!(length(n) > 1e-6 * scale) || !d.ok) { // A degenerate point (a pole): the normal from a step inside. let h = 1e-4 * (1.0 + max(abs(vert.u), abs(vert.v))); - let steps = array, 4>(vec2(h, h), vec2(-h, h), vec2(h, -h), vec2(-h, -h)); + // Steps (h, h), (-h, h), (h, -h), (-h, -h), without an array indexed by k (FXC). for (var k = 0u; k < 4u; k = k + 1u) { - let e = derivs(s, vert.u + steps[k].x, vert.v + steps[k].y); + let du = select(h, -h, (k & 1u) != 0u); + let dv = select(h, -h, (k & 2u) != 0u); + let e = derivs(s, vert.u + du, vert.v + dv); let m = oriented(s, cross(e.du, e.dv)); if (e.ok && length(m) > 0.0) { n = m; diff --git a/crates/gpu/src/tests.rs b/crates/gpu/src/tests.rs index 4254c0c7..4d12aa2b 100644 --- a/crates/gpu/src/tests.rs +++ b/crates/gpu/src/tests.rs @@ -4,12 +4,17 @@ use solidcraft_kernel::{Body, BoolOp}; use super::*; /// The GPU, or `None` with a note when the machine has no adapter (CI runners without a GPU). +/// A device that cannot build the pipeline fails the test: that is a shader the device rejects +/// (as D3D12's FXC once did), not a missing GPU. `SOLIDCRAFT_REQUIRE_GPU=1` (the Windows CI job, +/// on the WARP software adapter) fails without an adapter too. fn gpu() -> Option { match TessGpu::headless() { Ok(g) => { eprintln!("GPU: {}", g.adapter_name()); Some(g) } + Err(e @ GpuError::Pipeline(_)) => panic!("{e}"), + Err(e) if std::env::var_os("SOLIDCRAFT_REQUIRE_GPU").is_some() => panic!("a GPU is required here: {e}"), Err(e) => { eprintln!("skipping the GPU comparison: {e}"); None @@ -17,6 +22,80 @@ fn gpu() -> Option { } } +/// A shader the device rejects when the pipeline is created (its binding 1 is writable, the +/// layout's read-only), standing in for one a device's compiler cannot build. +const REJECTED: &str = " +@group(0) @binding(1) var surfaces: array; +@group(0) @binding(4) var out: array>; +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + surfaces[gid.x] = 1u; + out[gid.x] = vec4(0.0); +} +"; + +/// A pipeline the device cannot build is an error the caller can fall back from (to the CPU +/// mesh), not a panic in the device's uncaptured-error handler; and the device stays usable. +#[test] +fn a_pipeline_the_device_rejects_is_an_error_not_a_crash() { + let (adapter, device, queue) = match open_headless() { + Ok(d) => d, + Err(e) => { + assert!(std::env::var_os("SOLIDCRAFT_REQUIRE_GPU").is_none(), "a GPU is required here: {e}"); + eprintln!("skipping: {e}"); + return; + } + }; + let name = adapter.get_info().name; + match TessGpu::with_shader(name.clone(), device.clone(), queue.clone(), REJECTED) { + Err(GpuError::Pipeline(e)) => eprintln!("rejected as expected: {e}"), + Err(e) => panic!("unexpected error: {e}"), + Ok(_) => panic!("the rejected shader built a pipeline"), + } + // Not wgsl at all: also an error, not a panic. + assert!(matches!(TessGpu::with_shader(name.clone(), device.clone(), queue.clone(), "fn main( {"), Err(GpuError::Pipeline(_)))); + // The same device then builds the real pipeline and evaluates on it. + let gpu = ok(TessGpu::new(name, device, queue)); + let body = ok(solidcraft_kernel::torus(Vec3::ZERO, 6.0, 1.5)); + let plan = ok(body.display_plan(0.01)); + assert_eq!(ok(gpu.evaluate(&plan)).positions.len(), plan.evaluate().positions.len()); +} + +/// The shader as the D3D12 backend translates it (naga's HLSL for shader model 5.1, the FXC +/// target, with wgpu's options) avoids what FXC cannot compile without unrolling a loop: +/// arrays inside structs and two-dimensional arrays written by a loop counter. Runs everywhere; +/// the Windows CI job compiles the real thing with FXC. +#[test] +fn the_hlsl_for_fxc_has_no_arrays_fxc_must_unroll_loops_for() { + use naga::back::hlsl; + let module = ok(naga::front::wgsl::parse_str(SHADER)); + let info = ok(naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::default()).validate(&module)); + let options = hlsl::Options { + shader_model: hlsl::ShaderModel::V5_1, + force_loop_bounding: true, + restrict_indexing: true, + fake_missing_bindings: true, + ..Default::default() + }; + let pipeline = hlsl::PipelineOptions::default(); + let mut out = String::new(); + ok(hlsl::Writer::new(&mut out, &options, &pipeline).write(&module, &info, None)); + assert!(out.contains("void main("), "{out}"); + // Struct members that are arrays: `float n0_[8];` inside a `struct ... { };`. + let mut in_struct = false; + for line in out.lines() { + let l = line.trim(); + if l.starts_with("struct ") && !l.ends_with(';') { + in_struct = true; + } else if l.starts_with("};") { + in_struct = false; + } else if in_struct { + assert!(!l.contains('['), "an array inside a struct: {l}"); + } + assert!(!l.contains("]["), "a two-dimensional array: {l}"); + } +} + fn ok(r: std::result::Result) -> T { match r { Ok(v) => v, diff --git a/crates/ui-egui/src/gpu.rs b/crates/ui-egui/src/gpu.rs index 0578ac16..19f350c3 100644 --- a/crates/ui-egui/src/gpu.rs +++ b/crates/ui-egui/src/gpu.rs @@ -736,14 +736,30 @@ impl Batches { } } +/// The tessellation pipeline (`shader`: its source), where display plans are evaluated: on a +/// device that runs compute shaders with the storage buffers the evaluation binds (not WebGL2) +/// and builds the pipeline. Elsewhere `None`, and bodies are built on the CPU: a device whose +/// shader compiler rejects the shader (D3D12's FXC once did) only loses GPU evaluation. +fn tessellation(adapter: &wgpu::Adapter, device: &wgpu::Device, queue: &wgpu::Queue, shader: &str) -> Option { + let compute = adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) + && device.limits().max_storage_buffers_per_shader_stage >= 4; + if !compute { + return None; + } + let info = adapter.get_info(); + match solidcraft_gpu::TessGpu::with_shader(info.name.clone(), device.clone(), queue.clone(), shader) { + Ok(t) => Some(t), + Err(e) => { + log::warn!("bodies are built on the CPU on {} ({:?}): {e}", info.name, info.backend); + None + } + } +} + /// Create the pipelines. Call once with the app's render state and its depth/stencil/MSAA /// settings. Without a stencil the section cap does not hide the inside of cut bodies. pub fn install(rs: &egui_wgpu::RenderState, depth_bits: u8, stencil_bits: u8, samples: u32) -> GpuTarget { - // Display plans are evaluated where the device runs compute shaders with the storage - // buffers the evaluation binds (not WebGL2): elsewhere bodies are built on the CPU. - let compute = rs.adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) - && rs.device.limits().max_storage_buffers_per_shader_stage >= 4; - let tess = compute.then(|| solidcraft_gpu::TessGpu::new(rs.adapter.get_info().name, rs.device.clone(), rs.queue.clone())); + let tess = tessellation(&rs.adapter, &rs.device, &rs.queue, solidcraft_gpu::SHADER); let t = GpuTarget { format: rs.target_format, depth: egui_wgpu::depth_format_from_bits(depth_bits, stencil_bits), samples, tess: tess.is_some() }; let res = Resources::new(&rs.device, t, tess); rs.renderer.write().callback_resources.insert(res); @@ -912,8 +928,29 @@ impl egui_wgpu::CallbackTrait for ViewportCallback { /// Render `scene` offscreen on a device of its own, through the viewport's real resources and /// frame callback (`prepare`, then `paint`), and read the pixels back: straight RGBA, `size` × -/// `size`. `None` without a GPU adapter. With `tess`, planned bodies are evaluated on the GPU. +/// `size`. `None` without a GPU adapter. With `tess` (the tessellation shader), planned bodies +/// are evaluated on the GPU where the device builds that pipeline, as [`install`] sets it up. /// `highlight` is drawn as the viewport's highlight scene (hover and selection). +#[cfg(all(test, not(target_arch = "wasm32")))] +fn offscreen_device() -> Option<(wgpu::Adapter, wgpu::Device, wgpu::Queue)> { + let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle_from_env()); + let adapter = pollster::block_on(instance.request_adapter(&Default::default())).ok()?; + let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { + required_limits: wgpu::Limits::downlevel_defaults().using_resolution(adapter.limits()), + ..Default::default() + })) + .ok()?; + Some((adapter, device, queue)) +} + +/// Whether [`install`] would evaluate plans on the GPU with `shader`, on an offscreen device +/// (`None` without a GPU adapter). +#[cfg(all(test, not(target_arch = "wasm32")))] +pub(crate) fn offscreen_tessellates(shader: &str) -> Option { + let (adapter, device, queue) = offscreen_device()?; + Some(tessellation(&adapter, &device, &queue, shader).is_some()) +} + #[cfg(all(test, not(target_arch = "wasm32")))] pub(crate) fn render_offscreen( scene: GpuScene, @@ -922,18 +959,12 @@ pub(crate) fn render_offscreen( back: [f32; 3], eye: [f32; 4], size: u32, - tess: bool, + tess: Option<&str>, ) -> Option> { - let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle()); - let adapter = pollster::block_on(instance.request_adapter(&Default::default())).ok()?; - let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor { - required_limits: wgpu::Limits::downlevel_defaults().using_resolution(adapter.limits()), - ..Default::default() - })) - .ok()?; + let (adapter, device, queue) = offscreen_device()?; let (format, depth) = (wgpu::TextureFormat::Rgba8Unorm, wgpu::TextureFormat::Depth24PlusStencil8); - let target = GpuTarget { format, depth: Some(depth), samples: 1, tess }; - let compute = tess.then(|| solidcraft_gpu::TessGpu::new(adapter.get_info().name, device.clone(), queue.clone())); + let compute = tess.and_then(|shader| tessellation(&adapter, &device, &queue, shader)); + let target = GpuTarget { format, depth: Some(depth), samples: 1, tess: compute.is_some() }; let mut resources = egui_wgpu::CallbackResources::default(); resources.insert(Resources::new(&device, target, compute)); let none = || Arc::new(Mutex::new(None)); diff --git a/crates/ui-egui/src/viewport.rs b/crates/ui-egui/src/viewport.rs index f5ca73ba..707d0ea1 100644 --- a/crates/ui-egui/src/viewport.rs +++ b/crates/ui-egui/src/viewport.rs @@ -3316,8 +3316,8 @@ mod tests { let proj = projection(&app, rect); let e = proj.cam.eye().to_f32(); let eye = [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }]; - let planned = scene.planned.len(); - (planned, scene.tris.len(), crate::gpu::render_offscreen(scene, None, proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, tess)) + let (planned, shader) = (scene.planned.len(), tess.then_some(solidcraft_gpu::SHADER)); + (planned, scene.tris.len(), crate::gpu::render_offscreen(scene, None, proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, shader)) }; let (planned, gpu_tris, gpu) = render(true); let (none, cpu_tris, cpu) = render(false); @@ -3335,6 +3335,45 @@ mod tests { assert!(differ * 100 < drawn, "{differ} of {drawn} model pixels differ"); } + /// A device that cannot build the tessellation pipeline (D3D12's FXC rejected the shader + /// at startup and the app panicked) still shows the model: the failure is caught, logged, + /// and bodies are built on the CPU and drawn. A shader the device rejects stands in for FXC. + /// Skipped without a GPU adapter. + #[cfg(not(target_arch = "wasm32"))] + #[test] + fn bodies_are_built_on_the_cpu_when_the_tessellation_pipeline_fails() { + const REJECTED: &str = " + @group(0) @binding(1) var surfaces: array; + @compute @workgroup_size(64) + fn main(@builtin(global_invocation_id) gid: vec3) { surfaces[gid.x] = 1u; } + "; + let Some(tess) = crate::gpu::offscreen_tessellates(REJECTED) else { + eprintln!("skipping: no GPU adapter"); + return; + }; + assert!(!tess, "the rejected shader built a pipeline"); + let mut s = solidcraft_engine::Session::default(); + s.run_script(&solidcraft_engine::sample::script()).unwrap(); + let mut app = SolveApp::new(s, crate::Services::default()); + app.viewport.no_pixels = true; + app.fit_view(); + app.viewport.gpu = Some(GpuTarget { format: egui_wgpu::wgpu::TextureFormat::Rgba8Unorm, depth: None, samples: 1, tess }); + const SIZE: u32 = 256; + let rect = Rect::from_min_size(Pos2::ZERO, vec2(SIZE as f32, SIZE as f32)); + let scene = build_scene(&app); + assert!(scene.planned.is_empty(), "no body waits for the GPU evaluation"); + assert!(!scene.tris.is_empty(), "the bodies are meshed on the CPU"); + let proj = projection(&app, rect); + let e = proj.cam.eye().to_f32(); + let eye = [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }]; + let Some(img) = crate::gpu::render_offscreen(scene, None, proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, Some(REJECTED)) else { + return; + }; + let n = (SIZE * SIZE) as usize; + let drawn = img.chunks(4).filter(|p| p[..3] != [255, 255, 255]).count(); + assert!(drawn > n / 20, "the model is drawn: {drawn} pixels"); + } + /// A hovered face lights solidly over a body drawn from its display plan on the GPU: the /// highlight uses the plan's triangles, not the reference mesh's other triangulation of the /// same face, which fought the model at equal depth and shimmered in wedges. Every face of a @@ -3362,7 +3401,8 @@ mod tests { let proj = projection(&app, rect); let e = proj.cam.eye().to_f32(); let eye = [e[0], e[1], e[2], if proj.cam.fov > 0.0 { 1.0 } else { 0.0 }]; - crate::gpu::render_offscreen(scene, Some(highlight), proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, tess) + let shader = tess.then_some(solidcraft_gpu::SHADER); + crate::gpu::render_offscreen(scene, Some(highlight), proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, shader) }; let lit = |plain: &[u8], hovered: &[u8]| -> Vec { plain.chunks(4).zip(hovered.chunks(4)).map(|(a, b)| a != b).collect() }; let (Some(gpu_plain), Some(cpu_plain)) = (render(true, None), render(false, None)) else {