Compare commits
1
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
069ba1223f |
No files matched your search
@@ -286,31 +286,9 @@ 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, windows-fxc]
|
||||
needs: [changes, checks, clippy, build-tests, tests, doc-tests, oracle]
|
||||
if: always()
|
||||
runs-on: ubuntu-24.04
|
||||
steps:
|
||||
@@ -324,7 +302,6 @@ 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
|
||||
@@ -334,7 +311,6 @@ jobs:
|
||||
test "$TESTS" = success
|
||||
test "$DOCS" = success
|
||||
test "$ORACLE" = success
|
||||
test "$WINDOWS_FXC" = success
|
||||
fi
|
||||
|
||||
wasm:
|
||||
|
||||
Generated
-1
@@ -3351,7 +3351,6 @@ name = "solidcraft-gpu"
|
||||
version = "0.10.1"
|
||||
dependencies = [
|
||||
"log",
|
||||
"naga",
|
||||
"pollster 0.4.0",
|
||||
"solidcraft-geom",
|
||||
"solidcraft-kernel",
|
||||
|
||||
@@ -61,7 +61,6 @@ 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"
|
||||
|
||||
|
||||
@@ -21,5 +21,3 @@ 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"] }
|
||||
+26
-75
@@ -17,37 +17,6 @@ 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<Output = Option<wgpu::Error>>) -> Option<wgpu::Error> {
|
||||
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;
|
||||
@@ -60,8 +29,6 @@ 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),
|
||||
}
|
||||
@@ -111,47 +78,11 @@ pub struct TessGpu {
|
||||
}
|
||||
|
||||
impl TessGpu {
|
||||
/// 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<String>, device: wgpu::Device, queue: wgpu::Queue) -> Result<Self> {
|
||||
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<String>, device: wgpu::Device, queue: wgpu::Queue, shader: &str) -> Result<Self> {
|
||||
// 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) {
|
||||
/// The pipeline on an existing device (`adapter`: its name, for reports).
|
||||
pub fn new(adapter: impl Into<String>, device: wgpu::Device, queue: wgpu::Queue) -> Self {
|
||||
let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
|
||||
label: Some("solidcraft tessellation"),
|
||||
source: wgpu::ShaderSource::Wgsl(shader.into()),
|
||||
source: wgpu::ShaderSource::Wgsl(include_str!("tess.wgsl").into()),
|
||||
});
|
||||
let entry = |binding: u32, ty: wgpu::BufferBindingType| wgpu::BindGroupLayoutEntry {
|
||||
binding,
|
||||
@@ -183,14 +114,34 @@ impl TessGpu {
|
||||
compilation_options: Default::default(),
|
||||
cache: None,
|
||||
});
|
||||
(pipeline, layout)
|
||||
// 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,
|
||||
}
|
||||
}
|
||||
|
||||
/// A device of its own, without a window (native only).
|
||||
#[cfg(not(target_arch = "wasm32"))]
|
||||
pub fn headless() -> Result<Self> {
|
||||
let (adapter, device, queue) = open_headless()?;
|
||||
Self::new(adapter.get_info().name, device, queue)
|
||||
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))
|
||||
}
|
||||
|
||||
/// The adapter's name, for reports.
|
||||
|
||||
+40
-56
@@ -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.2 basis functions, A2.3 first
|
||||
// the same algorithms (The NURBS Book A2.1 span search, A2.3 basis functions and first
|
||||
// derivatives, the rational quotient rule), in f32.
|
||||
|
||||
const PLANE: u32 = 0u;
|
||||
@@ -60,6 +60,11 @@ struct Derivs {
|
||||
ok: bool,
|
||||
};
|
||||
|
||||
struct Basis {
|
||||
n0: array<f32, 8>,
|
||||
n1: array<f32, 8>,
|
||||
};
|
||||
|
||||
fn vec3_at(o: u32) -> vec3<f32> {
|
||||
return vec3<f32>(data[o], data[o + 1u], data[o + 2u]);
|
||||
}
|
||||
@@ -96,58 +101,42 @@ fn span_of(k0: u32, p: u32, n: u32, t: f32) -> u32 {
|
||||
return lo;
|
||||
}
|
||||
|
||||
// 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<f32, W> {
|
||||
var n: array<f32, W>;
|
||||
n[0] = 1.0;
|
||||
// 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<array<f32, 8>, 8>;
|
||||
var left: array<f32, 8>;
|
||||
var right: array<f32, 8>;
|
||||
ndu[0][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) {
|
||||
// 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;
|
||||
ndu[j][r] = right[r + 1u] + left[j - r];
|
||||
var temp = 0.0;
|
||||
if (den != 0.0) {
|
||||
temp = n[r] / den;
|
||||
if (ndu[j][r] != 0.0) {
|
||||
temp = ndu[r][j - 1u] / ndu[j][r];
|
||||
}
|
||||
n[r] = saved + right * temp;
|
||||
saved = left * temp;
|
||||
ndu[r][j] = saved + right[r + 1u] * temp;
|
||||
saved = left[j - r] * temp;
|
||||
}
|
||||
n[j] = saved;
|
||||
ndu[j][j] = saved;
|
||||
}
|
||||
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<f32, W> {
|
||||
var d: array<f32, W>;
|
||||
if (p == 0u) {
|
||||
return d;
|
||||
var b: Basis;
|
||||
for (var j = 0u; j <= p; j = j + 1u) {
|
||||
b.n0[j] = ndu[j][p];
|
||||
}
|
||||
let lower = basis_funs(k0, span, p - 1u, t);
|
||||
for (var r = 0u; r <= p; r = 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;
|
||||
}
|
||||
var d = 0.0;
|
||||
if (r >= 1u && ndu[p][r - 1u] != 0.0) {
|
||||
d = d + ndu[r - 1u][p - 1u] / ndu[p][r - 1u];
|
||||
}
|
||||
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;
|
||||
}
|
||||
if (r < p && ndu[p][r] != 0.0) {
|
||||
d = d - ndu[r][p - 1u] / ndu[p][r];
|
||||
}
|
||||
d[r] = x * f32(p);
|
||||
b.n1[r] = d * f32(p);
|
||||
}
|
||||
return d;
|
||||
return b;
|
||||
}
|
||||
|
||||
// Point and derivative of a rational function from its homogeneous value and derivative.
|
||||
@@ -163,19 +152,17 @@ 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 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);
|
||||
let bu = basis(s.knots_u, su, s.pu, u);
|
||||
let bv = basis(s.knots_v, sv, s.pv, v);
|
||||
var a = vec4<f32>(0.0);
|
||||
var au = vec4<f32>(0.0);
|
||||
var av = vec4<f32>(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 + bu0[i] * bv0[j] * c;
|
||||
au = au + bu1[i] * bv0[j] * c;
|
||||
av = av + bu0[i] * bv1[j] * c;
|
||||
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;
|
||||
}
|
||||
}
|
||||
r.ok = abs(a.w) > 0.0;
|
||||
@@ -202,14 +189,13 @@ 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 b0 = basis_funs(s.knots_u, span, s.pu, t);
|
||||
let b1 = basis_ders(s.knots_u, span, s.pu, t);
|
||||
let b = basis(s.knots_u, span, s.pu, t);
|
||||
var c = vec4<f32>(0.0);
|
||||
var dc = vec4<f32>(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 + b0[i] * h;
|
||||
dc = dc + b1[i] * h;
|
||||
c = c + b.n0[i] * h;
|
||||
dc = dc + b.n1[i] * h;
|
||||
}
|
||||
r.ok = abs(c.w) > 0.0;
|
||||
if (!r.ok) {
|
||||
@@ -277,11 +263,9 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
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)));
|
||||
// Steps (h, h), (-h, h), (h, -h), (-h, -h), without an array indexed by k (FXC).
|
||||
let steps = array<vec2<f32>, 4>(vec2<f32>(h, h), vec2<f32>(-h, h), vec2<f32>(h, -h), vec2<f32>(-h, -h));
|
||||
for (var k = 0u; k < 4u; k = k + 1u) {
|
||||
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 e = derivs(s, vert.u + steps[k].x, vert.v + steps[k].y);
|
||||
let m = oriented(s, cross(e.du, e.dv));
|
||||
if (e.ok && length(m) > 0.0) {
|
||||
n = m;
|
||||
|
||||
@@ -4,17 +4,12 @@ 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<TessGpu> {
|
||||
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
|
||||
@@ -22,80 +17,6 @@ fn gpu() -> Option<TessGpu> {
|
||||
}
|
||||
}
|
||||
|
||||
/// 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<storage, read_write> surfaces: array<u32>;
|
||||
@group(0) @binding(4) var<storage, read_write> out: array<vec4<f32>>;
|
||||
@compute @workgroup_size(64)
|
||||
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
|
||||
surfaces[gid.x] = 1u;
|
||||
out[gid.x] = vec4<f32>(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<T, E: std::fmt::Debug>(r: std::result::Result<T, E>) -> T {
|
||||
match r {
|
||||
Ok(v) => v,
|
||||
|
||||
+16
-47
@@ -736,30 +736,14 @@ 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<solidcraft_gpu::TessGpu> {
|
||||
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 {
|
||||
let tess = tessellation(&rs.adapter, &rs.device, &rs.queue, solidcraft_gpu::SHADER);
|
||||
// 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 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);
|
||||
@@ -928,29 +912,8 @@ 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` (the tessellation shader), planned bodies
|
||||
/// are evaluated on the GPU where the device builds that pipeline, as [`install`] sets it up.
|
||||
/// `size`. `None` without a GPU adapter. With `tess`, planned bodies are evaluated on the GPU.
|
||||
/// `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<bool> {
|
||||
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,
|
||||
@@ -959,12 +922,18 @@ pub(crate) fn render_offscreen(
|
||||
back: [f32; 3],
|
||||
eye: [f32; 4],
|
||||
size: u32,
|
||||
tess: Option<&str>,
|
||||
tess: bool,
|
||||
) -> Option<Vec<u8>> {
|
||||
let (adapter, device, queue) = offscreen_device()?;
|
||||
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 (format, depth) = (wgpu::TextureFormat::Rgba8Unorm, wgpu::TextureFormat::Depth24PlusStencil8);
|
||||
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 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 mut resources = egui_wgpu::CallbackResources::default();
|
||||
resources.insert(Resources::new(&device, target, compute));
|
||||
let none = || Arc::new(Mutex::new(None));
|
||||
|
||||
@@ -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, 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 = 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, gpu_tris, gpu) = render(true);
|
||||
let (none, cpu_tris, cpu) = render(false);
|
||||
@@ -3335,45 +3335,6 @@ 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<storage, read_write> surfaces: array<u32>;
|
||||
@compute @workgroup_size(64)
|
||||
fn main(@builtin(global_invocation_id) gid: vec3<u32>) { 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
|
||||
@@ -3401,8 +3362,7 @@ 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 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)
|
||||
crate::gpu::render_offscreen(scene, Some(highlight), proj.vp.to_f32(), proj.cam.back().to_f32(), eye, SIZE, tess)
|
||||
};
|
||||
let lit = |plain: &[u8], hovered: &[u8]| -> Vec<bool> { 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 {
|
||||
|
||||
Reference in new issue
Block a user