GPU-accelerated UI toolkit (Vulkan)
git clone https://git.lucas.co/cce-ui.git
W4a: compute jobs run in the browser
What a compute job is moves out of vk into the portable crate::compute:
Kernel, Binding, BindKind, workgroups, MAX_BINDINGS, the rules a job
is held to (check_job, the ping-pong slot_for / result_slot) and the
kernel's parse — naga's WGSL frontend, now a dependency everywhere,
validates a kernel and reads its @workgroup_size before any device sees
it. vk::compute re-exports every item at its old path.
web::ComputeDevice runs the same jobs on WebGPU and answers them the
same way; readback is a promise there, so its run and siblings are
async. Read-only storage is its own binding type in WebGPU's layouts
(the module says which binding is), and the device asks the adapter
for its own storage-buffer and workgroup limits, WebGPU's default being
eight storage buffers a stage. What WebGPU rejects is caught in a
validation error scope and returned as an Err.
examples/compute_probe/jobs.rs is the check, run natively by
compute_native and in headless Chromium by scripts/web-probe/compute: a
map, a uniform, 33- and 34-pass ping-pong solves, a 2D dispatch, ten
bindings, a bad kernel and a missing entry, each exact against a CPU
reference. Vulkan (lavapipe) and WebGPU (SwiftShader) print identical
output, digests included.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01WjL3pejMNY95NHv9BcmXaZ
CLAUDE.md | 28 +++-
Cargo.toml | 19 +++
examples/compute_native.rs | 19 +++
examples/compute_probe/jobs.rs | 151 +++++++++++++++++++
examples/compute_web.rs | 22 +++
scripts/check-wasm | 2 +-
scripts/web-probe/compute | 21 +++
scripts/web-probe/compute.html | 10 ++
scripts/web-probe/compute.mjs | 16 ++
src/compute.rs | 253 ++++++++++++++++++++++++++++++++
src/lib.rs | 1 +
src/vk/compute.rs | 157 ++------------------
src/web/compute.rs | 321 +++++++++++++++++++++++++++++++++++++++++
src/web/mod.rs | 44 +++++-
src/web/renderer.rs | 17 +--
15 files changed, 915 insertions(+), 166 deletions(-)
diff --git a/CLAUDE.md b/CLAUDE.md
index 111663f..0cce10c 100644
--- a/CLAUDE.md
+++ b/CLAUDE.md
@@ -115,6 +115,27 @@ Not there yet: the clipboard, IME composition, drag and drop, file dialogs, and
text is not the display list's (`display_list_text` false — it stages its own through the
native-only `stage_renderer`, so draws no text here).
+**Compute jobs run in the browser too** (`web::ComputeDevice`, since 2026-10-05). What a job
+IS moved out of `vk` into the portable `crate::compute` — `Kernel`, `Binding`, `BindKind`,
+`workgroups`, `MAX_BINDINGS`, and the rules a job is held to before any device sees it
+(`check_job`, the ping-pong `slot_for` / `result_slot`, `parse_kernel`: naga's WGSL
+frontend, now a dependency on every target, validates a kernel and reads its
+`@workgroup_size` — WebGPU can report neither) — and `vk::compute` re-exports every one at
+its old path. The browser device takes the same jobs and answers them the same way, with
+one difference the platform makes: readback is a promise, so its `run`, `run_over`,
+`run_passes`, `run_passes_over` and `workgroup_size` are `async`. Two things WebGPU does
+differently underneath: its layouts tell read-only storage from read-write (the module
+says which, `ParsedKernel::read_only_storage`), and a device starts at the spec's default
+of eight storage buffers a stage, so it asks for the adapter's own (SwiftShader offers
+ten; a job past the adapter's ceiling is an `Err` naming the limit). What WebGPU rejects
+is caught in a validation error scope and returned. `examples/compute_probe/jobs.rs` is
+the check — a map, a uniform, a 33- and a 34-pass ping-pong, a 2D dispatch, ten bindings,
+a bad kernel and a missing entry, each exact against a CPU reference in f32 — run by
+`compute_native` and by `scripts/web-probe/compute`: the two outputs are identical to the
+bit (lavapipe vs SwiftShader, 2026-10-05). A reference written for a length that is not
+a multiple of four floats must know that `arrayLength` counts the 16-byte padding, on
+both devices.
+
**The reference app runs on both, through one input script.** `examples/demo_web.rs` is
`src/main.rs`'s `DemoApp` (included by `#[path]`, hence `pub(crate)`) in a page;
`scripts/web-probe/demo <dir>` builds it, serves it with the machine's fonts and replays the
@@ -1169,6 +1190,9 @@ cce-system-interface) to confirm behavior, not just the test suite.
- `config.rs` — KDL loading and `kdl_to_json` conversion (see workspace `CLAUDE.md` for paths).
- `context.rs` — `UiContext`: the retained widget tree, event routing, spatial grid, dirty
tracking, hit-testing.
+- `compute.rs` — what a compute job is, apart from the device that runs it: `Kernel`,
+ `Binding`, the job rules and naga's parse (see "Compute jobs run in the browser too").
+ `vk::ComputeDevice` and `web::ComputeDevice` run them.
- `history.rs` — `History<T>`: the undo/redo snapshot stack (cap, gestures, grouped runs).
The toolkit defines the stack and the routing, never the step — see the trait section.
- `widget/` — `container/` (vbox/hbox/scroll/menu/treelist/…), `input/` (button/slider/text_box/
@@ -1224,7 +1248,9 @@ cce-system-interface) to confirm behavior, not just the test suite.
Vulkan path: an sRGB VIEW of the canvas's unorm format, the parameter block as a
dynamic-offset uniform, a 1x1 backdrop, the blur snapshot as end-pass / copy / resume,
every frame drawn whole. And `shell.rs`, the browser shell: `run`, `Fonts`, `Sizing`,
- `capture` (see "And an `Application` runs in a page" above).
+ `capture` (see "And an `Application` runs in a page" above); `compute.rs`, the async
+ `ComputeDevice`; and `request_device`, the adapter and device every one of them asks
+ for (with the limits a caller names raised to the adapter's).
- `protocol.rs` — inline-generated Wayland protocol bindings.
- `ipc.rs` — the `/tmp/<prefix>-<WAYLAND_DISPLAY>.sock` helpers (`socket_path`, `send_command`,
the bounded `read_request_line`, `focus_window`), and `ipc::instance`: single-instance
diff --git a/Cargo.toml b/Cargo.toml
index b76fe95..c2ad963 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -20,6 +20,10 @@ doc_editor = []
cce-vault = { git = "https://github.com/lsgalante/cce-vault.git", rev = "41ca52bd51ea3655a791c58005be9fce939e4333", optional = true }
glam = "0.29"
bytemuck = { version = "1", features = ["derive"] }
+# WGSL parsing and validation everywhere — a compute kernel is checked, and
+# its `@workgroup_size` read, before any device sees it; native adds the
+# SPIR-V writer below.
+naga = { version = "24", features = ["wgsl-in"] }
# Text shaping. This was reached through glyphon until the wgpu path was retired;
# every `glyphon::` item cce-ui ever used was a cosmic-text re-export (the sole
# exception, `TextBounds`, is now defined in backend/text.rs), so depending
@@ -160,6 +164,16 @@ web-sys = { version = "0.3", features = [
"gpu_texture_usage",
"gpu_shader_stage",
"gpu_map_mode",
+ "GpuDeviceDescriptor",
+ "GpuSupportedLimits",
+ "GpuAdapterInfo",
+ "GpuComputePipeline",
+ "GpuComputePipelineDescriptor",
+ "GpuProgrammableStage",
+ "GpuComputePassEncoder",
+ "GpuComputePassDescriptor",
+ "GpuErrorFilter",
+ "GpuError",
] }
# The renderer probe: one scene through both renderers (examples/probe/).
@@ -168,6 +182,11 @@ web-sys = { version = "0.3", features = [
name = "probe_web"
crate-type = ["cdylib"]
+# The compute probe's browser half (examples/compute_probe/); empty natively.
+[[example]]
+name = "compute_web"
+crate-type = ["cdylib"]
+
# The reference app (src/main.rs) in a browser, on the browser shell; a
# cdylib for wasm-bindgen, empty natively.
[[example]]
diff --git a/examples/compute_native.rs b/examples/compute_native.rs
new file mode 100644
index 0000000..5496faa
--- /dev/null
+++ b/examples/compute_native.rs
@@ -0,0 +1,19 @@
+//! The compute probe on Vulkan (`vk::ComputeDevice`): runs the shared jobs
+//! (`compute_probe/jobs.rs`) and prints one line each. Its browser half is
+//! `compute_web`; `scripts/web-probe/compute` diffs the two.
+
+#[path = "compute_probe/jobs.rs"]
+mod jobs;
+
+fn main() {
+ let mut dev = match cce_ui::vk::ComputeDevice::new() {
+ Ok(d) => d,
+ Err(e) => {
+ eprintln!("{e}");
+ std::process::exit(2);
+ }
+ };
+ for line in compute_jobs!(dev,) {
+ println!("{line}");
+ }
+}
diff --git a/examples/compute_probe/jobs.rs b/examples/compute_probe/jobs.rs
new file mode 100644
index 0000000..f3d18b1
--- /dev/null
+++ b/examples/compute_probe/jobs.rs
@@ -0,0 +1,151 @@
+//! The compute probe's jobs, shared by both halves: each runs on a device
+//! and is held to a CPU reference, and prints a digest of its result bytes,
+//! so the native run (`compute_native`, Vulkan) and the browser's
+//! (`compute_web`, WebGPU) can be compared line for line. The arithmetic is
+//! exact in f32 (small integers, power-of-two weights), so a conformant
+//! device must match the reference — and the other device — to the bit.
+
+use cce_ui::compute::{Binding, Kernel};
+
+pub const DOUBLE: &str = "
+@group(0) @binding(0) var<storage, read> a: array<f32>;
+@group(0) @binding(1) var<storage, read_write> b: array<f32>;
+@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+ if id.x < arrayLength(&b) { b[id.x] = a[id.x] * 2.0; }
+}";
+
+pub const SAXPY: &str = "
+struct P { a: f32, n: u32 }
+@group(0) @binding(0) var<uniform> p: P;
+@group(0) @binding(1) var<storage, read> x: array<f32>;
+@group(0) @binding(2) var<storage, read_write> y: array<f32>;
+@compute @workgroup_size(32) fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+ if id.x < p.n { y[id.x] = p.a * x[id.x] + y[id.x]; }
+}";
+
+pub const BLUR: &str = "
+@group(0) @binding(0) var<storage, read> src: array<f32>;
+@group(0) @binding(1) var<storage, read_write> dst: array<f32>;
+@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+ let n = arrayLength(&dst);
+ let i = id.x;
+ if i >= n { return; }
+ let l = src[max(i, 1u) - 1u];
+ let r = src[min(i + 1u, n - 1u)];
+ dst[i] = 0.25 * l + 0.5 * src[i] + 0.25 * r;
+}";
+
+pub const GRID: &str = "
+@group(0) @binding(0) var<storage, read_write> g: array<f32>;
+@compute @workgroup_size(8, 8) fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+ if id.x < 32u && id.y < 24u { g[id.y * 32u + id.x] = f32(id.x) + 1000.0 * f32(id.y); }
+}";
+
+/// A kernel over `n` bindings: `n - 1` inputs summed into the last.
+pub fn many_source(n: usize) -> String {
+ let mut s = String::new();
+ for i in 0..n - 1 {
+ s += &format!("@group(0) @binding({i}) var<storage, read> in{i}: array<f32>;\n");
+ }
+ s += &format!("@group(0) @binding({}) var<storage, read_write> out: array<f32>;\n", n - 1);
+ s += "@compute @workgroup_size(16) fn main(@builtin(global_invocation_id) id: vec3<u32>) {\n";
+ s += " let i = id.x; if i >= arrayLength(&out) { return; }\n var t = 0.0;\n";
+ for i in 0..n - 1 {
+ s += &format!(" t += in{i}[i];\n");
+ }
+ s += " out[i] = t;\n}\n";
+ s
+}
+
+/// FNV-1a over a result's bytes: what two devices are compared by.
+pub fn digest(v: &[f32]) -> String {
+ let mut h: u64 = 0xcbf29ce484222325;
+ for b in bytemuck::cast_slice::<f32, u8>(v) {
+ h ^= *b as u64;
+ h = h.wrapping_mul(0x100000001b3);
+ }
+ format!("{h:016x}")
+}
+
+pub fn report(name: &str, got: &[f32], want: &[f32]) -> String {
+ let worst = got.iter().zip(want).map(|(g, w)| (g - w).abs()).fold(0.0f32, f32::max);
+ let ok = got.len() == want.len() && got.iter().zip(want).all(|(g, w)| g.to_bits() == w.to_bits());
+ format!("{name}: {} n={} max|d|={worst} digest={}", if ok { "exact" } else { "DIFFERS" }, got.len(), digest(got))
+}
+
+/// The jobs, written once over whichever device runs them. A macro rather
+/// than a generic function: one device's `run` is synchronous and the
+/// other's async, and `$await` is the one word between them.
+#[macro_export]
+macro_rules! compute_jobs {
+ ($dev:expr, $($await:tt)*) => {{
+ use cce_ui::compute::{Binding, Kernel};
+ use jobs::*;
+ let mut lines: Vec<String> = Vec::new();
+
+ // A map over 1000 elements, sized from the entry's @workgroup_size.
+ let a: Vec<f32> = (0..1000).map(|i| i as f32 - 500.0).collect();
+ let mut b = vec![0.0f32; 1000];
+ let r = $dev.run_over(&Kernel::new(DOUBLE, "main"), &mut [Binding::input(&a), Binding::rw(&mut b)], 1000)$($await)*;
+ let want: Vec<f32> = a.iter().map(|v| v * 2.0).collect();
+ lines.push(match r { Ok(()) => report("double", &b, &want), Err(e) => format!("double: ERR {e}") });
+
+ // A uniform block beside storage.
+ #[repr(C)]
+ #[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
+ struct P { a: f32, n: u32, _pad: [u32; 2] }
+ let x: Vec<f32> = (0..777).map(|i| (i % 13) as f32).collect();
+ let mut y: Vec<f32> = (0..777).map(|i| (i % 7) as f32).collect();
+ let want: Vec<f32> = x.iter().zip(&y).map(|(x, y)| 2.5 * x + y).collect();
+ let p = P { a: 2.5, n: 777, _pad: [0; 2] };
+ let r = $dev.run_over(&Kernel::new(SAXPY, "main"), &mut [Binding::uniform(&p), Binding::input(&x), Binding::rw(&mut y)], 777)$($await)*;
+ lines.push(match r { Ok(()) => report("saxpy", &y, &want), Err(e) => format!("saxpy: ERR {e}") });
+
+ // Ping-pong passes, odd and even counts: the result lands in the
+ // output binding either way. 4096 elements: a length the 16-byte
+ // padding leaves alone, since `arrayLength` counts the padding.
+ for passes in [33u32, 34] {
+ let src: Vec<f32> = (0..4096).map(|i| if i % 512 == 256 { 4096.0 } else { 0.0 }).collect();
+ let mut dst = vec![0.0f32; 4096];
+ let mut want = src.clone();
+ for _ in 0..passes {
+ let n = want.len();
+ want = (0..n).map(|i| 0.25 * want[i.saturating_sub(1)] + 0.5 * want[i] + 0.25 * want[(i + 1).min(n - 1)]).collect();
+ }
+ let r = $dev
+ .run_passes_over(&Kernel::new(BLUR, "main"), &mut [Binding::input(&src), Binding::rw(&mut dst)], 4096, passes, Some((0, 1)))
+ $($await)*;
+ lines.push(match r { Ok(()) => report(&format!("blur x{passes}"), &dst, &want), Err(e) => format!("blur x{passes}: ERR {e}") });
+ }
+
+ // A 2D dispatch.
+ let mut g = vec![-1.0f32; 32 * 24];
+ let r = $dev.run(&Kernel::new(GRID, "main"), &mut [Binding::rw(&mut g)], [4, 3, 1])$($await)*;
+ let want: Vec<f32> = (0..32 * 24).map(|i| (i % 32) as f32 + 1000.0 * (i / 32) as f32).collect();
+ lines.push(match r { Ok(()) => report("grid", &g, &want), Err(e) => format!("grid: ERR {e}") });
+
+ // Ten bindings: past WebGPU's default of eight storage buffers a
+ // stage (the browser device asks the adapter for its own limit),
+ // within what every adapter offers — SwiftShader's is ten.
+ let ins: Vec<Vec<f32>> = (0..9).map(|k| (0..100).map(|i| (k * 100 + i) as f32).collect()).collect();
+ let mut out = vec![0.0f32; 100];
+ let want: Vec<f32> = (0..100).map(|i| ins.iter().map(|v| v[i]).sum()).collect();
+ let mut binds: Vec<Binding> = ins.iter().map(|v| Binding::input(v)).collect();
+ binds.push(Binding::rw(&mut out));
+ let src = many_source(10);
+ let r = $dev.run_over(&Kernel::new(src.as_str(), "main"), &mut binds, 100)$($await)*;
+ drop(binds);
+ lines.push(match r { Ok(()) => report("ten bindings", &out, &want), Err(e) => format!("ten bindings: ERR {e}") });
+
+ // A user's bad kernel is an Err, with naga's word for what is wrong.
+ let mut z = vec![0.0f32; 4];
+ let r = $dev.run(&Kernel::new("fn main( {", "main"), &mut [Binding::rw(&mut z)], [1, 1, 1])$($await)*;
+ lines.push(format!("bad wgsl: {}", match r { Err(e) if e.contains("parse error") => "Err(parse error)".to_string(), other => format!("{:?}", other) }));
+ let r = $dev.run(&Kernel::new(DOUBLE, "nope"), &mut [Binding::input(&a), Binding::rw(&mut z)], [1, 1, 1])$($await)*;
+ lines.push(format!("no entry: {}", match r { Err(e) => e, Ok(()) => "Ok".into() }));
+ lines
+ }};
+}
+
+#[allow(dead_code)]
+fn _uses(_: Binding, _: Kernel) {}
diff --git a/examples/compute_web.rs b/examples/compute_web.rs
new file mode 100644
index 0000000..583463f
--- /dev/null
+++ b/examples/compute_web.rs
@@ -0,0 +1,22 @@
+//! The compute probe on WebGPU (`web::ComputeDevice`), in a browser: the
+//! shared jobs (`compute_probe/jobs.rs`), one line each, returned to the
+//! page. A cdylib for wasm-bindgen; natively it is empty.
+
+#[cfg(target_arch = "wasm32")]
+#[path = "compute_probe/jobs.rs"]
+mod jobs;
+
+#[cfg(target_arch = "wasm32")]
+mod web {
+ use super::jobs;
+ use wasm_bindgen::prelude::*;
+
+ /// The device's name and every job's line, newline-separated.
+ #[wasm_bindgen]
+ pub async fn run_jobs() -> Result<String, JsValue> {
+ std::panic::set_hook(Box::new(|info| web_sys::console::error_1(&info.to_string().into())));
+ let mut dev = cce_ui::web::ComputeDevice::new().await?;
+ web_sys::console::log_1(&dev.device_name().into());
+ Ok(crate::compute_jobs!(dev, .await).join("\n"))
+ }
+}
diff --git a/scripts/check-wasm b/scripts/check-wasm
index 8696f84..7bb42a6 100755
--- a/scripts/check-wasm
+++ b/scripts/check-wasm
@@ -18,5 +18,5 @@ if ! rustup target list --installed 2>/dev/null | grep -qx wasm32-unknown-unknow
fi
cargo check --lib --target wasm32-unknown-unknown
cargo check --lib --target wasm32-unknown-unknown --features markdown,doc_editor
-cargo check --example probe_web --example demo_web --target wasm32-unknown-unknown
+cargo check --example probe_web --example demo_web --example compute_web --target wasm32-unknown-unknown
echo "check-wasm: ok"
diff --git a/scripts/web-probe/compute b/scripts/web-probe/compute
new file mode 100755
index 0000000..e670698
--- /dev/null
+++ b/scripts/web-probe/compute
@@ -0,0 +1,21 @@
+#!/usr/bin/env bash
+# web-probe/compute — run the compute probe's jobs (examples/compute_probe/)
+# on WebGPU in headless Chromium and print one line each, the format the
+# native half prints:
+#
+# cargo run --example compute_native > native.txt
+# scripts/web-probe/compute > web.txt && diff native.txt web.txt
+#
+# Every job's arithmetic is exact in f32, so each line says `exact` and the
+# digests of the two runs agree to the bit. Needs what `run` needs, but no
+# fonts.
+set -euo pipefail
+here=$(cd "$(dirname "$0")" && pwd)
+cd "$here/../.."
+cargo build --release --target wasm32-unknown-unknown --example compute_web >&2
+site=$(mktemp -d)
+trap 'rm -rf "$site"' EXIT
+wasm-bindgen --target web --no-typescript --out-dir "$site" \
+ target/wasm32-unknown-unknown/release/examples/compute_web.wasm
+cp "$here/compute.html" "$site/compute.html"
+node "$here/compute.mjs" "$site"
diff --git a/scripts/web-probe/compute.html b/scripts/web-probe/compute.html
new file mode 100644
index 0000000..83b33ef
--- /dev/null
+++ b/scripts/web-probe/compute.html
@@ -0,0 +1,10 @@
+<!doctype html>
+<meta charset="utf-8">
+<title>cce-ui compute probe</title>
+<!-- The compute probe's browser half (examples/compute_web.rs). -->
+<script type="module">
+import init, { run_jobs } from './compute_web.js';
+await init();
+window.runJobs = () => run_jobs();
+window.probeReady = true;
+</script>
diff --git a/scripts/web-probe/compute.mjs b/scripts/web-probe/compute.mjs
new file mode 100644
index 0000000..29e2cfb
--- /dev/null
+++ b/scripts/web-probe/compute.mjs
@@ -0,0 +1,16 @@
+// compute.mjs <site-dir>: open the compute probe page (browser.mjs), run
+// the shared jobs on WebGPU, and print one line each — the format
+// `cargo run --example compute_native` prints, so the two can be diffed.
+import { open } from './browser.mjs';
+
+const [root] = process.argv.slice(2);
+if (!root) { console.error('usage: compute.mjs <site-dir>'); process.exit(2); }
+const { page, logs, close } = await open(root, 'compute.html', 320, 200);
+await page.waitForFunction(() => window.probeReady === true);
+let ok = true;
+try {
+ console.log(await page.evaluate(() => window.runJobs()));
+} catch (e) { ok = false; logs.push('[probe] ' + String(e)); }
+for (const l of logs) console.error(l);
+await close();
+process.exit(ok ? 0 : 1);
diff --git a/src/compute.rs b/src/compute.rs
new file mode 100644
index 0000000..3eff0b4
--- /dev/null
+++ b/src/compute.rs
@@ -0,0 +1,253 @@
+//! What a compute job is, apart from the device that runs it: a [`Kernel`]
+//! (WGSL source and an entry point), its [`Binding`]s in `@binding(i)` order,
+//! and the rules a job is held to before anything reaches a GPU — the
+//! binding count, the dispatch, the ping-pong pair, the kernel parsed and
+//! validated by naga with its own diagnostics.
+//!
+//! Two devices run jobs: `vk::ComputeDevice` (a headless Vulkan device,
+//! synchronous) and, in the browser, `web::ComputeDevice` (WebGPU, whose
+//! readback is a promise, so its `run` is async). Both take these types and
+//! answer a job the same way, so a kernel and its bindings are written once.
+
+/// A compute shader: WGSL source and the `@compute` entry point to run.
+#[derive(Clone, Debug, PartialEq, Eq, Hash)]
+pub struct Kernel {
+ pub source: String,
+ pub entry: String,
+}
+
+impl Kernel {
+ pub fn new(source: impl Into<String>, entry: impl Into<String>) -> Self {
+ Kernel { source: source.into(), entry: entry.into() }
+ }
+}
+
+/// How a binding is declared to the shader, in `@binding(i)` order.
+#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
+pub enum BindKind {
+ /// `var<storage, read>` or `var<storage, read_write>`.
+ Storage,
+ /// `var<uniform>`: a small parameter block, 16-byte layout rules apply.
+ Uniform,
+}
+
+/// One buffer of a job, bound at `@group(0) @binding(i)` for its index in
+/// the list handed to a device's `run`.
+pub enum Binding<'a> {
+ /// Read-write storage: uploaded before the dispatch and READ BACK into
+ /// the same slice after it.
+ Storage(&'a mut [u8]),
+ /// Read-only storage: uploaded, never read back.
+ Input(&'a [u8]),
+ /// A uniform block: uploaded, never read back.
+ Uniform(&'a [u8]),
+}
+
+impl<'a> Binding<'a> {
+ /// A read-write binding over a typed slice (`&mut [f32]`, `&mut [[f32; 3]]`, …).
+ pub fn rw<T: bytemuck::Pod>(data: &'a mut [T]) -> Self {
+ Binding::Storage(bytemuck::cast_slice_mut(data))
+ }
+
+ /// A read-only storage binding over a typed slice.
+ pub fn input<T: bytemuck::Pod>(data: &'a [T]) -> Self {
+ Binding::Input(bytemuck::cast_slice(data))
+ }
+
+ /// A uniform binding over one `Pod` struct.
+ pub fn uniform<T: bytemuck::Pod>(value: &'a T) -> Self {
+ Binding::Uniform(bytemuck::bytes_of(value))
+ }
+
+ pub(crate) fn kind(&self) -> BindKind {
+ match self {
+ Binding::Storage(_) | Binding::Input(_) => BindKind::Storage,
+ Binding::Uniform(_) => BindKind::Uniform,
+ }
+ }
+
+ pub(crate) fn bytes(&self) -> &[u8] {
+ match self {
+ Binding::Storage(b) => b,
+ Binding::Input(b) => b,
+ Binding::Uniform(b) => b,
+ }
+ }
+}
+
+/// Workgroups needed to cover `items` at `per_group` invocations each — the
+/// `@workgroup_size` of the entry point, which a device's `run_over` reads
+/// for you.
+pub fn workgroups(items: u32, per_group: u32) -> u32 {
+ items.div_ceil(per_group.max(1)).max(1)
+}
+
+/// The most bindings one job may carry.
+pub const MAX_BINDINGS: usize = 16;
+
+/// Storage bindings are bound whole, so a buffer's size has to be a multiple
+/// of the widest element stride a shader may declare; 16 covers `vec4<f32>`.
+pub(crate) const BUFFER_ALIGN: usize = 16;
+
+/// A binding of `len` bytes as the buffer it is uploaded into: at least one
+/// alignment unit, rounded up to one. The padding is uploaded as zeros.
+pub(crate) fn padded_len(len: usize) -> usize {
+ len.max(BUFFER_ALIGN).div_ceil(BUFFER_ALIGN) * BUFFER_ALIGN
+}
+
+/// Hold a job to the rules before any device is asked: the binding count,
+/// a dispatch with no zero in it, at least one pass, and a ping-pong pair
+/// that names a read-only input and a read-write output of one length.
+pub(crate) fn check_job(
+ bindings: &[Binding<'_>],
+ groups: [u32; 3],
+ passes: u32,
+ ping_pong: Option<(usize, usize)>,
+) -> Result<(), String> {
+ if bindings.len() > MAX_BINDINGS {
+ return Err(format!("{} bindings; a job may carry at most {MAX_BINDINGS}", bindings.len()));
+ }
+ if groups.iter().any(|&g| g == 0) {
+ return Err(format!("workgroup count {groups:?} has a zero"));
+ }
+ if passes == 0 {
+ return Err("a job needs at least one pass".to_string());
+ }
+ if let Some((a, b)) = ping_pong {
+ if a == b || a >= bindings.len() || b >= bindings.len() {
+ return Err(format!("ping-pong pair ({a}, {b}) does not name two distinct bindings of {}", bindings.len()));
+ }
+ if !matches!(bindings[a], Binding::Input(_)) {
+ return Err(format!("ping-pong binding {a} must be a read-only Input: it is where the first pass reads"));
+ }
+ if !matches!(bindings[b], Binding::Storage(_)) {
+ return Err(format!("ping-pong binding {b} must be a read-write Storage: it is where the result lands"));
+ }
+ if bindings[a].bytes().len() != bindings[b].bytes().len() {
+ return Err(format!(
+ "ping-pong bindings {a} and {b} differ in length ({} vs {} bytes)",
+ bindings[a].bytes().len(),
+ bindings[b].bytes().len()
+ ));
+ }
+ }
+ Ok(())
+}
+
+/// The buffer bound at `binding` in a pass: a ping-pong pass with `swapped`
+/// set (every second one) binds the pair's two buffers the other way round.
+pub(crate) fn slot_for(binding: usize, swapped: bool, ping_pong: Option<(usize, usize)>) -> usize {
+ match ping_pong {
+ Some((a, b)) if swapped && binding == a => b,
+ Some((a, b)) if swapped && binding == b => a,
+ _ => binding,
+ }
+}
+
+/// The buffer a read-write binding is read back from: the ping-pong output
+/// from whichever buffer the LAST pass wrote, every other one from its own.
+pub(crate) fn result_slot(binding: usize, passes: u32, ping_pong: Option<(usize, usize)>) -> usize {
+ match ping_pong {
+ Some((a, b)) if binding == b && passes % 2 == 0 => a,
+ _ => binding,
+ }
+}
+
+/// A kernel parsed and validated, with what a device needs to know of it.
+pub(crate) struct ParsedKernel {
+ pub module: naga::Module,
+ #[cfg_attr(target_arch = "wasm32", allow(dead_code))]
+ pub info: naga::valid::ModuleInfo,
+ pub workgroup_size: [u32; 3],
+}
+
+impl ParsedKernel {
+ /// Whether the module declares `@group(0) @binding(i)` as read-only
+ /// storage (`var<storage, read>`). WebGPU's layouts tell read-only from
+ /// read-write storage, where Vulkan's do not.
+ #[cfg_attr(not(target_arch = "wasm32"), allow(dead_code))]
+ pub fn read_only_storage(&self, binding: u32) -> bool {
+ self.module.global_variables.iter().any(|(_, g)| {
+ g.binding.as_ref().is_some_and(|b| b.group == 0 && b.binding == binding)
+ && matches!(g.space, naga::AddressSpace::Storage { access } if !access.contains(naga::StorageAccess::STORE))
+ })
+ }
+}
+
+/// Parse and validate a kernel, every failure reported with naga's own
+/// diagnostic — a kernel may be a user's, so a bad one is an `Err`, never a
+/// panic — and find its `@compute` entry point, naming what the module does
+/// offer when there is none of that name.
+pub(crate) fn parse_kernel(kernel: &Kernel) -> Result<ParsedKernel, String> {
+ let module = naga::front::wgsl::parse_str(&kernel.source)
+ .map_err(|e| format!("WGSL parse error: {}", e.emit_to_string(&kernel.source).trim_end()))?;
+ let entry = module
+ .entry_points
+ .iter()
+ .find(|ep| ep.name == kernel.entry && ep.stage == naga::ShaderStage::Compute)
+ .ok_or_else(|| {
+ let offered: Vec<&str> = module
+ .entry_points
+ .iter()
+ .filter(|ep| ep.stage == naga::ShaderStage::Compute)
+ .map(|ep| ep.name.as_str())
+ .collect();
+ format!(
+ "no @compute entry point named `{}`; the module offers {}",
+ kernel.entry,
+ if offered.is_empty() { "none".to_string() } else { offered.join(", ") }
+ )
+ })?;
+ let workgroup_size = entry.workgroup_size;
+ let info = naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::empty())
+ .validate(&module)
+ .map_err(|e| format!("WGSL validation error: {}", e.emit_to_string(&kernel.source).trim_end()))?;
+ Ok(ParsedKernel { module, info, workgroup_size })
+}
+
+#[cfg(test)]
+mod tests {
+ use super::*;
+
+ #[test]
+ fn a_job_is_held_to_the_rules_before_a_device_sees_it() {
+ let mut out = [0.0f32; 4];
+ let inp = [0.0f32; 4];
+ let short = [0.0f32; 2];
+ assert!(check_job(&[Binding::rw(&mut out)], [1, 1, 1], 1, None).is_ok());
+ assert!(check_job(&[Binding::rw(&mut out)], [1, 0, 1], 1, None).unwrap_err().contains("zero"));
+ assert!(check_job(&[Binding::rw(&mut out)], [1, 1, 1], 0, None).unwrap_err().contains("pass"));
+ assert!(check_job(&[Binding::input(&inp), Binding::rw(&mut out)], [1, 1, 1], 2, Some((0, 1))).is_ok());
+ assert!(check_job(&[Binding::input(&inp), Binding::rw(&mut out)], [1, 1, 1], 2, Some((1, 0)))
+ .unwrap_err()
+ .contains("read-only Input"));
+ assert!(check_job(&[Binding::input(&short), Binding::rw(&mut out)], [1, 1, 1], 2, Some((0, 1)))
+ .unwrap_err()
+ .contains("differ in length"));
+ }
+
+ #[test]
+ fn the_ping_pong_result_is_wherever_the_last_pass_wrote() {
+ let pp = Some((0, 1));
+ assert_eq!((slot_for(0, true, pp), slot_for(1, true, pp), slot_for(2, true, pp)), (1, 0, 2));
+ assert_eq!(slot_for(0, false, pp), 0);
+ assert_eq!(result_slot(1, 3, pp), 1);
+ assert_eq!(result_slot(1, 4, pp), 0);
+ assert_eq!(result_slot(1, 4, None), 1);
+ }
+
+ #[test]
+ fn a_kernel_reports_its_workgroup_size_its_read_only_bindings_and_its_errors() {
+ let src = "@group(0) @binding(0) var<storage, read> a: array<f32>;
+ @group(0) @binding(1) var<storage, read_write> b: array<f32>;
+ @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+ if id.x < arrayLength(&b) { b[id.x] = a[id.x] * 2.0; }
+ }";
+ let k = parse_kernel(&Kernel::new(src, "main")).unwrap();
+ assert_eq!(k.workgroup_size, [64, 1, 1]);
+ assert!(k.read_only_storage(0));
+ assert!(!k.read_only_storage(1));
+ assert!(parse_kernel(&Kernel::new(src, "nope")).err().unwrap().contains("offers main"));
+ assert!(parse_kernel(&Kernel::new("fn (", "main")).err().unwrap().contains("parse error"));
+ }
+}
diff --git a/src/lib.rs b/src/lib.rs
index 5773daf..c9de488 100644
--- a/src/lib.rs
+++ b/src/lib.rs
@@ -3,6 +3,7 @@
// else builds for the browser too: `cargo check --lib --target
// wasm32-unknown-unknown` is the check.
pub mod color;
+pub mod compute;
pub mod widget;
pub mod config;
pub mod input;
diff --git a/src/vk/compute.rs b/src/vk/compute.rs
index c67113e..e96efd5 100644
--- a/src/vk/compute.rs
+++ b/src/vk/compute.rs
@@ -35,78 +35,8 @@ use ash::vk;
use std::collections::HashMap;
use std::ffi::CString;
-/// A compute shader: WGSL source and the `@compute` entry point to run.
-#[derive(Clone, Debug, PartialEq, Eq, Hash)]
-pub struct Kernel {
- pub source: String,
- pub entry: String,
-}
-
-impl Kernel {
- pub fn new(source: impl Into<String>, entry: impl Into<String>) -> Self {
- Kernel { source: source.into(), entry: entry.into() }
- }
-}
-
-/// How a binding is declared to the shader, in `@binding(i)` order.
-#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
-pub enum BindKind {
- /// `var<storage, read>` or `var<storage, read_write>`.
- Storage,
- /// `var<uniform>`: a small parameter block, 16-byte layout rules apply.
- Uniform,
-}
-
-/// One buffer of a job, bound at `@group(0) @binding(i)` for its index in
-/// the list handed to [`ComputeDevice::run`].
-pub enum Binding<'a> {
- /// Read-write storage: uploaded before the dispatch and READ BACK into
- /// the same slice after it.
- Storage(&'a mut [u8]),
- /// Read-only storage: uploaded, never read back.
- Input(&'a [u8]),
- /// A uniform block: uploaded, never read back.
- Uniform(&'a [u8]),
-}
-
-impl<'a> Binding<'a> {
- /// A read-write binding over a typed slice (`&mut [f32]`, `&mut [[f32; 3]]`, …).
- pub fn rw<T: bytemuck::Pod>(data: &'a mut [T]) -> Self {
- Binding::Storage(bytemuck::cast_slice_mut(data))
- }
-
- /// A read-only storage binding over a typed slice.
- pub fn input<T: bytemuck::Pod>(data: &'a [T]) -> Self {
- Binding::Input(bytemuck::cast_slice(data))
- }
-
- /// A uniform binding over one `Pod` struct.
- pub fn uniform<T: bytemuck::Pod>(value: &'a T) -> Self {
- Binding::Uniform(bytemuck::bytes_of(value))
- }
-
- fn kind(&self) -> BindKind {
- match self {
- Binding::Storage(_) | Binding::Input(_) => BindKind::Storage,
- Binding::Uniform(_) => BindKind::Uniform,
- }
- }
-
- fn bytes(&self) -> &[u8] {
- match self {
- Binding::Storage(b) => b,
- Binding::Input(b) => b,
- Binding::Uniform(b) => b,
- }
- }
-}
-
-/// Workgroups needed to cover `items` at `per_group` invocations each — the
-/// `@workgroup_size` of the entry point, which [`ComputeDevice::run_over`]
-/// reads for you.
-pub fn workgroups(items: u32, per_group: u32) -> u32 {
- items.div_ceil(per_group.max(1)).max(1)
-}
+pub use crate::compute::{workgroups, BindKind, Binding, Kernel, MAX_BINDINGS};
+use crate::compute::{check_job, padded_len, parse_kernel, result_slot, slot_for};
#[derive(Clone, PartialEq, Eq, Hash)]
struct PipelineKey {
@@ -122,10 +52,6 @@ struct Pipeline {
workgroup_size: [u32; 3],
}
-/// Storage bindings are bound whole, so a buffer's size has to be a multiple
-/// of the widest element stride a shader may declare; 16 covers `vec4<f32>`.
-const BUFFER_ALIGN: usize = 16;
-
/// A headless device that runs compute jobs. See the module docs.
pub struct ComputeDevice {
pipelines: HashMap<PipelineKey, Pipeline>,
@@ -138,9 +64,6 @@ pub struct ComputeDevice {
core: VkCore,
}
-/// The most bindings one job may carry (the descriptor pool is sized to it).
-pub const MAX_BINDINGS: usize = 16;
-
impl ComputeDevice {
/// A device on the machine's preferred GPU (`CCE_VK_DEVICE` steers it,
/// as for every renderer). `Err` when there is no usable Vulkan at all.
@@ -272,33 +195,7 @@ impl ComputeDevice {
passes: u32,
ping_pong: Option<(usize, usize)>,
) -> Result<(), String> {
- if bindings.len() > MAX_BINDINGS {
- return Err(format!("{} bindings; a job may carry at most {MAX_BINDINGS}", bindings.len()));
- }
- if groups.iter().any(|&g| g == 0) {
- return Err(format!("workgroup count {groups:?} has a zero"));
- }
- if passes == 0 {
- return Err("a job needs at least one pass".to_string());
- }
- if let Some((a, b)) = ping_pong {
- if a == b || a >= bindings.len() || b >= bindings.len() {
- return Err(format!("ping-pong pair ({a}, {b}) does not name two distinct bindings of {}", bindings.len()));
- }
- if !matches!(bindings[a], Binding::Input(_)) {
- return Err(format!("ping-pong binding {a} must be a read-only Input: it is where the first pass reads"));
- }
- if !matches!(bindings[b], Binding::Storage(_)) {
- return Err(format!("ping-pong binding {b} must be a read-write Storage: it is where the result lands"));
- }
- if bindings[a].bytes().len() != bindings[b].bytes().len() {
- return Err(format!(
- "ping-pong bindings {a} and {b} differ in length ({} vs {} bytes)",
- bindings[a].bytes().len(),
- bindings[b].bytes().len()
- ));
- }
- }
+ check_job(bindings, groups, passes, ping_pong)?;
let kinds: Vec<BindKind> = bindings.iter().map(Binding::kind).collect();
let key = PipelineKey { kernel: kernel.clone(), kinds };
let (pipeline, layout, set_layout) = {
@@ -312,7 +209,7 @@ impl ComputeDevice {
let mut sizes = Vec::with_capacity(bindings.len());
for (i, b) in bindings.iter().enumerate() {
let bytes = b.bytes();
- let padded = bytes.len().max(BUFFER_ALIGN).div_ceil(BUFFER_ALIGN) * BUFFER_ALIGN;
+ let padded = padded_len(bytes.len());
if b.kind() == BindKind::Uniform {
let cap = unsafe {
self.core.instance.get_physical_device_properties(self.core.physical_device).limits.max_uniform_buffer_range
@@ -366,17 +263,10 @@ impl ComputeDevice {
.set_layouts(&set_layouts),
)
.map_err(|e| format!("descriptor set: {e}"))?;
- let slot_for = |binding: usize, swapped: bool| -> usize {
- match ping_pong {
- Some((a, b)) if swapped && binding == a => b,
- Some((a, b)) if swapped && binding == b => a,
- _ => binding,
- }
- };
let mut infos: Vec<[vk::DescriptorBufferInfo; 1]> = Vec::with_capacity(set_count * bindings.len());
for (si, _) in sets.iter().enumerate() {
for i in 0..bindings.len() {
- let slot = slot_for(i, si == 1);
+ let slot = slot_for(i, si == 1, ping_pong);
infos.push([vk::DescriptorBufferInfo::default()
.buffer(self.slots[slot].buffer)
.offset(0)
@@ -452,15 +342,9 @@ impl ComputeDevice {
// Read back the read-write bindings — the ping-pong output from
// whichever buffer the last pass wrote.
- let last_written = |i: usize| -> usize {
- match ping_pong {
- Some((a, b)) if i == b && passes % 2 == 0 => a,
- _ => i,
- }
- };
for (i, b) in bindings.iter_mut().enumerate() {
if let Binding::Storage(out) = b {
- let mapped = self.slots[last_written(i)]
+ let mapped = self.slots[result_slot(i, passes, ping_pong)]
.allocation
.as_ref()
.and_then(|a| a.mapped_slice())
@@ -570,36 +454,15 @@ impl Drop for ComputeDevice {
/// workgroup size. `compile_wgsl` in the renderer panics on a bad shader,
/// which is right for the toolkit's own; a kernel here may be a user's.
fn compile_kernel(kernel: &Kernel) -> Result<(Vec<u32>, [u32; 3]), String> {
- let module = naga::front::wgsl::parse_str(&kernel.source)
- .map_err(|e| format!("WGSL parse error: {}", e.emit_to_string(&kernel.source).trim_end()))?;
- let entry = module
- .entry_points
- .iter()
- .find(|ep| ep.name == kernel.entry && ep.stage == naga::ShaderStage::Compute)
- .ok_or_else(|| {
- let offered: Vec<&str> = module
- .entry_points
- .iter()
- .filter(|ep| ep.stage == naga::ShaderStage::Compute)
- .map(|ep| ep.name.as_str())
- .collect();
- format!(
- "no @compute entry point named `{}`; the module offers {}",
- kernel.entry,
- if offered.is_empty() { "none".to_string() } else { offered.join(", ") }
- )
- })?;
- let workgroup_size = entry.workgroup_size;
- let info = naga::valid::Validator::new(naga::valid::ValidationFlags::all(), naga::valid::Capabilities::empty())
- .validate(&module)
- .map_err(|e| format!("WGSL validation error: {}", e.emit_to_string(&kernel.source).trim_end()))?;
+ let parsed = parse_kernel(kernel)?;
let options = naga::back::spv::Options {
lang_version: (1, 0),
flags: naga::back::spv::WriterFlags::LABEL_VARYINGS,
..Default::default()
};
- let spirv = naga::back::spv::write_vec(&module, &info, &options, None).map_err(|e| format!("SPIR-V: {e}"))?;
- Ok((spirv, workgroup_size))
+ let spirv =
+ naga::back::spv::write_vec(&parsed.module, &parsed.info, &options, None).map_err(|e| format!("SPIR-V: {e}"))?;
+ Ok((spirv, parsed.workgroup_size))
}
#[cfg(test)]
diff --git a/src/web/compute.rs b/src/web/compute.rs
new file mode 100644
index 0000000..36ca9ad
--- /dev/null
+++ b/src/web/compute.rs
@@ -0,0 +1,321 @@
+//! `crate::compute` jobs on WebGPU: the browser's `vk::ComputeDevice`.
+//!
+//! The same kernels, bindings and rules (`crate::compute`), answered the same
+//! way — a ping-pong job's result lands in its output binding, every
+//! read-write binding is read back into its slice — with one difference
+//! the platform makes: reading a buffer back is a promise, so [`run`] and its
+//! siblings are `async`. A job is still one submission, its passes in one
+//! compute pass; WebGPU orders the dispatches of a pass by itself.
+//!
+//! A kernel is parsed and validated by naga before WebGPU sees it, so a
+//! user's bad WGSL is an `Err` with naga's diagnostic, as natively; what
+//! WebGPU still rejects (a limit, a layout) is caught in a validation error
+//! scope and returned too. The device asks for the adapter's own limits on
+//! storage buffers and workgroups: WebGPU's defaults allow eight storage
+//! buffers a stage, where a job may carry sixteen. What the adapter offers
+//! is the ceiling — SwiftShader's is ten — and a job past it is an `Err`
+//! naming the limit. (cce-designer's kernels bind at most seven.)
+//!
+//! [`run`]: ComputeDevice::run
+
+use std::collections::HashMap;
+
+use wasm_bindgen::{JsCast, JsValue};
+use web_sys::{
+ gpu_buffer_usage as buffer_usage, gpu_map_mode as map_mode, gpu_shader_stage as shader_stage, GpuBindGroupDescriptor,
+ GpuBindGroupEntry, GpuBindGroupLayout, GpuBindGroupLayoutDescriptor, GpuBindGroupLayoutEntry, GpuBuffer,
+ GpuBufferBinding, GpuBufferBindingLayout, GpuBufferBindingType, GpuBufferDescriptor, GpuComputePassDescriptor,
+ GpuComputePipeline, GpuComputePipelineDescriptor, GpuDevice, GpuErrorFilter, GpuPipelineLayoutDescriptor,
+ GpuProgrammableStage, GpuQueue, GpuShaderModuleDescriptor,
+};
+
+use crate::compute::{
+ check_job, padded_len, parse_kernel, result_slot, slot_for, workgroups, BindKind, Binding, Kernel,
+};
+
+/// The limits a compute device asks the adapter for in full.
+const LIMITS: &[&str] = &[
+ "maxStorageBuffersPerShaderStage",
+ "maxUniformBuffersPerShaderStage",
+ "maxStorageBufferBindingSize",
+ "maxUniformBufferBindingSize",
+ "maxBufferSize",
+ "maxComputeWorkgroupStorageSize",
+ "maxComputeInvocationsPerWorkgroup",
+ "maxComputeWorkgroupSizeX",
+ "maxComputeWorkgroupSizeY",
+ "maxComputeWorkgroupSizeZ",
+ "maxComputeWorkgroupsPerDimension",
+];
+
+#[derive(Clone, PartialEq, Eq, Hash)]
+struct PipelineKey {
+ kernel: Kernel,
+ kinds: Vec<BindKind>,
+}
+
+struct Pipeline {
+ pipeline: GpuComputePipeline,
+ layout: GpuBindGroupLayout,
+ workgroup_size: [u32; 3],
+}
+
+/// A buffer and its size in bytes.
+struct Slot {
+ buffer: GpuBuffer,
+ size: usize,
+}
+
+/// A WebGPU device that runs compute jobs. See the module docs.
+pub struct ComputeDevice {
+ _gpu: web_sys::Gpu,
+ adapter: web_sys::GpuAdapter,
+ device: GpuDevice,
+ queue: GpuQueue,
+ pipelines: HashMap<PipelineKey, Pipeline>,
+ /// One buffer per binding index, grown when a job needs more room.
+ slots: Vec<Option<Slot>>,
+ /// The mappable copies read-write bindings are read back through.
+ readback: Vec<Option<Slot>>,
+}
+
+fn js_err(e: JsValue) -> String {
+ e.as_string()
+ .or_else(|| js_sys::Reflect::get(&e, &"message".into()).ok().and_then(|m| m.as_string()))
+ .unwrap_or_else(|| format!("{e:?}"))
+}
+
+impl ComputeDevice {
+ /// A device on the browser's WebGPU adapter. `Err` when the browser
+ /// offers none.
+ pub async fn new() -> Result<Self, String> {
+ let (gpu, adapter, device) =
+ super::request_device(LIMITS).await.map_err(|e| format!("no WebGPU compute device: {}", js_err(e)))?;
+ let queue = device.queue();
+ Ok(Self { _gpu: gpu, adapter, device, queue, pipelines: HashMap::new(), slots: Vec::new(), readback: Vec::new() })
+ }
+
+ /// The adapter as the browser describes it, for a log line or a status
+ /// readout. A browser may say little: it guards what fingerprints a machine.
+ pub fn device_name(&self) -> String {
+ let info = self.adapter.info();
+ let parts: Vec<String> =
+ [info.vendor(), info.architecture(), info.device(), info.description()].into_iter().filter(|s| !s.is_empty()).collect();
+ if parts.is_empty() { "WebGPU".into() } else { format!("WebGPU {}", parts.join(" ")) }
+ }
+
+ /// The entry point's `@workgroup_size`, compiling the kernel if needed.
+ pub async fn workgroup_size(&mut self, kernel: &Kernel, kinds: &[BindKind]) -> Result<[u32; 3], String> {
+ let key = PipelineKey { kernel: kernel.clone(), kinds: kinds.to_vec() };
+ Ok(self.pipeline(&key).await?.workgroup_size)
+ }
+
+ /// Run the kernel over `items` invocations along x, the workgroup count
+ /// from the entry point's own `@workgroup_size` (see `vk::ComputeDevice::run_over`).
+ pub async fn run_over(&mut self, kernel: &Kernel, bindings: &mut [Binding<'_>], items: u32) -> Result<(), String> {
+ let kinds: Vec<BindKind> = bindings.iter().map(Binding::kind).collect();
+ let wg = self.workgroup_size(kernel, &kinds).await?;
+ self.execute(kernel, bindings, [workgroups(items, wg[0]), 1, 1], 1, None).await
+ }
+
+ /// Upload every binding, dispatch `groups` workgroups of the kernel, and
+ /// read every [`Binding::Storage`] back into its slice once the GPU is done.
+ pub async fn run(&mut self, kernel: &Kernel, bindings: &mut [Binding<'_>], groups: [u32; 3]) -> Result<(), String> {
+ self.execute(kernel, bindings, groups, 1, None).await
+ }
+
+ /// [`run_passes`](Self::run_passes) sized over `items`, like [`run_over`](Self::run_over).
+ pub async fn run_passes_over(
+ &mut self,
+ kernel: &Kernel,
+ bindings: &mut [Binding<'_>],
+ items: u32,
+ passes: u32,
+ ping_pong: Option<(usize, usize)>,
+ ) -> Result<(), String> {
+ let kinds: Vec<BindKind> = bindings.iter().map(Binding::kind).collect();
+ let wg = self.workgroup_size(kernel, &kinds).await?;
+ self.execute(kernel, bindings, [workgroups(items, wg[0]), 1, 1], passes, ping_pong).await
+ }
+
+ /// `passes` dispatches of the kernel in one submission, with an optional
+ /// ping-pong pair — the semantics of `vk::ComputeDevice::run_passes`.
+ pub async fn run_passes(
+ &mut self,
+ kernel: &Kernel,
+ bindings: &mut [Binding<'_>],
+ groups: [u32; 3],
+ passes: u32,
+ ping_pong: Option<(usize, usize)>,
+ ) -> Result<(), String> {
+ self.execute(kernel, bindings, groups, passes, ping_pong).await
+ }
+
+ async fn execute(
+ &mut self,
+ kernel: &Kernel,
+ bindings: &mut [Binding<'_>],
+ groups: [u32; 3],
+ passes: u32,
+ ping_pong: Option<(usize, usize)>,
+ ) -> Result<(), String> {
+ check_job(bindings, groups, passes, ping_pong)?;
+ let kinds: Vec<BindKind> = bindings.iter().map(Binding::kind).collect();
+ let key = PipelineKey { kernel: kernel.clone(), kinds };
+ let (pipeline, layout) = {
+ let p = self.pipeline(&key).await?;
+ (p.pipeline.clone(), p.layout.clone())
+ };
+
+ // Buffers: one per binding index, reused when big enough. The
+ // padding is uploaded too, as zeros, so an `arrayLength` that counts
+ // it reads what it would natively.
+ let uniform_cap = self.device.limits().max_uniform_buffer_binding_size() as usize;
+ let mut sizes = Vec::with_capacity(bindings.len());
+ for (i, b) in bindings.iter().enumerate() {
+ let bytes = b.bytes();
+ let padded = padded_len(bytes.len());
+ if b.kind() == BindKind::Uniform && padded > uniform_cap {
+ return Err(format!("binding {i}: a uniform block of {} bytes exceeds the device's {uniform_cap}", bytes.len()));
+ }
+ let usage = buffer_usage::STORAGE | buffer_usage::UNIFORM | buffer_usage::COPY_SRC | buffer_usage::COPY_DST;
+ ensure(&self.device, &mut self.slots, i, padded, usage, "compute-binding")?;
+ let mut upload = bytes.to_vec();
+ upload.resize(padded, 0);
+ let slot = self.slots[i].as_ref().unwrap();
+ self.queue.write_buffer_with_u32_and_u8_slice(&slot.buffer, 0, &upload).map_err(js_err)?;
+ sizes.push(padded);
+ }
+
+ self.device.push_error_scope(GpuErrorFilter::Validation);
+ // One bind group, or two with the ping-pong pair swapped in the
+ // second, so alternate passes bind the buffers the other way round.
+ let group_count = if ping_pong.is_some() { 2 } else { 1 };
+ let mut groups_bound = Vec::with_capacity(group_count);
+ for g in 0..group_count {
+ let entries: Vec<GpuBindGroupEntry> = (0..bindings.len())
+ .map(|i| {
+ let slot = slot_for(i, g == 1, ping_pong);
+ let binding = GpuBufferBinding::new(&self.slots[slot].as_ref().unwrap().buffer);
+ binding.set_size(sizes[slot] as u32);
+ GpuBindGroupEntry::new_with_gpu_buffer_binding(i as u32, &binding)
+ })
+ .collect();
+ groups_bound.push(self.device.create_bind_group(&GpuBindGroupDescriptor::new(&entries, &layout)));
+ }
+ let encoder = self.device.create_command_encoder();
+ let pass = encoder.begin_compute_pass_with_descriptor(&GpuComputePassDescriptor::new());
+ pass.set_pipeline(&pipeline);
+ for p in 0..passes {
+ let g = if ping_pong.is_some() && p % 2 == 1 { 1 } else { 0 };
+ pass.set_bind_group(0, Some(&groups_bound[g]));
+ pass.dispatch_workgroups_with_workgroup_count_y_and_workgroup_count_z(groups[0], groups[1], groups[2]);
+ }
+ pass.end();
+
+ // Copy every read-write binding's result into a mappable buffer.
+ let mut reads = Vec::new();
+ for (i, b) in bindings.iter().enumerate() {
+ if let Binding::Storage(out) = b {
+ let from = result_slot(i, passes, ping_pong);
+ let len = padded_len(out.len());
+ ensure(&self.device, &mut self.readback, i, len, buffer_usage::MAP_READ | buffer_usage::COPY_DST, "compute-readback")?;
+ encoder
+ .copy_buffer_to_buffer_with_u32_and_u32_and_u32(
+ &self.slots[from].as_ref().unwrap().buffer,
+ 0,
+ &self.readback[i].as_ref().unwrap().buffer,
+ 0,
+ len as u32,
+ )
+ .map_err(js_err)?;
+ reads.push(i);
+ }
+ }
+ self.queue.submit(&[encoder.finish()]);
+ if let Some(err) = wasm_bindgen_futures::JsFuture::from(self.device.pop_error_scope()).await.map_err(js_err)?.dyn_ref::<web_sys::GpuError>()
+ {
+ return Err(format!("WebGPU rejected the job: {}", err.message()));
+ }
+
+ for i in reads {
+ let Binding::Storage(out) = &mut bindings[i] else { unreachable!() };
+ let buffer = &self.readback[i].as_ref().unwrap().buffer;
+ wasm_bindgen_futures::JsFuture::from(buffer.map_async_with_u32_and_u32(map_mode::READ, 0, padded_len(out.len()) as u32))
+ .await
+ .map_err(|e| format!("binding {i}: readback: {}", js_err(e)))?;
+ let mapped = js_sys::Uint8Array::new(&buffer.get_mapped_range().map_err(js_err)?.into());
+ mapped.subarray(0, out.len() as u32).copy_to(out);
+ buffer.unmap();
+ }
+ Ok(())
+ }
+
+ async fn pipeline(&mut self, key: &PipelineKey) -> Result<&Pipeline, String> {
+ if !self.pipelines.contains_key(key) {
+ let built = self.build_pipeline(key).await?;
+ self.pipelines.insert(key.clone(), built);
+ }
+ Ok(&self.pipelines[key])
+ }
+
+ async fn build_pipeline(&self, key: &PipelineKey) -> Result<Pipeline, String> {
+ let parsed = parse_kernel(&key.kernel)?;
+ // Read-only storage is its own binding type here (WebGPU tells it
+ // from read-write, which Vulkan does not): the module says which.
+ let entries: Vec<GpuBindGroupLayoutEntry> = key
+ .kinds
+ .iter()
+ .enumerate()
+ .map(|(i, k)| {
+ let ty = match k {
+ BindKind::Uniform => GpuBufferBindingType::Uniform,
+ BindKind::Storage if parsed.read_only_storage(i as u32) => GpuBufferBindingType::ReadOnlyStorage,
+ BindKind::Storage => GpuBufferBindingType::Storage,
+ };
+ let buffer = GpuBufferBindingLayout::new();
+ buffer.set_type(ty);
+ let entry = GpuBindGroupLayoutEntry::new(i as u32, shader_stage::COMPUTE);
+ entry.set_buffer(&buffer);
+ entry
+ })
+ .collect();
+ self.device.push_error_scope(GpuErrorFilter::Validation);
+ let layout = self.device.create_bind_group_layout(&GpuBindGroupLayoutDescriptor::new(&entries)).map_err(js_err)?;
+ let pipeline_layout = self.device.create_pipeline_layout(&GpuPipelineLayoutDescriptor::new(&[js_sys::JsOption::wrap(layout.clone())]));
+ let module = self.device.create_shader_module(&GpuShaderModuleDescriptor::new(&key.kernel.source));
+ let stage = GpuProgrammableStage::new(&module);
+ stage.set_entry_point(&key.kernel.entry);
+ let pipeline = self.device.create_compute_pipeline(&GpuComputePipelineDescriptor::new(&pipeline_layout, &stage));
+ if let Some(err) = wasm_bindgen_futures::JsFuture::from(self.device.pop_error_scope()).await.map_err(js_err)?.dyn_ref::<web_sys::GpuError>()
+ {
+ return Err(format!("compute pipeline: {}", err.message()));
+ }
+ Ok(Pipeline { pipeline, layout, workgroup_size: parsed.workgroup_size })
+ }
+}
+
+/// `slots[i]` holds a buffer of at least `size` bytes with `usage`.
+fn ensure(device: &GpuDevice, slots: &mut Vec<Option<Slot>>, i: usize, size: usize, usage: u32, label: &str) -> Result<(), String> {
+ if slots.len() <= i {
+ slots.resize_with(i + 1, || None);
+ }
+ if slots[i].as_ref().is_some_and(|s| s.size >= size) {
+ return Ok(());
+ }
+ if let Some(old) = slots[i].take() {
+ old.buffer.destroy();
+ }
+ let desc = GpuBufferDescriptor::new(size as u32, usage);
+ desc.set_label(label);
+ slots[i] = Some(Slot { buffer: device.create_buffer(&desc).map_err(js_err)?, size });
+ Ok(())
+}
+
+impl Drop for ComputeDevice {
+ fn drop(&mut self) {
+ for s in self.slots.iter().chain(self.readback.iter()).flatten() {
+ s.buffer.destroy();
+ }
+ }
+}
diff --git a/src/web/mod.rs b/src/web/mod.rs
index 6644aff..4f375d5 100644
--- a/src/web/mod.rs
+++ b/src/web/mod.rs
@@ -3,13 +3,55 @@
//! shaders (`draw::shaders`), atlas (`draw::glyphs`) and image queue
//! (`draw::images`), and the shell that runs an `Application` on a canvas
//! with it ([`run`]) — the page's events in through the shared `Driver`,
-//! `Pacer::turn` on animation frames.
+//! `Pacer::turn` on animation frames — and a [`ComputeDevice`] that runs
+//! `crate::compute` jobs on WebGPU.
//!
//! wasm32 only, and it needs web-sys's WebGPU bindings switched on
//! (`--cfg=web_sys_unstable_apis`, set for this crate in `.cargo/config.toml`).
+mod compute;
mod renderer;
mod shell;
+pub use compute::ComputeDevice;
pub use renderer::{Capture, PendingCapture, WebRenderer};
pub use shell::{capture, run, Fonts, Sizing};
+
+use wasm_bindgen::JsValue;
+
+/// Ask the browser for a WebGPU device, with `limits` (WebGPU limit names)
+/// raised to what the adapter offers — a device is created at the spec's
+/// defaults otherwise, which are below what most adapters can do. What the
+/// device rejects later, and why it was lost, goes to the console: a WebGPU
+/// validation error is otherwise silent.
+pub(crate) async fn request_device(
+ limits: &[&str],
+) -> Result<(web_sys::Gpu, web_sys::GpuAdapter, web_sys::GpuDevice), JsValue> {
+ let window = web_sys::window().ok_or("no window")?;
+ let gpu = window.navigator().gpu();
+ let adapter = gpu
+ .request_adapter()
+ .await?
+ .into_option()
+ .ok_or("this browser offers no WebGPU adapter")?;
+ let desc = web_sys::GpuDeviceDescriptor::new();
+ if !limits.is_empty() {
+ let offered = adapter.limits();
+ let required = js_sys::Object::new();
+ for name in limits {
+ let v = js_sys::Reflect::get(&offered, &JsValue::from_str(name))?;
+ if !v.is_undefined() {
+ js_sys::Reflect::set(&required, &JsValue::from_str(name), &v)?;
+ }
+ }
+ js_sys::Reflect::set(&desc, &JsValue::from_str("requiredLimits"), &required)?;
+ }
+ let device: web_sys::GpuDevice = adapter.request_device_with_descriptor(&desc).await?;
+ js_sys::Function::new_with_args(
+ "d",
+ "d.onuncapturederror = (e) => console.error('cce-ui WebGPU:', e.error.message); \
+ d.lost.then((i) => console.error('cce-ui WebGPU device lost:', i.reason, i.message));",
+ )
+ .call1(&JsValue::NULL, &device)?;
+ Ok((gpu, adapter, device))
+}
diff --git a/src/web/renderer.rs b/src/web/renderer.rs
index d9354fc..2bd2e8c 100644
--- a/src/web/renderer.rs
+++ b/src/web/renderer.rs
@@ -287,22 +287,7 @@ fn uniform_binding(buffer: &GpuBuffer, size: u32) -> GpuBufferBinding {
impl WebRenderer {
/// Ask the browser for a WebGPU device and set `canvas` up to draw into.
pub async fn new(canvas: HtmlCanvasElement) -> Result<Self, JsValue> {
- let window = web_sys::window().ok_or("no window")?;
- let gpu = window.navigator().gpu();
- let adapter = gpu
- .request_adapter()
- .await?
- .into_option()
- .ok_or("this browser offers no WebGPU adapter")?;
- let device: GpuDevice = adapter.request_device().await?;
- // Report what the device rejects and why it was lost, on the console:
- // a WebGPU validation error is otherwise silent.
- js_sys::Function::new_with_args(
- "d",
- "d.onuncapturederror = (e) => console.error('cce-ui WebGPU:', e.error.message); \
- d.lost.then((i) => console.error('cce-ui WebGPU device lost:', i.reason, i.message));",
- )
- .call1(&JsValue::NULL, &device)?;
+ let (gpu, adapter, device) = super::request_device(&[]).await?;
let queue = device.queue();
let context: GpuCanvasContext = canvas
.get_context("webgpu")?