Transaction-Science/open-standards
Watch
1
0
Fork
You've already forked open-standards
 0
Code Issues Pull requests Projects Releases 15 Packages Wiki Activity Actions 11

wai-quantum 0.4.2: tile runs certified on their own, and the WebGPU member in a browser #225

Merged
dcharlot merged 5 commits from wai-quantum-webgpu-tile into main 2026-10-12 15:27:57 -04:00
Conversation 0 Commits 5 Files changed 30 +4079 -240
30 changed files with 4079 additions and 240 deletions
Download patch file Download diff file Expand all files Collapse all files

27
.github/workflows/wai-conformance.yml vendored
View file

@ -501,7 +501,10 @@ jobs:
quantum-wasip2:
name: quantum core and corpus under wasm32-wasip2 (wasmtime)
runs-on: linux
timeout-minutes: 60
# The fabric's planner test runs fabric-v1 under every planner configuration in a
# debug wasm build, several minutes on its own; the timeout leaves it room rather
# than the test being skipped on wasm32.
timeout-minutes: 90
steps:
- uses: actions/checkout@v5
@ -562,10 +565,10 @@ jobs:
done
# The fabric's groundwork for accelerator members (certification one case at
# a time, the fabric-v1 corpus and its pins, the 32-bit-lane words and their
# range guard, the program a member runs, the WGSL kernel's host layout,
# member ids and mutants), each by name: a module filter is satisfied by any
# one test of the module, so a test dropped or compiled out on wasm32 would
# otherwise go unnoticed.
# range guard, the program a member runs in every planner configuration, the
# WGSL kernels' host layouts, member ids and mutants), each by name: a module
# filter is satisfied by any one test of the module, so a test dropped or
# compiled out on wasm32 would otherwise go unnoticed.
for t in quantum_fabric::tests::certification_is_certify \
quantum_fabric::tests::fxmul_wide_is_the_exact_product \
quantum_fabric::tests::apply_words_is_apply \
@ -580,6 +583,12 @@ jobs:
quantum_fabric::tests::fabric_v1_covers_every_dispatch_shape \
quantum_fabric::tests::gate_words_are_the_table \
quantum_fabric::plan::tests::execute_is_the_native_simulator \
quantum_fabric::plan::tests::every_configuration_is_the_native_simulator \
quantum_fabric::plan::tests::long_circuits_tile_as_they_run_gate_by_gate \
quantum_fabric::plan::tests::tiles_are_exact_on_any_words \
quantum_fabric::plan::tests::tiles_cover_every_path \
quantum_fabric::plan::tests::the_v1_tail_phases_are_not_skipped \
quantum_fabric::plan::tests::malformed_programs_are_refused \
quantum_fabric::wgsl::tests::implementation_id_is_pinned \
quantum_fabric::wgsl::tests::gate_record_layout \
quantum_fabric::wgsl::tests::grids_cover_every_item_once \
@ -587,7 +596,13 @@ jobs:
quantum_fabric::wgsl::tests::member_ids_are_pinned \
quantum_fabric::wgsl::tests::the_cpu_reference_is_placed_first_when_nothing_is_metered \
quantum_fabric::wgsl::tests::mutants_apply_once \
quantum_fabric::wgsl::tests::mutant_witnesses_exist; do
quantum_fabric::wgsl::tests::mutant_witnesses_exist \
quantum_fabric::wgsl::tests::placement_breaks_an_implementation_tie_by_name \
quantum_fabric::wgsl::tests::tile_implementation_id_is_pinned \
quantum_fabric::wgsl::tests::tile_records_layout \
quantum_fabric::wgsl::tests::tile_grids_cover_every_tile_once \
quantum_fabric::wgsl::tests::tile_mutants_apply_once \
quantum_fabric::wgsl::tests::tile_mutant_witnesses_exist; do
grep -q "^test $t \.\.\. ok$" "$RUNNER_TEMP/wasip2-lib.log" \
|| { echo "::error::$t did not run"; exit 1; }
done

88
.github/workflows/wai-quantum-webgpu.yml vendored
View file

@ -1,7 +1,10 @@
name: wai-quantum WebGPU member
# The fabric's WebGPU member: one WGSL gate kernel, compiled for every adapter a
# runner has, certified against the CPU reference's pinned corpus digests.
# runner has, certified against the CPU reference's pinned corpus digests. A
# second implementation, tile runs (the same kernel followed by apply_tile, which
# applies a run of gates tile by tile in one dispatch), is certified on its own
# against the same digests, on every adapter and in the browser.
#
# Each member job declares the adapters its runner must have in
# WAI_QUANTUM_WEBGPU_EXPECT. The device test fails if the adapters found differ,
@ -14,6 +17,12 @@ name: wai-quantum WebGPU member
# --no-default-features: its default feature, `codecs`, links system codec
# libraries that neither GPU host has, and the corpus needs none of them.
#
# The browser job runs the same kernel through a browser's WebGPU, compiled by the
# browser's own WGSL compiler: wai-web's `quantum_webgpu` bundle in a headless page
# (wai/quantum-conformance/browser_webgpu.mjs), which must certify a GPU member with
# the four pinned digests, reproduce every vector and refuse every reject file with
# its class. It is not a required status until it has been green three runs in a row.
#
# There is no windows job while the windows runner is offline: a job queued on an
# offline runner never settles. Add it back from the comment at the end when the
# runner has survived a reboot.
@ -35,7 +44,11 @@ on:
- 'wai/wai-rs/Cargo.lock'
- 'wai/wai-rs/src/lib.rs'
- 'rust-toolchain.toml'
# The checks and browser jobs build wai-web with its feature, --locked.
- 'wai/wai-web/Cargo.toml'
- 'wai/wai-web/Cargo.lock'
- 'wai/wai-web/src/quantum_webgpu.rs'
- 'wai/wai-web/src/lib.rs'
- 'wai/wai-web/demo/quantum_webgpu*'
- 'scripts/wai-quantum-webgpu-member.sh'
- '.github/workflows/wai-quantum-webgpu.yml'
@ -50,7 +63,11 @@ on:
- 'wai/wai-rs/Cargo.lock'
- 'wai/wai-rs/src/lib.rs'
- 'rust-toolchain.toml'
# The checks and browser jobs build wai-web with its feature, --locked.
- 'wai/wai-web/Cargo.toml'
- 'wai/wai-web/Cargo.lock'
- 'wai/wai-web/src/quantum_webgpu.rs'
- 'wai/wai-web/src/lib.rs'
- 'wai/wai-web/demo/quantum_webgpu*'
- 'scripts/wai-quantum-webgpu-member.sh'
- '.github/workflows/wai-quantum-webgpu.yml'
@ -80,7 +97,7 @@ jobs:
SUDO=""; [ "$(id -u)" -ne 0 ] && SUDO="sudo"
$SUDO apt-get update -qq
$SUDO apt-get install -y -qq --no-install-recommends spirv-tools
- name: The shader under naga 30.0.1
- name: The shaders under naga 30.0.1
shell: bash
working-directory: wai/tools/quantum-wgsl-check
run: |
@ -88,9 +105,13 @@ jobs:
cargo run --locked --release -- "$RUNNER_TEMP" 2>&1 | tee "$RUNNER_TEMP/wgsl-check.log"
grep -qx "quantum-wgsl-check: ok" "$RUNNER_TEMP/wgsl-check.log"
spirv-val --target-env vulkan1.1 "$RUNNER_TEMP/lanes32.spv"
spirv-val --target-env vulkan1.1 "$RUNNER_TEMP/lanes32_tile.spv"
- name: The member builds for a browser
working-directory: wai/wai-quantum
run: cargo build --locked --lib --target wasm32-unknown-unknown --no-default-features --features quantum_fabric_webgpu
- name: The browser bundle builds with the member
working-directory: wai/wai-web
run: cargo build --locked --release --target wasm32-unknown-unknown --features quantum_webgpu
- name: On wasm32-wasip2 the feature is inert
shell: bash
working-directory: wai/wai-quantum
@ -110,6 +131,7 @@ jobs:
set -euo pipefail
cargo package --locked --list --allow-dirty > "$RUNNER_TEMP/package.txt"
grep -qx 'src/quantum_fabric/lanes32.wgsl' "$RUNNER_TEMP/package.txt"
grep -qx 'src/quantum_fabric/lanes32_tile.wgsl' "$RUNNER_TEMP/package.txt"
- name: MSRV 1.97.0
working-directory: wai/wai-quantum
run: cargo +1.97.0 check --locked --lib --features quantum_fabric_webgpu
@ -210,6 +232,68 @@ jobs:
working-directory: src/wai/wai-rs
run: CARGO_TARGET_DIR="$HOME/hw-target/wai-rs" nice -n 10 cargo run --locked --no-default-features --features quantum_fabric_webgpu --bin wai_quantum_vectors -- verify ../quantum-conformance --webgpu
browser:
name: member · browser WebGPU (macos)
runs-on: macos
timeout-minutes: 60
steps:
- uses: actions/checkout@v5
# launchd's PATH omits cargo on the macOS host runner.
- name: Put cargo on PATH
shell: bash
run: echo "$HOME/.cargo/bin" >> "$GITHUB_PATH"
# The page's file is held to the crate by wai-quantum's tests/fabric_expected.rs:
# exactly the four certificates of the pinned corpora (gateset, corpus, seed,
# cases, wide circuits and the digest the CPU computes) and both implementation
# ids. Editing the file alone, say giving fabric-v1 random-v1's digests so that
# no circuit wider than one tile runs, fails here.
- name: The page's pins are the crate's
shell: bash
working-directory: wai/wai-quantum
run: |
set -euo pipefail
nice -n 10 cargo test --locked --features quantum_fabric --test fabric_expected 2>&1 | tee "$RUNNER_TEMP/fabric-expected.log"
grep -q '^test the_browser_legs_pins_are_the_crates \.\.\. ok$' "$RUNNER_TEMP/fabric-expected.log" \
|| { echo "::error::the pins test did not run"; exit 1; }
# Pinned to the wasm-bindgen version in wai-web/Cargo.lock: the CLI and the crate
# must match exactly.
- name: wasm-bindgen CLI 0.2.127
shell: bash
run: |
set -euo pipefail
rustup target add wasm32-unknown-unknown
cargo install --locked wasm-bindgen-cli --version 0.2.127
- name: The member's bundle
shell: bash
working-directory: wai/wai-web
run: |
set -euo pipefail
nice -n 10 cargo build --locked --release --target wasm32-unknown-unknown --features quantum_webgpu
wasm-bindgen --target web --out-dir demo/pkg-webgpu --out-name wai_quantum_webgpu \
target/wasm32-unknown-unknown/release/wai_web.wasm
# Playwright and its headless browser go into the job's temporary directory, not
# the host's caches. `chromium` is Playwright's name for its engine.
- name: Playwright 1.60.0 and its headless browser
shell: bash
run: |
set -euo pipefail
node --version
pw="$RUNNER_TEMP/playwright"
mkdir -p "$pw"
npm install --prefix "$pw" --no-save --no-audit --no-fund [email protected]
PLAYWRIGHT_BROWSERS_PATH="$pw/browsers" "$pw/node_modules/.bin/playwright" install --only-shell chromium
echo "PLAYWRIGHT_HOST=$pw/" >> "$GITHUB_ENV"
echo "PLAYWRIGHT_BROWSERS_PATH=$pw/browsers" >> "$GITHUB_ENV"
- name: The member in a browser
shell: bash
run: |
set -euo pipefail
WAI_QUANTUM_WEBGPU_REQUIRE=1 nice -n 10 node wai/quantum-conformance/browser_webgpu.mjs 2>&1 \
| tee "$RUNNER_TEMP/browser-webgpu.log"
grep -qE '^browser webgpu: member=local/gpu-webgpu-[0-9a-f]{12} class=gpu certified bit-exact 4 of 4; tiles certified bit-exact 4 of 4; ' \
"$RUNNER_TEMP/browser-webgpu.log" \
|| { echo "::error::no GPU member certified bit-exact 4 of 4, per-gate and in tile runs, in the browser"; exit 1; }
# Windows, when its runner is back (DX12, translator naga+dxbc, and Vulkan, as separate members):
# member-windows:
# runs-on: windows

26
scripts/wai-quantum-webgpu-member.sh
View file

@ -2,11 +2,13 @@
#
# wai-quantum-webgpu-member.sh
#
# Runs the WebGPU member's device tests and fails unless (1) all of them ran and
# passed, (2) the adapters found are exactly WAI_QUANTUM_WEBGPU_EXPECT, (3) each
# declared adapter earned one bit-exact certificate per pinned corpus digest, and
# (4) each declared adapter killed every shader mutant. These checks are what make
# a filtered, skipped or adapter-less run fail instead of pass.
# Runs the WebGPU member's device tests and fails unless (1) all 13 of them ran
# and passed, (2) the adapters found are exactly WAI_QUANTUM_WEBGPU_EXPECT, (3) each
# declared adapter earned one bit-exact certificate per pinned corpus digest under
# each implementation (the per-gate kernel, wgsl-lanes32/, and tile runs,
# wgsl-lanes32-tile/), and (4) each declared adapter killed every shader mutant of
# both. These checks are what make a filtered, skipped or adapter-less run fail
# instead of pass.
#
# WAI_QUANTUM_WEBGPU_EXPECT=<key[,key...]|none> [WAI_QUANTUM_WEBGPU_WIDE=1] \
# scripts/wai-quantum-webgpu-member.sh
@ -21,8 +23,8 @@ export WAI_QUANTUM_WEBGPU_REPORT="$dir/webgpu-member.report"
cargo test --locked --features quantum_fabric_webgpu --test webgpu_member \
-- --nocapture --test-threads=1 2>&1 | tee "$log"
grep -q "test result: ok. 12 passed; 0 failed; 0 ignored; 0 measured; 0 filtered out" "$log" \
|| { echo "::error::the member's 12 tests did not all run and pass"; exit 1; }
grep -q "test result: ok. 13 passed; 0 failed; 0 ignored; 0 measured; 0 filtered out" "$log" \
|| { echo "::error::the member's 13 tests did not all run and pass"; exit 1; }
grep -qxF "webgpu adapters: as declared ($WAI_QUANTUM_WEBGPU_EXPECT)" "$WAI_QUANTUM_WEBGPU_REPORT" \
|| { echo "::error::the adapters found are not the declared set"; exit 1; }
@ -37,13 +39,19 @@ for key in $(tr ',' '\n' <<<"$WAI_QUANTUM_WEBGPU_EXPECT" | sort -u); do
want=$(tr ',' '\n' <<<"$WAI_QUANTUM_WEBGPU_EXPECT" | grep -cxF "$key")
for p in "${pins[@]}"; do
set -- $p
n=$(grep -cE "^webgpu certified: key=$key member=local/[a-z]+-[a-z0-9]+-[0-9a-f]{12} kernel=$1 corpus=$2 implementation=wgsl-lanes32/[0-9a-f]{16} parity=bit-exact cases=[0-9]+ digest=$3$" "$WAI_QUANTUM_WEBGPU_REPORT" || true)
[ "$n" -eq "$want" ] || { echo "::error::$key: $n bit-exact $2 certificates for $1, $want expected"; exit 1; }
for implementation in wgsl-lanes32 wgsl-lanes32-tile; do
n=$(grep -cE "^webgpu certified: key=$key member=local/[a-z]+-[a-z0-9]+-[0-9a-f]{12} kernel=$1 corpus=$2 implementation=$implementation/[0-9a-f]{16} parity=bit-exact cases=[0-9]+ digest=$3$" "$WAI_QUANTUM_WEBGPU_REPORT" || true)
[ "$n" -eq "$want" ] || { echo "::error::$key: $n bit-exact $2 certificates for $1 under $implementation, $want expected"; exit 1; }
done
done
n=$(grep -cxF "webgpu mutants: key=$key killed 7 of 7" "$WAI_QUANTUM_WEBGPU_REPORT" || true)
[ "$n" -eq "$want" ] || { echo "::error::$key: a mutant survived, or the mutant test did not run"; exit 1; }
n=$(grep -cxF "webgpu tile mutants: key=$key killed 3 of 3" "$WAI_QUANTUM_WEBGPU_REPORT" || true)
[ "$n" -eq "$want" ] || { echo "::error::$key: a tile mutant survived, or the mutant test did not run"; exit 1; }
if [ "${WAI_QUANTUM_WEBGPU_WIDE:-0}" = 1 ] && [ "${key%%-*}" = gpu ]; then
grep -qE "^webgpu wide: key=$key n=2[56] matches native$" "$WAI_QUANTUM_WEBGPU_REPORT" \
|| { echo "::error::$key: no wide run at the member's full size"; exit 1; }
grep -qE "^webgpu tile wide: key=$key n=2[56] matches native$" "$WAI_QUANTUM_WEBGPU_REPORT" \
|| { echo "::error::$key: no wide tile run at the member's full size"; exit 1; }
fi
done

2
wai/.gitignore vendored
View file

@ -51,6 +51,8 @@ wai-web/demo/scpkg/
# wasm-pack build outputs (rebuilt by wasm-pack / demo/verify_int_synth.sh).
wai-web/demo/pkg/
wai-web/demo/pkg-node/
# The WebGPU member's bundle (demo/quantum_webgpu.html; built with --features quantum_webgpu).
wai-web/demo/pkg-webgpu/
# int_hyper conditional-codec demo artifacts (PTQ model + bitstream, large —
# regenerated by tools/wai_sota_int_synth_export.py into demo/int_hyper/).
wai-web/demo/int_hyper/

43
wai/CHANGELOG.md
View file

@ -6,6 +6,49 @@ wire format may change and such changes are recorded here.
## v0.4.0 — current
### wai-quantum 0.4.2: tile runs, and the WebGPU member in a browser
Additive. The default features, `full`, the per-gate implementation and every existing pin
and digest are unchanged.
- **Tile runs: a second implementation, `wgsl-lanes32-tile/54e443520c020013`.**
`plan::Config` (`PER_GATE`, `TILED`; `tile_size`, `run_cap`), `plan::Step::Tile`,
`compile_with`, `schedule`, `execute_from`, `PlannedGate::{new, is_identity,
is_well_formed}` and `TILE_BITS_MAX`; `PLAN_REVISION` is 2. A tiled plan applies each
maximal run of consecutive gates that act inside a tile (diagonal gates on any target,
others with their target inside the tile) to tiles of up to 1024 amplitudes, in
op-log order, and leaves out exact identities only (circuit/1's `P(19..=32)` is
`2^30 − 1`, not the identity, and stays). The CPU executor mirrors a tile kernel's
index arithmetic and equals the per-gate program on any words; every configuration
certifies with the lanes32 pins. A program whose gate's `diag` disagrees with its
words, or whose op index is `u32::MAX`, is refused by the executor and by the member.
`wgsl::TILE_SHADER` (`lanes32.wgsl` followed by `lanes32_tile.wgsl`, entry point
`apply_tile`: one workgroup of 256 per tile, the tile updated in place in the state
planes, a storage barrier between gates, the tile size in the run record, never in
the text; a run record whose tiles or grid do not cover the state exactly once is
refused) has its own id (`tile_implementation_id`), records (`tile_gate_record`,
`tile_run_record`), grid (`tile_grid`) and mutants (`TILE_MUTANTS`, with the
detectors `Race` and `Malformed`, `race_witness`, `malformed_tile_run`).
`WebGpuMember::tiled()` gives the same device running it, probed and certified on its
own; the member splits a run into dispatches of at most `SUBMIT_BUDGET` pair-updates.
Discovery opens per-gate members, which are the fallback. Given both implementations'
entries for one member and no energy figures, `place` breaks the tie by
implementation name and places the tile implementation (`wgsl-lanes32-tile/…` sorts
first). The checker holds `TILE_SHADER` to its own rules, barrier placement and
barrier uniformity included, and its SPIR-V to the validator; the member's device
tests, the member script and the browser leg require the four pinned digests under
both implementations and every tile mutant caught. `tests/fabric_expected.rs` (not
in the package) holds `quantum-conformance/fabric/expected.json` to the crate.
- **The member in a browser.** wai-web's feature `quantum_webgpu` binds the member for JS
(`quantum_webgpu_discover`, `QuantumWebGpu`: `certify`, `statevector_hash_hex`) in a
bundle of its own (`--out-name wai_quantum_webgpu`); the default build and its bundle
are unchanged. `demo/quantum_webgpu.html` runs the member in a worker, its kernel
compiled by the browser's own WGSL compiler. `quantum-conformance/browser_webgpu.mjs`
drives the page headless and requires the four certificates bit-exact with the pinned
digests (`quantum-conformance/fabric/expected.json`), every vector's statevector hash,
every reject file refused with its class, and a circuit wider than the member holds
refused. The `checks` job builds the bundle, and a `browser` job runs the page on macOS.
### wai-quantum 0.4.1: the fabric's WebGPU member
Additive. The default features, `full`, and every existing pin and digest are unchanged.

19
wai/quantum-conformance/README.md
View file

@ -57,6 +57,25 @@ matching manifest. It is run by hand, before the website deploys a bundle. Note
the IDE's `statevector_hash_hex()` is empty until `run()` is called — the circuit is
built eagerly, the state is not.
## Browser, on the fabric's WebGPU member
```sh
WAI_QUANTUM_WEBGPU_REQUIRE=1 node wai/quantum-conformance/browser_webgpu.mjs
```
Serves `wai-web/demo` and this directory on 127.0.0.1 and opens
`quantum_webgpu.html` in Playwright's headless browser, with the `quantum_webgpu`
bundle built into `wai-web/demo/pkg-webgpu` (`wai-web/demo/README.md`). The gate
kernel runs on the page's WebGPU adapter, compiled by the browser's own WGSL
compiler. The page must certify the kernel bit-exact for both gatesets on random-v1
and fabric-v1 with the digests `fabric/expected.json` pins (the CPU reference's), and
run every vector of both trees to its statevector hash, on the per-gate implementation
and again on the tile implementation `fabric/expected.json` names (the browser's
compiler holds its barrier to WGSL's uniformity rules); refuse every reject file with
its listed class; and refuse a circuit wider than the member holds. Without
`WAI_QUANTUM_WEBGPU_REQUIRE=1`, a page offered no adapter is reported and passes. CI
runs it on macOS (`wai-quantum-webgpu.yml`, job `browser`).
## Referee
```sh

242
wai/quantum-conformance/browser_webgpu.mjs Normal file
View file

@ -0,0 +1,242 @@
// The fabric's WebGPU member in a browser, held to the CPU reference's pins.
//
// The native members are checked by scripts/wai-quantum-webgpu-member.sh. This
// checks the browser's: the same WGSL kernel compiled by the browser's own WGSL
// compiler, run in a worker by wai-web/demo/quantum_webgpu.html with the bundle
// built by
//
// cd wai/wai-web
// cargo build --locked --release --target wasm32-unknown-unknown --features quantum_webgpu
// wasm-bindgen --target web --out-dir demo/pkg-webgpu --out-name wai_quantum_webgpu \
// target/wasm32-unknown-unknown/release/wai_web.wasm
//
// It serves wai-web/demo and quantum-conformance on 127.0.0.1, opens the page in
// Playwright's headless browser and requires:
//
// - an adapter opened as a member, its arithmetic probed in the page, running the
// implementation fabric/expected.json names;
// - the four certificates fabric/expected.json lists (both gatesets, random-v1 and
// fabric-v1), each bit-exact, on that member, for that implementation, over the
// listed number of cases, with the pinned corpus digest;
// - every vector of vectors/ and v2/vectors/ run on the member to its manifest's
// statevector hash;
// - the same adapter running the tile implementation fabric/expected.json names
// (tile runs: the browser's WGSL compiler holds its barrier to the uniformity
// rules, which the native translator does not check), with the same four
// certificates and every vector;
// - every file of reject/ refused with its listed class (`WQC parse (<class>): ...`);
// - a circuit one qubit wider than the member holds refused: as too-wide, or as
// too-many-qubits where the member holds the reference's maximum; and, with the
// same adapter offering only the WebGPU default buffer limits, a member holding
// 25 qubits that refuses 26 as too-wide.
//
// On success it prints
//
// browser webgpu: member=<id> class=<c> certified bit-exact 4 of 4; tiles certified bit-exact 4 of 4; vectors <n> of <n>; tile vectors <n> of <n>; rejects <m> of <m>
//
// With no adapter it prints `browser webgpu: no adapter (<why>)`, which is a
// failure when WAI_QUANTUM_WEBGPU_REQUIRE=1 and not otherwise.
//
// [WAI_QUANTUM_WEBGPU_REQUIRE=1] [PLAYWRIGHT_HOST=<dir with node_modules/playwright>] \
// node wai/quantum-conformance/browser_webgpu.mjs
//
// Headless, a page is offered no hardware adapter unless the browser draws through
// the platform's GPU interface: on macOS that is `--use-angle=metal`, kept below.
// WAI_QUANTUM_WEBGPU_BROWSER_ARGS (space-separated) replaces the launch arguments.
import http from "node:http";
import fs from "node:fs";
import path from "node:path";
import { fileURLToPath } from "node:url";
import { createRequire } from "node:module";
const HERE = path.dirname(fileURLToPath(import.meta.url));
const DEMO = path.resolve(HERE, "../wai-web/demo");
const PLAYWRIGHT_HOST = process.env.PLAYWRIGHT_HOST ||
"/Users/dcharlot/data-share/vibe-coding/joulesperbit/";
const REQUIRE = process.env.WAI_QUANTUM_WEBGPU_REQUIRE === "1";
const PORT = Number(process.env.PORT || 8795);
// The certificates take the longest: a few hundred circuits per gateset, each run
// on the device and read back.
const TIMEOUT_MS = Number(process.env.WAI_QUANTUM_WEBGPU_TIMEOUT_MS || 30 * 60 * 1000);
const ARGS = process.env.WAI_QUANTUM_WEBGPU_BROWSER_ARGS !== undefined
? process.env.WAI_QUANTUM_WEBGPU_BROWSER_ARGS.split(/\s+/).filter(Boolean)
: process.platform === "darwin" ? ["--use-angle=metal"] : [];
const { chromium } = createRequire(PLAYWRIGHT_HOST)("playwright");
const readJson = (p) => JSON.parse(fs.readFileSync(p, "utf8"));
// What the page must report, read from the committed corpus.
const EXPECTED = readJson(path.join(HERE, "fabric", "expected.json"));
const VECTORS = {};
for (const prefix of ["", "v2/"]) {
for (const [file, e] of Object.entries(readJson(path.join(HERE, prefix, "expected_hashes.json")))) {
VECTORS[prefix + file] = e.statevector_hash;
}
}
const REJECTS = Object.fromEntries(
Object.entries(readJson(path.join(HERE, "reject", "reject.json"))).map(([file, kind]) => ["reject/" + file, kind]));
const MAX_QUBITS = 26; // wai::quantum::MAX_QUBITS
const DEFAULT_LIMITS_QUBITS = 25; // two planes of 2^25 u32 in a 128 MiB binding
// Two roots, at the paths they have under wai/, so the page's default corpus
// location (../../quantum-conformance/) holds here as it does under any static
// server run from wai/.
const ROOTS = { "/wai-web/demo/": DEMO, "/quantum-conformance/": HERE };
const MIME = { ".js": "text/javascript", ".mjs": "text/javascript", ".wasm": "application/wasm", ".html": "text/html", ".json": "application/json" };
const server = http.createServer((req, res) => {
const url = decodeURIComponent(req.url.split("?")[0]);
const prefix = Object.keys(ROOTS).find((p) => url.startsWith(p));
const root = prefix && ROOTS[prefix];
const f = root && path.resolve(root, "." + url.slice(prefix.length - 1));
if (!f || !f.startsWith(root + path.sep) || !fs.existsSync(f) || !fs.statSync(f).isFile()) {
res.writeHead(404);
return res.end("not found");
}
res.writeHead(200, { "content-type": MIME[path.extname(f)] || "application/octet-stream" });
fs.createReadStream(f).pipe(res);
});
let browser;
let fail = 0;
const report = (label, problem) => {
if (problem) {
fail++;
console.log(`FAIL ${label}: ${problem}`);
} else {
console.log(`PASS ${label}`);
}
};
process.on("unhandledRejection", (e) => { console.error("UNHANDLED:", e); process.exitCode = 1; });
try {
if (!fs.existsSync(path.join(DEMO, "pkg-webgpu", "wai_quantum_webgpu.js"))) {
console.error(`no pkg-webgpu/wai_quantum_webgpu.js under ${DEMO}; build it as this file's header says`);
process.exit(2);
}
await new Promise((r) => server.listen(PORT, "127.0.0.1", r));
browser = await chromium.launch({ args: ARGS });
const page = await browser.newPage();
const errs = [];
page.on("console", (m) => { if (m.type() === "error") errs.push(m.text()); });
page.on("pageerror", (e) => errs.push(String(e)));
// domcontentloaded, never networkidle: the page's work runs in a worker after
// load, and networkidle stalls on some hosts.
await page.goto(`http://127.0.0.1:${PORT}/wai-web/demo/quantum_webgpu.html`, { waitUntil: "domcontentloaded", timeout: 30000 });
await page.waitForFunction(() => window.quantumWebGpu && window.quantumWebGpu.done === true, null, { timeout: TIMEOUT_MS, polling: 1000 });
const messages = await page.evaluate(() => window.quantumWebGpu.messages);
const of = (type, limits) => messages.filter((m) => m.type === type && (limits === undefined || m.limits === limits));
for (const m of messages.filter((m) => m.type === "error" || m.type === "refused")) {
report(`the page (${m.type})`, m.error || m.why);
}
const none = of("no-adapter")[0];
const member = of("member", "offered")[0];
if (!member) {
const why = none ? none.why : "the page reported no member and no reason";
console.log(`browser webgpu: no adapter (${why})`);
if (REQUIRE || !none || fail) process.exitCode = 1;
} else {
// The member.
console.log(`member ${member.id}: ${member.detail}`);
report(`member ${member.id} (class ${member.class})`,
!/^local\/(gpu|swgpu)-webgpu-[0-9a-f]{12}$/.test(member.id) ? "not a WebGPU member id"
: (member.class === "cpu") !== member.id.startsWith("local/swgpu-") ? "its class and id disagree"
: null);
report("the probe passed in the page", member.detail.includes("probe passed in this process") ? null : `detail: ${member.detail}`);
report(`implementation ${member.implementation}`,
member.implementation === EXPECTED.implementation ? null : `fabric/expected.json pins ${EXPECTED.implementation}`);
// The tile member: the same adapter, the pinned tile implementation.
const tiled = of("tiled")[0];
report(`the tile member runs ${EXPECTED.tile_implementation}`,
!tiled ? "not reported" : tiled.error ? `refused: ${tiled.error}`
: tiled.id !== member.id ? `member ${tiled.id}`
: tiled.implementation !== EXPECTED.tile_implementation ? `implementation ${tiled.implementation}`
: null);
// The four certificates, per implementation.
const certify = (tile, implementation) => {
let certified = 0;
for (const want of EXPECTED.certificates) {
const label = `${tile ? "tile " : ""}${want.kernel} ${want.corpus}`;
const got = of("certificate").find((m) => m.label === label && !!m.tile === tile);
const c = got && got.got;
const generator = ` seed=${want.seed} cases=${want.cases}` + (want.corpus === "fabric-v1" ? ` wide=${want.wide}` : "");
const problem = !got ? "not reported"
: got.error ? `refused: ${got.error}`
: c.parity !== "bit-exact" ? `parity ${c.parity}`
: c.kernel !== want.kernel ? `kernel ${c.kernel}`
: c.implementation !== implementation ? `implementation ${c.implementation}`
: c.member !== member.id ? `issued to ${c.member}`
: c.cases !== want.cases_compared ? `${c.cases} cases, not ${want.cases_compared}`
: !c.generator.endsWith(generator) ? `generator ${c.generator}`
: c.corpus_hex !== want.corpus_hex ? `corpus ${c.corpus_hex}, pinned ${want.corpus_hex}`
: null;
report(`certificate ${label} bit-exact with the pinned digest`, problem);
if (!problem) certified++;
}
return certified;
};
const certified = certify(false, member.implementation);
const tileCertified = certify(true, EXPECTED.tile_implementation);
// Every vector, per implementation.
const run = (tile) => {
let vectors = 0;
for (const [label, want] of Object.entries(VECTORS)) {
const got = of("vector").find((m) => m.label === label && !!m.tile === tile);
const problem = !got ? "not reported" : got.error ? `refused: ${got.error}` : got.got !== want ? `browser=${got.got} corpus=${want}` : null;
report(`${tile ? "tile " : ""}${label}`, problem);
if (!problem) vectors++;
}
return vectors;
};
const vectors = run(false);
const tileVectors = run(true);
// Every reject file.
let rejects = 0;
for (const [label, kind] of Object.entries(REJECTS)) {
const got = of("reject").find((m) => m.label === label);
const m = got && got.error && /^WQC parse \(([a-z-]+)\): /.exec(got.error);
const problem = !got ? "not reported" : !got.error ? "accepted" : !m || m[1] !== kind ? `refused as: ${got.error}` : null;
report(`${label} refused as ${kind}`, problem);
if (!problem) rejects++;
}
// One qubit too wide, on the limits offered and on the WebGPU defaults.
const wide = of("wide", "offered")[0];
const wantClass = member.max_qubits + 1 > MAX_QUBITS ? "too-many-qubits" : "too-wide";
report(`${member.max_qubits + 1} qubits on the member refused as ${wantClass}`,
!wide ? "not reported" : !wide.error ? "accepted" : !wide.error.startsWith(`simulate (${wantClass}): `) ? `refused as: ${wide.error}` : null);
const capped = of("member", "webgpu-defaults")[0];
const cappedWide = of("wide", "webgpu-defaults")[0];
if (capped) console.log(`member ${capped.id} (the same adapter on the WebGPU default limits): ${capped.detail}`);
report(`on the WebGPU default limits the member holds ${DEFAULT_LIMITS_QUBITS} qubits and refuses ${DEFAULT_LIMITS_QUBITS + 1} as too-wide`,
!capped ? "no member on the default limits"
: capped.max_qubits !== DEFAULT_LIMITS_QUBITS ? `it holds ${capped.max_qubits}`
: !cappedWide ? "not reported"
: !cappedWide.error ? "accepted"
: !cappedWide.error.startsWith(`simulate (too-wide): ${DEFAULT_LIMITS_QUBITS + 1} qubits; this member holds at most ${DEFAULT_LIMITS_QUBITS}`) ? `refused as: ${cappedWide.error}`
: null);
report("console errors", errs.length ? `${errs.length}\n ${errs.join("\n ")}` : null);
const n = Object.keys(VECTORS).length;
const r = Object.keys(REJECTS).length;
if (fail === 0) {
const k = EXPECTED.certificates.length;
console.log(`browser webgpu: member=${member.id} class=${member.class} certified bit-exact ${certified} of ${k}; tiles certified bit-exact ${tileCertified} of ${k}; vectors ${vectors} of ${n}; tile vectors ${tileVectors} of ${n}; rejects ${rejects} of ${r}`);
} else {
console.log(`browser webgpu: ${fail} checks failed`);
}
process.exitCode = fail ? 1 : 0;
}
} catch (e) {
console.error("HARNESS ERROR:", e);
process.exitCode = 1;
} finally {
if (browser) await browser.close();
server.close();
}

46
wai/quantum-conformance/fabric/expected.json Normal file
View file

@ -0,0 +1,46 @@
{
"implementation": "wgsl-lanes32/dc8e19f61e0e87ab",
"tile_implementation": "wgsl-lanes32-tile/54e443520c020013",
"certificates": [
{
"gateset": 1,
"kernel": "wai.quantum.circuit/apply",
"corpus": "random-v1",
"seed": 7,
"cases": 200,
"wide": 0,
"cases_compared": 200,
"corpus_hex": "d18454c639acfd9c7b66450f4ded1aac4ab46cc10a7bedbcf4958a96e204fb36"
},
{
"gateset": 2,
"kernel": "wai.quantum.circuit2/apply",
"corpus": "random-v1",
"seed": 7,
"cases": 200,
"wide": 0,
"cases_compared": 200,
"corpus_hex": "3dbde276c1095bf827faec71d5980e050f8a55a5ba7d97cd3bbdaec3f43a13b2"
},
{
"gateset": 1,
"kernel": "wai.quantum.circuit/apply",
"corpus": "fabric-v1",
"seed": 7,
"cases": 200,
"wide": 24,
"cases_compared": 224,
"corpus_hex": "41c726a0ac80f7defef92e15ecf893bdb7a3076b5e7f670fa0b7142e5c033df7"
},
{
"gateset": 2,
"kernel": "wai.quantum.circuit2/apply",
"corpus": "fabric-v1",
"seed": 7,
"cases": 200,
"wide": 24,
"cases_compared": 224,
"corpus_hex": "9a84b55af68f1a8b07390b234cf562dd8d85cef3185a697ab85690c8790a7809"
}
]
}

2
wai/tools/quantum-wgsl-check/Cargo.lock generated
View file

@ -599,7 +599,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
[[package]]
name = "wai-quantum"
version = "0.4.1"
version = "0.4.2"
dependencies = [
"blake3",
"ed25519-dalek",

600
wai/tools/quantum-wgsl-check/src/main.rs
View file

@ -1,63 +1,131 @@
//! `quantum-wgsl-check`: holds the WGSL gate kernel of `wai-quantum` to the
//! rules its bit-exactness argument rests on (the module documentation of
//! `wai_quantum::quantum_fabric::wgsl`), on naga's IR, and translates it to
//! `quantum-wgsl-check`: holds the WGSL gate kernels of `wai-quantum` to the
//! rules their bit-exactness arguments rest on (the module documentation of
//! `wai_quantum::quantum_fabric::wgsl`), on naga's IR, and translates them to
//! MSL, SPIR-V and HLSL.
//!
//! ```text
//! quantum-wgsl-check [--shader <file>] [<directory>]
//! quantum-wgsl-check [--shader <file>] [--tile-shader <file>] [<directory>]
//! ```
//!
//! Without `--shader`, the text checked is
//! `wai_quantum::quantum_fabric::wgsl::SHADER`, the exact constant the crate
//! compiles. `--shader <file>` checks that file's text instead, so that a
//! planted violation can be shown to fail; the mutants are then applied to that
//! text. With a directory, the SPIR-V is written there as `lanes32.spv`, for
//! `spirv-val`.
//! Two texts are checked, each under its own rules: `SHADER`, the per-gate
//! kernel, and `TILE_SHADER`, the same kernel followed by the tile runs. By
//! default they are `wai_quantum::quantum_fabric::wgsl::SHADER` and
//! `TILE_SHADER`, the exact constants the crate compiles. `--shader <file>` or
//! `--tile-shader <file>` checks that file's text in place of one, so that a
//! planted violation can be shown to fail; its mutants are then applied to
//! that text. With a directory, the SPIR-V is written there as `lanes32.spv`
//! and `lanes32_tile.spv`, for `spirv-val`.
//!
//! The checks, in order; the first failure ends the run with exit status 1:
//! The checks, in order, for each text; the first failure ends the run with
//! exit status 1:
//!
//! 1. The text parses, and validates with every validation flag and no
//! capability.
//! 2. The entry points are exactly `apply` and `probe`, both compute, both with
//! workgroup size `[64, 1, 1]`.
//! 3. The globals are bindings 0 to 5 of group 0: storage read-write at 0, 1, 2
//! and 4, uniform at 3 and 5. No workgroup variable.
//! 2. The entry points are exactly the expected ones, all compute: `apply` and
//! `probe` with workgroup size `[64, 1, 1]`, and in `TILE_SHADER`
//! `apply_tile` with `[256, 1, 1]`.
//! 3. The bound globals are group 0: storage read-write at 0, 1, 2 and 4,
//! uniform at 3 and 5, and in `TILE_SHADER` storage read-only at 6 and
//! uniform at 7. Neither has a workgroup variable (the SPIR-V naga writes
//! for a workgroup array or structure carries layout decorations a Vulkan
//! validator rejects; were one allowed, `workgroup_bytes` bounds it).
//! 4. Every scalar type, in the type arena and in the type of every
//! expression, is `bool` or a 32-bit unsigned integer.
//! 5. No conversion or bitcast, no division or remainder, no built-in math
//! function; every shift amount is a `u32` literal from 1 to 31, or a
//! constant initialised with one.
//! 6. Every atomic operation is `atomicOr` or `atomicMin`; no barrier.
//! 6. Every atomic operation is `atomicOr` or `atomicMin`. `SHADER` has no
//! barrier. In `TILE_SHADER`, every barrier is a `storageBarrier()` in
//! `apply_tile`, placed where the uniformity argument puts it: at the top
//! level of the entry point, or at the top level of the body of a `for`
//! loop at the top level, a loop with no `break` but its condition's and no
//! `continue` or `return`. And every barrier is in uniform control flow:
//! no condition of an `if`, `switch` or loop that encloses it, and no
//! `return`, `break` or `continue` on a path to it under such a condition,
//! depends on the invocation (its local or global id, or anything but the
//! workgroup id, the number of workgroups and the workgroup size among the
//! built-ins) or on memory other invocations write (read-write storage,
//! workgroup memory, atomics), directly or through a variable or a call.
//! Uniform buffers, read-only storage and buffer lengths are uniform. (Naga
//! does not check barrier uniformity; this is the same analysis in the
//! small, and a browser's WGSL compiler checks it too.)
//! 7. The module translates to MSL 2.4, to SPIR-V, and to HLSL shader model
//! 5.1, the model the DX12 member compiles for.
//! 8. Every mutant of `wgsl::MUTANTS` changes the text, parses and validates.
//! 9. The implementation id is printed (for `SHADER` only).
//! 8. Every mutant (`wgsl::MUTANTS` for `SHADER`, `wgsl::TILE_MUTANTS` for
//! `TILE_SHADER`) changes the text, parses and validates.
//! 9. The implementation id is printed (for the crate's own text only).
//!
//! The last line of a passing run is `quantum-wgsl-check: ok`.
use std::collections::BTreeMap;
use std::collections::{BTreeMap, BTreeSet, HashSet};
use std::path::{Path, PathBuf};
use std::process::ExitCode;
use naga::back::{hlsl, msl, spv};
use naga::valid::{Capabilities, ModuleInfo, ValidationFlags, Validator};
use naga::{AddressSpace, Arena, AtomicFunction, BinaryOperator, Block, Expression, Handle, Literal, Module, Scalar, ScalarKind, ShaderStage, Statement, StorageAccess, TypeInner};
use wai_quantum::quantum_fabric::wgsl;
use naga::{AddressSpace, Arena, AtomicFunction, Barrier, BinaryOperator, Binding, Block, BuiltIn, Expression, Function, Handle, Literal, Module, Scalar, ScalarKind, ShaderStage, Statement, StorageAccess, TypeInner};
use wai_quantum::quantum_fabric::wgsl::{self, Mutant};
const NAME: &str = "quantum-wgsl-check";
type Check<T = ()> = Result<T, String>;
const READ_WRITE: &str = "storage read-write";
const READ: &str = "storage read";
const UNIFORM: &str = "uniform";
/// What one text is held to.
struct Rules {
/// The constant's name, as the output says it.
label: &'static str,
/// The SPIR-V file written for `spirv-val`.
spv: &'static str,
/// The entry points and their workgroup sizes along x.
entry_points: &'static [(&'static str, u32)],
/// The bound globals of group 0.
bindings: &'static [(u32, &'static str)],
/// Bytes of workgroup memory allowed; zero forbids workgroup variables.
workgroup_bytes: u32,
/// The one entry point that may hold barriers, if any, and the one kind
/// of barrier it may hold.
barriers_in: Option<(&'static str, Barrier)>,
mutants: &'static [Mutant],
implementation: fn() -> String,
}
const PER_GATE: Rules = Rules {
label: "SHADER",
spv: "lanes32.spv",
entry_points: &[("apply", wgsl::WORKGROUP), ("probe", wgsl::WORKGROUP)],
bindings: &[(0, READ_WRITE), (1, READ_WRITE), (2, READ_WRITE), (3, UNIFORM), (4, READ_WRITE), (5, UNIFORM)],
workgroup_bytes: 0,
barriers_in: None,
mutants: &wgsl::MUTANTS,
implementation: wgsl::implementation_id,
};
const TILE: Rules = Rules {
label: "TILE_SHADER",
spv: "lanes32_tile.spv",
entry_points: &[("apply", wgsl::WORKGROUP), ("apply_tile", wgsl::TILE_WORKGROUP), ("probe", wgsl::WORKGROUP)],
bindings: &[(0, READ_WRITE), (1, READ_WRITE), (2, READ_WRITE), (3, UNIFORM), (4, READ_WRITE), (5, UNIFORM), (6, READ), (7, UNIFORM)],
workgroup_bytes: 0,
barriers_in: Some(("apply_tile", Barrier::STORAGE)),
mutants: &wgsl::TILE_MUTANTS,
implementation: wgsl::tile_implementation_id,
};
fn main() -> ExitCode {
let args: Vec<String> = std::env::args().skip(1).collect();
let (shader, out) = match parse_args(&args) {
let (shader, tile_shader, out) = match parse_args(&args) {
Ok(a) => a,
Err(e) => {
eprintln!("{NAME}: {e}\nusage: {NAME} [--shader <file>] [<directory>]");
eprintln!("{NAME}: {e}\nusage: {NAME} [--shader <file>] [--tile-shader <file>] [<directory>]");
return ExitCode::from(2);
}
};
match run(shader.as_deref(), out.as_deref()) {
let result = run(&PER_GATE, wgsl::SHADER, shader.as_deref(), out.as_deref()).and_then(|()| run(&TILE, wgsl::TILE_SHADER, tile_shader.as_deref(), out.as_deref()));
match result {
Ok(()) => {
println!("{NAME}: ok");
ExitCode::SUCCESS
@ -69,12 +137,16 @@ fn main() -> ExitCode {
}
}
fn parse_args(args: &[String]) -> Check<(Option<PathBuf>, Option<PathBuf>)> {
let (mut shader, mut out) = (None, None);
type Args = (Option<PathBuf>, Option<PathBuf>, Option<PathBuf>);
fn parse_args(args: &[String]) -> Check<Args> {
let (mut shader, mut tile_shader, mut out) = (None, None, None);
let mut it = args.iter();
while let Some(a) = it.next() {
if a == "--shader" {
shader = Some(PathBuf::from(it.next().ok_or("--shader needs a file")?));
} else if a == "--tile-shader" {
tile_shader = Some(PathBuf::from(it.next().ok_or("--tile-shader needs a file")?));
} else if a.starts_with("--") {
return Err(format!("unknown option {a}"));
} else if out.is_none() {
@ -83,39 +155,52 @@ fn parse_args(args: &[String]) -> Check<(Option<PathBuf>, Option<PathBuf>)> {
return Err(format!("unexpected argument {a}"));
}
}
Ok((shader, out))
Ok((shader, tile_shader, out))
}
fn run(shader: Option<&Path>, out: Option<&Path>) -> Check {
let text = match shader {
fn run(rules: &Rules, own: &str, file: Option<&Path>, out: Option<&Path>) -> Check {
let label = rules.label;
let text = match file {
None => {
println!("{NAME}: checking wai_quantum::quantum_fabric::wgsl::SHADER");
wgsl::SHADER.to_string()
println!("{NAME}: {label}: checking wai_quantum::quantum_fabric::wgsl::{label}");
own.to_string()
}
Some(p) => {
println!("{NAME}: checking {} in place of the compiled SHADER", p.display());
println!("{NAME}: {label}: checking {} in place of the compiled {label}", p.display());
std::fs::read_to_string(p).map_err(|e| format!("reading {}: {e}", p.display()))?
}
};
let (module, info) = compile(&text)?;
println!("{NAME}: 1. parsed and validated (every validation flag, no capability)");
entry_points(&module)?;
println!("{NAME}: 2. entry points apply and probe: compute, workgroup size [64, 1, 1]");
globals(&module)?;
println!("{NAME}: 3. group 0: storage read-write at 0, 1, 2, 4; uniform at 3, 5; no workgroup variable");
types(&module, &info)?;
println!("{NAME}: 4. every scalar is bool or u32");
expressions(&module)?;
println!("{NAME}: 5. no conversion, division, remainder or math built-in; every shift amount a literal in 1..=31");
let at = |e: String| format!("{label}: {e}");
let (module, info) = compile(&text).map_err(at)?;
println!("{NAME}: {label}: 1. parsed and validated (every validation flag, no capability)");
entry_points(rules, &module).map_err(at)?;
let eps: Vec<String> = rules.entry_points.iter().map(|(n, x)| format!("{n} [{x}, 1, 1]")).collect();
println!("{NAME}: {label}: 2. compute entry points {}", eps.join(", "));
let workgroup = globals(rules, &module).map_err(at)?;
let bound: Vec<String> = rules.bindings.iter().map(|(b, s)| format!("{b} {s}")).collect();
println!("{NAME}: {label}: 3. group 0: {}; workgroup memory {workgroup} bytes (at most {})", bound.join(", "), rules.workgroup_bytes);
types(&module, &info).map_err(at)?;
println!("{NAME}: {label}: 4. every scalar is bool or u32");
expressions(&module).map_err(at)?;
println!("{NAME}: {label}: 5. no conversion, division, remainder or math built-in; every shift amount a literal in 1..=31");
let mut barriers = 0;
for (name, f) in functions(&module) {
statements(&f.body, name)?;
let allowed = rules.barriers_in.filter(|(ep, _)| *ep == name).map(|(_, kind)| kind);
statements(&f.body, name, allowed, Depth::Top, &mut barriers).map_err(at)?;
}
println!("{NAME}: 6. atomics are atomicOr and atomicMin only; no barrier");
translate(&module, &info, out)?;
mutants(&text, shader.is_some())?;
println!("{NAME}: 8. every mutant changes the text, parses and validates");
if shader.is_none() {
println!("{NAME}: implementation {}", wgsl::implementation_id());
match rules.barriers_in {
None => println!("{NAME}: {label}: 6. atomics are atomicOr and atomicMin only; no barrier"),
Some((ep, kind)) => {
let e = module.entry_points.iter().find(|e| e.name == ep).ok_or_else(|| at(format!("6. no entry point {ep}")))?;
uniform_barriers(&module, &e.function, ep).map_err(at)?;
println!("{NAME}: {label}: 6. atomics are atomicOr and atomicMin only; {barriers} barrier(s) of {kind:?}, all in {ep} at the top level of its body or of a top-level for loop without break, continue or return, and in uniform control flow");
}
}
translate(rules, &module, &info, out).map_err(at)?;
mutants(rules, own, &text, file.is_some()).map_err(at)?;
println!("{NAME}: {label}: 8. every mutant changes the text, parses and validates");
if file.is_none() {
println!("{NAME}: {label}: implementation {}", (rules.implementation)());
}
Ok(())
}
@ -128,34 +213,46 @@ fn compile(text: &str) -> Check<(Module, ModuleInfo)> {
}
/// Step 2.
fn entry_points(module: &Module) -> Check {
fn entry_points(rules: &Rules, module: &Module) -> Check {
let mut names: Vec<&str> = module.entry_points.iter().map(|e| e.name.as_str()).collect();
names.sort_unstable();
if names != ["apply", "probe"] {
return Err(format!("2. the entry points are {names:?}, not [\"apply\", \"probe\"]"));
let mut want: Vec<&str> = rules.entry_points.iter().map(|(n, _)| *n).collect();
want.sort_unstable();
if names != want {
return Err(format!("2. the entry points are {names:?}, not {want:?}"));
}
for e in &module.entry_points {
if e.stage != ShaderStage::Compute {
return Err(format!("2. entry point `{}` is a {:?} stage", e.name, e.stage));
}
if e.workgroup_size != [wgsl::WORKGROUP, 1, 1] || e.workgroup_size_overrides.is_some() {
return Err(format!("2. entry point `{}` has workgroup size {:?} (overrides {:?})", e.name, e.workgroup_size, e.workgroup_size_overrides));
let x = rules.entry_points.iter().find(|(n, _)| *n == e.name).map(|(_, x)| *x).unwrap_or(0);
if e.workgroup_size != [x, 1, 1] || e.workgroup_size_overrides.is_some() {
return Err(format!("2. entry point `{}` has workgroup size {:?} (overrides {:?}), not [{x}, 1, 1]", e.name, e.workgroup_size, e.workgroup_size_overrides));
}
}
Ok(())
}
/// Step 3.
fn globals(module: &Module) -> Check {
const READ_WRITE: &str = "storage read-write";
const UNIFORM: &str = "uniform";
let want = BTreeMap::from([(0, READ_WRITE), (1, READ_WRITE), (2, READ_WRITE), (3, UNIFORM), (4, READ_WRITE), (5, UNIFORM)]);
/// Step 3. Returns the bytes of workgroup memory.
fn globals(rules: &Rules, module: &Module) -> Check<u32> {
let want: BTreeMap<u32, &str> = rules.bindings.iter().copied().collect();
let mut found = BTreeMap::new();
let mut workgroup = 0u32;
for (_, v) in module.global_variables.iter() {
let name = v.name.as_deref().unwrap_or("(unnamed)");
let space = match v.space {
AddressSpace::WorkGroup => return Err(format!("3. `var<workgroup> {name}`: the kernel holds no workgroup memory")),
AddressSpace::WorkGroup => {
if rules.workgroup_bytes == 0 {
return Err(format!("3. `var<workgroup> {name}`: the kernel holds no workgroup memory"));
}
if v.binding.is_some() {
return Err(format!("3. workgroup variable `{name}` has a binding"));
}
workgroup += module.types[v.ty].inner.size(module.to_ctx());
continue;
}
AddressSpace::Storage { access } if access.contains(StorageAccess::LOAD | StorageAccess::STORE) => READ_WRITE,
AddressSpace::Storage { access } if access == StorageAccess::LOAD => READ,
AddressSpace::Uniform => UNIFORM,
other => return Err(format!("3. global `{name}` is in address space {other:?}")),
};
@ -170,7 +267,10 @@ fn globals(module: &Module) -> Check {
if found != want {
return Err(format!("3. the bindings are {found:?}, not {want:?}"));
}
Ok(())
if workgroup > rules.workgroup_bytes {
return Err(format!("3. {workgroup} bytes of workgroup memory, more than {}", rules.workgroup_bytes));
}
Ok(workgroup)
}
fn scalar_ok(s: Scalar) -> bool {
@ -253,8 +353,66 @@ fn expressions(module: &Module) -> Check {
Ok(())
}
/// Where a statement sits, for the barrier rule of step 6.
#[derive(Clone, Copy, PartialEq, Eq)]
enum Depth {
/// The top level of a function's body.
Top,
/// The top level of the body of a `for` loop that is at the top level.
LoopBody,
/// Anywhere else.
Nested,
}
/// Whether a block holds a statement that leaves a loop or the function
/// (`break`, `continue`, `return`), outside nested loops for `break` and
/// `continue`.
fn exits(block: &Block, in_nested_loop: bool) -> Option<&'static str> {
for s in block.iter() {
let found = match s {
Statement::Return { .. } => Some("return"),
Statement::Break if !in_nested_loop => Some("break"),
Statement::Continue if !in_nested_loop => Some("continue"),
Statement::Block(b) => exits(b, in_nested_loop),
Statement::If { accept, reject, .. } => exits(accept, in_nested_loop).or_else(|| exits(reject, in_nested_loop)),
// A `break` in a switch leaves the switch only.
Statement::Switch { cases, .. } => cases.iter().find_map(|c| exits(&c.body, in_nested_loop).filter(|e| *e != "break")),
Statement::Loop { body, continuing, .. } => exits(body, true).or_else(|| exits(continuing, true)).filter(|e| *e == "return"),
_ => None,
};
if found.is_some() {
return found;
}
}
None
}
/// Whether a block holds a barrier anywhere.
fn has_barrier(block: &Block) -> bool {
block.iter().any(|s| match s {
Statement::ControlBarrier(_) | Statement::MemoryBarrier(_) => true,
Statement::Block(b) => has_barrier(b),
Statement::If { accept, reject, .. } => has_barrier(accept) || has_barrier(reject),
Statement::Switch { cases, .. } => cases.iter().any(|c| has_barrier(&c.body)),
Statement::Loop { body, continuing, .. } => has_barrier(body) || has_barrier(continuing),
_ => false,
})
}
/// The body of a `for` loop as naga lowers it: emits, then
/// `if (cond) {} else { break; }`, then the loop's own block. Returns that
/// block when the loop has this form.
fn for_body(body: &Block) -> Option<&Block> {
let rest: Vec<&Statement> = body.iter().skip_while(|s| matches!(s, Statement::Emit(_))).collect();
match rest.as_slice() {
[Statement::If { accept, reject, .. }, Statement::Block(user)] if accept.is_empty() && reject.len() == 1 && matches!(reject.iter().next(), Some(Statement::Break)) => Some(user),
_ => None,
}
}
/// Step 6, recursively through every block.
fn statements(block: &Block, name: &str) -> Check {
fn statements(block: &Block, name: &str, barrier_kind: Option<Barrier>, depth: Depth, barriers: &mut u32) -> Check {
let barriers_allowed = barrier_kind.is_some();
for s in block.iter() {
match s {
Statement::Atomic { fun, .. } => {
@ -262,20 +420,49 @@ fn statements(block: &Block, name: &str) -> Check {
return Err(format!("6. in `{name}`: the atomic operation {fun:?}"));
}
}
Statement::ControlBarrier(_) | Statement::MemoryBarrier(_) => return Err(format!("6. in `{name}`: a barrier")),
Statement::Block(b) => statements(b, name)?,
Statement::ControlBarrier(b) => {
if !barriers_allowed {
return Err(format!("6. in `{name}`: a barrier"));
}
if Some(*b) != barrier_kind {
return Err(format!("6. in `{name}`: a barrier of {b:?}, not of {barrier_kind:?}"));
}
if depth == Depth::Nested {
return Err(format!("6. in `{name}`: a barrier nested below the top level of the body or of a top-level for loop"));
}
*barriers += 1;
}
Statement::MemoryBarrier(_) => return Err(format!("6. in `{name}`: a memory barrier")),
Statement::Block(b) => statements(b, name, barrier_kind, if depth == Depth::Top { Depth::Top } else { Depth::Nested }, barriers)?,
Statement::If { accept, reject, .. } => {
statements(accept, name)?;
statements(reject, name)?;
statements(accept, name, barrier_kind, Depth::Nested, barriers)?;
statements(reject, name, barrier_kind, Depth::Nested, barriers)?;
}
Statement::Switch { cases, .. } => {
for c in cases {
statements(&c.body, name)?;
statements(&c.body, name, barrier_kind, Depth::Nested, barriers)?;
}
}
Statement::Loop { body, continuing, .. } => {
statements(body, name)?;
statements(continuing, name)?;
Statement::Loop { body, continuing, break_if } => {
// A barrier in a loop is allowed only in a top-level `for`
// loop, at the top level of its own block, and only when
// nothing but the condition's `break` leaves the loop.
let user = for_body(body).filter(|_| depth == Depth::Top && break_if.is_none());
if barriers_allowed && (has_barrier(body) || has_barrier(continuing)) {
let Some(user) = user else {
return Err(format!("6. in `{name}`: a barrier in a loop that is not a top-level for loop"));
};
if let Some(e) = exits(user, false).or_else(|| exits(continuing, false)) {
return Err(format!("6. in `{name}`: a `{e}` in a loop that holds a barrier"));
}
if has_barrier(continuing) {
return Err(format!("6. in `{name}`: a barrier in a loop's continuing block"));
}
statements(user, name, barrier_kind, Depth::LoopBody, barriers)?;
} else {
statements(body, name, barrier_kind, Depth::Nested, barriers)?;
statements(continuing, name, barrier_kind, Depth::Nested, barriers)?;
}
}
_ => {}
}
@ -283,40 +470,279 @@ fn statements(block: &Block, name: &str) -> Check {
Ok(())
}
/// Whether a global's contents can differ between invocations of one
/// workgroup during a dispatch: memory some invocation may write.
fn varying_global(module: &Module, g: Handle<naga::GlobalVariable>) -> bool {
match module.global_variables[g].space {
AddressSpace::Uniform | AddressSpace::Handle | AddressSpace::Immediate => false,
AddressSpace::Storage { access } => access.contains(StorageAccess::STORE),
_ => true,
}
}
/// Built-ins with one value for a whole workgroup.
fn uniform_builtin(b: &Binding) -> bool {
matches!(b, Binding::BuiltIn(BuiltIn::WorkGroupId | BuiltIn::NumWorkGroups | BuiltIn::WorkGroupSize))
}
/// The functions whose result can vary between invocations whatever their
/// arguments: they read varying memory, use an atomic, or call one that does.
fn varying_functions(module: &Module) -> HashSet<Handle<Function>> {
let mut varying = HashSet::new();
loop {
let before = varying.len();
for (h, f) in module.functions.iter() {
let reads = f.expressions.iter().any(|(_, e)| match *e {
Expression::GlobalVariable(g) => varying_global(module, g),
Expression::AtomicResult { .. } => true,
_ => false,
});
if reads || calls(&f.body).iter().any(|c| varying.contains(c)) {
varying.insert(h);
}
}
if varying.len() == before {
return varying;
}
}
}
/// The functions a block calls.
fn calls(block: &Block) -> Vec<Handle<Function>> {
let mut out = Vec::new();
for s in block.iter() {
match s {
Statement::Call { function, .. } => out.push(*function),
Statement::Block(b) => out.extend(calls(b)),
Statement::If { accept, reject, .. } => {
out.extend(calls(accept));
out.extend(calls(reject));
}
Statement::Switch { cases, .. } => cases.iter().for_each(|c| out.extend(calls(&c.body))),
Statement::Loop { body, continuing, .. } => {
out.extend(calls(body));
out.extend(calls(continuing));
}
_ => {}
}
}
out
}
/// How control leaves a block under a varying condition.
#[derive(Clone, Copy, Default)]
struct Exits {
/// A `return` some invocations take and others do not.
returns: bool,
/// A `break` or `continue` some invocations take and others do not.
breaks: bool,
}
/// The uniformity analysis of one function, in the small: which values and
/// which control flow can differ between the invocations of a workgroup.
struct Uniformity<'a> {
module: &'a Module,
f: &'a Function,
name: &'a str,
varying_fns: HashSet<Handle<Function>>,
/// Varying arguments (by position).
args: Vec<bool>,
/// Varying local variables (by handle index): stored a varying value, or
/// stored under varying control.
vars: Vec<bool>,
/// Call results that vary.
results: HashSet<Handle<Expression>>,
changed: bool,
report: bool,
errors: BTreeSet<String>,
}
impl Uniformity<'_> {
fn varies(&self, h: Handle<Expression>) -> bool {
match self.f.expressions[h] {
Expression::Literal(_) | Expression::Constant(_) | Expression::Override(_) | Expression::ZeroValue(_) => false,
// A buffer's length is the same for every invocation.
Expression::ArrayLength(_) => false,
Expression::FunctionArgument(i) => self.args[i as usize],
// Pointers: what they point at matters when it is loaded.
Expression::GlobalVariable(_) | Expression::LocalVariable(_) => false,
Expression::Load { pointer } => self.varies_at(pointer),
Expression::Access { base, index } => self.varies(base) || self.varies(index),
Expression::AccessIndex { base, .. } => self.varies(base),
Expression::Splat { value, .. } => self.varies(value),
Expression::Swizzle { vector, .. } => self.varies(vector),
Expression::Compose { ref components, .. } => components.iter().any(|c| self.varies(*c)),
Expression::Unary { expr, .. } => self.varies(expr),
Expression::Binary { left, right, .. } => self.varies(left) || self.varies(right),
Expression::Select { condition, accept, reject } => self.varies(condition) || self.varies(accept) || self.varies(reject),
Expression::Relational { argument, .. } => self.varies(argument),
Expression::CallResult(_) => self.results.contains(&h),
// Atomics, and anything this analysis does not know: assumed to vary.
_ => true,
}
}
/// Whether the memory a pointer expression names can vary: its root's
/// contents, or any index on the way.
fn varies_at(&self, pointer: Handle<Expression>) -> bool {
match self.f.expressions[pointer] {
Expression::GlobalVariable(g) => varying_global(self.module, g),
Expression::LocalVariable(v) => self.vars[v.index()],
Expression::Access { base, index } => self.varies(index) || self.varies_at(base),
Expression::AccessIndex { base, .. } => self.varies_at(base),
Expression::FunctionArgument(i) => self.args[i as usize],
_ => true,
}
}
/// The local variable a pointer expression stores into, if any.
fn local_root(&self, pointer: Handle<Expression>) -> Option<usize> {
match self.f.expressions[pointer] {
Expression::LocalVariable(v) => Some(v.index()),
Expression::Access { base, .. } | Expression::AccessIndex { base, .. } => self.local_root(base),
_ => None,
}
}
fn mark_var(&mut self, v: usize) {
if !self.vars[v] {
self.vars[v] = true;
self.changed = true;
}
}
/// Walks `block` with its control flow varying or not; returns how it
/// leaves under varying conditions.
fn block(&mut self, block: &Block, mut varying: bool) -> Exits {
let mut exits = Exits::default();
for s in block.iter() {
match s {
Statement::Store { pointer, value } => {
if let Some(v) = self.local_root(*pointer)
&& (varying || self.varies(*value) || self.varies_at(*pointer))
{
self.mark_var(v);
}
}
Statement::Call { function, arguments, result: Some(r) } => {
if (self.varying_fns.contains(function) || arguments.iter().any(|a| self.varies(*a))) && self.results.insert(*r) {
self.changed = true;
}
}
Statement::ControlBarrier(_) | Statement::MemoryBarrier(_) => {
if varying && self.report {
self.errors.insert(format!("6. in `{}`: a barrier in control flow that can differ between the invocations of a workgroup", self.name));
}
}
Statement::If { condition, accept, reject } => {
let c = varying || self.varies(*condition);
let (a, r) = (self.block(accept, c), self.block(reject, c));
exits.returns |= a.returns || r.returns;
exits.breaks |= a.breaks || r.breaks;
varying |= a.returns || a.breaks || r.returns || r.breaks;
}
Statement::Switch { selector, cases } => {
let c = varying || self.varies(*selector);
let mut any = false;
for case in cases {
let e = self.block(&case.body, c);
// A `break` leaves the switch; after it, control may still differ.
exits.returns |= e.returns;
any |= e.returns || e.breaks || (c && !case.body.is_empty());
}
varying |= any && c;
}
Statement::Loop { body, continuing, break_if } => {
let pass = |u: &mut Self, v: bool| {
let (b, c) = (u.block(body, v), u.block(continuing, v));
let leaves = break_if.is_some_and(|e| u.varies(e));
(b.returns || c.returns, b.breaks || c.breaks || leaves)
};
let (mut returns, mut leaves) = pass(self, varying);
if (returns || leaves) && !varying {
// Some invocations leave the loop, or an iteration,
// where others do not: the whole loop varies.
(returns, leaves) = pass(self, true);
}
exits.returns |= returns;
varying |= returns || leaves;
}
Statement::Block(b) => {
let e = self.block(b, varying);
exits.returns |= e.returns;
exits.breaks |= e.breaks;
varying |= e.returns || e.breaks;
}
Statement::Return { .. } => exits.returns |= varying,
Statement::Break | Statement::Continue => exits.breaks |= varying,
_ => {}
}
}
exits
}
}
/// Step 6, uniformity: every barrier of the entry point `name` is in control
/// flow that is the same for every invocation of a workgroup.
fn uniform_barriers(module: &Module, f: &Function, name: &str) -> Check {
let args = f.arguments.iter().map(|a| !a.binding.as_ref().is_some_and(uniform_builtin)).collect();
let mut u = Uniformity { module, f, name, varying_fns: varying_functions(module), args, vars: vec![false; f.local_variables.len()], results: HashSet::new(), changed: true, report: false, errors: BTreeSet::new() };
for (h, v) in f.local_variables.iter() {
if v.init.is_some_and(|e| u.varies(e)) {
u.vars[h.index()] = true;
}
}
// Varying is monotone: walk until nothing new varies, then once more to report.
while u.changed {
u.changed = false;
u.block(&f.body, false);
}
u.report = true;
u.block(&f.body, false);
match u.errors.into_iter().next() {
Some(e) => Err(e),
None => Ok(()),
}
}
/// Step 7.
fn translate(module: &Module, info: &ModuleInfo, out: Option<&Path>) -> Check {
fn translate(rules: &Rules, module: &Module, info: &ModuleInfo, out: Option<&Path>) -> Check {
let label = rules.label;
let options = msl::Options { lang_version: (2, 4), ..Default::default() };
msl::write_string(module, info, &options, &msl::PipelineOptions::default()).map_err(|e| format!("7. MSL: {e}"))?;
let words = spv::write_vec(module, info, &spv::Options::default(), None).map_err(|e| format!("7. SPIR-V: {e}"))?;
let mut hlsl_text = String::new();
// Shader model 5.1 (the default), the model the DX12 member compiles for.
hlsl::Writer::new(&mut hlsl_text, &hlsl::Options::default(), &hlsl::PipelineOptions::default()).write(module, info, None).map_err(|e| format!("7. HLSL: {e}"))?;
println!("{NAME}: 7. translated to MSL 2.4, SPIR-V and HLSL shader model 5.1");
println!("{NAME}: {label}: 7. translated to MSL 2.4, SPIR-V and HLSL shader model 5.1");
match out {
Some(dir) => {
let path = dir.join("lanes32.spv");
let path = dir.join(rules.spv);
let bytes: Vec<u8> = words.iter().flat_map(|w| w.to_le_bytes()).collect();
std::fs::write(&path, bytes).map_err(|e| format!("7. writing {}: {e}", path.display()))?;
println!("{NAME}: 7. SPIR-V written to {}", path.display());
println!("{NAME}: {label}: 7. SPIR-V written to {}", path.display());
}
None => println!("{NAME}: 7. SPIR-V not written: no directory given"),
None => println!("{NAME}: {label}: 7. SPIR-V not written: no directory given"),
}
Ok(())
}
/// Step 8. On `SHADER` the mutants are `wgsl::mutate`'s; on a file given with
/// `--shader`, the same edits applied to its text, each `find` exactly once.
fn mutants(text: &str, file: bool) -> Check {
for m in &wgsl::MUTANTS {
let mutated = if file {
let n = text.matches(m.find).count();
if n != 1 {
return Err(format!("8. mutant {}: `{}` occurs {n} times", m.name, m.find));
/// Step 8. Each mutant's `find` must occur exactly once in the text checked
/// (the crate's own, or a file given in its place).
fn mutants(rules: &Rules, own: &str, text: &str, file: bool) -> Check {
for m in rules.mutants {
let n = text.matches(m.find).count();
if n != 1 {
return Err(format!("8. mutant {}: `{}` occurs {n} times", m.name, m.find));
}
let mutated = text.replacen(m.find, m.replace, 1);
if !file {
// The crate's own mutant text is this one.
let crate_text = if rules.label == PER_GATE.label { wgsl::mutate(m) } else { wgsl::mutate_tile(m) };
if crate_text != mutated || text != own {
return Err(format!("8. mutant {}: the crate's mutant text is not this edit of {}", m.name, rules.label));
}
text.replacen(m.find, m.replace, 1)
} else {
wgsl::mutate(m)
};
}
if mutated == text {
return Err(format!("8. mutant {} does not change the text", m.name));
}

2
wai/wai-pack/Cargo.lock generated
View file

@ -698,7 +698,7 @@ dependencies = [
[[package]]
name = "wai-quantum"
version = "0.4.1"
version = "0.4.2"
dependencies = [
"blake3",
"ed25519-dalek",

2
wai/wai-quantum-serve/Cargo.lock generated
View file

@ -589,7 +589,7 @@ dependencies = [
[[package]]
name = "wai-quantum"
version = "0.4.1"
version = "0.4.2"
dependencies = [
"blake3",
"ed25519-dalek",

2
wai/wai-quantum/Cargo.lock generated
View file

@ -995,7 +995,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
[[package]]
name = "wai-quantum"
version = "0.4.1"
version = "0.4.2"
dependencies = [
"blake3",
"ed25519-dalek",

9
wai/wai-quantum/Cargo.toml
View file

@ -1,6 +1,6 @@
[package]
name = "wai-quantum"
version = "0.4.1"
version = "0.4.2"
edition = "2024"
rust-version = "1.97"
license = "Apache-2.0"
@ -11,6 +11,9 @@ documentation = "https://docs.rs/wai-quantum"
keywords = ["quantum", "simulation", "wasm", "deterministic", "qec"]
categories = ["science", "simulation", "wasm", "mathematics"]
readme = "README.md"
# Holds a file outside this package (quantum-conformance/fabric/expected.json) to
# the crate; it runs in the repository only.
exclude = ["tests/fabric_expected.rs"]
# docs.rs builds with DEFAULT features only, which would document 3 of 21
# modules and hide the entire operational + ML layers. Document everything.
@ -159,6 +162,10 @@ full = [
name = "webgpu_member"
required-features = ["quantum_fabric_webgpu"]
[[test]]
name = "fabric_expected"
required-features = ["quantum_fabric"]
[[example]]
name = "qio_gset"
required-features = ["quantum_qio"]

2
wai/wai-quantum/README.md
View file

@ -94,7 +94,7 @@ never creates information.
| `quantum_channel` | **transport itself**: Kraus channels (bit flip, dephasing, depolarizing, amplitude damping), and linear-inversion tomography to see what actually arrived — with the law that a channel can only *lose* information, never create it |
| `quantum_qbom` | Quantum Bill of Materials — stage receipts into one manifest |
| `quantum_vml`, `quantum_phasor`, `quantum_qfhrr`, `quantum_kernel`, `quantum_qdata`, `quantum_phasor_meter` | the interference / phasor ML layer |
| `quantum_fabric_webgpu` | **the fabric's WebGPU member**, opt-in and not in `full`: the 32-bit-lane gate kernel in WGSL, one dispatch per gate, through Metal, Vulkan or DX12 natively and through a browser's WebGPU on `wasm32-unknown-unknown`. Each adapter is probed against the CPU's arithmetic in the process that uses it and certified bit-exact by reproducing the CPU reference's pinned corpus digests; a result that would leave 32 bits is refused, never returned. The one feature that reaches a platform interface |
| `quantum_fabric_webgpu` | **the fabric's WebGPU member**, opt-in and not in `full`: the 32-bit-lane gate kernel in WGSL, one dispatch per gate, through Metal, Vulkan or DX12 natively and through a browser's WebGPU on `wasm32-unknown-unknown`; a second implementation applies runs of gates tile by tile, one dispatch per run, and is certified on its own. Each adapter is probed against the CPU's arithmetic in the process that uses it and certified bit-exact by reproducing the CPU reference's pinned corpus digests; a result that would leave 32 bits is refused, never returned. The one feature that reaches a platform interface |
| `full` | everything |
Default features are `quantum`, `quantum_toolchain`, `quantum_receipt`.

50
wai/wai-quantum/src/quantum_fabric.rs
View file

@ -37,10 +37,10 @@
//! a 32-bit member must check each result, as [`lanes::simulate_checked`]
//! does. [`lanes`] is that kernel written with only the operations such a
//! fabric has, and its certificate against the native simulator is bit-exact.
//! [`plan`] is the program an accelerator member runs for a circuit, with a
//! CPU executor of exactly that program, and [`wgsl`] is the same kernel in
//! WGSL for a WebGPU member, with its host layout and the argument that it is
//! bit-exact. The
//! [`plan`] is the program an accelerator member runs for a circuit, one step
//! per gate or in tile runs, with a CPU executor of exactly that program, and
//! [`wgsl`] is the same kernel in WGSL for a WebGPU member, per gate and in
//! tile runs, with its host layouts and the argument that it is bit-exact. The
//! floating-point modules can earn bit-exactness only on a fabric with 64-bit
//! floats and the same order of operations; elsewhere their parity is a
//! measured bound.
@ -1218,20 +1218,12 @@ pub mod lanes {
if tbit == 0 || !tbit.is_power_of_two() || tbit > len / 2 || ctrl & tbit != 0 || ctrl & !(len - 1) != 0 {
return None;
}
let e = |k: usize| w.m[k] as i32;
let mut fits = true;
for i in 0..len {
if i & tbit == 0 && i & ctrl == ctrl {
let j = i | tbit;
let (a0, a1) = (amps[i], amps[j]);
let row = |r: usize| -> [i64; 2] {
let (mr, mi, nr, ni) = (e(4 * r), e(4 * r + 1), e(4 * r + 2), e(4 * r + 3));
[
fx_exact(mr, a0[0]) - fx_exact(mi, a0[1]) + fx_exact(nr, a1[0]) - fx_exact(ni, a1[1]),
fx_exact(mr, a0[1]) + fx_exact(mi, a0[0]) + fx_exact(nr, a1[1]) + fx_exact(ni, a1[0]),
]
};
let (o0, o1) = (row(0), row(1));
let (o0, o1) = (exact_row(w, 0, a0, a1), exact_row(w, 1, a0, a1));
fits &= o0.iter().chain(&o1).all(|v| i32::try_from(*v).is_ok());
// `as` keeps the low 32 bits.
amps[i] = [o0[0] as i32, o0[1] as i32];
@ -1241,6 +1233,18 @@ pub mod lanes {
Some(fits)
}
/// Row `r` (0 or 1) of a gate's words on the pair `(a0, a1)`, exactly:
/// `m_r0·a0 + m_r1·a1`, every product the exact `fx` and the sums in 64
/// bits. The low 32 bits of each part are the lane's.
pub(super) fn exact_row(w: &GateWords, r: usize, a0: Lane, a1: Lane) -> [i64; 2] {
let e = |k: usize| w.m[k] as i32;
let (mr, mi, nr, ni) = (e(4 * r), e(4 * r + 1), e(4 * r + 2), e(4 * r + 3));
[
fx_exact(mr, a0[0]) - fx_exact(mi, a0[1]) + fx_exact(nr, a1[0]) - fx_exact(ni, a1[1]),
fx_exact(mr, a0[1]) + fx_exact(mi, a0[0]) + fx_exact(nr, a1[1]) + fx_exact(ni, a1[0]),
]
}
/// Gates in order from `state`, as a 32-bit-lane member runs them: the
/// lanes, and the lowest position in `gates` of a gate whose exact result
/// left 32 bits. `None` if any gate's words are malformed for this state
@ -1852,6 +1856,21 @@ mod tests {
}
}
/// The planner configurations every test of a program runs under: one
/// step per gate, the tiled default, and small tiles with short and long
/// runs, the identity skipped or kept.
pub(super) fn configurations() -> Vec<plan::Config> {
let mut all = vec![plan::Config::PER_GATE, plan::Config::TILED];
for tile_bits in [1, 3] {
for run_max in [1, 7] {
for skip_identity in [true, false] {
all.push(plan::Config { tile_bits, run_max, skip_identity });
}
}
}
all
}
#[test]
fn a_rounded_product_is_never_negated() {
let v2 = [[759250126, 0], [1, 0], [759250126, 0], [1, 0]];
@ -1863,6 +1882,11 @@ mod tests {
assert_eq!(lanes::simulate(&c).unwrap(), want);
let run = lanes::simulate_checked(&c, 0).unwrap();
assert_eq!((run.amps.as_slice(), run.out_of_range), (want.as_slice(), None));
// Every planner configuration, tile runs included, keeps the tie.
for cfg in configurations() {
let run = plan::execute(&plan::compile_with(&c, 0, &cfg).unwrap());
assert_eq!((run.amps.as_slice(), run.out_of_range), (want.as_slice(), None), "{cfg:?}");
}
// H with `m11·a_j` computed as `−(m01·a_j)`: one rounded product,
// negated, where the table rounds `−h·a_j` on its own.
let h = gs.inv_sqrt2() as i32;

123
wai/wai-quantum/src/quantum_fabric/lanes32_tile.wgsl Normal file
View file

@ -0,0 +1,123 @@
// Tile runs: one workgroup owns one tile of run.size consecutive amplitudes of
// the state and applies a run of gates to it in place, in op-log order, with a
// storage barrier between gates. Each gate either has its target inside the
// tile (tbit < size) or is diagonal (m01 = m10 = 0) and acts per amplitude.
// Uses fx64, add64, sub64, fits, row and the bindings re, im, report of
// lanes32.wgsl, which this file is appended to. No shift amount here is
// anything but a literal; sizes and indices come multiplied, not shifted. No
// workgroup memory: the tile stays in re and im, which no other workgroup of
// the dispatch touches.
//
// Uniformity: the early returns depend only on run (a uniform buffer),
// workgroup_id, num_workgroups and the state planes' lengths. The barrier sits
// at the top level of the r loop, whose bound run.count is a uniform-buffer
// value; no break, continue or return occurs inside that loop, and no branch
// that depends on a gate or an invocation holds a barrier.
struct TileGate {
row0: vec4<u32>, // m00.re, m00.im, m01.re, m01.im
row1: vec4<u32>, // m10.re, m10.im, m11.re, m11.im
tbit: u32, // 1 << target
ctrl: u32, // the controls' bits, OR-ed
diag: u32, // 1: m01 = m10 = (0, 0)
index: u32, // the gate's position in the op-log
}
struct TileRun {
first: u32, // the run's first record in tile_gates
count: u32, // gates in the run, at least 1
size: u32, // amplitudes per tile: a power of two, at most TILE_MAX
x_groups: u32, // workgroups along x
tiles: u32, // tiles in the state: 2^n / size
pad0: u32,
pad1: u32,
pad2: u32,
}
@group(0) @binding(6) var<storage, read> tile_gates: array<TileGate>;
@group(0) @binding(7) var<uniform> run: TileRun;
const TILE_MAX: u32 = 1024u;
const TILE_LANES: u32 = 256u;
const NONE: u32 = 0xffffffffu;
const ZERO2: vec2<u32> = vec2<u32>(0u, 0u);
@compute @workgroup_size(256)
fn apply_tile(@builtin(workgroup_id) wid: vec3<u32>,
@builtin(num_workgroups) groups: vec3<u32>,
@builtin(local_invocation_index) lid: u32) {
let size = run.size;
let span = run.tiles * size; // 2^n
// A malformed run record, or one whose tiles or grid do not cover the state
// exactly once, is reported and never run (uniform: run is a uniform buffer,
// num_workgroups a uniform built-in, and a buffer's length is the same for
// every invocation).
if (size == 0u || size > TILE_MAX || (size & (size - 1u)) != 0u || run.count == 0u
|| run.tiles == 0u || span != arrayLength(&re) || span != arrayLength(&im)
|| run.x_groups == 0u || groups.x != run.x_groups || groups.z != 1u
|| groups.x * groups.y < run.tiles) {
atomicOr(&report[0], 4u);
return;
}
let tile = wid.y * run.x_groups + wid.x;
if (tile >= run.tiles) { return; } // workgroup-uniform
let base = tile * size;
// A tile outside the state is reported, and nothing of it is touched.
let inside = base + size <= arrayLength(&re) && base + size <= arrayLength(&im);
if (!inside) { atomicOr(&report[0], 2u); }
var bad = NONE;
var malformed = false;
for (var r = 0u; r < run.count; r++) {
// A record outside tile_gates reads as malformed, never out of bounds.
let at = run.first + r;
let held = at < arrayLength(&tile_gates);
let h = tile_gates[select(0u, at, held)];
let shape = held && h.tbit != 0u && (h.tbit & (h.tbit - 1u)) == 0u && h.tbit < span
&& (h.ctrl & h.tbit) == 0u && (h.ctrl & ~(span - 1u)) == 0u;
var ok = true;
if (shape && h.diag == 1u) {
for (var k = 0u; k < 4u; k++) {
let e = lid + TILE_LANES * k;
let gi = base + e;
if (inside && e < size && (gi & h.ctrl) == h.ctrl) {
let a = vec2<u32>(re[gi], im[gi]);
var o: vec4<u32>;
if ((gi & h.tbit) == 0u) { o = row(h.row0, a, ZERO2); } else { o = row(h.row1, ZERO2, a); }
re[gi] = o.y;
im[gi] = o.w;
if (!(fits(o.xy) && fits(o.zw))) { ok = false; }
}
}
} else if (shape && h.tbit < size) {
for (var k = 0u; k < 2u; k++) {
let q = lid + TILE_LANES * k;
if (inside && q < (size >> 1u)) {
let low = q & (h.tbit - 1u);
let gi = base + (((q ^ low) << 1u) | low); // q with a 0 inserted at the target bit
let gj = gi | h.tbit;
if ((gi & h.ctrl) == h.ctrl) {
let a0 = vec2<u32>(re[gi], im[gi]);
let a1 = vec2<u32>(re[gj], im[gj]);
let o0 = row(h.row0, a0, a1);
let o1 = row(h.row1, a0, a1);
re[gi] = o0.y;
im[gi] = o0.w;
re[gj] = o1.y;
im[gj] = o1.w;
if (!(fits(o0.xy) && fits(o0.zw) && fits(o1.xy) && fits(o1.zw))) { ok = false; }
}
}
}
} else {
malformed = true;
}
if (!ok) { bad = select(bad, h.index, h.index < bad); }
storageBarrier(); // this gate's writes before the next gate's reads
}
if (malformed) { atomicOr(&report[0], 4u); }
if (bad != NONE) {
atomicOr(&report[0], 1u);
atomicMin(&report[1], bad);
}
}

597
wai/wai-quantum/src/quantum_fabric/plan.rs
View file

@ -4,20 +4,107 @@
//! configuration.
//!
//! A [`Program`] holds the circuit's gates as [`GateWords`], in op-log order,
//! and the steps that apply them. Each step here is one gate over the whole
//! state. The arithmetic is [`lanes::apply_words`]: the stored lanes are the
//! low 32 bits of the exact values, and a run reports the first op whose exact
//! result left 32 bits, which no valid circuit of at most [`lanes::horizon`]
//! gates can do.
//! and the steps that apply them. A [`Step::Gate`] is one gate over the whole
//! state. A [`Step::Tile`] is a run of consecutive gates applied tile by tile:
//! the state is cut into tiles of `size` consecutive amplitudes, and each tile
//! goes through every gate of the run, in op-log order, on its own, so that a
//! member can apply a whole run in one dispatch, one group of invocations per
//! tile. A gate joins a run when it acts inside a tile: a gate whose matrix is
//! diagonal (`m01 = m10 = 0`) acts on each amplitude alone, whatever its
//! target, and any other gate acts on pairs that lie inside one tile when its
//! target's bit is below the tile's size. Every other gate is a step of its
//! own. [`Config`] says how a
//! circuit is planned: [`Config::PER_GATE`], one step per gate, is what
//! [`compile`] makes; [`Config::TILED`] makes tile runs of up to 1024
//! amplitudes.
//!
//! The arithmetic is [`lanes::apply_words`]'s in every configuration: the
//! stored lanes are the low 32 bits of the exact values, and a run reports the
//! first op whose exact result left 32 bits, which no valid circuit of at most
//! [`lanes::horizon`] gates can do.
//!
//! # Why every configuration gives the same bits
//!
//! A planner may only move data exactly; it never changes what is computed.
//!
//! - **A tile run is its gates, one after another.** Each gate of a run acts
//! on amplitudes of one tile only, so the tiles evolve independently, and
//! within a tile the gates apply in op-log order to the values the previous
//! gate left. Every pair or amplitude a gate updates is computed from the
//! same inputs by the same exact rows as in [`lanes::apply_words`], and the
//! same pairs are updated (a control is never the target, so both members of
//! a pair pass or fail the control test together). Splitting a run in two is
//! exact for the same reason.
//! - **A diagonal gate acts per amplitude.** `apply_words` computes
//! `a_i' = m00·a_i + m01·a_j` with `m01 = 0`, and the exact `fx(0, x)` is 0
//! for every `x`, so the value is that of `m00·a_i + 0·0`; likewise
//! `a_j' = 0·0 + m11·a_j`. The range check covers the same four parts.
//! - **The identity is skipped** only when its words are exactly
//! `[2^30, 0, 0, 0, 0, 0, 2^30, 0]`: `fx(2^30, x) = x` and `fx(0, x) = 0`
//! for every 32-bit `x`, so the gate changes nothing and cannot leave 32
//! bits. Circuit/1's `P(19..=32)` is `(2^30 − 1, 0)`, not the identity, and
//! is never skipped.
//! - **The report is the lowest op index** whose exact result left 32 bits.
//! The lanes after a gate are the same in every configuration, so the set of
//! gates whose exact result leaves 32 bits is the same, and so is its least
//! element.
//!
//! What no configuration does, because each changes bits: negate a rounded
//! product (`fx(−m, x)` differs from `−fx(m, x)` at a tie), multiply two
//! gates' matrices into one, reorder gates (even ones on disjoint qubits), or
//! skip a gate that is not exactly the identity.
use super::lanes::{self, GateWords, Lane};
use crate::quantum::{Circuit, MAX_QUBITS, ONE, QuantumError};
/// Bumped whenever the meaning of a [`Program`] changes. An implementation id
/// that depends on more than [`Step::Gate`] includes it; the per-gate WGSL
/// kernel's id ([`super::wgsl::implementation_id`]) versions the one step it
/// runs with [`super::wgsl::HOST_LAYOUT`] instead.
pub const PLAN_REVISION: u32 = 1;
/// that runs more than [`Step::Gate`] includes it
/// ([`super::wgsl::tile_implementation_id`]); the per-gate WGSL kernel's id
/// ([`super::wgsl::implementation_id`]) versions the one step it runs with
/// [`super::wgsl::HOST_LAYOUT`] instead. Revision 2 added [`Step::Tile`].
pub const PLAN_REVISION: u32 = 2;
/// The largest tile, as a power of two: 1024 amplitudes, four for each
/// invocation of a tile kernel's workgroup of 256.
pub const TILE_BITS_MAX: u8 = 10;
/// The identity's words: `m00 = m11 = 2^30`, every other entry zero.
const IDENTITY: [u32; 8] = [ONE as u32, 0, 0, 0, 0, 0, ONE as u32, 0];
/// How a circuit is planned.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub struct Config {
/// Amplitudes per tile: `2^min(n, tile_bits)`. Zero plans no tiles (one
/// step per gate). At most [`TILE_BITS_MAX`].
pub tile_bits: u8,
/// Gates per tile run, at most; further capped at `2^30 / 2^n` (at least
/// one), so that one run's work stays bounded on the widest states.
pub run_max: u32,
/// Leave out every gate whose words are exactly the identity's.
pub skip_identity: bool,
}
impl Config {
/// One [`Step::Gate`] per op, nothing skipped: what [`compile`] makes.
pub const PER_GATE: Config = Config { tile_bits: 0, run_max: 1, skip_identity: false };
/// Tiles of up to 1024 amplitudes, runs of up to 1024 gates, the identity
/// skipped.
pub const TILED: Config = Config { tile_bits: TILE_BITS_MAX, run_max: 1024, skip_identity: true };
/// Amplitudes per tile on `n_qubits`: `2^min(n, tile_bits)`, or 0 when
/// the configuration plans no tiles or the state has no qubits.
pub fn tile_size(&self, n_qubits: u8) -> u32 {
if self.tile_bits == 0 || n_qubits == 0 { 0 } else { 1 << n_qubits.min(self.tile_bits).min(TILE_BITS_MAX) }
}
/// Gates per tile run on `n_qubits`, at most: `run_max`, capped at
/// `2^30 / 2^n` and at least one. A member may split a run into several
/// dispatches to bound each one's work (splitting a run is exact).
pub fn run_cap(&self, n_qubits: u8) -> u32 {
let budget = if n_qubits >= 30 { 1 } else { (1u32 << 30) >> n_qubits };
self.run_max.min(budget).max(1)
}
}
/// One gate of a program.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
@ -29,12 +116,45 @@ pub struct PlannedGate {
pub diag: bool,
}
impl PlannedGate {
/// The gate with these words at op-log position `index`; `diag` is read
/// from the words.
pub fn new(words: GateWords, index: u32) -> PlannedGate {
PlannedGate { words, index, diag: words.m[2..6] == [0; 4] }
}
/// Whether a member may run the gate: `diag` says what the words say (a
/// tile kernel trusts it to act per amplitude), and the index is not
/// `u32::MAX`, which the tile kernel's report keeps for "no op". Every
/// gate [`PlannedGate::new`] and [`compile`] make is well formed:
/// `validate` holds an op-log's length, and so every index, below
/// `u32::MAX`.
pub fn is_well_formed(&self) -> bool {
self.diag == (self.words.m[2..6] == [0; 4]) && self.index != u32::MAX
}
/// Whether the words are exactly the identity's (whatever the target and
/// controls): the gate changes nothing and cannot leave 32 bits.
pub fn is_identity(&self) -> bool {
self.words.m == IDENTITY
}
/// Whether the gate acts inside tiles of `size` amplitudes: diagonal, or
/// its target's bit below `size`.
fn fits_tile(&self, size: u32) -> bool {
self.diag || self.words.tbit < size
}
}
/// One step of a program.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
#[non_exhaustive]
pub enum Step {
/// `gates[gate]` over the whole state: one dispatch.
Gate { gate: u32 },
/// `gates[first..first + count]`, in order, on each tile of `size`
/// consecutive amplitudes (a power of two): one dispatch.
Tile { first: u32, count: u32, size: u32 },
}
/// What a member runs for one circuit from one basis state.
@ -49,24 +169,60 @@ pub struct Program {
/// Validates exactly as [`Circuit::simulate_from`] does, with the same errors
/// in the same order (`validate()`, then [`QuantumError::TooManyQubits`]),
/// then plans one [`Step::Gate`] per op, in op-log order.
/// then plans one [`Step::Gate`] per op, in op-log order:
/// [`compile_with`]`(c, start, &Config::PER_GATE)`.
pub fn compile(c: &Circuit, start: usize) -> Result<Program, QuantumError> {
compile_with(c, start, &Config::PER_GATE)
}
/// Validates as [`compile`] does, then plans the ops under `cfg`
/// ([`schedule`]).
///
/// Panics if `cfg.tile_bits` is above [`TILE_BITS_MAX`] or `cfg.run_max` is
/// zero.
pub fn compile_with(c: &Circuit, start: usize, cfg: &Config) -> Result<Program, QuantumError> {
c.validate()?;
if c.n_qubits > MAX_QUBITS {
return Err(QuantumError::TooManyQubits(c.n_qubits));
}
let gates: Vec<PlannedGate> = c
.ops
.iter()
.enumerate()
.map(|(k, g)| {
let words = lanes::gate_words(g, c.gateset);
// `validate` holds the op count to `u32`.
PlannedGate { words, index: k as u32, diag: words.m[2..6] == [0; 4] }
})
.collect();
let steps = (0..gates.len() as u32).map(|gate| Step::Gate { gate }).collect();
Ok(Program { n_qubits: c.n_qubits, start: start % (1usize << c.n_qubits), gates, steps })
// `validate` holds the op count to `u32`.
let gates = c.ops.iter().enumerate().map(|(k, g)| PlannedGate::new(lanes::gate_words(g, c.gateset), k as u32)).collect();
Ok(schedule(c.n_qubits, start, gates, cfg))
}
/// Plans `gates`, in op-log order, on a state of `n_qubits` from the basis
/// state `start % 2^n`: with `cfg.skip_identity`, every gate whose words are
/// exactly the identity's (and whose target and controls fit the state) is
/// left out; then each maximal run of consecutive gates that act inside a
/// tile of [`Config::tile_size`] amplitudes, split every [`Config::run_cap`]
/// gates, is a [`Step::Tile`], and every other gate a [`Step::Gate`].
///
/// Panics if `n_qubits` is above [`MAX_QUBITS`], `cfg.tile_bits` above
/// [`TILE_BITS_MAX`], or `cfg.run_max` zero.
pub fn schedule(n_qubits: u8, start: usize, gates: Vec<PlannedGate>, cfg: &Config) -> Program {
assert!(n_qubits <= MAX_QUBITS, "a program holds at most {MAX_QUBITS} qubits, not {n_qubits}");
assert!(cfg.tile_bits <= TILE_BITS_MAX, "tiles hold at most 2^{TILE_BITS_MAX} amplitudes, not 2^{}", cfg.tile_bits);
assert!(cfg.run_max >= 1, "a tile run holds at least one gate");
// An identity whose target or controls a kernel refuses stays, to be refused.
let len = 1usize << n_qubits;
let gates: Vec<PlannedGate> = if cfg.skip_identity { gates.into_iter().filter(|g| !(g.is_identity() && tile_shape(&g.words, len))).collect() } else { gates };
let (size, cap) = (cfg.tile_size(n_qubits), cfg.run_cap(n_qubits) as usize);
let mut steps = Vec::new();
let mut k = 0;
while k < gates.len() {
if size != 0 && gates[k].fits_tile(size) {
let first = k;
while k < gates.len() && k - first < cap && gates[k].fits_tile(size) {
k += 1;
}
// `validate` holds the op count, and so `first` and `k`, to `u32`.
steps.push(Step::Tile { first: first as u32, count: (k - first) as u32, size });
} else {
steps.push(Step::Gate { gate: k as u32 });
k += 1;
}
}
Program { n_qubits, start: start % len, gates, steps }
}
/// A run in lanes.
@ -79,32 +235,120 @@ pub struct LaneRun {
pub out_of_range: Option<u32>,
}
/// The program on the CPU: `amps[start] = (2^30, 0)`, then each step with
/// [`lanes::apply_words`]. Panics only if the program is malformed, which
/// [`compile`] never makes.
/// The program on the CPU: `amps[start] = (2^30, 0)`, then each step as
/// [`execute_from`] runs it. Panics only if the program is malformed, which
/// [`compile`], [`compile_with`] and [`schedule`] never make.
pub fn execute(p: &Program) -> LaneRun {
let mut amps = vec![[0i32; 2]; 1usize << p.n_qubits];
amps[p.start] = [ONE as i32, 0];
execute_from(p, &amps).expect("a planned program's words fit its state")
}
/// The program's steps on the CPU from `state` (which replaces the basis
/// state `p.start`): each [`Step::Gate`] with [`lanes::apply_words`], each
/// [`Step::Tile`] tile by tile with a tile kernel's own index arithmetic.
/// `None` if `state` is not `2^n` lanes or the program is malformed: a gate
/// that is not [`PlannedGate::is_well_formed`], a step naming gates it does
/// not hold, a tile size that is not a power of two of at most
/// `2^TILE_BITS_MAX` and the state's length, an empty run, or words a kernel
/// refuses (those [`lanes::apply_words`] refuses, and in a run a gate that
/// neither is diagonal nor has its target's bit below the tile's size). A
/// member refuses the same programs.
pub fn execute_from(p: &Program, state: &[Lane]) -> Option<LaneRun> {
if p.n_qubits > MAX_QUBITS || state.len() != 1usize << p.n_qubits || !p.gates.iter().all(PlannedGate::is_well_formed) {
return None;
}
let mut amps = state.to_vec();
let mut out_of_range: Option<u32> = None;
let mut flag = |index: u32| out_of_range = Some(out_of_range.map_or(index, |k| k.min(index)));
for step in &p.steps {
match *step {
Step::Gate { gate } => {
let g = &p.gates[gate as usize];
let fits = lanes::apply_words(&mut amps, &g.words).expect("a compiled program's words fit its state");
if !fits {
out_of_range = Some(out_of_range.map_or(g.index, |k| k.min(g.index)));
let g = p.gates.get(gate as usize)?;
if !lanes::apply_words(&mut amps, &g.words)? {
flag(g.index);
}
}
Step::Tile { first, count, size } => {
let run = p.gates.get(first as usize..(first as usize).checked_add(count as usize)?)?;
let size = size as usize;
if run.is_empty() || !size.is_power_of_two() || size > 1 << TILE_BITS_MAX || size > amps.len() {
return None;
}
for g in run {
if !tile_shape(&g.words, amps.len()) || !(g.diag || (g.words.tbit as usize) < size) {
return None;
}
}
for base in (0..amps.len()).step_by(size) {
for g in run {
if !tile_gate(&mut amps[base..base + size], base, g) {
flag(g.index);
}
}
}
}
}
}
LaneRun { amps, out_of_range }
Some(LaneRun { amps, out_of_range })
}
/// The predicate both kernels refuse a record by, for a state of `len` lanes:
/// `tbit` a power of two below `len`, `ctrl` clear of `tbit` and of every bit
/// at or above `len`.
fn tile_shape(w: &GateWords, len: usize) -> bool {
let (tbit, ctrl) = (w.tbit as usize, w.ctrl as usize);
tbit.is_power_of_two() && tbit < len && ctrl & tbit == 0 && ctrl & !(len - 1) == 0
}
/// The low 32 bits of an exact value, and whether it is a 32-bit one.
fn lane(v: [i64; 2]) -> (Lane, bool) {
// `as` keeps the low 32 bits.
([v[0] as i32, v[1] as i32], i32::try_from(v[0]).is_ok() && i32::try_from(v[1]).is_ok())
}
/// One gate of a run on the tile starting at amplitude `base`, as a tile
/// kernel computes it: a diagonal gate per amplitude, row 0 with a zero
/// partner where the target's bit is clear and row 1 where it is set; any
/// other gate per pair, the pair `q` being `q` with a zero inserted at the
/// target's bit and that index with the bit set. The control test reads the
/// state-wide index. Returns whether every exact result fits in 32 bits.
fn tile_gate(tile: &mut [Lane], base: usize, g: &PlannedGate) -> bool {
let (t, c) = (g.words.tbit as usize, g.words.ctrl as usize);
let mut fits = true;
if g.diag {
for (e, a) in tile.iter_mut().enumerate() {
let gi = base + e;
if gi & c == c {
let row = if gi & t == 0 { lanes::exact_row(&g.words, 0, *a, [0, 0]) } else { lanes::exact_row(&g.words, 1, [0, 0], *a) };
let (v, ok) = lane(row);
*a = v;
fits &= ok;
}
}
} else {
for q in 0..tile.len() / 2 {
let low = q & (t - 1);
let li = ((q ^ low) << 1) | low;
let lj = li | t;
if (base + li) & c == c {
let (a0, a1) = (tile[li], tile[lj]);
let (v0, ok0) = lane(lanes::exact_row(&g.words, 0, a0, a1));
let (v1, ok1) = lane(lanes::exact_row(&g.words, 1, a0, a1));
(tile[li], tile[lj]) = (v0, v1);
fits &= ok0 && ok1;
}
}
}
fits
}
#[cfg(test)]
mod tests {
use super::*;
use crate::quantum::{BaseGate, Gate, Gateset};
use crate::quantum_fabric::{Corpus, FABRIC_DEFAULT};
use crate::quantum_fabric::tests::configurations;
use crate::quantum_fabric::{Certification, Corpus, FABRIC_DEFAULT, Parity, apply_kernel, wgsl};
fn widened(lanes: &[Lane]) -> Vec<[i64; 2]> {
lanes.iter().map(|a| [i64::from(a[0]), i64::from(a[1])]).collect()
@ -114,6 +358,18 @@ mod tests {
c.simulate_from(start).unwrap().amps.iter().map(|a| [a.re, a.im]).collect()
}
fn hex(b: &[u8]) -> String {
b.iter().map(|x| format!("{x:02x}")).collect()
}
fn next(s: &mut u64) -> u64 {
*s = s.wrapping_add(0x9E37_79B9_7F4A_7C15);
let mut z = *s;
z = (z ^ (z >> 30)).wrapping_mul(0xBF58_476D_1CE4_E5B9);
z = (z ^ (z >> 27)).wrapping_mul(0x94D0_49BB_1331_11EB);
z ^ (z >> 31)
}
#[test]
fn execute_is_the_native_simulator() {
for gs in Gateset::ALL {
@ -141,17 +397,23 @@ mod tests {
assert_eq!(p.steps, (0..6).map(|gate| Step::Gate { gate }).collect::<Vec<_>>());
assert_eq!(p.gates.iter().map(|g| g.index).collect::<Vec<_>>(), [0, 1, 2, 3, 4, 5]);
assert_eq!(p.gates.iter().map(|g| g.diag).collect::<Vec<_>>(), [false, true, false, true, false, true]);
assert_eq!(compile_with(&c, 11, &Config::PER_GATE).unwrap(), p);
// No qubits: no gates, and the one amplitude.
let empty = compile(&Circuit::with_gateset(0, Gateset::V2), 9).unwrap();
assert_eq!((empty.start, empty.gates.len()), (0, 0));
assert_eq!(execute(&empty), LaneRun { amps: vec![[ONE as i32, 0]], out_of_range: None });
// It refuses what the native simulator refuses, with the same error.
for c in crate::quantum_fabric::tests::refused_circuits() {
assert_eq!(compile(&c, 0).unwrap_err(), c.simulate_from(0).unwrap_err(), "{c:?}");
assert_eq!(compile_with(&Circuit::with_gateset(0, Gateset::V2), 9, &Config::TILED).unwrap(), empty);
// It refuses what the native simulator refuses, with the same error,
// in every configuration.
for cfg in configurations() {
for c in crate::quantum_fabric::tests::refused_circuits() {
assert_eq!(compile_with(&c, 0, &cfg).unwrap_err(), c.simulate_from(0).unwrap_err(), "{c:?}");
}
}
let mut wide = Circuit::new(MAX_QUBITS + 1);
wide.ops.push(Gate { base: BaseGate::X, controls: vec![], target: 0, param: 0 });
assert_eq!(compile(&wide, 0).unwrap_err(), QuantumError::TooManyQubits(MAX_QUBITS + 1));
assert_eq!(compile_with(&wide, 0, &Config::TILED).unwrap_err(), QuantumError::TooManyQubits(MAX_QUBITS + 1));
assert_eq!(compile(&Circuit::new(MAX_QUBITS + 1), 0).unwrap_err(), QuantumError::TooManyQubits(MAX_QUBITS + 1));
// Validation comes first, as in the native simulator.
let mut both = Circuit::new(MAX_QUBITS + 1);
@ -159,4 +421,271 @@ mod tests {
assert_eq!(compile(&both, 0).unwrap_err(), both.simulate_from(0).unwrap_err());
assert!(matches!(compile(&both, 0), Err(QuantumError::Invalid(_))));
}
/// The lanes32 digests, random-v1 and fabric-v1, per gateset.
const PINS: [(Gateset, &str, &str); 2] = [
(Gateset::V1, "d18454c639acfd9c7b66450f4ded1aac4ab46cc10a7bedbcf4958a96e204fb36", "41c726a0ac80f7defef92e15ecf893bdb7a3076b5e7f670fa0b7142e5c033df7"),
(Gateset::V2, "3dbde276c1095bf827faec71d5980e050f8a55a5ba7d97cd3bbdaec3f43a13b2", "9a84b55af68f1a8b07390b234cf562dd8d85cef3185a697ab85690c8790a7809"),
];
/// Every configuration's program equals the native simulator over
/// random-v1 (300 circuits) and fabric-v1, both gatesets, and certifies
/// with the lanes32 pins: a certificate of one configuration
/// (implementation `lanes32-plan`) has the pinned corpus digest.
#[test]
fn every_configuration_is_the_native_simulator() {
let configs = configurations();
assert_eq!(configs.len(), 10);
for (gs, random_pin, fabric_pin) in PINS {
// fabric-v1 is random-v1's first 200 circuits, then the wide ones.
let fabric = FABRIC_DEFAULT.circuits(gs);
let more = Corpus::Random { seed: 7, cases: 300 }.circuits(gs);
assert_eq!(fabric[..200], more[..200]);
let circuits: Vec<&Circuit> = fabric.iter().chain(&more[200..]).collect();
let reference: Vec<_> = circuits.iter().map(|c| lanes::reference_output(c).unwrap()).collect();
for cfg in &configs {
let (mut random_v1, mut fabric_v1) = (Certification::new(), Certification::new());
for (k, (c, want)) in circuits.iter().zip(&reference).enumerate() {
let run = execute(&compile_with(c, 0, cfg).unwrap());
assert_eq!(run.out_of_range, None, "{cfg:?} {gs:?} case {k}");
let out = lanes::output(&run.amps);
assert_eq!(&out, want, "{cfg:?} {gs:?} case {k}");
let input = lanes::circuit_bytes(c);
if k < 200 {
random_v1.add(&input, want, &out);
}
if k < fabric.len() {
fabric_v1.add(&input, want, &out);
}
}
// A start other than zero, on random-v1's first circuits.
for c in &circuits[..40] {
assert_eq!(widened(&execute(&compile_with(c, 5, cfg).unwrap()).amps), native(c, 5), "{cfg:?} {gs:?} start 5");
}
let random_v1 = random_v1.finish(apply_kernel(gs), "lanes32-plan", "local/test", &Corpus::Random { seed: 7, cases: 200 }.generator(gs));
let fabric_v1 = fabric_v1.finish(apply_kernel(gs), "lanes32-plan", "local/test", &FABRIC_DEFAULT.generator(gs));
for (cert, cases, pin) in [(random_v1, 200, random_pin), (fabric_v1, 224, fabric_pin)] {
assert_eq!((cert.parity, cert.cases), (Parity::BitExact, cases), "{cfg:?} {gs:?}");
assert_eq!(hex(&cert.corpus), pin, "{cfg:?} {gs:?}");
}
}
}
}
/// A random circuit of `ops` gates on `n` qubits: every base gate, up to
/// three distinct controls, `P(1..=32)`.
fn random_circuit(n: u8, ops: usize, seed: u64, gs: Gateset) -> Circuit {
let bases = [BaseGate::I, BaseGate::X, BaseGate::Y, BaseGate::Z, BaseGate::H, BaseGate::S, BaseGate::Sdg, BaseGate::T, BaseGate::Tdg, BaseGate::P];
let mut s = seed;
let mut c = Circuit::with_gateset(n, gs);
for _ in 0..ops {
let base = bases[(next(&mut s) % bases.len() as u64) as usize];
let target = (next(&mut s) % u64::from(n)) as u8;
let mut controls = Vec::new();
for _ in 0..next(&mut s) % 4 {
let q = (next(&mut s) % u64::from(n)) as u8;
if q != target && !controls.contains(&q) {
controls.push(q);
}
}
let param = if base == BaseGate::P { 1 + (next(&mut s) % 32) as u16 } else { 0 };
c.ops.push(Gate { base, controls, target, param });
}
c
}
/// Long circuits: the tiled program equals the per-gate one and the
/// native simulator, with runs long enough to be split at their cap.
#[test]
fn long_circuits_tile_as_they_run_gate_by_gate() {
let short_runs = Config { tile_bits: 3, run_max: 7, skip_identity: true };
for gs in Gateset::ALL {
for (n, ops, seed) in [(4u8, 3000usize, 0x6c6f_6e67u64), (11, 600, 0x7469_6c65), (12, 400, 0x7275_6e73)] {
let c = random_circuit(n, ops, seed + n as u64, gs);
assert!(c.ops.len() as u64 <= lanes::horizon(n, gs));
let per_gate = execute(&compile(&c, 0).unwrap());
assert_eq!((per_gate.out_of_range, widened(&per_gate.amps)), (None, native(&c, 0)), "{gs:?} n = {n}");
for cfg in [Config::TILED, short_runs] {
let p = compile_with(&c, 0, &cfg).unwrap();
let longest = p.steps.iter().filter_map(|s| match *s { Step::Tile { count, .. } => Some(count), _ => None }).max().unwrap();
assert!(longest > 1, "{cfg:?} {gs:?} n = {n}");
assert!(longest <= cfg.run_cap(n));
assert_eq!(execute(&p), per_gate, "{cfg:?} {gs:?} n = {n}");
}
}
// A 4-qubit circuit is one tile, and every gate fits it: the
// runs are cut only at the cap.
let c = random_circuit(4, 3000, 3, gs);
let p = compile_with(&c, 0, &Config::TILED).unwrap();
let counts: Vec<u32> = p.steps.iter().map(|s| match *s { Step::Tile { count, size, .. } => { assert_eq!(size, 16); count } other => panic!("{other:?}") }).collect();
assert_eq!(counts.iter().sum::<u32>() as usize, p.gates.len());
assert!(counts[..counts.len() - 1].iter().all(|&k| k == 1024), "{counts:?}");
}
}
/// Any words, whatever their range: the tiled program equals the gates
/// one by one ([`lanes::run_words`]), lanes and report, with diagonal
/// gates on every target, controls above and below the tile's bits, and
/// runs cut by gates that do not fit a tile.
#[test]
fn tiles_are_exact_on_any_words() {
let mut s = 0x7469_6c65_7321u64;
let (mut flagged, mut clean, mut tiled) = (0, 0, 0);
for case in 0..1500 {
let n = 1 + (next(&mut s) % 12) as u8;
let full = case % 2 == 0;
let value = |s: &mut u64| -> i32 { if full { next(s) as i32 } else { (next(s) % (1 << 23)) as i32 - (1 << 22) } };
let state: Vec<Lane> = (0..1usize << n).map(|_| [value(&mut s), value(&mut s)]).collect();
let gates: Vec<GateWords> = (0..1 + next(&mut s) % 24)
.map(|_| {
let t = (next(&mut s) % u64::from(n)) as u32;
let ctrl = (0..u32::from(n)).filter(|&b| b != t && next(&mut s) % 4 == 0).fold(0u32, |acc, b| acc | 1 << b);
let mut m = [0u32; 8];
let entry = |s: &mut u64| -> u32 { if full { next(s) as u32 } else { ((next(s) % (1 << 30)) as i64 - (1 << 29)) as i32 as u32 } };
m.iter_mut().for_each(|x| *x = entry(&mut s));
match next(&mut s) % 4 {
0 | 1 => m[2..6].fill(0),
2 => m = IDENTITY,
_ => {}
}
GateWords { m, tbit: 1 << t, ctrl }
})
.collect();
let want = lanes::run_words(&state, &gates).expect("valid shapes");
let planned: Vec<PlannedGate> = gates.iter().enumerate().map(|(k, w)| PlannedGate::new(*w, k as u32)).collect();
for cfg in [Config::TILED, Config { tile_bits: 2, run_max: 3, skip_identity: false }] {
let p = schedule(n, 0, planned.clone(), &cfg);
if p.steps.iter().any(|s| matches!(s, Step::Tile { count, .. } if *count > 1)) {
tiled += 1;
}
assert_eq!(execute_from(&p, &state), Some(want.clone()), "case {case}, n = {n}, {cfg:?}");
}
if want.out_of_range.is_some() {
flagged += 1;
} else {
clean += 1;
}
}
assert!(flagged >= 100 && clean >= 100 && tiled >= 1000, "{flagged} flagged, {clean} clean, {tiled} with a run of two gates or more");
}
#[test]
fn tiles_cover_every_path() {
for gs in Gateset::ALL {
let mut count = std::collections::BTreeMap::<&str, usize>::new();
let mut add = |what: &'static str, k: usize| *count.entry(what).or_default() += k;
for c in FABRIC_DEFAULT.circuits(gs) {
let p = compile_with(&c, 0, &Config::TILED).unwrap();
add("identity skipped", c.ops.len() - p.gates.len());
for step in &p.steps {
match *step {
Step::Gate { .. } => add("per-gate step", 1),
Step::Tile { first, count, size } => {
add("tile run of two gates or more", usize::from(count > 1));
add("tile of 1024 amplitudes", usize::from(size == 1024));
add("tile below 1024 amplitudes", usize::from(size < 1024));
let tiles = (1u32 << c.n_qubits) / size;
add("tile grid of more than one row", usize::from(wgsl::tile_grid(tiles, wgsl::TILE_X_GROUPS_MAX).y > 1));
for g in &p.gates[first as usize..(first + count) as usize] {
add("diagonal gate on a target outside the tile", usize::from(g.diag && g.words.tbit >= size));
add("non-diagonal gate in a run", usize::from(!g.diag));
add("control above the tile's bits", usize::from(g.words.ctrl >= size));
add("control inside the tile", usize::from(g.words.ctrl & (size - 1) != 0));
}
}
}
}
}
assert_eq!(count.len(), 10, "{gs:?}: {count:?}");
for (what, k) in &count {
assert!(*k > 0, "{gs:?}: no {what}");
}
}
// The cap: 1024 gates per run up to 20 qubits, then half as many for
// every qubit more, 16 at 26.
for (n, cap) in [(1, 1024), (20, 1024), (21, 512), (24, 64), (26, 16)] {
assert_eq!(Config::TILED.run_cap(n), cap, "n = {n}");
}
assert_eq!((Config::TILED.tile_size(0), Config::TILED.tile_size(3), Config::TILED.tile_size(10), Config::TILED.tile_size(26)), (0, 8, 1024, 1024));
assert_eq!((Config::PER_GATE.tile_size(12), Config::PER_GATE.run_cap(12)), (0, 1));
}
#[test]
fn the_v1_tail_phases_are_not_skipped() {
// P(19..=32) of circuit/1 is (2^30 − 1, 0): not the identity.
let p19 = lanes::gate_words(&Gate { base: BaseGate::P, controls: vec![], target: 0, param: 19 }, Gateset::V1);
assert_eq!(p19.m, [ONE as u32, 0, 0, 0, 0, 0, (ONE - 1) as u32, 0]);
assert!(!PlannedGate::new(p19, 0).is_identity() && PlannedGate::new(p19, 0).diag);
let mut c = Circuit::with_gateset(1, Gateset::V1);
c.h(0).p(19, 0);
let p = compile_with(&c, 0, &Config::TILED).unwrap();
assert_eq!(p.gates.len(), 2);
assert_eq!(p.steps, [Step::Tile { first: 0, count: 2, size: 2 }]);
let run = execute(&p);
assert_eq!(widened(&run.amps), native(&c, 0));
let h = Gateset::V1.inv_sqrt2() as i32;
assert_eq!(run.amps, [[h, 0], [h - 1, 0]]);
// It takes 2^29 + 1 to 2^29.
let tail = schedule(1, 0, vec![PlannedGate::new(p19, 0)], &Config::TILED);
assert_eq!(tail.steps, [Step::Tile { first: 0, count: 1, size: 2 }]);
let run = execute_from(&tail, &[[0, 0], [(1 << 29) + 1, 0]]).unwrap();
assert_eq!(run.amps, [[0, 0], [1 << 29, 0]]);
// The identity itself is skipped, and an identity left in changes
// nothing.
let mut c = Circuit::with_gateset(2, Gateset::V1);
c.h(0);
c.ops.push(Gate { base: BaseGate::I, controls: vec![1], target: 0, param: 0 });
c.h(1);
let skipped = compile_with(&c, 0, &Config::TILED).unwrap();
assert_eq!(skipped.gates.iter().map(|g| g.index).collect::<Vec<_>>(), [0, 2]);
let kept = compile_with(&c, 0, &Config { skip_identity: false, ..Config::TILED }).unwrap();
assert_eq!(kept.gates.len(), 3);
assert_eq!(execute(&skipped), execute(&kept));
}
/// Programs a tile kernel refuses are refused here too.
#[test]
fn malformed_programs_are_refused() {
let x2 = PlannedGate::new(lanes::gate_words(&Gate { base: BaseGate::X, controls: vec![], target: 2, param: 0 }, Gateset::V2), 0);
let z0 = PlannedGate::new(lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 0, param: 0 }, Gateset::V2), 0);
let state = vec![[0, 0]; 8];
let program = |gates: Vec<PlannedGate>, steps: Vec<Step>| Program { n_qubits: 3, start: 0, gates, steps };
// A non-diagonal gate whose target is the tile's size: its pairs leave the tile.
assert_eq!(execute_from(&program(vec![x2], vec![Step::Tile { first: 0, count: 1, size: 4 }]), &state), None);
assert!(execute_from(&program(vec![x2], vec![Step::Tile { first: 0, count: 1, size: 8 }]), &state).is_some());
// Sizes and ranges a run cannot have.
for (first, count, size) in [(0, 0, 4), (0, 2, 4), (1, 1, 4), (0, 1, 3), (0, 1, 16), (0, 1, 0)] {
assert_eq!(execute_from(&program(vec![z0], vec![Step::Tile { first, count, size }]), &state), None, "{first} {count} {size}");
}
assert_eq!(execute_from(&program(vec![z0], vec![Step::Gate { gate: 1 }]), &state), None);
// Words a kernel refuses, inside a run.
for w in [GateWords { tbit: 3, ..z0.words }, GateWords { tbit: 8, ..z0.words }, GateWords { ctrl: 1, ..z0.words }, GateWords { ctrl: 8, ..z0.words }] {
assert_eq!(execute_from(&program(vec![PlannedGate { words: w, ..z0 }], vec![Step::Tile { first: 0, count: 1, size: 8 }]), &state), None, "{w:?}");
}
// An identity a kernel would refuse is not skipped, so it is refused.
let refused_identity = PlannedGate::new(GateWords { m: IDENTITY, tbit: 3, ctrl: 0 }, 0);
let p = schedule(3, 0, vec![refused_identity], &Config::TILED);
assert_eq!(p.gates.len(), 1);
assert_eq!(execute_from(&p, &state), None);
// A gate whose `diag` disagrees with its words: X planned as diagonal
// would act per amplitude, and wrongly; a diagonal gate planned as not
// diagonal is refused alike. In a run and alone.
let x0 = PlannedGate::new(lanes::gate_words(&Gate { base: BaseGate::X, controls: vec![], target: 0, param: 0 }, Gateset::V2), 0);
for g in [PlannedGate { diag: true, ..x0 }, PlannedGate { diag: false, ..z0 }] {
assert!(!g.is_well_formed());
for step in [Step::Tile { first: 0, count: 1, size: 8 }, Step::Gate { gate: 0 }] {
assert_eq!(execute_from(&program(vec![g], vec![step]), &state), None, "{g:?} {step:?}");
}
}
assert!(x0.is_well_formed() && z0.is_well_formed());
// The index the tile kernel keeps for "no op".
let last = PlannedGate { index: u32::MAX, ..z0 };
assert!(!last.is_well_formed());
for step in [Step::Tile { first: 0, count: 1, size: 8 }, Step::Gate { gate: 0 }] {
assert_eq!(execute_from(&program(vec![last], vec![step]), &state), None);
}
assert!(execute_from(&program(vec![PlannedGate { index: u32::MAX - 1, ..z0 }], vec![Step::Tile { first: 0, count: 1, size: 8 }]), &state).is_some());
// A state of the wrong length.
assert_eq!(execute_from(&program(vec![z0], vec![Step::Gate { gate: 0 }]), &state[..4]), None);
}
}

401
wai/wai-quantum/src/quantum_fabric/wgsl.rs
View file

@ -90,14 +90,73 @@
//! to. The rules are checked on the translator's own reading of the text by
//! the unpublished crate `wai/tools/quantum-wgsl-check`.
//!
//! # Tile runs: a second implementation
//!
//! [`TILE_SHADER`] is [`SHADER`] followed by `lanes32_tile.wgsl`, whose entry
//! point `apply_tile` runs a [`Step::Tile`](super::plan::Step::Tile) of a
//! program planned with [`Config::TILED`](super::plan::Config::TILED): one
//! workgroup of [`TILE_WORKGROUP`] invocations per tile of up to 1024
//! consecutive amplitudes, which it alone touches for the whole dispatch while
//! it applies every gate of the run to the tile in place, in op-log order. A
//! run of gates is one dispatch instead of one per gate. The steps of a tiled
//! program that are single gates run on the same module's `apply`. It is its
//! own implementation ([`tile_implementation_id`]), certified on its own. A
//! member runs the per-gate kernel unless its caller asks for tile runs, and
//! falls back to it; [`place`](super::place), with both implementations'
//! entries for one member and no energy figures, breaks the tie by
//! implementation name and places the tile implementation
//! (`wgsl-lanes32-tile/…` sorts before `wgsl-lanes32/…`).
//!
//! - **The same function.** `apply_tile` computes every output with `row`, the
//! per-gate kernel's exact 64-bit row, and stores its low words: for a pair,
//! the per-gate kernel's computation; for a diagonal gate, the row with a
//! zero partner, whose exact value is the per-gate one because `fx64(0, x)`
//! is exactly 0 for every `x`. A gate is flagged when any exact output does
//! not fit, and the report keeps the lowest op index flagged over every
//! tile and gate. The tiled program on the CPU ([`plan::execute`]) mirrors
//! this index arithmetic tile by tile, and equals the per-gate program on
//! any words ([`plan`](super::plan) documents why it may).
//! - **No race.** Within a gate each amplitude (diagonal) or pair (otherwise)
//! belongs to one invocation; a storage barrier (workgroup execution scope,
//! storage memory) separates each gate from the next, so a gate reads what
//! the gate before wrote; tiles are disjoint, one workgroup each, and the
//! dispatch ends before the next step reads the state. A run record is
//! refused (bit 4) and nothing is touched unless its size is a power of two
//! of at most 1024, its run holds a gate, its tiles cover each state plane
//! exactly, and its grid (`num_workgroups`, one deep) holds every tile once
//! with `x_groups` workgroups to a row. A tile outside the state is reported
//! (bit 2) and not touched; a record outside the gate buffer reads as
//! malformed, never out of bounds.
//! - **A uniform barrier.** The early returns depend only on the run record (a
//! uniform buffer), the workgroup id, the number of workgroups and the
//! state planes' lengths. The barrier sits at the top level
//! of the loop over the run's gates, whose bound is in the uniform record;
//! that loop holds no `break`, `continue` or `return`, and no branch that
//! depends on a gate or an invocation holds a barrier. Naga does not check
//! barrier uniformity; a browser's WGSL compiler, which does, compiles this
//! text in the browser leg of the member's checks.
//! - **The size travels in a uniform.** The tile size, the gate range and the
//! grid are fields of the run record, never substituted into the text, so
//! the implementation id names the exact text every device compiles.
//!
//! The tile file keeps the per-gate rules (unsigned 32-bit values only, literal
//! shift amounts, no conversion, division or float, no workgroup memory) and
//! adds exactly one barrier, `storageBarrier()`. It holds no workgroup memory
//! because the SPIR-V the native translator writes for a workgroup array or
//! structure carries explicit layout decorations, which a Vulkan SPIR-V
//! validator rejects for the workgroup storage class; the tile stays in the
//! state's own planes instead. A device must be caught running each of
//! [`TILE_MUTANTS`].
//!
//! # What the argument does not cover
//!
//! Whether a device and its shader compiler conform is not proved here. A
//! member earns that per adapter: by reproducing the pinned corpus digests of
//! [`FABRIC_DEFAULT`](super::FABRIC_DEFAULT), and by being caught running each
//! of [`MUTANTS`]. Any edit to [`SHADER`] or to the host layout is a new
//! implementation, with a new [`implementation_id`], whose certificates must be
//! earned again.
//! of [`MUTANTS`] (and, for tile runs, [`TILE_MUTANTS`]). Any edit to
//! [`SHADER`], to the tile file, to the host layout or to the meaning of a
//! program is a new implementation, with a new [`implementation_id`] or
//! [`tile_implementation_id`], whose certificates must be earned again.
//!
//! [`lanes`]: super::lanes
//! [`lanes::apply`]: super::lanes::apply
@ -105,7 +164,8 @@
//! [`lanes::horizon`]: super::lanes::horizon
//! [`plan::execute`]: super::plan::execute
use super::lanes::GateWords;
use super::lanes::{GateWords, Lane};
use super::plan::{PLAN_REVISION, PlannedGate, Program, Step};
use super::{FabricClass, put};
use crate::quantum::MAX_QUBITS;
@ -116,8 +176,9 @@ pub const SHADER: &str = include_str!("lanes32.wgsl");
pub const IMPLEMENTATION: &str = "wgsl-lanes32";
/// Version of the host-side layout: planar `u32` planes, the 64-byte gate
/// record, the grid rule. Bumped on any change to them; it is part of the
/// implementation id.
/// record, the grid rule, and for tile runs the 48-byte tile gate record, the
/// 32-byte run record and the tile grid rule. Bumped on any change to them;
/// it is part of both implementation ids.
pub const HOST_LAYOUT: u32 = 1;
/// Invocations per workgroup, as the kernel declares.
@ -192,6 +253,84 @@ pub fn gate_grid(n_qubits: u8) -> Grid {
grid(1 << (n_qubits - 1), X_GROUPS_MAX)
}
/// The kernel with tile runs, `lanes32.wgsl` followed by `lanes32_tile.wgsl`,
/// exactly as it is compiled: the entry points `apply` and `probe` of
/// [`SHADER`], and `apply_tile`.
pub const TILE_SHADER: &str = concat!(include_str!("lanes32.wgsl"), include_str!("lanes32_tile.wgsl"));
/// The tile implementation's name, which its id starts with.
pub const TILE_IMPLEMENTATION: &str = "wgsl-lanes32-tile";
/// Invocations per workgroup of `apply_tile`, as it declares: a tile of up to
/// 1024 amplitudes is four per invocation, or two pairs.
pub const TILE_WORKGROUP: u32 = 256;
/// Workgroups along x of a tile step, at most: at 1024 amplitudes a tile, two
/// rows or more from 14 qubits on, inside the certification corpus.
pub const TILE_X_GROUPS_MAX: u32 = 8;
/// Bytes of one tile gate record, the tile kernel's `TileGate`.
pub const TILE_GATE_RECORD_BYTES: usize = 48;
/// Bytes of one run record, the tile kernel's `TileRun` uniform.
pub const TILE_RUN_RECORD_BYTES: usize = 32;
/// `"wgsl-lanes32-tile/<16 hex>"`: the first 8 bytes of
/// BLAKE3(`"wai:quantum-fabric-wgsl\x01"` ‖ put([`TILE_SHADER`]) ‖
/// [`HOST_LAYOUT`] ‖ [`PLAN_REVISION`], both big-endian `u32`), in hex. The
/// tile size is a field of the run record, never part of the text, so this
/// names the exact text every device compiles.
pub fn tile_implementation_id() -> String {
let mut o = DOMAIN_IMPLEMENTATION.to_vec();
put(&mut o, TILE_SHADER.as_bytes());
o.extend_from_slice(&HOST_LAYOUT.to_be_bytes());
o.extend_from_slice(&PLAN_REVISION.to_be_bytes());
format!("{TILE_IMPLEMENTATION}/{}", hex(&blake3::hash(&o).as_bytes()[..8]))
}
/// The 48-byte little-endian storage record of one gate of a tile run:
/// `m[0..8]`, `tbit`, `ctrl`, `diag` (1 or 0), `index`.
pub fn tile_gate_record(g: &PlannedGate) -> [u8; TILE_GATE_RECORD_BYTES] {
let mut r = [0u8; TILE_GATE_RECORD_BYTES];
let w = &g.words;
let words = w.m.into_iter().chain([w.tbit, w.ctrl, u32::from(g.diag), g.index]);
for (k, v) in words.enumerate() {
r[4 * k..4 * k + 4].copy_from_slice(&v.to_le_bytes());
}
r
}
/// The 32-byte little-endian uniform record of one tile run: `first` (its
/// first record in the tile gate buffer), `count`, `size`, `x_groups`,
/// `tiles`, then three zero words.
pub fn tile_run_record(first: u32, count: u32, size: u32, x_groups: u32, tiles: u32) -> [u8; TILE_RUN_RECORD_BYTES] {
let mut r = [0u8; TILE_RUN_RECORD_BYTES];
for (k, v) in [first, count, size, x_groups, tiles, 0, 0, 0].into_iter().enumerate() {
r[4 * k..4 * k + 4].copy_from_slice(&v.to_le_bytes());
}
r
}
/// A tile step's dispatch: `x` by `y` workgroups, where the workgroup at
/// `(wx, wy)` holds tile `wy·x + wx` if it is below `tiles`.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub struct TileGrid {
pub x: u32,
pub y: u32,
pub tiles: u32,
}
/// Workgroups covering `tiles` tiles: `x = min(tiles, max_x)`,
/// `y = ⌈tiles / x⌉`. In production `max_x` is [`TILE_X_GROUPS_MAX`]; both
/// are powers of two, so `y = tiles / x` exactly.
///
/// Panics if `tiles` or `max_x` is zero.
pub fn tile_grid(tiles: u32, max_x: u32) -> TileGrid {
assert!(tiles >= 1 && max_x >= 1, "a tile grid covers at least one tile, with at least one workgroup along x");
let x = tiles.min(max_x);
TileGrid { x, y: tiles.div_ceil(x), tiles }
}
/// What a WebGPU member is, as its id hashes it. The device module fills it
/// from the adapter's information and the granted limits; the fields are plain
/// values here so that the id is testable without a device.
@ -272,6 +411,13 @@ pub enum Detector {
Certificate,
/// A result that left 32 bits is not reported.
Range,
/// The lanes of [`race_witness`], a tile run in which invocations read
/// amplitudes other invocations wrote for the gate before, differ from
/// the CPU's.
Race,
/// [`malformed_tile_run`], a run holding a gate the tile kernel must
/// refuse, is not refused as a host defect.
Malformed,
}
/// A textual edit of [`SHADER`] that a conforming device must be caught
@ -310,6 +456,72 @@ pub fn mutate(m: &Mutant) -> String {
SHADER.replacen(m.find, m.replace, 1)
}
/// The mutants of the tile kernel. Each `find` occurs exactly once in
/// [`TILE_SHADER`], in its tile part.
pub const TILE_MUTANTS: [Mutant; 3] = [
// The barrier between gates is gone: a gate may read what the gate before has not yet written.
Mutant {
name: "T1",
find: "storageBarrier(); // this gate's writes before the next gate's reads",
replace: "// no barrier between gates",
detector: Detector::Race,
},
// The report never keeps the op index of a result that left 32 bits.
Mutant { name: "T2", find: "select(bad, h.index, h.index < bad)", replace: "bad", detector: Detector::Range },
// A gate whose pairs leave the tile is run instead of refused.
Mutant { name: "T3", find: "h.tbit < size", replace: "h.tbit <= size", detector: Detector::Malformed },
];
/// [`TILE_SHADER`] with `m` applied. Panics unless `m.find` occurs exactly
/// once.
pub fn mutate_tile(m: &Mutant) -> String {
let n = TILE_SHADER.matches(m.find).count();
assert_eq!(n, 1, "mutant {}: `{}` occurs {n} times in the tile kernel", m.name, m.find);
TILE_SHADER.replacen(m.find, m.replace, 1)
}
/// A state of `2^n_qubits` lanes (10 or more qubits, at most 20) and a run of
/// 1024 gates that a tiled plan makes one tile step per tile, built so that a
/// tile kernel without a barrier between gates goes wrong: diagonal gates
/// whose controls (qubits 5, 6 and 7) leave most invocations of a workgroup
/// idle alternate with gates on qubits 0 to 4, whose pairs are held by other
/// invocations than the amplitudes the diagonal gates updated. Idle
/// invocations run ahead into the next gate and read amplitudes not yet
/// written. Every value spans the whole 32-bit range, so a stale read shows in
/// the lanes. Panics unless `10 ≤ n_qubits ≤ 20`.
pub fn race_witness(n_qubits: u8) -> (Vec<Lane>, Vec<GateWords>) {
assert!((10..=20).contains(&n_qubits), "the race witness takes 10 to 20 qubits, not {n_qubits}");
let mut s = 0x7261_6365u64 + u64::from(n_qubits);
let mut next = || {
s = s.wrapping_add(0x9E37_79B9_7F4A_7C15);
let mut z = s;
z = (z ^ (z >> 30)).wrapping_mul(0xBF58_476D_1CE4_E5B9);
z = (z ^ (z >> 27)).wrapping_mul(0x94D0_49BB_1331_11EB);
z ^ (z >> 31)
};
let state: Vec<Lane> = (0..1usize << n_qubits).map(|_| [next() as i32, next() as i32]).collect();
let gates = (0..1024u32)
.map(|r| {
let mut m = [0u32; 8];
m.iter_mut().for_each(|x| *x = next() as u32);
if r % 2 == 0 {
m[2..6].fill(0);
GateWords { m, tbit: 1 << (8 + (r / 2) % 2), ctrl: 0b1110_0000 }
} else {
GateWords { m, tbit: 1 << ((r / 2) % 5), ctrl: 0 }
}
})
.collect();
(state, gates)
}
/// A program a tile kernel must refuse as a host defect: on 3 qubits, a tile
/// step of 4 amplitudes holding `X` on qubit 2, whose pairs leave the tile.
pub fn malformed_tile_run() -> Program {
let x = GateWords { m: [0, 0, 1 << 30, 0, 1 << 30, 0, 0, 0], tbit: 4, ctrl: 0 };
Program { n_qubits: 3, start: 0, gates: vec![PlannedGate::new(x, 0)], steps: vec![Step::Tile { first: 0, count: 1, size: 4 }] }
}
#[cfg(test)]
mod tests {
use super::*;
@ -517,6 +729,30 @@ mod tests {
}
}
/// With both implementations' entries for one member and no energy
/// figures, placement breaks the tie by implementation name: the tile
/// implementation sorts first and is placed, as the documentation says.
#[test]
fn placement_breaks_an_implementation_tie_by_name() {
assert!(tile_implementation_id() < implementation_id());
let id = hardware().member_id();
for gs in Gateset::ALL {
let kernel = apply_kernel(gs);
let entry = |implementation: String| {
let certificate = Certificate { kernel: kernel.into(), implementation: implementation.clone(), member: id.clone(), generator: FABRIC_DEFAULT.generator(gs), parity: Parity::BitExact, cases: 224, corpus: [7; 32] };
KernelEntry { kernel: kernel.into(), implementation, certificate, energy: None }
};
for kernels in [vec![entry(implementation_id()), entry(tile_implementation_id())], vec![entry(tile_implementation_id()), entry(implementation_id())]] {
let m = Member { id: id.clone(), class: FabricClass::Gpu, presence: Presence::Detected, kernels, detail: String::new() };
let (_, k) = place(std::slice::from_ref(&m), kernel, &Parity::BitExact).unwrap();
assert_eq!(k.implementation, tile_implementation_id(), "{gs:?}");
}
// Offered alone, the per-gate entry is placed.
let m = Member { id: id.clone(), class: FabricClass::Gpu, presence: Presence::Detected, kernels: vec![entry(implementation_id())], detail: String::new() };
assert_eq!(place(std::slice::from_ref(&m), kernel, &Parity::BitExact).unwrap().1.implementation, implementation_id());
}
}
#[test]
fn mutants_apply_once() {
let names: Vec<&str> = MUTANTS.iter().map(|m| m.name).collect();
@ -642,4 +878,157 @@ mod tests {
let z = lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 0, param: 0 }, Gateset::V2);
assert_eq!(lanes::apply_words(&mut [[0, 0], [i32::MIN, 0]], &z), Some(false));
}
/// Changes only with a reviewed edit to the tile kernel or to the per-gate
/// kernel it is appended to, to the argument in this module's
/// documentation, to the host layout ([`HOST_LAYOUT`]) or to the meaning
/// of a program ([`PLAN_REVISION`]): every certificate of the old id must
/// then be earned again.
const PINNED_TILE_IMPLEMENTATION: &str = "wgsl-lanes32-tile/54e443520c020013";
#[test]
fn tile_implementation_id_is_pinned() {
let id = tile_implementation_id();
assert_eq!(id, PINNED_TILE_IMPLEMENTATION);
let digest = id.strip_prefix("wgsl-lanes32-tile/").expect("the implementation name");
assert_eq!(digest.len(), 16);
assert!(digest.bytes().all(|b| b.is_ascii_digit() || (b'a'..=b'f').contains(&b)), "{id}");
assert_ne!(id.split('/').nth(1), implementation_id().split('/').nth(1));
// The tile kernel is appended to the per-gate one, both with LF line
// endings, and names no tile size.
assert!(TILE_SHADER.starts_with(SHADER) && TILE_SHADER.len() > SHADER.len());
assert!(!TILE_SHADER.contains('\r'));
assert_eq!(TILE_SHADER.matches("@compute").count(), 3);
assert_eq!(TILE_SHADER.matches("storageBarrier()").count(), 1);
assert!(!TILE_SHADER.contains("workgroupBarrier") && !TILE_SHADER.contains("var<workgroup>"));
assert!(!SHADER.contains("Barrier"));
}
#[test]
fn tile_records_layout() {
let g = PlannedGate { words: GateWords { m: [1, 2, 3, 4, 5, 6, 7, 8], tbit: 4, ctrl: 3 }, index: 9, diag: true };
let r = tile_gate_record(&g);
let word = |r: &[u8], k: usize| u32::from_le_bytes(r[4 * k..4 * k + 4].try_into().unwrap());
for k in 0..8 {
assert_eq!(word(&r, k), k as u32 + 1, "m[{k}]");
}
assert_eq!((word(&r, 8), word(&r, 9), word(&r, 10), word(&r, 11)), (4, 3, 1, 9));
assert_eq!(word(&tile_gate_record(&PlannedGate { diag: false, ..g }), 10), 0);
let r = tile_run_record(5, 6, 1024, 8, 16);
assert_eq!((0..5).map(|k| word(&r, k)).collect::<Vec<_>>(), [5, 6, 1024, 8, 16]);
assert_eq!(r[20..32], [0; 12]);
// Little-endian: the low byte first.
assert_eq!(tile_run_record(0x0102_0304, 1, 2, 1, 1)[0..4], [0x04, 0x03, 0x02, 0x01]);
}
#[test]
fn tile_grids_cover_every_tile_once() {
for n in 0..=16u32 {
let tiles = 1u32 << n;
for max_x in [1, 3, TILE_X_GROUPS_MAX] {
let g = tile_grid(tiles, max_x);
assert!(g.x >= 1 && g.x <= max_x && g.tiles == tiles, "{g:?}");
let mut hits = vec![0u32; tiles as usize];
for wy in 0..g.y {
for wx in 0..g.x {
let tile = wy * g.x + wx;
if tile < tiles {
hits[tile as usize] += 1;
}
}
}
assert!(hits.iter().all(|&h| h == 1), "{g:?}");
assert!(u64::from(g.x) * u64::from(g.y - 1) < u64::from(tiles), "{g:?}");
}
// In production both are powers of two: an exact grid.
let g = tile_grid(tiles, TILE_X_GROUPS_MAX);
assert_eq!(g.x * g.y, tiles);
}
// 26 qubits in tiles of 1024: (8, 8192).
assert_eq!(tile_grid(1 << 16, TILE_X_GROUPS_MAX), TileGrid { x: 8, y: 8192, tiles: 1 << 16 });
}
#[test]
fn tile_mutants_apply_once() {
let names: Vec<&str> = TILE_MUTANTS.iter().map(|m| m.name).collect();
assert_eq!(names, ["T1", "T2", "T3"]);
for m in &TILE_MUTANTS {
let text = mutate_tile(m);
assert_ne!(text, TILE_SHADER, "{}", m.name);
// The edit is in the tile part: the per-gate kernel is untouched.
assert!(text.starts_with(SHADER), "{}", m.name);
assert!(!SHADER.contains(m.find), "{}", m.name);
}
let detectors: Vec<Detector> = TILE_MUTANTS.iter().map(|m| m.detector).collect();
assert_eq!(detectors, [Detector::Race, Detector::Range, Detector::Malformed]);
// T1 leaves no barrier.
assert!(!mutate_tile(&TILE_MUTANTS[0]).contains("storageBarrier()"));
// The per-gate mutants apply to the tile kernel's copy of the per-gate one as well.
for m in &MUTANTS {
assert_eq!(TILE_SHADER.matches(m.find).count(), 1, "{}", m.name);
}
}
/// A tile run with no barrier between its gates, in one schedule a device
/// may choose: each of the [`TILE_WORKGROUP`] invocations runs every gate
/// of the run before the next invocation starts. With the barrier, the
/// only order is gate by gate.
fn run_invocation_by_invocation(tile: &mut [Lane], base: usize, gates: &[PlannedGate]) {
let size = tile.len();
let lanes_of = |v: [i64; 2]| [v[0] as i32, v[1] as i32];
for lid in 0..TILE_WORKGROUP as usize {
for g in gates {
let (t, c) = (g.words.tbit as usize, g.words.ctrl as usize);
if g.diag {
for e in (0..4).map(|k| lid + TILE_WORKGROUP as usize * k).filter(|&e| e < size && (base + e) & c == c) {
let a = tile[e];
tile[e] = lanes_of(if (base + e) & t == 0 { lanes::exact_row(&g.words, 0, a, [0, 0]) } else { lanes::exact_row(&g.words, 1, [0, 0], a) });
}
} else {
for q in (0..2).map(|k| lid + TILE_WORKGROUP as usize * k).filter(|&q| q < size / 2) {
let low = q & (t - 1);
let li = ((q ^ low) << 1) | low;
let lj = li | t;
if (base + li) & c == c {
let (a0, a1) = (tile[li], tile[lj]);
(tile[li], tile[lj]) = (lanes_of(lanes::exact_row(&g.words, 0, a0, a1)), lanes_of(lanes::exact_row(&g.words, 1, a0, a1)));
}
}
}
}
}
}
#[test]
fn tile_mutant_witnesses_exist() {
use crate::quantum_fabric::plan::{self, Config};
// T1: the race witness is one tile step per tile, gate by gate it is
// the CPU's, and invocation by invocation, which a kernel without the
// barrier between gates may do, it is not.
for n in [10u8, 11] {
let (state, words) = race_witness(n);
let planned: Vec<PlannedGate> = words.iter().enumerate().map(|(k, w)| PlannedGate::new(*w, k as u32)).collect();
let p = plan::schedule(n, 0, planned.clone(), &Config::TILED);
assert_eq!(p.steps, [Step::Tile { first: 0, count: 1024, size: 1024 }], "n = {n}");
let want = lanes::run_words(&state, &words).unwrap();
assert_eq!(plan::execute_from(&p, &state), Some(want.clone()));
let mut racy = state.clone();
for base in (0..racy.len()).step_by(1024) {
run_invocation_by_invocation(&mut racy[base..base + 1024], base, &planned);
}
let differing = racy.iter().zip(&want.amps).filter(|(a, b)| a != b).count();
assert!(differing > 100, "n = {n}: {differing} lanes differ");
}
// T2: a result that leaves 32 bits inside a tile step is reported.
let z = lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 0, param: 0 }, Gateset::V2);
let p = plan::schedule(1, 0, vec![PlannedGate::new(z, 0)], &Config::TILED);
assert_eq!(p.steps, [Step::Tile { first: 0, count: 1, size: 2 }]);
assert_eq!(plan::execute_from(&p, &[[0, 0], [i32::MIN, 0]]).unwrap().out_of_range, Some(0));
// T3: the malformed run is refused, and the same gate in a tile that
// holds its pairs is not.
let bad = malformed_tile_run();
assert_eq!(plan::execute_from(&bad, &[[0, 0]; 8]), None);
let held = Program { steps: vec![Step::Tile { first: 0, count: 1, size: 8 }], ..bad.clone() };
assert!(plan::execute_from(&held, &[[0, 0]; 8]).is_some());
}
}

320
wai/wai-quantum/src/quantum_fabric_webgpu.rs
View file

@ -17,6 +17,19 @@
//! the state comes back in [`READBACK_CHUNK`] pieces. The kernel reports
//! exactly when the reference's result leaves 32 bits; the member then returns
//! [`WebGpuError::OutOfRange`], never a state.
//! - **Tile runs** ([`WebGpuMember::tiled`]): the same device running the tile
//! implementation ([`wgsl::tile_implementation_id`]), which plans circuits
//! with [`Config::TILED`] and runs each run of gates that act inside a tile
//! as one dispatch, one workgroup per tile; its run records go in a second
//! uniform ring and its gates in a storage buffer, per submission. It is
//! probed and certified on its own. Discovery opens per-gate members; a tile
//! member exists only when a caller asks for one, and the per-gate kernel
//! is the one to fall back to. Placement has no preference of its own:
//! given both implementations' entries for one member and no energy
//! figures, [`place`](crate::quantum_fabric::place) breaks the tie by
//! implementation name, and `wgsl-lanes32-tile/…` sorts before
//! `wgsl-lanes32/…`, so the tile implementation is placed; a caller who
//! wants the per-gate kernel offers only its entries.
//! - **Certification** ([`WebGpuMember::certify`], [`WebGpuMember::member`]):
//! the member's own run of a corpus, folded case by case into a
//! [`Certification`]. The corpus digest covers inputs and outputs only, so a
@ -61,7 +74,7 @@
use crate::quantum::{Amp, Circuit, Gateset, MAX_QUBITS, ONE, QuantumError, StateVector};
use crate::quantum_fabric::lanes::{self, GateWords, Lane};
use crate::quantum_fabric::plan::{self, LaneRun, Program, Step};
use crate::quantum_fabric::plan::{self, Config, LaneRun, PlannedGate, Program, Step};
use crate::quantum_fabric::wgsl::{self, AdapterIdentity};
use crate::quantum_fabric::{Certificate, Certification, Corpus, FabricClass, KernelEntry, Member, Presence, apply_kernel};
use core::future::Future;
@ -253,13 +266,15 @@ struct Sink {
lost: Option<String>,
}
/// The kernel compiled on a device: both entry points, with their layouts.
/// The kernel compiled on a device: its entry points, with their layouts;
/// `apply_tile` only for the tile kernel.
#[derive(Clone)]
struct Pipelines {
apply_layout: wgpu::BindGroupLayout,
probe_layout: wgpu::BindGroupLayout,
apply: wgpu::ComputePipeline,
probe: wgpu::ComputePipeline,
tile: Option<(wgpu::BindGroupLayout, wgpu::ComputePipeline)>,
}
/// One WebGPU adapter as a fabric member: a device opened on it, the kernel
@ -277,7 +292,12 @@ pub struct WebGpuMember {
detail: String,
max_qubits: u8,
max_x: u32,
/// Gate records in the uniform ring: gates per submission, at most.
/// How circuits and lanes are planned: [`Config::PER_GATE`], or
/// [`Config::TILED`] for the tile implementation.
config: Config,
/// Records in each uniform ring: gate steps and tile steps per
/// submission, at most, and tile gates per submission unless one run
/// holds more.
ring_records: u32,
/// Bytes of state read back per copy.
readback_chunk: u64,
@ -289,7 +309,7 @@ pub struct WebGpuMember {
impl std::fmt::Debug for WebGpuMember {
fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
f.debug_struct("WebGpuMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).field("max_x", &self.max_x).field("ring_records", &self.ring_records).field("readback_chunk", &self.readback_chunk).field("detail", &self.detail).finish()
f.debug_struct("WebGpuMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).field("config", &self.config).field("max_x", &self.max_x).field("ring_records", &self.ring_records).field("readback_chunk", &self.readback_chunk).field("detail", &self.detail).finish()
}
}
@ -336,7 +356,7 @@ async fn open(adapter: wgpu::Adapter) -> Result<WebGpuMember, Refusal> {
}
});
// 5. The kernel.
let pipelines = match kernel_pipelines(&device, Cow::Borrowed(wgsl::SHADER)).await {
let pipelines = match kernel_pipelines(&device, Cow::Borrowed(wgsl::SHADER), false).await {
Ok(p) => p,
Err(e) => return Err(refuse(&identity, e)),
};
@ -361,6 +381,7 @@ async fn open(adapter: wgpu::Adapter) -> Result<WebGpuMember, Refusal> {
detail,
max_qubits,
max_x: wgsl::X_GROUPS_MAX,
config: Config::PER_GATE,
ring_records: RING_RECORDS,
readback_chunk: READBACK_CHUNK,
implementation: wgsl::implementation_id(),
@ -446,16 +467,20 @@ impl Scopes {
}
}
/// The kernel's two entry points with explicit layouts: `apply` over
/// `{re, im, report: storage; gate record: uniform at a dynamic offset}`, and
/// `probe` over `{cases: storage; probe grid: uniform}`.
async fn kernel_pipelines(device: &wgpu::Device, source: Cow<'_, str>) -> Result<Pipelines, WebGpuError> {
let storage = |binding: u32| wgpu::BindGroupLayoutEntry {
/// The kernel's entry points with explicit layouts: `apply` over
/// `{re, im, report: storage; gate record: uniform at a dynamic offset}`,
/// `probe` over `{cases: storage; probe grid: uniform}`, and with `tiles`
/// `apply_tile` over `{re, im, report: storage; tile gates: read-only storage;
/// run record: uniform at a dynamic offset}`, four storage buffers, the
/// downlevel limit exactly.
async fn kernel_pipelines(device: &wgpu::Device, source: Cow<'_, str>, tiles: bool) -> Result<Pipelines, WebGpuError> {
let buffer = |binding: u32, read_only: bool| wgpu::BindGroupLayoutEntry {
binding,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: false }, has_dynamic_offset: false, min_binding_size: None },
ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only }, has_dynamic_offset: false, min_binding_size: None },
count: None,
};
let storage = |binding: u32| buffer(binding, false);
let uniform = |binding: u32, dynamic: bool, size: u64| wgpu::BindGroupLayoutEntry {
binding,
visibility: wgpu::ShaderStages::COMPUTE,
@ -482,8 +507,16 @@ async fn kernel_pipelines(device: &wgpu::Device, source: Cow<'_, str>) -> Result
};
let apply = pipeline(&apply_layout, "apply");
let probe = pipeline(&probe_layout, "probe");
let tile = tiles.then(|| {
let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("wai-quantum apply_tile"),
entries: &[storage(0), storage(1), storage(2), buffer(6, true), uniform(7, true, wgsl::TILE_RUN_RECORD_BYTES as u64)],
});
let tile = pipeline(&layout, "apply_tile");
(layout, tile)
});
scopes.close().await?;
Ok(Pipelines { apply_layout, probe_layout, apply, probe })
Ok(Pipelines { apply_layout, probe_layout, apply, probe, tile })
}
/// Resolves when `map_async`'s callback has run. Natively it waits on its own
@ -676,9 +709,16 @@ impl WebGpuMember {
self.settled()
}
/// How this member plans circuits and lanes: [`Config::PER_GATE`], or
/// [`Config::TILED`] for the tile implementation ([`Self::tiled`]).
pub fn config(&self) -> Config {
self.config
}
/// The kernel on any state of `2^n` lanes (`n ≥ 1`) and any words, each
/// gate's index its position: what [`lanes::run_words`] gives, with words
/// it calls malformed refused as [`WebGpuError::Internal`].
/// gate's index its position, planned with [`Self::config`]: what
/// [`lanes::run_words`] gives, with words it calls malformed refused as
/// [`WebGpuError::Internal`].
pub async fn run_lanes(&self, state: &[Lane], gates: &[GateWords]) -> Result<LaneRun, WebGpuError> {
let len = state.len();
if len < 2 || !len.is_power_of_two() {
@ -688,12 +728,16 @@ impl WebGpuMember {
if n > u32::from(self.max_qubits) {
return Err(WebGpuError::TooWide { n_qubits: n.min(255) as u8, max: self.max_qubits });
}
let gates: Vec<(GateWords, u32)> = gates.iter().enumerate().map(|(k, w)| (*w, u32::try_from(k).unwrap_or(u32::MAX))).collect();
self.run(n as u8, Init::Lanes(state), &gates).await
let planned: Vec<PlannedGate> = gates.iter().enumerate().map(|(k, w)| PlannedGate::new(*w, u32::try_from(k).unwrap_or(u32::MAX))).collect();
let p = plan::schedule(n as u8, 0, planned, &self.config);
self.run(n as u8, Init::Lanes(state), &p.gates, &p.steps).await
}
/// A compiled program ([`plan::compile`]): the lanes and the first op
/// whose exact result left 32 bits, as [`plan::execute`] gives them.
/// A planned program ([`plan::compile`], [`plan::compile_with`],
/// [`plan::schedule`]): the lanes and the first op whose exact result left
/// 32 bits, as [`plan::execute`] gives them. Tile steps need the tile
/// implementation ([`Self::tiled`]); a member without it, or a step it
/// cannot lay out, is refused as [`WebGpuError::Internal`].
pub async fn run_program(&self, p: &Program) -> Result<LaneRun, WebGpuError> {
let n = p.n_qubits;
if n == 0 {
@ -705,24 +749,15 @@ impl WebGpuMember {
if n > MAX_QUBITS || p.start >= 1usize << n {
return Err(WebGpuError::Internal(2));
}
let mut gates = Vec::new();
gates.try_reserve_exact(p.steps.len()).map_err(|e| WebGpuError::Memory(e.to_string()))?;
for step in &p.steps {
match *step {
Step::Gate { gate } => {
let g = p.gates.get(gate as usize).ok_or(WebGpuError::Internal(4))?;
gates.push((g.words, g.index));
}
}
}
self.run(n, Init::Basis(p.start), &gates).await
self.run(n, Init::Basis(p.start), &p.gates, &p.steps).await
}
/// Equal to `c.simulate_from(start)`, or a typed refusal: what the CPU
/// reference refuses ([`WebGpuError::Circuit`]), more qubits than the
/// member holds, or a result that leaves 32 bits.
/// member holds, or a result that leaves 32 bits. Planned with
/// [`Self::config`].
pub async fn simulate_from(&self, c: &Circuit, start: usize) -> Result<StateVector, WebGpuError> {
let p = plan::compile(c, start).map_err(WebGpuError::Circuit)?;
let p = plan::compile_with(c, start, &self.config).map_err(WebGpuError::Circuit)?;
if p.n_qubits > self.max_qubits {
return Err(WebGpuError::TooWide { n_qubits: p.n_qubits, max: self.max_qubits });
}
@ -742,14 +777,15 @@ impl WebGpuMember {
}
/// The certificate of [`apply_kernel`]`(gs)` as [`Self::implementation`]
/// on this member, over `corpus`: each circuit run from `|0…0⟩` against the
/// native simulator. A range report inside a corpus is an error: its
/// circuits are valid and within the horizon.
/// on this member, over `corpus`: each circuit planned with
/// [`Self::config`] and run from `|0…0⟩`, against the native simulator. A
/// range report inside a corpus is an error: its circuits are valid and
/// within the horizon.
pub async fn certify(&self, gs: Gateset, corpus: Corpus) -> Result<Certificate, WebGpuError> {
let mut cert = Certification::new();
for c in corpus.circuits(gs) {
let reference = lanes::reference_output(&c).map_err(WebGpuError::Circuit)?;
let run = self.run_program(&plan::compile(&c, 0).map_err(WebGpuError::Circuit)?).await?;
let run = self.run_program(&plan::compile_with(&c, 0, &self.config).map_err(WebGpuError::Circuit)?).await?;
if let Some(gate) = run.out_of_range {
return Err(WebGpuError::OutOfRange { gate });
}
@ -794,14 +830,35 @@ impl WebGpuMember {
/// The same device running another kernel text, as implementation
/// `"mutant"`: for showing that a device running an edited kernel is
/// caught. Fails if the text does not compile.
/// caught. On a tile member ([`Self::tiled`]) the text must hold
/// `apply_tile` too. Fails if the text does not compile.
#[doc(hidden)]
pub async fn with_shader(&self, source: &str) -> Result<WebGpuMember, WebGpuError> {
self.alive()?;
let pipelines = kernel_pipelines(&self.device, Cow::Owned(source.to_string())).await?;
let pipelines = kernel_pipelines(&self.device, Cow::Owned(source.to_string()), self.pipelines.tile.is_some()).await?;
Ok(WebGpuMember { pipelines, implementation: "mutant".into(), ..self.clone() })
}
/// The same device running the tile implementation:
/// [`wgsl::TILE_SHADER`] compiled, its arithmetic probed in this process
/// as discovery probes the per-gate kernel's, circuits and lanes planned
/// with [`Config::TILED`], and implementation
/// [`wgsl::tile_implementation_id`]. A tile member earns its own
/// certificates ([`Self::certify`], [`Self::member`]); this member is
/// unchanged. See the module documentation for how placement orders the
/// two implementations.
pub async fn tiled(&self) -> Result<WebGpuMember, WebGpuError> {
self.alive()?;
let pipelines = kernel_pipelines(&self.device, Cow::Borrowed(wgsl::TILE_SHADER), true).await?;
let member = WebGpuMember { pipelines, implementation: wgsl::tile_implementation_id(), config: Config::TILED, ..self.clone() };
let pairs = probe_pairs(4096, 0x5eed);
let out = member.probe(&pairs).await?;
match first_mismatch(&pairs, &out) {
None => Ok(member),
Some((a, b)) => Err(WebGpuError::Probe { a, b }),
}
}
/// An error if the device has been lost.
fn alive(&self) -> Result<(), WebGpuError> {
match &lock(&self.sink).lost {
@ -823,6 +880,13 @@ impl WebGpuMember {
Err(WebGpuError::Device(std::mem::take(&mut s.messages).join("; ")))
}
/// The dispatch of a tile step over `tiles` tiles: at most
/// [`wgsl::TILE_X_GROUPS_MAX`] workgroups along x, fewer under a forced
/// geometry ([`Self::with_max_x`]).
fn tile_grid(&self, tiles: u32) -> wgsl::TileGrid {
wgsl::tile_grid(tiles, self.max_x.min(wgsl::TILE_X_GROUPS_MAX))
}
fn buffer(&self, label: &str, size: u64, usage: wgpu::BufferUsages) -> wgpu::Buffer {
self.device.create_buffer(&wgpu::BufferDescriptor { label: Some(label), size, usage, mapped_at_creation: false })
}
@ -871,14 +935,68 @@ impl WebGpuMember {
Ok(())
}
/// Runs `gates` (words and op-log index) over `2^n` lanes from `init`:
/// one dispatch per gate, in place, in order.
async fn run(&self, n: u8, init: Init<'_>, gates: &[(GateWords, u32)]) -> Result<LaneRun, WebGpuError> {
/// Runs `steps` over `gates` on `2^n` lanes from `init`, in place, in
/// order: one dispatch per gate step, over the whole state, and one or
/// more per tile step, one workgroup per tile. A tile step whose gates
/// would make one dispatch of more than [`SUBMIT_BUDGET`] pair-updates is
/// split into dispatches of fewer gates each (splitting a run is exact),
/// so that every submission stays within the budget. A program the host
/// cannot lay out (a gate that is not
/// [`PlannedGate::is_well_formed`], a gate or range it does not hold, a
/// tile size of zero or above the state's length, a tile step on a member
/// without the tile kernel, a grid past the per-dimension limit) is
/// refused as [`WebGpuError::Internal`] before anything runs; the kernels
/// refuse the rest with report bit 4.
async fn run(&self, n: u8, init: Init<'_>, gates: &[PlannedGate], steps: &[Step]) -> Result<LaneRun, WebGpuError> {
self.alive()?;
let len = 1usize << n;
let plane = 4u64 << n;
let grid = wgsl::grid(1u32 << (n - 1), self.max_x);
let ring_size = u64::from(self.ring_records) * self.stride;
let pairs = 1u64 << (n - 1);
// What the host lays out: well-formed gates, every step in range, and
// the tile kernel for tile steps, each split to at most `per_dispatch`
// gates.
if !gates.iter().all(PlannedGate::is_well_formed) {
return Err(WebGpuError::Internal(4));
}
let per_dispatch = (SUBMIT_BUDGET / pairs).clamp(1, u64::from(u32::MAX)) as u32;
let mut laid: Vec<Step> = Vec::new();
laid.try_reserve(steps.len()).map_err(|e| WebGpuError::Memory(e.to_string()))?;
let mut longest = 0usize;
for step in steps {
match *step {
Step::Gate { gate } => {
if gate as usize >= gates.len() {
return Err(WebGpuError::Internal(4));
}
laid.push(*step);
}
Step::Tile { first, count, size } => {
let end = (first as usize).checked_add(count as usize);
let tiles = if size == 0 || size as usize > len { 0 } else { (len / size as usize) as u32 };
if self.pipelines.tile.is_none() || end.is_none_or(|e| e > gates.len()) || tiles == 0 || self.tile_grid(tiles).y > 65535 {
return Err(WebGpuError::Internal(4));
}
if count == 0 {
// The kernel refuses an empty run itself (bit 4).
laid.push(*step);
}
let mut at = first;
while at < first + count {
let part = per_dispatch.min(first + count - at);
laid.push(Step::Tile { first: at, count: part, size });
longest = longest.max(part as usize);
at += part;
}
}
}
}
let steps = laid.as_slice();
let tile_steps = steps.iter().any(|s| matches!(s, Step::Tile { .. }));
// Tile gates per submission: a ring's worth, or the longest run.
let tile_capacity = (self.ring_records as usize).max(longest);
// Buffers and the starting state.
let scopes = Scopes::push(&self.device);
@ -917,31 +1035,127 @@ impl WebGpuMember {
},
],
});
// The tile kernel's own: run records in a second uniform ring, the
// runs' gates in a storage buffer, rewritten per submission.
let tiles = match (&self.pipelines.tile, tile_steps) {
(Some((layout, pipeline)), true) => {
let runs = self.buffer("wai-quantum tile runs", ring_size, wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST);
let tile_gates = self.buffer("wai-quantum tile gates", (tile_capacity * wgsl::TILE_GATE_RECORD_BYTES) as u64, wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST);
let bind = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("wai-quantum apply_tile"),
layout,
entries: &[
wgpu::BindGroupEntry { binding: 0, resource: re.as_entire_binding() },
wgpu::BindGroupEntry { binding: 1, resource: im.as_entire_binding() },
wgpu::BindGroupEntry { binding: 2, resource: report.as_entire_binding() },
wgpu::BindGroupEntry { binding: 6, resource: tile_gates.as_entire_binding() },
wgpu::BindGroupEntry {
binding: 7,
resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding { buffer: &runs, offset: 0, size: NonZeroU64::new(wgsl::TILE_RUN_RECORD_BYTES as u64) }),
},
],
});
Some((pipeline, runs, tile_gates, bind))
}
_ => None,
};
scopes.close().await?;
upload?;
// The gates, a chunk per submission, each waited on natively before
// the next, so that POLL_TIMEOUT bounds one chunk's GPU time.
let per_submit = (SUBMIT_BUDGET >> (n - 1)).clamp(1, u64::from(self.ring_records)) as usize;
// The steps, a chunk per submission, each waited on natively before
// the next, so that POLL_TIMEOUT bounds one chunk's GPU time. A chunk
// holds at least one step (at most SUBMIT_BUDGET pair-updates, as laid
// out above), and more while its pair-updates stay within
// SUBMIT_BUDGET, each ring holds its records and the tile gate buffer
// holds the runs' gates.
let ring_records = self.ring_records as usize;
let stride = self.stride as usize;
for chunk in gates.chunks(per_submit) {
let mut records = vec![0u8; chunk.len() * stride];
for (k, (words, index)) in chunk.iter().enumerate() {
records[k * stride..k * stride + wgsl::GATE_RECORD_BYTES].copy_from_slice(&wgsl::gate_record(words, grid.items, grid.row_stride, *index));
let mut next = 0;
while next < steps.len() {
let first = next;
let (mut cost, mut gate_steps, mut tile_runs, mut tile_gates) = (0u64, 0usize, 0usize, 0usize);
while next < steps.len() {
let (work, is_tile, held) = match steps[next] {
Step::Gate { .. } => (pairs, false, 0),
Step::Tile { count, .. } => (pairs.saturating_mul(u64::from(count)), true, count as usize),
};
let full = cost.saturating_add(work) > SUBMIT_BUDGET || if is_tile { tile_runs == ring_records || tile_gates + held > tile_capacity } else { gate_steps == ring_records };
if next > first && full {
break;
}
cost = cost.saturating_add(work);
if is_tile {
(tile_runs, tile_gates) = (tile_runs + 1, tile_gates + held);
} else {
gate_steps += 1;
}
next += 1;
}
let chunk = &steps[first..next];
let mut gate_records = vec![0u8; gate_steps * stride];
let mut run_records = vec![0u8; tile_runs * stride];
let mut tile_records = Vec::with_capacity(tile_gates * wgsl::TILE_GATE_RECORD_BYTES);
let (mut g, mut t) = (0usize, 0usize);
for step in chunk {
match *step {
Step::Gate { gate } => {
let pg = &gates[gate as usize];
gate_records[g * stride..g * stride + wgsl::GATE_RECORD_BYTES].copy_from_slice(&wgsl::gate_record(&pg.words, grid.items, grid.row_stride, pg.index));
g += 1;
}
Step::Tile { first, count, size } => {
let at = (tile_records.len() / wgsl::TILE_GATE_RECORD_BYTES) as u32;
for pg in &gates[first as usize..(first + count) as usize] {
tile_records.extend_from_slice(&wgsl::tile_gate_record(pg));
}
let tg = self.tile_grid((len / size as usize) as u32);
run_records[t * stride..t * stride + wgsl::TILE_RUN_RECORD_BYTES].copy_from_slice(&wgsl::tile_run_record(at, count, size, tg.x, tg.tiles));
t += 1;
}
}
}
let scopes = Scopes::push(&self.device);
// Ordered on the queue after the submissions that read the ring before.
self.queue.write_buffer(&ring, 0, &records);
// Ordered on the queue after the submissions that read them before.
if !gate_records.is_empty() {
self.queue.write_buffer(&ring, 0, &gate_records);
}
if let Some((_, runs, tile_gate_buffer, _)) = &tiles
&& !run_records.is_empty()
{
self.queue.write_buffer(runs, 0, &run_records);
self.queue.write_buffer(tile_gate_buffer, 0, &tile_records);
}
let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("wai-quantum gates") });
{
let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("wai-quantum gates"), timestamp_writes: None });
pass.set_pipeline(&self.pipelines.apply);
for k in 0..chunk.len() {
pass.set_bind_group(0, &bind, &[(k * stride) as u32]);
pass.dispatch_workgroups(grid.x, grid.y, 1);
let (mut g, mut t, mut on_tiles) = (0usize, 0usize, None);
for step in chunk {
match (*step, &tiles) {
(Step::Gate { .. }, _) => {
if on_tiles != Some(false) {
pass.set_pipeline(&self.pipelines.apply);
on_tiles = Some(false);
}
pass.set_bind_group(0, &bind, &[(g * stride) as u32]);
pass.dispatch_workgroups(grid.x, grid.y, 1);
g += 1;
}
(Step::Tile { size, .. }, Some((pipeline, _, _, tile_bind))) => {
if on_tiles != Some(true) {
pass.set_pipeline(pipeline);
on_tiles = Some(true);
}
let tg = self.tile_grid((len / size as usize) as u32);
pass.set_bind_group(0, tile_bind, &[(t * stride) as u32]);
pass.dispatch_workgroups(tg.x, tg.y, 1);
t += 1;
}
(Step::Tile { .. }, None) => unreachable!("tile steps were refused above without the tile kernel"),
}
}
}
let _index = self.queue.submit([encoder.finish()]);
drop((gate_records, run_records, tile_records));
scopes.close().await?;
#[cfg(not(target_arch = "wasm32"))]
wait(&self.device, &self.sink, _index)?;

61
wai/wai-quantum/tests/fabric_expected.rs Normal file
View file

@ -0,0 +1,61 @@
//! `quantum-conformance/fabric/expected.json` is what the browser leg of the
//! fabric's WebGPU member holds a page to (`browser_webgpu.mjs`). This holds
//! the file to the crate: it must be exactly the four certificates of the
//! fabric's pinned corpora (both gatesets; random-v1 and fabric-v1 with their
//! seeds, sizes and wide circuits; the corpus digests the CPU computes) and
//! the two implementation ids, so that editing the file alone cannot weaken
//! the leg, for instance by giving fabric-v1 random-v1's digest and so never
//! running a circuit wider than one tile. A missing or unreadable file fails.
//!
//! The file lives outside this package, so this test is excluded from it.
use std::path::Path;
use serde_json::{Value, json};
use wai_quantum::quantum::Gateset;
use wai_quantum::quantum_fabric::{Corpus, FABRIC_DEFAULT, apply_kernel, lanes, wgsl};
fn hex(b: &[u8]) -> String {
b.iter().map(|x| format!("{x:02x}")).collect()
}
#[test]
fn the_browser_legs_pins_are_the_crates() {
let path = Path::new(env!("CARGO_MANIFEST_DIR")).join("../quantum-conformance/fabric/expected.json");
let text = std::fs::read_to_string(&path).unwrap_or_else(|e| panic!("{}: {e}", path.display()));
let file: Value = serde_json::from_str(&text).unwrap_or_else(|e| panic!("{}: {e}", path.display()));
let Corpus::Fabric { seed, cases, wide } = FABRIC_DEFAULT else { panic!("FABRIC_DEFAULT is a fabric corpus") };
let random = Corpus::Random { seed, cases };
let mut certificates = Vec::new();
for (corpus, wide) in [(random, 0), (FABRIC_DEFAULT, wide)] {
for (number, gs) in [(1, Gateset::V1), (2, Gateset::V2)] {
let cert = lanes::certify_corpus(gs, "local/test", corpus);
certificates.push(json!({
"gateset": number,
"kernel": apply_kernel(gs),
"corpus": corpus.name(),
"seed": seed,
"cases": cases,
"wide": wide,
"cases_compared": cert.cases,
"corpus_hex": hex(&cert.corpus),
}));
}
}
let want = json!({
"implementation": wgsl::implementation_id(),
"tile_implementation": wgsl::tile_implementation_id(),
"certificates": certificates,
});
assert_eq!(file, want, "{} is not the crate's pins; the crate's are:\n{}", path.display(), serde_json::to_string_pretty(&want).unwrap());
// The four digests are the pinned ones, and fabric-v1's are not random-v1's.
let digests: Vec<&str> = file["certificates"].as_array().unwrap().iter().map(|c| c["corpus_hex"].as_str().unwrap()).collect();
assert_eq!(
digests,
[
"d18454c639acfd9c7b66450f4ded1aac4ab46cc10a7bedbcf4958a96e204fb36",
"3dbde276c1095bf827faec71d5980e050f8a55a5ba7d97cd3bbdaec3f43a13b2",
"41c726a0ac80f7defef92e15ecf893bdb7a3076b5e7f670fa0b7142e5c033df7",
"9a84b55af68f1a8b07390b234cf562dd8d85cef3185a697ab85690c8790a7809",
]
);
}

302
wai/wai-quantum/tests/webgpu_member.rs
View file

@ -20,11 +20,12 @@ use std::future::Future;
use std::io::Write;
use std::path::Path;
use std::sync::OnceLock;
use std::time::Instant;
use wai_quantum::ed25519_dalek::SigningKey;
use wai_quantum::quantum::{BaseGate, Circuit, Gate, Gateset, MAX_QUBITS, QuantumError, from_wqc};
use wai_quantum::quantum_fabric::lanes::{self, GateWords, Lane};
use wai_quantum::quantum_fabric::plan::Program;
use wai_quantum::quantum_fabric::wgsl::{self, Detector, MUTANTS};
use wai_quantum::quantum_fabric::plan::{self, Config, PlannedGate, Program, Step};
use wai_quantum::quantum_fabric::wgsl::{self, Detector, MUTANTS, Mutant, TILE_MUTANTS};
use wai_quantum::quantum_fabric::{Corpus, FABRIC_DEFAULT, FabricClass, FabricReceipt, Parity, Presence, REFERENCE, apply_kernel, detect_local, place};
use wai_quantum::quantum_fabric_webgpu::{Discovery, WebGpuError, WebGpuMember, blocking, probe_pairs};
@ -451,36 +452,81 @@ fn g7_refusals_are_typed() {
}
}
/// Whether a member running `mutant`'s text is caught by the mutant's
/// detector.
fn caught(mm: &WebGpuMember, mutant: &Mutant, key: &str) -> bool {
match mutant.detector {
Detector::Probe => {
let edges = probe_pairs(0, 0);
let out = run(mm.probe(&edges));
let differing = edges.iter().zip(&out).filter(|&(&(a, b), o)| *o != (lanes::mul_wide(a, b), lanes::fxmul_wide(a, b))).count();
println!("webgpu mutant {} on {key}: {differing} of 484 edge pairs differ", mutant.name);
differing >= 1
}
Detector::Certificate => Gateset::ALL.into_iter().all(|gs| match attempt(mm.certify(gs, RANDOM_V1)) {
Err(e) => {
println!("webgpu mutant {} on {key}: {gs:?} certification refused: {e}", mutant.name);
true
}
Ok(c) => c.parity != Parity::BitExact,
}),
Detector::Range => attempt(mm.run_lanes(&Z_STATE, &[z_words()])).map(|r| r.out_of_range) != Ok(Some(0)),
Detector::Race => {
// A race need not show on every run: three tries per size.
let mut differing = Vec::new();
for n in [10u8, 11] {
let (state, words) = wgsl::race_witness(n);
let want = lanes::run_words(&state, &words).expect("valid shapes");
for _ in 0..3 {
match attempt(mm.run_lanes(&state, &words)) {
Ok(got) => differing.push(got.amps.iter().zip(&want.amps).filter(|(a, b)| a != b).count() + usize::from(got.out_of_range != want.out_of_range)),
Err(e) => {
println!("webgpu mutant {} on {key}: n = {n} refused: {e}", mutant.name);
differing.push(usize::MAX);
}
}
}
}
println!("webgpu mutant {} on {key}: lanes differing per run of the race witness: {differing:?}", mutant.name);
differing.iter().any(|&d| d > 0)
}
Detector::Malformed => attempt(mm.run_program(&wgsl::malformed_tile_run())) != Err(WebGpuError::Internal(4)),
}
}
#[test]
fn g8_mutants_are_killed() {
let edges = probe_pairs(0, 0);
assert_eq!(edges.len(), 484);
assert_eq!(probe_pairs(0, 0).len(), 484);
for m in members() {
let mut killed = 0;
for mutant in &MUTANTS {
let mm = attempt(m.with_shader(&wgsl::mutate(mutant))).unwrap_or_else(|e| panic!("{}: mutant {} does not compile: {e}", m.key(), mutant.name));
assert_eq!(mm.implementation(), "mutant");
let caught = match mutant.detector {
Detector::Probe => {
let out = run(mm.probe(&edges));
let differing = edges.iter().zip(&out).filter(|&(&(a, b), o)| *o != (lanes::mul_wide(a, b), lanes::fxmul_wide(a, b))).count();
println!("webgpu mutant {} on {}: {differing} of 484 edge pairs differ", mutant.name, m.key());
differing >= 1
}
Detector::Certificate => Gateset::ALL.into_iter().all(|gs| match attempt(mm.certify(gs, RANDOM_V1)) {
Err(e) => {
println!("webgpu mutant {} on {}: {gs:?} certification refused: {e}", mutant.name, m.key());
true
}
Ok(c) => c.parity != Parity::BitExact,
}),
Detector::Range => attempt(mm.run_lanes(&Z_STATE, &[z_words()])).map(|r| r.out_of_range) != Ok(Some(0)),
};
assert!(caught, "{}: mutant {} survived", m.key(), mutant.name);
assert!(caught(&mm, mutant, m.key()), "{}: mutant {} survived", m.key(), mutant.name);
killed += 1;
}
assert_eq!(killed, MUTANTS.len());
report(&format!("webgpu mutants: key={} killed {killed} of {}", m.key(), MUTANTS.len()));
// The tile kernel's mutants, on the tile member. First what the
// detectors rely on: the kernel as written reports the range, refuses
// the malformed run and runs the race witness as the CPU does.
let tiled = run(m.tiled());
assert_eq!(run(tiled.run_lanes(&Z_STATE, &[z_words()])).out_of_range, Some(0), "{}", m.key());
assert_eq!(attempt(tiled.run_program(&wgsl::malformed_tile_run())).unwrap_err(), WebGpuError::Internal(4), "{}", m.key());
for n in [10u8, 11] {
let (state, words) = wgsl::race_witness(n);
assert_eq!(run(tiled.run_lanes(&state, &words)), lanes::run_words(&state, &words).unwrap(), "{} race witness n = {n}", m.key());
}
let mut killed = 0;
for mutant in &TILE_MUTANTS {
let mm = attempt(tiled.with_shader(&wgsl::mutate_tile(mutant))).unwrap_or_else(|e| panic!("{}: tile mutant {} does not compile: {e}", m.key(), mutant.name));
assert_eq!((mm.implementation(), mm.config()), ("mutant", Config::TILED));
assert!(caught(&mm, mutant, m.key()), "{}: tile mutant {} survived", m.key(), mutant.name);
killed += 1;
}
assert_eq!(killed, TILE_MUTANTS.len());
report(&format!("webgpu tile mutants: key={} killed {killed} of {}", m.key(), TILE_MUTANTS.len()));
}
}
@ -541,3 +587,217 @@ fn g11_native_futures_resolve_in_one_poll() {
// Every other test reached its results through `blocking::resolve`.
report("webgpu futures: one poll each");
}
/// The milliseconds `f` takes, at best over `tries`, and its last result.
fn timed<T>(tries: usize, mut f: impl FnMut() -> T) -> (f64, T) {
let mut best = f64::INFINITY;
let mut last = None;
for _ in 0..tries {
let t = Instant::now();
let v = f();
best = best.min(t.elapsed().as_secs_f64() * 1e3);
last = Some(v);
}
(best, last.expect("at least one try"))
}
#[test]
fn g12_tiles_match_per_gate() {
for m in members() {
let tiled = run(m.tiled());
assert_eq!((tiled.id(), tiled.implementation(), tiled.config()), (m.id(), wgsl::tile_implementation_id().as_str(), Config::TILED));
assert!(tiled.implementation().starts_with("wgsl-lanes32-tile/"));
// Any state and any words, mapped through tile steps: the tile member
// equals the per-gate member and the CPU, lanes and report.
let mut s = 0x7469_6c65_6432_u64;
let (mut flagged, mut clean, mut runs) = (0, 0, 0);
for case in 0..2000 {
let n = 1 + (next(&mut s) % 12) as u32;
let full = case % 2 == 0;
let value = |s: &mut u64| -> i32 { if full { next(s) as i32 } else { within(s, 1 << 22) } };
let entry = |s: &mut u64| -> u32 { if full { next(s) as u32 } else { within(s, 1 << 29) as u32 } };
let state: Vec<Lane> = (0..1usize << n).map(|_| [value(&mut s), value(&mut s)]).collect();
let gates: Vec<GateWords> = (0..1 + next(&mut s) % 24)
.map(|_| {
let t = (next(&mut s) % u64::from(n)) as u32;
let ctrl = (0..n).filter(|&b| b != t && next(&mut s) % 4 == 0).fold(0u32, |acc, b| acc | 1 << b);
let mut w = [0u32; 8];
w.iter_mut().for_each(|x| *x = entry(&mut s));
// Half the gates diagonal, on every target.
if next(&mut s) % 2 == 0 {
w[2..6].fill(0);
}
GateWords { m: w, tbit: 1 << t, ctrl }
})
.collect();
let planned: Vec<PlannedGate> = gates.iter().enumerate().map(|(k, w)| PlannedGate::new(*w, k as u32)).collect();
if plan::schedule(n as u8, 0, planned, &Config::TILED).steps.iter().any(|st| matches!(st, Step::Tile { count, .. } if *count > 1)) {
runs += 1;
}
let want = lanes::run_words(&state, &gates).expect("valid shapes");
let got = run(tiled.run_lanes(&state, &gates));
assert_eq!(got.out_of_range, want.out_of_range, "{}: case {case}, n = {n}", m.key());
assert!(got.amps == want.amps, "{}: case {case}, n = {n}: the lanes differ", m.key());
assert!(run(m.run_lanes(&state, &gates)) == got, "{}: case {case}, n = {n}: per-gate and tiled differ", m.key());
if want.out_of_range.is_some() {
flagged += 1;
} else {
clean += 1;
}
}
assert!(flagged >= 100 && clean >= 100 && runs >= 1000, "{}: {flagged} flagged, {clean} clean, {runs} with a run of two gates or more", m.key());
report(&format!("webgpu tile kernel: key={} ok", m.key()));
// The four pinned certificates, under the tile implementation.
for (gs, random_pin, fabric_pin) in PINS {
for (corpus, cases, pin) in [(RANDOM_V1, 200, random_pin), (FABRIC_DEFAULT, 224, fabric_pin)] {
let c = blocking::certify(&tiled, gs, corpus).unwrap_or_else(|e| panic!("{} tiled {gs:?} {}: {e}", m.key(), corpus.name()));
assert_eq!((c.parity, c.cases), (Parity::BitExact, cases), "{} tiled {gs:?} {}", m.key(), corpus.name());
assert_eq!(hex(&c.corpus), pin, "{} tiled {gs:?} {}", m.key(), corpus.name());
assert_eq!((c.kernel.as_str(), c.implementation.as_str(), c.member.as_str()), (apply_kernel(gs), tiled.implementation(), m.id()));
assert_eq!(c.generator, corpus.generator(gs));
report(&format!(
"webgpu certified: key={} member={} kernel={} corpus={} implementation={} parity=bit-exact cases={} digest={}",
m.key(),
m.id(),
c.kernel,
corpus.name(),
c.implementation,
c.cases,
hex(&c.corpus)
));
}
}
// fabric-v1 (tiles of 1024, up to 128 of them) on forced geometries: one
// workgroup to a row, and three (a ragged last row), with rings of two
// records and the planes read back in pieces.
for max_x in [1, 3] {
let forced = tiled.with_max_x(max_x).with_ring_records(2).with_readback_chunk(12 * 1024 + 4);
for (gs, _, pin) in PINS {
let c = run(forced.certify(gs, FABRIC_DEFAULT));
assert_eq!((c.parity, hex(&c.corpus)), (Parity::BitExact, pin.to_string()), "{} tiled max_x {max_x} {gs:?}", m.key());
}
}
report(&format!("webgpu tile geometry: key={} ok", m.key()));
// The conformance vectors.
let root = Path::new(env!("CARGO_MANIFEST_DIR")).join("../quantum-conformance");
let mut total = 0;
for (dir, manifest) in [("vectors", "expected_hashes.json"), ("v2/vectors", "v2/expected_hashes.json")] {
let manifest: serde_json::Value = serde_json::from_str(&std::fs::read_to_string(root.join(manifest)).expect("the manifest reads")).expect("the manifest parses");
let mut files: Vec<_> = std::fs::read_dir(root.join(dir)).expect("the vectors list").map(|e| e.expect("an entry").path()).filter(|p| p.extension().is_some_and(|x| x == "wqc")).collect();
files.sort();
assert!(!files.is_empty(), "no vectors in {dir}");
for path in files {
let name = path.file_name().unwrap().to_str().unwrap().to_string();
let (c, _) = from_wqc(&std::fs::read(&path).unwrap()).unwrap_or_else(|e| panic!("{dir}/{name}: {e}"));
let state = run(tiled.simulate(&c));
assert_eq!(Some(hex(&state.statevector_hash()).as_str()), manifest[&name]["statevector_hash"].as_str(), "{} tiled {dir}/{name}", m.key());
total += 1;
}
}
report(&format!("webgpu tile vectors: key={} {total} ok", m.key()));
// Repeated runs agree.
let c = random_circuit(18, 200, 0x7265_7065_6174, Gateset::V1);
let native = hex(&c.simulate().unwrap().statevector_hash());
let hashes: Vec<String> = (0..5).map(|_| hex(&blocking::simulate_from(&tiled, &c, 0).unwrap().statevector_hash())).collect();
assert!(hashes.iter().all(|h| *h == native), "{}: {hashes:?}, native {native}", m.key());
report(&format!("webgpu tile repeated: key={} 5 runs, one hash", m.key()));
// A start other than |0…0⟩, reduced modulo 2^n as the native simulator does.
for n in [12u8, 16] {
let c = random_circuit(n, 200, 0x7374_6172 + u64::from(n), Gateset::V2);
for start in [5usize, (1 << n) + 5] {
let native = c.simulate_from(start).unwrap().statevector_hash();
assert_eq!(hex(&blocking::simulate_from(&tiled, &c, start).unwrap().statevector_hash()), hex(&native), "{} tiled n = {n} start {start}", m.key());
}
}
// The first op whose result leaves 32 bits sits in a later submission
// than the first, and a later op leaves them again: the report names
// the first op's place in the op-log. The filler is circuit/1's P(19),
// which is not the identity, so the tile planner keeps every gate; on
// two qubits it touches only amplitudes that are zero.
let p19 = lanes::gate_words(&Gate { base: BaseGate::P, controls: vec![], target: 0, param: 19 }, Gateset::V1);
let z1 = lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 1, param: 0 }, Gateset::V2);
let state: [Lane; 4] = [[0, 0], [0, 0], [i32::MIN, 0], [0, 0]];
let gates: Vec<GateWords> = (0..5000).map(|k| if k == 4100 || k == 4500 { z1 } else { p19 }).collect();
let want = lanes::run_words(&state, &gates).expect("valid shapes");
assert_eq!(want.out_of_range, Some(4100));
for member in [tiled.clone(), tiled.with_ring_records(3)] {
assert_eq!(run(member.run_lanes(&state, &gates)), want, "{} tiled, first out of range at 4100, {member:?}", m.key());
}
// A run longer than one dispatch may hold: at 20 qubits a dispatch of
// SUBMIT_BUDGET pair-updates is 512 gates, and a run may hold 1024; the
// member splits it, and the state is the per-gate member's.
let mut s = 0x7370_6c69u64;
let diagonal = [BaseGate::Z, BaseGate::S, BaseGate::Sdg, BaseGate::T, BaseGate::Tdg, BaseGate::P];
let mut long_run = Circuit::with_gateset(20, Gateset::V2);
for _ in 0..600 {
let base = diagonal[(next(&mut s) % 6) as usize];
let target = (next(&mut s) % 20) as u8;
let param = if base == BaseGate::P { 1 + (next(&mut s) % 32) as u16 } else { 0 };
long_run.ops.push(Gate { base, controls: vec![], target, param });
}
let planned = plan::compile_with(&long_run, 0, &Config::TILED).unwrap();
assert_eq!(planned.steps, [Step::Tile { first: 0, count: 600, size: 1024 }]);
assert!(600 > (wai_quantum::quantum_fabric_webgpu::SUBMIT_BUDGET >> 19));
let per_gate = blocking::simulate_from(m, &long_run, 0).unwrap().statevector_hash();
assert_eq!(hex(&blocking::simulate_from(&tiled, &long_run, 0).unwrap().statevector_hash()), hex(&per_gate), "{} a split run", m.key());
// Planned gates a member may not run are refused before anything
// runs, by both implementations: a `diag` that disagrees with the
// words, and the op index the tile kernel keeps for "no op".
let x0 = PlannedGate::new(lanes::gate_words(&Gate { base: BaseGate::X, controls: vec![], target: 0, param: 0 }, Gateset::V2), 0);
for g in [PlannedGate { diag: true, ..x0 }, PlannedGate { index: u32::MAX, ..x0 }] {
for (member, step) in [(m, Step::Gate { gate: 0 }), (&tiled, Step::Gate { gate: 0 }), (&tiled, Step::Tile { first: 0, count: 1, size: 4 })] {
let p = Program { n_qubits: 2, start: 0, gates: vec![g], steps: vec![step] };
assert_eq!(attempt(member.run_program(&p)).unwrap_err(), WebGpuError::Internal(4), "{} {g:?} {step:?}", m.key());
}
}
report(&format!("webgpu tile edges: key={} ok", m.key()));
// The widest states, as g6 runs them per gate: a run cap of 16 gates
// and a grid of (8, 8192) tiles at 26 qubits.
if std::env::var("WAI_QUANTUM_WEBGPU_WIDE").is_ok_and(|v| v == "1") && m.class() == FabricClass::Gpu {
let mut sizes = vec![24u8];
if m.max_qubits() > 24 {
sizes.push(m.max_qubits());
}
for n in sizes {
let qft = Circuit::qft(n);
let random = random_circuit(n, 200, 0x77_6964_65 + u64::from(n), Gateset::V2);
for c in [&qft, &random] {
let native = c.simulate().unwrap().statevector_hash();
let got = blocking::simulate_from(&tiled, c, 0).unwrap_or_else(|e| panic!("{} tiled n = {n}: {e}", m.key()));
assert_eq!(hex(&got.statevector_hash()), hex(&native), "{} tiled n = {n}", m.key());
}
report(&format!("webgpu tile wide: key={} n={n} matches native", m.key()));
}
}
// Timing, per-gate and tiled on the same circuits: printed, never
// asserted. The hashes must agree.
// On 10 qubits every gate fits one tile: runs of up to 1024 gates.
let long = random_circuit(10, 4000, 0x6c6f_6e67, Gateset::V2);
for (what, c) in [
("qft", Circuit::qft(16)),
("random", random_circuit(16, 400, 0x7469_6d65, Gateset::V2)),
("random", random_circuit(20, 400, 0x7469_6d65 + 1, Gateset::V2)),
("random", random_circuit(22, 200, 0x7469_6d65 + 2, Gateset::V2)),
("long", long),
] {
let steps = plan::compile_with(&c, 0, &Config::TILED).unwrap().steps.len();
let (per_gate_ms, a) = timed(3, || blocking::simulate_from(m, &c, 0).unwrap().statevector_hash());
let (tiled_ms, b) = timed(3, || blocking::simulate_from(&tiled, &c, 0).unwrap().statevector_hash());
assert_eq!(hex(&a), hex(&b), "{} {what} n = {}", m.key(), c.n_qubits);
report(&format!(
"webgpu timing: key={} {what} n={} gates={}: per-gate {per_gate_ms:.1} ms ({} dispatches), tiled {tiled_ms:.1} ms ({steps} dispatches)",
m.key(),
c.n_qubits,
c.ops.len(),
c.ops.len()
));
}
}
}

2
wai/wai-rs/Cargo.lock generated
View file

@ -4459,7 +4459,7 @@ dependencies = [
[[package]]
name = "wai-quantum"
version = "0.4.1"
version = "0.4.2"
dependencies = [
"blake3",
"ed25519-dalek",

764
wai/wai-web/Cargo.lock generated
View file

@ -2,12 +2,36 @@
# It is not intended for manual editing.
version = 4
[[package]]
name = "allocator-api2"
version = "0.2.21"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "683d7910e743518b0e34f1186f92494becacb047c7b6bf616c96772180fef923"
[[package]]
name = "android_system_properties"
version = "0.1.6"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ae221649c9976a6f6c56ae1facf410f3ddb33cc661c4b7b61020a912d4237fbc"
dependencies = [
"libc",
]
[[package]]
name = "arrayvec"
version = "0.7.8"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d3fb67a6e08acf24fdeccbac2cb6ac4305825bd1f117462e0e6f2f193345ad56"
[[package]]
name = "ash"
version = "0.38.0+1.3.281"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f"
dependencies = [
"libloading",
]
[[package]]
name = "async-trait"
version = "0.1.92"
@ -31,6 +55,21 @@ version = "1.8.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "2af50177e190e07a26ab74f8b1efbfe2ef87da2116221318cb1c2e82baf7de06"
[[package]]
name = "bit-set"
version = "0.10.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "09ec2f926cc3060f09db9ebc5b52823d85268d24bb917e472c0c4bea35780a7d"
dependencies = [
"bit-vec",
]
[[package]]
name = "bit-vec"
version = "0.9.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b71798fca2c1fe1086445a7258a4bc81e6e49dcd24c8d0dd9a1e57395b603f51"
[[package]]
name = "bitflags"
version = "2.13.1"
@ -59,12 +98,41 @@ dependencies = [
"generic-array",
]
[[package]]
name = "block2"
version = "0.6.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "cdeb9d870516001442e364c5220d3574d2da8dc765554b4a617230d33fa58ef5"
dependencies = [
"objc2",
]
[[package]]
name = "bumpalo"
version = "3.20.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "72f5acc6cb2ba439de613abc23857ec3d78374d8ed5ac84e9d11336e87da8649"
[[package]]
name = "bytemuck"
version = "1.25.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "95832e849adfb21180ccb6826a99da14e5d266ae5c2e668e1602cf234f153797"
dependencies = [
"bytemuck_derive",
]
[[package]]
name = "bytemuck_derive"
version = "1.12.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "6a1f896587b6f2c069c73d2f0913e2d590c3990285cd2f0b6aa02b786b4c679c"
dependencies = [
"proc-macro2",
"quote",
"syn 3.0.4",
]
[[package]]
name = "cast"
version = "0.3.0"
@ -105,6 +173,23 @@ version = "1.0.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801"
[[package]]
name = "cfg_aliases"
version = "0.2.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f079e83a288787bcd14a6aea84cee5c87a67c5a3e660c30f557a3d24761b3527"
[[package]]
name = "codespan-reporting"
version = "0.13.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "af491d569909a7e4dee0ad7db7f5341fef5c614d5b8ec8cf765732aba3cff681"
dependencies = [
"serde",
"termcolor",
"unicode-width",
]
[[package]]
name = "const-oid"
version = "0.9.6"
@ -135,6 +220,12 @@ dependencies = [
"libc",
]
[[package]]
name = "crunchy"
version = "0.2.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5"
[[package]]
name = "crypto-common"
version = "0.1.7"
@ -192,6 +283,25 @@ dependencies = [
"crypto-common",
]
[[package]]
name = "dispatch2"
version = "0.3.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "1e0e367e4e7da84520dedcac1901e4da967309406d1e51017ae1abfb97adbd38"
dependencies = [
"bitflags",
"objc2",
]
[[package]]
name = "document-features"
version = "0.2.12"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d4b8a88685455ed29a21542a33abd9cb6510b6b129abadabdcef0f4c55bc8f61"
dependencies = [
"litrs",
]
[[package]]
name = "ed25519"
version = "2.2.3"
@ -250,6 +360,12 @@ version = "0.1.11"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d45db016d36b838f563236e9193d0ee6ce38f3f68b6c94e914b4929c96bbb890"
[[package]]
name = "foldhash"
version = "0.2.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "77ce24cb58228fbb8aa041425bb1050850ac19177686ea6e0f41a70416f56fdb"
[[package]]
name = "futures-core"
version = "0.3.34"
@ -306,11 +422,51 @@ dependencies = [
"r-efi",
]
[[package]]
name = "gpu-allocator"
version = "0.28.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "51255ea7cfaadb6c5f1528d43e92a82acb2b96c43365989a28b2d44ee38f8795"
dependencies = [
"ash",
"hashbrown 0.16.1",
"log",
"presser",
"thiserror",
"windows",
]
[[package]]
name = "half"
version = "2.7.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b"
dependencies = [
"cfg-if",
"crunchy",
"num-traits",
"zerocopy",
]
[[package]]
name = "hashbrown"
version = "0.16.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "841d1cc9bed7f9236f321df977030373f4a4163ae1a7dbfe1a51a2c1a51d9100"
dependencies = [
"allocator-api2",
"equivalent",
"foldhash",
]
[[package]]
name = "hashbrown"
version = "0.17.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ed5909b6e89a2db4456e54cd5f673791d7eca6732202bbf2a9cc504fe2f9b84a"
dependencies = [
"foldhash",
]
[[package]]
name = "heck"
@ -325,7 +481,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d466e9454f08e4a911e14806c24e16fba1b4c121d1ea474396f396069cf949d9"
dependencies = [
"equivalent",
"hashbrown",
"hashbrown 0.17.1",
]
[[package]]
@ -351,6 +507,16 @@ version = "0.2.189"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3eaf3ede3fee6db1a4c2ee091bf8a8b4dccdc6d17f656fb07896ee72867612f2"
[[package]]
name = "libloading"
version = "0.8.9"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55"
dependencies = [
"cfg-if",
"windows-link",
]
[[package]]
name = "libm"
version = "0.2.16"
@ -363,6 +529,21 @@ version = "0.12.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "32a66949e030da00e8c7d4434b251670a91556f4144941d37452769c25d58a53"
[[package]]
name = "litrs"
version = "1.0.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "11d3d7f243d5c5a8b9bb5d6dd2b1602c0cb0b9db1621bafc7ed66e35ff9fe092"
[[package]]
name = "lock_api"
version = "0.4.14"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "224399e74b87b5f3557511d98dff8b14089b3dadafcab6bb93eab67d3aace965"
dependencies = [
"scopeguard",
]
[[package]]
name = "log"
version = "0.4.34"
@ -385,6 +566,44 @@ dependencies = [
"walkdir",
]
[[package]]
name = "naga"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a616d2fb8c89516ac2723a581f69d6c18576046bed761bd6b305e5618e6ae130"
dependencies = [
"arrayvec",
"bit-set",
"bitflags",
"cfg-if",
"cfg_aliases",
"codespan-reporting",
"half",
"hashbrown 0.17.1",
"indexmap",
"libm",
"log",
"naga-types",
"num-traits",
"once_cell",
"rustc-hash",
"spirv",
"thiserror",
"unicode-ident",
]
[[package]]
name = "naga-types"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "590afbf58a6f4f62873cd5cff4468061844bafa1cdf399cc954537c22d768d49"
dependencies = [
"hashbrown 0.17.1",
"indexmap",
"rustc-hash",
"thiserror",
]
[[package]]
name = "nu-ansi-term"
version = "0.50.3"
@ -423,6 +642,92 @@ dependencies = [
"libm",
]
[[package]]
name = "objc2"
version = "0.6.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "08849bbd4767dfae9457696856ae1c84fe4e0281bbe4a7abff2d0e06fb7981f8"
dependencies = [
"objc2-encode",
]
[[package]]
name = "objc2-core-foundation"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "2a180dd8642fa45cdb7dd721cd4c11b1cadd4929ce112ebd8b9f5803cc79d536"
dependencies = [
"bitflags",
"dispatch2",
"objc2",
]
[[package]]
name = "objc2-core-graphics"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e022c9d066895efa1345f8e33e584b9f958da2fd4cd116792e15e07e4720a807"
dependencies = [
"dispatch2",
"objc2",
"objc2-core-foundation",
"objc2-io-surface",
]
[[package]]
name = "objc2-encode"
version = "4.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ef25abbcd74fb2609453eb695bd2f860d389e457f67dc17cafc8b8cbc89d0c33"
[[package]]
name = "objc2-foundation"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e3e0adef53c21f888deb4fa59fc59f7eb17404926ee8a6f59f5df0fd7f9f3272"
dependencies = [
"bitflags",
"objc2",
"objc2-core-foundation",
]
[[package]]
name = "objc2-io-surface"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "180788110936d59bab6bd83b6060ffdfffb3b922ba1396b312ae795e1de9d81d"
dependencies = [
"bitflags",
"objc2",
"objc2-core-foundation",
]
[[package]]
name = "objc2-metal"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a0125f776a10d00af4152d74616409f0d4a2053a6f57fa5b7d6aa2854ac04794"
dependencies = [
"bitflags",
"block2",
"objc2",
"objc2-foundation",
]
[[package]]
name = "objc2-quartz-core"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "96c1358452b371bf9f104e21ec536d37a650eb10f7ee379fff67d2e08d537f1f"
dependencies = [
"bitflags",
"objc2",
"objc2-core-foundation",
"objc2-core-graphics",
"objc2-foundation",
"objc2-metal",
]
[[package]]
name = "once_cell"
version = "1.21.4"
@ -435,6 +740,38 @@ version = "11.1.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d6790f58c7ff633d8771f42965289203411a5e5c68388703c06e14f24770b41e"
[[package]]
name = "ordered-float"
version = "5.5.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "8c7c9e0d9b23589f26070720bac724174bfec1083e82f7854cdd0267518343c0"
dependencies = [
"num-traits",
]
[[package]]
name = "parking_lot"
version = "0.12.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "93857453250e3077bd71ff98b6a65ea6621a19bb0f559a85248955ac12c45a1a"
dependencies = [
"lock_api",
"parking_lot_core",
]
[[package]]
name = "parking_lot_core"
version = "0.9.12"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "2621685985a2ebf1c516881c026032ac7deafcda1a2c9b7850dc81e3dfcb64c1"
dependencies = [
"cfg-if",
"libc",
"redox_syscall",
"smallvec",
"windows-link",
]
[[package]]
name = "pin-project-lite"
version = "0.2.17"
@ -457,6 +794,27 @@ version = "0.3.34"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f6b464fbc74e149a392436b17d523f769e057cb6877f6a5c4618bc6f11800548"
[[package]]
name = "portable-atomic"
version = "1.15.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "05c8b63e8d9609db387f0324918f81d68fe27748f084ef092fb35954d0539a85"
[[package]]
name = "portable-atomic-util"
version = "0.2.8"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "10ab3eb7f3becc3a1cbc4f2c6f20267996cfc1a6467a873763411b136a122715"
dependencies = [
"portable-atomic",
]
[[package]]
name = "presser"
version = "0.3.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e8cf8e6a8aa66ce33f63993ffc4ea4271eb5b0530a9002db8455ea6050c77bfa"
[[package]]
name = "proc-macro2"
version = "1.0.107"
@ -466,6 +824,12 @@ dependencies = [
"unicode-ident",
]
[[package]]
name = "profiling"
version = "1.0.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3d595e54a326bc53c1c197b32d295e14b169e3cfeaa8dc82b529f947fba6bcf5"
[[package]]
name = "quote"
version = "1.0.47"
@ -490,6 +854,51 @@ dependencies = [
"getrandom 0.2.17",
]
[[package]]
name = "range-alloc"
version = "0.1.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ca45419789ae5a7899559e9512e58ca889e41f04f1f2445e9f4b290ceccd1d08"
[[package]]
name = "raw-window-handle"
version = "0.6.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "20675572f6f24e9e76ef639bc5552774ed45f1c30e2951e1e99c59888861c539"
[[package]]
name = "raw-window-metal"
version = "1.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "40d213455a5f1dc59214213c7330e074ddf8114c9a42411eb890c767357ce135"
dependencies = [
"objc2",
"objc2-core-foundation",
"objc2-foundation",
"objc2-quartz-core",
]
[[package]]
name = "redox_syscall"
version = "0.5.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ed2bf2547551a7053d6fdfafda3f938979645c44812fbfcda098faae3f1a362d"
dependencies = [
"bitflags",
]
[[package]]
name = "renderdoc-sys"
version = "1.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "19b30a45b0cd0bcca8037f3d0dc3421eaf95327a17cad11964fb8179b4fc4832"
[[package]]
name = "rustc-hash"
version = "1.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2"
[[package]]
name = "rustc_version"
version = "0.4.1"
@ -527,6 +936,12 @@ dependencies = [
"winapi-util",
]
[[package]]
name = "scopeguard"
version = "1.2.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "94143f37725109f92c262ed2cf5e59bce7498c01bcc1502d7b9afe439a4e9f49"
[[package]]
name = "semver"
version = "1.0.28"
@ -629,6 +1044,21 @@ version = "0.4.12"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0c790de23124f9ab44544d7ac05d60440adc586479ce501c1d6d7da3cd8c9cf5"
[[package]]
name = "smallvec"
version = "1.16.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "5b3dc8af474f516a851ff4bd12db780f948b9250ad37211e4eec0bccea54e01b"
[[package]]
name = "spirv"
version = "0.4.0+sdk-1.4.341.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d9571ea910ebd84c86af4b3ed27f9dbdc6ad06f17c5f96146b2b671e2976744f"
dependencies = [
"bitflags",
]
[[package]]
name = "spki"
version = "0.7.3"
@ -639,6 +1069,12 @@ dependencies = [
"der",
]
[[package]]
name = "static_assertions"
version = "1.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f"
[[package]]
name = "subtle"
version = "2.6.1"
@ -680,6 +1116,35 @@ dependencies = [
"windows-sys",
]
[[package]]
name = "termcolor"
version = "1.4.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "06794f8f6c5c898b3275aebefa6b8a1cb24cd2c6c79397ab15774837a0bc5755"
dependencies = [
"winapi-util",
]
[[package]]
name = "thiserror"
version = "2.0.21"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "09e52cb86a36cede5cb101bf8908837b3e4c6e5e59fe7fd85c23fb56200d189e"
dependencies = [
"thiserror-impl",
]
[[package]]
name = "thiserror-impl"
version = "2.0.21"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "fe5197923287db20a58125f0bc85c062f7f2c892de97b18c356f9efb14b28524"
dependencies = [
"proc-macro2",
"quote",
"syn 3.0.4",
]
[[package]]
name = "toml"
version = "0.9.12+spec-1.1.0"
@ -731,6 +1196,12 @@ version = "1.0.24"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75"
[[package]]
name = "unicode-width"
version = "0.2.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b4ac048d71ede7ee76d585517add45da530660ef4390e49b098733c6e897f254"
[[package]]
name = "version_check"
version = "0.9.5"
@ -755,11 +1226,12 @@ dependencies = [
[[package]]
name = "wai-quantum"
version = "0.4.1"
version = "0.4.2"
dependencies = [
"blake3",
"ed25519-dalek",
"serde_json",
"wgpu",
]
[[package]]
@ -774,6 +1246,7 @@ dependencies = [
"serde_json",
"wai",
"wasm-bindgen",
"wasm-bindgen-futures",
"wasm-bindgen-test",
]
@ -889,6 +1362,169 @@ version = "0.2.127"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "33ff1c1b360982e93b6d8ea9c04836f71dba0817a16f91e229cf3a51bdd9d987"
[[package]]
name = "web-sys"
version = "0.3.104"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c435338968042f4f59a557f690a253676d47ce13ceb55d70100e7facf6620a30"
dependencies = [
"js-sys",
"wasm-bindgen",
]
[[package]]
name = "wgpu"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "527ccdf43dd5b2e8676eed9984ce00e2bbb0a1b85b70c1969dcb6cd2eb55ab9e"
dependencies = [
"arrayvec",
"bitflags",
"bytemuck",
"cfg-if",
"cfg_aliases",
"document-features",
"hashbrown 0.17.1",
"js-sys",
"log",
"naga",
"portable-atomic",
"profiling",
"raw-window-handle",
"smallvec",
"static_assertions",
"wasm-bindgen",
"wasm-bindgen-futures",
"web-sys",
"wgpu-core",
"wgpu-hal",
"wgpu-types",
]
[[package]]
name = "wgpu-core"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "14c018fce9b6270aa203c2fdd56f3cce996713534bd757e4ea58c8560b121f14"
dependencies = [
"arrayvec",
"bit-set",
"bit-vec",
"bitflags",
"bytemuck",
"cfg_aliases",
"document-features",
"hashbrown 0.17.1",
"indexmap",
"log",
"naga",
"naga-types",
"once_cell",
"parking_lot",
"portable-atomic",
"profiling",
"raw-window-handle",
"rustc-hash",
"smallvec",
"thiserror",
"wgpu-core-deps-apple",
"wgpu-core-deps-windows-linux-android",
"wgpu-hal",
"wgpu-naga-bridge",
"wgpu-types",
]
[[package]]
name = "wgpu-core-deps-apple"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "061f3d319a40d39d00b1ecc2c33b89fe21d4e6fe01859df3500a3a8ecccd6b68"
dependencies = [
"wgpu-hal",
]
[[package]]
name = "wgpu-core-deps-windows-linux-android"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7586165fd5f6d881cb9ce4bb71f40d6caab2c0f1837e3fc1d9788a197fb6004f"
dependencies = [
"wgpu-hal",
]
[[package]]
name = "wgpu-hal"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b6b7fb58561a792bc237628ba0792e332de418fefe145f13b5ed8201e6d52f58"
dependencies = [
"android_system_properties",
"arrayvec",
"ash",
"bit-set",
"bitflags",
"block2",
"bytemuck",
"cfg-if",
"cfg_aliases",
"gpu-allocator",
"hashbrown 0.17.1",
"libc",
"libloading",
"log",
"naga",
"naga-types",
"objc2",
"objc2-core-foundation",
"objc2-core-graphics",
"objc2-foundation",
"objc2-metal",
"objc2-quartz-core",
"once_cell",
"ordered-float",
"parking_lot",
"portable-atomic",
"portable-atomic-util",
"profiling",
"range-alloc",
"raw-window-handle",
"raw-window-metal",
"renderdoc-sys",
"smallvec",
"static_assertions",
"thiserror",
"wgpu-naga-bridge",
"wgpu-types",
"windows",
"windows-core",
]
[[package]]
name = "wgpu-naga-bridge"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d2f62e73117bb7a62bfd9c5a5841438a823f6566c6442a808ee269d2d055c081"
dependencies = [
"naga",
"wgpu-types",
]
[[package]]
name = "wgpu-types"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "99dad6f1fbdbbdb4c278a6508b059d44688f5cebddf78d005a46a31340269286"
dependencies = [
"bitflags",
"bytemuck",
"js-sys",
"log",
"naga-types",
"raw-window-handle",
"static_assertions",
"web-sys",
]
[[package]]
name = "winapi-util"
version = "0.1.11"
@ -898,12 +1534,107 @@ dependencies = [
"windows-sys",
]
[[package]]
name = "windows"
version = "0.62.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "527fadee13e0c05939a6a05d5bd6eec6cd2e3dbd648b9f8e447c6518133d8580"
dependencies = [
"windows-collections",
"windows-core",
"windows-future",
"windows-numerics",
]
[[package]]
name = "windows-collections"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "23b2d95af1a8a14a3c7367e1ed4fc9c20e0a26e79551b1454d72583c97cc6610"
dependencies = [
"windows-core",
]
[[package]]
name = "windows-core"
version = "0.62.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b8e83a14d34d0623b51dce9581199302a221863196a1dde71a7663a4c2be9deb"
dependencies = [
"windows-implement",
"windows-interface",
"windows-link",
"windows-result",
"windows-strings",
]
[[package]]
name = "windows-future"
version = "0.3.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e1d6f90251fe18a279739e78025bd6ddc52a7e22f921070ccdc67dde84c605cb"
dependencies = [
"windows-core",
"windows-link",
"windows-threading",
]
[[package]]
name = "windows-implement"
version = "0.60.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "053e2e040ab57b9dc951b72c264860db7eb3b0200ba345b4e4c3b14f67855ddf"
dependencies = [
"proc-macro2",
"quote",
"syn 2.0.119",
]
[[package]]
name = "windows-interface"
version = "0.59.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3f316c4a2570ba26bbec722032c4099d8c8bc095efccdc15688708623367e358"
dependencies = [
"proc-macro2",
"quote",
"syn 2.0.119",
]
[[package]]
name = "windows-link"
version = "0.2.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5"
[[package]]
name = "windows-numerics"
version = "0.3.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "6e2e40844ac143cdb44aead537bbf727de9b044e107a0f1220392177d15b0f26"
dependencies = [
"windows-core",
"windows-link",
]
[[package]]
name = "windows-result"
version = "0.4.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7781fa89eaf60850ac3d2da7af8e5242a5ea78d1a11c49bf2910bb5a73853eb5"
dependencies = [
"windows-link",
]
[[package]]
name = "windows-strings"
version = "0.5.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7837d08f69c77cf6b07689544538e017c1bfcf57e34b4c0ff58e6c2cd3b37091"
dependencies = [
"windows-link",
]
[[package]]
name = "windows-sys"
version = "0.61.2"
@ -913,6 +1644,15 @@ dependencies = [
"windows-link",
]
[[package]]
name = "windows-threading"
version = "0.2.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3949bd5b99cafdf1c7ca86b43ca564028dfe27d66958f2470940f73d86d75b37"
dependencies = [
"windows-link",
]
[[package]]
name = "winnow"
version = "0.7.15"
@ -925,6 +1665,26 @@ version = "1.0.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "23b97319f7b8343df12cc98938e5c3eb436064524c8d2b4e30a1d3a36eecdf81"
[[package]]
name = "zerocopy"
version = "0.8.62"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "86502bf56ac7c77571a32e2647bb2a15894565e981fb2a48d7bde2d91c965a9d"
dependencies = [
"zerocopy-derive",
]
[[package]]
name = "zerocopy-derive"
version = "0.8.62"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "5457206954b06561e2608c7e19cf58b1926586d999c246eebe4502f7e2039d1a"
dependencies = [
"proc-macro2",
"quote",
"syn 2.0.119",
]
[[package]]
name = "zeroize"
version = "1.9.0"

11
wai/wai-web/Cargo.toml
View file

@ -22,6 +22,17 @@ js-sys = "0.3"
wai = { path = "../wai-rs", default-features = false, features = ["world", "audio_scene", "haptic", "splat4d", "provenance", "score", "film", "neural_int", "generate", "video_int", "audio_int", "world_model", "world_int", "splat_int", "quantum", "quantum_ops", "quantum_cal", "quantum_mitigate", "quantum_qec", "quantum_control", "quantum_noise", "quantum_compile", "quantum_atom", "quantum_qbom", "quantum_receipt", "quantum_toolchain", "quantum_mps", "quantum_stabilizer", "quantum_pauli", "quantum_vml", "quantum_phasor", "quantum_qfhrr", "quantum_kernel", "quantum_phasor_meter", "quantum_qdata", "video_workflow", "staged", "stage_measure", "receipt_claims"] }
ed25519-dalek = "2" # actor keys for browser-side op authoring (same version as wai-rs)
blake3 = "1" # in-page receipt check: hash a decoded image against a pinned digest
# Awaits the WebGPU member's futures behind the promises of src/quantum_webgpu.rs;
# only the `quantum_webgpu` feature compiles it.
wasm-bindgen-futures = { version = "0.4.77", optional = true }
[features]
# The fabric's WebGPU member in a page (src/quantum_webgpu.rs): the gate kernel's
# WGSL compiled by the browser's own WebGPU, certified in the page against the CPU
# reference's pinned corpus digests. Built as its own bundle with its own
# --out-name (demo/quantum_webgpu.html); off by default, so the main bundle is
# unchanged.
quantum_webgpu = ["wai/quantum_fabric", "wai/quantum_fabric_webgpu", "dep:wasm-bindgen-futures"]
[dev-dependencies]
# Pinned to the wasm-bindgen release Cargo.lock holds, so the test runner and the

32
wai/wai-web/demo/README.md
View file

@ -115,3 +115,35 @@ give the wasm32 build's verdicts on every session it sealed:
# from wai-web/ (claim_verify.py needs the `cryptography` package)
bash demo/verify_measure.sh
```
## `quantum_webgpu.html` — the fabric's WebGPU member in a page
`src/quantum_webgpu.rs` (feature `quantum_webgpu`, its own bundle) binds
`wai::quantum_fabric_webgpu` for JS: `quantum_webgpu_discover()` opens the adapter the
page is offered as a fabric member and probes its arithmetic against the CPU's, and the
member it returns certifies the gate kernel against the CPU reference on a corpus
(`certify`) and runs a `WQC` container to its statevector hash (`statevector_hash_hex`).
The kernel is the same WGSL text the native member runs, compiled by the page's own
WebGPU. A bit-exact member reproduces the CPU reference's pinned corpus digests,
`quantum-conformance/fabric/expected.json`.
The page runs `quantum_webgpu_worker.js` in a module worker: the four pinned
certificates (both gatesets, random-v1 and fabric-v1) and every vector of the quantum
corpus, on the member and again on the same adapter running the tile implementation
(`tiled()`: runs of gates applied tile by tile, one dispatch per run, whose barrier the
page's WGSL compiler holds to the uniformity rules); then every reject file (each refused
with its listed class), and a circuit one qubit wider than the member holds. A second worker repeats the last check with the adapter
offering only the WebGPU default buffer limits, where the member holds 25 qubits and
refuses 26 as too wide.
```
# from wai-web/: the bundle has its own --out-name, so its imports never meet the main bundle's
cargo build --release --target wasm32-unknown-unknown --features quantum_webgpu
wasm-bindgen --target web --out-dir demo/pkg-webgpu --out-name wai_quantum_webgpu \
target/wasm32-unknown-unknown/release/wai_web.wasm
# serve wai/: the page reads ../../quantum-conformance/
cd .. && python3 -m http.server 8741 # → http://localhost:8741/wai-web/demo/quantum_webgpu.html
```
`quantum-conformance/browser_webgpu.mjs` drives the page headless and checks every
result against the committed corpus.

126
wai/wai-web/demo/quantum_webgpu.html Normal file
View file

@ -0,0 +1,126 @@
<!DOCTYPE html>
<html lang="en">
<head>
<meta charset="UTF-8">
<meta name="viewport" content="width=device-width, initial-scale=1.0">
<title>WAI — the fabric's WebGPU member</title>
<style>
:root { --bg:#06080c; --bg2:#0d1117; --border:#1e2736; --text:#e6edf5; --text2:#8b99ad; --text3:#4a5568; --g:#00e5a0; --w:#f0b429; --r:#ff6b6b; --mono:'JetBrains Mono','SF Mono',Menlo,monospace; }
* { box-sizing:border-box; margin:0; padding:0; }
body { background:var(--bg); color:var(--text); font-family:var(--mono); padding:24px 16px; }
h1 { font-size:15px; letter-spacing:.04em; margin-bottom:4px; } h1 em { font-style:normal; color:var(--g); }
.sub { font-size:11px; color:var(--text2); margin-bottom:14px; max-width:880px; line-height:1.5; }
code { color:var(--w); }
#status { font-size:12px; margin-bottom:10px; color:var(--text2); }
#log { list-style:none; font-size:11px; line-height:1.6; max-width:1100px; background:var(--bg2); border:1px solid var(--border); border-radius:10px; padding:10px 14px; overflow-wrap:anywhere; }
#log li.ok::before { content:"pass "; color:var(--g); }
#log li.bad::before { content:"FAIL "; color:var(--r); }
#log li.info::before { content:" "; white-space:pre; }
#log li.info { color:var(--text2); }
</style>
</head>
<body>
<h1>WAI <em>wai.quantum.fabric</em> — the WebGPU member in this page</h1>
<p class="sub">
The gate kernel of <code>wai.quantum.circuit</code> and <code>wai.quantum.circuit2</code>, written in
WGSL with unsigned 32-bit arithmetic only, is compiled here by this page's own WebGPU and run in a
worker. The adapter's arithmetic is probed against the CPU's first; the kernel is then certified
against the CPU reference on two corpora per gateset. A bit-exact member reproduces the reference's
pinned corpus digests (<code>quantum-conformance/fabric/expected.json</code>); every vector of the
quantum corpus must reproduce its statevector hash, and every reject file must be refused with its class.
The same adapter then runs the tile implementation (runs of gates applied tile by tile, one dispatch
per run), certified against the same digests.
</p>
<div id="status">starting</div>
<ol id="log"></ol>
<script type="module">
// Results for a driver: every message the workers posted, and whether both finished.
const state = (window.quantumWebGpu = { messages: [], done: false });
const log = document.getElementById("log");
const status = document.getElementById("status");
const params = new URLSearchParams(location.search);
function line(kind, text) {
const li = document.createElement("li");
li.className = kind;
li.textContent = text;
log.appendChild(li);
}
function show(m) {
switch (m.type) {
case "member":
line("info", `member ${m.id} (class ${m.class}, key ${m.key}) on the ${m.limits === "offered" ? "limits the adapter offers" : "WebGPU default limits"}: up to ${m.max_qubits} qubits; implementation ${m.implementation}`);
line("info", `detail: ${m.detail}`);
break;
case "certificate":
if (m.error) line("bad", `certificate ${m.label}: ${m.error}`);
else line(m.got.parity === "bit-exact" && m.got.corpus_hex === m.pin ? "ok" : "bad",
`certificate ${m.label}: ${m.got.parity}, ${m.got.cases} cases, corpus ${m.got.corpus_hex}${m.got.corpus_hex === m.pin ? " (the pin)" : `, pinned ${m.pin}`}`);
break;
case "vector":
line(m.got === m.want ? "ok" : "bad", `${m.tile ? "tile " : ""}${m.label}: ${m.error ? `refused: ${m.error}` : `statevector ${m.got}${m.got === m.want ? "" : `, manifest ${m.want}`}`}`);
break;
case "tiled":
if (m.error) line("bad", `tile member: ${m.error}`);
else line("info", `tile member ${m.id}: implementation ${m.implementation}`);
break;
case "reject": {
const ok = !!m.error && m.error.startsWith(`WQC parse (${m.want}): `);
line(ok ? "ok" : "bad", `${m.label} refused as ${m.want}${ok ? "" : `: ${m.error ? `refused as ${m.error}` : "accepted"}`}`);
break;
}
case "wide": {
const ok = !!m.error && m.error.startsWith(`simulate (${m.want}): `);
line(ok ? "ok" : "bad", `${m.n} qubits on a member holding ${m.max_qubits}: ${m.error ? `refused, ${m.error}` : "accepted"}`);
break;
}
case "no-adapter":
line("info", `no adapter: ${m.why}`);
break;
case "done":
break;
default:
line("bad", `${m.type}: ${m.why || m.error || JSON.stringify(m)}`);
}
}
// One worker to the end; its messages, recorded and shown.
function runWorker(query) {
return new Promise((resolve) => {
const worker = new Worker(`quantum_webgpu_worker.js?${query}`, { type: "module" });
worker.onmessage = (e) => {
state.messages.push(e.data);
show(e.data);
if (e.data.type === "done") {
worker.terminate();
resolve(e.data);
}
};
const fail = (why) => {
const m = { type: "error", error: why };
state.messages.push(m);
show(m);
worker.terminate();
resolve(m);
};
worker.onerror = (e) => fail(`worker: ${e.message || "failed to load"}`);
worker.onmessageerror = () => fail("worker: a message could not be read");
});
}
const corpus = params.get("corpus");
const base = corpus ? `corpus=${encodeURIComponent(corpus)}` : "";
status.textContent = "the member on the limits the adapter offers";
await runWorker(base);
// The too-wide check on the WebGPU default limits, unless the page was offered no adapter.
if (state.messages.some((m) => m.type === "member")) {
status.textContent = "the same adapter on the WebGPU default limits";
await runWorker([base, "limits=webgpu-defaults"].filter(Boolean).join("&"));
}
const failed = log.querySelectorAll("li.bad").length;
status.textContent = failed ? `done: ${failed} failed` : "done";
state.done = true;
</script>
</body>
</html>

207
wai/wai-web/demo/quantum_webgpu_worker.js Normal file
View file

@ -0,0 +1,207 @@
// The fabric's WebGPU member, run in a module worker (quantum_webgpu.html starts
// it). It opens the adapter this worker is offered as a member, probing its
// arithmetic against the CPU's, then:
//
// - certifies the gate kernel of both gatesets on random-v1 and fabric-v1, the
// four certificates quantum-conformance/fabric/expected.json pins;
// - runs every vector of quantum-conformance/vectors/ and v2/vectors/ on the
// member to its statevector hash;
// - does both again on the same adapter running the tile implementation
// (member.tiled()), whose kernel this worker's WGSL compiler also holds to the
// uniformity rules (messages with tile: true);
// - offers it every file of quantum-conformance/reject/, each of which must be
// refused with its listed class;
// - offers it a circuit one qubit wider than it holds: refused as too-wide, or as
// too-many-qubits where it already holds the reference's maximum.
//
// Each result is posted to the page as it arrives; the last message is
// {type: "done"}. Query parameters:
//
// corpus=<url> the quantum-conformance directory
// (default ../../quantum-conformance/, as served from wai/)
// limits=webgpu-defaults the adapter offers only the WebGPU default buffer limits
// (128 MiB per storage binding, 256 MiB per buffer), as many
// pages are offered; the worker then runs the too-wide check
// alone. An adapter that holds the reference's maximum can
// never be offered a valid circuit too wide for it; this
// shows the refusal on the same device.
import init, { quantum_webgpu_discover } from "./pkg-webgpu/wai_quantum_webgpu.js";
const params = new URL(self.location.href).searchParams;
const CORPUS = new URL(params.get("corpus") || "../../quantum-conformance/", self.location.href);
const LIMITS = params.get("limits");
// The widest state the reference simulates (wai::quantum::MAX_QUBITS).
const MAX_QUBITS = 26;
// The WebGPU default limits, in bytes.
const DEFAULT_STORAGE_BINDING = 134217728;
const DEFAULT_BUFFER = 268435456;
const post = (m) => self.postMessage(m);
const message = (e) => String((e && e.message) || e);
async function fetchBytes(url) {
const r = await fetch(url);
if (!r.ok) throw new Error(`${url}: HTTP ${r.status}`);
return new Uint8Array(await r.arrayBuffer());
}
async function fetchJson(url) {
const r = await fetch(url);
if (!r.ok) throw new Error(`${url}: HTTP ${r.status}`);
return r.json();
}
// The same container with its contract's n_qubits set to `n`: a valid circuit
// whose ops touch only the qubits they did, now `n` qubits wide. Sections keep
// their order; their offsets are recomputed.
function widen(wqc, n) {
const view = new DataView(wqc.buffer, wqc.byteOffset, wqc.byteLength);
const count = view.getUint16(4, true);
const blob = 6 + 9 * count;
const sections = [];
for (let k = 0; k < count; k++) {
const at = 6 + 9 * k;
const off = view.getUint32(at + 1, true);
const len = view.getUint32(at + 5, true);
sections.push({ kind: view.getUint8(at), bytes: wqc.subarray(blob + off, blob + off + len) });
}
const contract = sections.find((s) => s.kind === 1);
if (!contract) throw new Error("the container has no contract section");
const text = new TextDecoder().decode(contract.bytes);
const wider = text.replace(/"n_qubits":\d+/, `"n_qubits":${n}`);
if (wider === text) throw new Error("the contract names no n_qubits");
contract.bytes = new TextEncoder().encode(wider);
const out = new Uint8Array(blob + sections.reduce((t, s) => t + s.bytes.length, 0));
const w = new DataView(out.buffer);
out.set(wqc.subarray(0, 6));
let off = 0;
sections.forEach((s, k) => {
const at = 6 + 9 * k;
w.setUint8(at, s.kind);
w.setUint32(at + 1, off, true);
w.setUint32(at + 5, s.bytes.length, true);
out.set(s.bytes, blob + off);
off += s.bytes.length;
});
return out;
}
// Every adapter this worker is then offered reports the WebGPU default buffer
// limits; the member requests what the adapter reports, so its device is granted
// those.
function offerDefaultLimits() {
const gpu = navigator.gpu;
const request = gpu.requestAdapter.bind(gpu);
gpu.requestAdapter = async (options) => {
const adapter = await request(options);
if (adapter) {
Object.defineProperty(adapter.limits, "maxStorageBufferBindingSize", { value: DEFAULT_STORAGE_BINDING });
Object.defineProperty(adapter.limits, "maxBufferSize", { value: DEFAULT_BUFFER });
}
return adapter;
};
}
// qft5 of circuit2, widened to one qubit more than the member holds.
async function tooWide(member, limits) {
const n = member.max_qubits() + 1;
const bytes = widen(await fetchBytes(new URL("v2/vectors/qft5.wqc", CORPUS)), n);
const want = n > MAX_QUBITS ? "too-many-qubits" : "too-wide";
try {
const hash = await member.statevector_hash_hex(bytes);
post({ type: "wide", limits, n, max_qubits: member.max_qubits(), want, accepted: hash });
} catch (e) {
post({ type: "wide", limits, n, max_qubits: member.max_qubits(), want, error: message(e) });
}
}
// The pinned certificates, then every vector of both trees, on `member`.
async function certifyAndRun(member, expected, tile) {
for (const c of expected.certificates) {
const label = `${tile ? "tile " : ""}${c.kernel} ${c.corpus}`;
try {
const got = JSON.parse(await member.certify(c.gateset, c.corpus, BigInt(c.seed), c.cases, c.wide));
post({ type: "certificate", tile, label, gateset: c.gateset, corpus: c.corpus, pin: c.corpus_hex, got });
} catch (e) {
post({ type: "certificate", tile, label, gateset: c.gateset, corpus: c.corpus, pin: c.corpus_hex, error: message(e) });
}
}
for (const prefix of ["", "v2/"]) {
const manifest = await fetchJson(new URL(`${prefix}expected_hashes.json`, CORPUS));
for (const [file, entry] of Object.entries(manifest)) {
const label = prefix + file;
try {
const bytes = await fetchBytes(new URL(`${prefix}vectors/${file}`, CORPUS));
post({ type: "vector", tile, label, want: entry.statevector_hash, got: await member.statevector_hash_hex(bytes) });
} catch (e) {
post({ type: "vector", tile, label, want: entry.statevector_hash, error: message(e) });
}
}
}
}
async function run() {
if (!navigator.gpu) {
post({ type: "no-adapter", why: "navigator.gpu is absent in this worker" });
return;
}
if (LIMITS === "webgpu-defaults") offerDefaultLimits();
else if (LIMITS) throw new Error(`limits=${LIMITS}: the one value is webgpu-defaults`);
await init();
let member;
try {
member = await quantum_webgpu_discover();
} catch (e) {
const why = message(e);
post(why.startsWith("no adapter: ") ? { type: "no-adapter", why: why.slice("no adapter: ".length) } : { type: "refused", why });
return;
}
post({
type: "member",
limits: LIMITS || "offered",
id: member.member_id(),
key: member.key(),
class: member.class(),
detail: member.detail(),
max_qubits: member.max_qubits(),
implementation: member.implementation(),
});
if (LIMITS) {
await tooWide(member, LIMITS);
return;
}
// The four certificates and every vector of both trees, on the per-gate
// member and then on the tile member.
const expected = await fetchJson(new URL("fabric/expected.json", CORPUS));
await certifyAndRun(member, expected, false);
let tiled;
try {
tiled = await member.tiled();
post({ type: "tiled", id: tiled.member_id(), implementation: tiled.implementation() });
} catch (e) {
post({ type: "tiled", error: message(e) });
}
if (tiled) await certifyAndRun(tiled, expected, true);
// Every reject file.
const listing = await fetchJson(new URL("reject/reject.json", CORPUS));
for (const [file, kind] of Object.entries(listing)) {
const label = "reject/" + file;
try {
const bytes = await fetchBytes(new URL(`reject/${file}`, CORPUS));
post({ type: "reject", label, want: kind, accepted: await member.statevector_hash_hex(bytes) });
} catch (e) {
post({ type: "reject", label, want: kind, error: message(e) });
}
}
await tooWide(member, "offered");
}
run()
.catch((e) => post({ type: "error", error: message(e) }))
.finally(() => post({ type: "done" }));

7
wai/wai-web/src/lib.rs
View file

@ -153,6 +153,13 @@ pub mod quantum_phasor;
// a total-energy roll-up, and show a swapped stage break the bill.
pub mod quantum_qbom;
// The fabric's WebGPU member in a page (wai.quantum.fabric): the browser's adapter
// opened and probed, the gate kernel certified on it against the CPU reference's
// pinned corpus digests, and WQC containers run to their statevector hash. Only in
// the separate `quantum_webgpu` bundle; the main bundle does not carry it.
#[cfg(all(feature = "quantum_webgpu", target_arch = "wasm32"))]
mod quantum_webgpu;
// Staged delivery (extensions/staged-delivery.md): wai::staged::StageSink, the sink a
// native player runs, bound for JS — bases added, refinement layers offered, holds
// expired and bases removed in any order, every unit's stage, producer, tier, output

204
wai/wai-web/src/quantum_webgpu.rs Normal file
View file

@ -0,0 +1,204 @@
//! The fabric's WebGPU member in a page (`wai::quantum_fabric_webgpu`), bound
//! for JS in the separate `quantum_webgpu` bundle.
//!
//! [`quantum_webgpu_discover`] opens the adapter the page is offered as a
//! member and probes its arithmetic against the CPU's in this page. The member
//! then certifies the gate kernel against the CPU reference on a corpus
//! ([`QuantumWebGpu::certify`]) and runs `WQC` containers to their statevector
//! hash ([`QuantumWebGpu::statevector_hash_hex`]).
//!
//! The kernel is the same WGSL text the native member runs, compiled here by
//! the browser's own WGSL compiler, so a bit-exact certificate from a page is
//! evidence from an independent translator. [`QuantumWebGpu::tiled`] gives the
//! same adapter running the tile implementation, whose barrier the browser's
//! compiler holds to WGSL's uniformity rules, which the native translator does
//! not check. The corpus digest covers inputs
//! and outputs only, so a bit-exact member reproduces the CPU reference's
//! pinned digests. Certificates are earned in the page that runs the work.
//!
//! Every method that uses the device returns a promise over a clone of the
//! member's handle, so no borrow of the object is held across an await.
use std::rc::Rc;
use wai::quantum::{from_wqc, Gateset, QuantumError};
use wai::quantum_fabric::{Corpus, Parity};
use wai::quantum_fabric_webgpu::{discover, WebGpuError, WebGpuMember};
use wasm_bindgen::prelude::*;
use wasm_bindgen_futures::future_to_promise;
fn hex(b: &[u8]) -> String {
b.iter().map(|x| format!("{x:02x}")).collect()
}
/// The class word of a refusal: the reference's own class for a circuit it
/// refuses (`malformed`, `invalid`, `contract`, `too-many-qubits`, as
/// `reject.json` lists them), else the member's reason.
fn class(e: &WebGpuError) -> &'static str {
match e {
WebGpuError::Circuit(q) => q.kind(),
WebGpuError::TooWide { .. } => "too-wide",
WebGpuError::OutOfRange { .. } => "out-of-range",
WebGpuError::NoAdapter(_) => "no-adapter",
WebGpuError::Unsuitable { .. } => "unsuitable",
WebGpuError::Device(_) => "device",
WebGpuError::Memory(_) => "memory",
WebGpuError::Timeout => "timeout",
WebGpuError::Pending => "pending",
WebGpuError::Internal(_) => "internal",
WebGpuError::Probe { .. } => "probe",
_ => "member",
}
}
/// A promise's rejection: an `Error` whose message is
/// `<stage> (<class>): <reason>`, the form the quantum sink throws.
fn refusal(stage: &str, class: &str, reason: &dyn std::fmt::Display) -> JsValue {
JsError::new(&format!("{stage} ({class}): {reason}")).into()
}
fn member_refusal(stage: &str, e: &WebGpuError) -> JsValue {
refusal(stage, class(e), e)
}
fn parse_refusal(e: &QuantumError) -> JsValue {
refusal("WQC parse", e.kind(), e)
}
fn parity_label(p: &Parity) -> &'static str {
match p {
Parity::BitExact => "bit-exact",
Parity::Within { .. } => "within",
Parity::Statistical { .. } => "statistical",
Parity::Uncertified => "uncertified",
}
}
/// The browser's adapter as a fabric member, its arithmetic probed in this page.
#[wasm_bindgen]
pub struct QuantumWebGpu {
member: Rc<WebGpuMember>,
}
/// The browser's adapter as a fabric member, opened and probed in this page.
/// Rejects with `no adapter: <why>` when the page is offered none, or with the
/// adapter's refusal (`adapter <key> (<name>) refused: <why>`).
#[wasm_bindgen]
pub async fn quantum_webgpu_discover() -> Result<QuantumWebGpu, JsError> {
let d = discover().await;
if let Some(member) = d.found.into_iter().next() {
return Ok(QuantumWebGpu { member: Rc::new(member) });
}
match d.refused.first() {
Some(r) => Err(JsError::new(&format!("adapter {} ({:?}) refused: {}", r.key, r.adapter, r.error))),
None => Err(JsError::new("no adapter: the page's WebGPU offered none")),
}
}
#[wasm_bindgen]
impl QuantumWebGpu {
/// The member id: `local/gpu-webgpu-<d12>`, or `local/swgpu-webgpu-<d12>`
/// for a software adapter.
pub fn member_id(&self) -> String {
self.member.id().to_string()
}
/// The key a runner declares for it.
pub fn key(&self) -> String {
self.member.key().to_string()
}
/// `"gpu"`, or `"cpu"` for a software adapter.
pub fn class(&self) -> String {
self.member.class().label().to_string()
}
/// What discovery found: adapter, translator, granted limits. Not signed. A
/// page's WebGPU may report no adapter description; the member id then rests
/// on the rest of the identity, and this says so.
pub fn detail(&self) -> String {
let detail = self.member.detail();
if self.member.identity().name.is_empty() {
format!("{detail}; adapter identity withheld")
} else {
detail.to_string()
}
}
/// The widest circuit this member's buffers hold.
pub fn max_qubits(&self) -> u8 {
self.member.max_qubits()
}
/// The kernel implementation the member runs (`wgsl-lanes32/<16 hex>`, or
/// `wgsl-lanes32-tile/<16 hex>` for a tile member).
pub fn implementation(&self) -> String {
self.member.implementation().to_string()
}
/// The same adapter running the tile implementation: the tile kernel
/// compiled by the page's WebGPU, its arithmetic probed in this page, and
/// circuits planned in tile runs. Resolves to a `QuantumWebGpu` that
/// certifies and runs as this one does; rejects with
/// `tiled (<class>): <reason>`.
pub fn tiled(&self) -> js_sys::Promise {
let member = Rc::clone(&self.member);
future_to_promise(async move {
let tiled = member.tiled().await.map_err(|e| member_refusal("tiled", &e))?;
Ok(QuantumWebGpu { member: Rc::new(tiled) }.into())
})
}
/// Certify the gate kernel of `gateset` (1 for `wai.quantum.circuit`, 2 for
/// `wai.quantum.circuit2`) on this member over a corpus: `"random-v1"`
/// (`seed`, `cases`; `wide` must be 0) or `"fabric-v1"` (`seed`, `cases`,
/// `wide`). `seed` is a `BigInt`. Resolves to the certificate as JSON:
/// `{kernel, implementation, member, generator, parity, cases, corpus_hex,
/// certificate_hash_hex}`; rejects with `certify (<class>): <reason>`.
pub fn certify(&self, gateset: u8, corpus: String, seed: u64, cases: u32, wide: u32) -> js_sys::Promise {
let member = Rc::clone(&self.member);
future_to_promise(async move {
let gs = match gateset {
1 => Gateset::V1,
2 => Gateset::V2,
g => return Err(refusal("certify", "argument", &format!("gateset {g}: 1 is wai.quantum.circuit, 2 is wai.quantum.circuit2"))),
};
let corpus = match corpus.as_str() {
"random-v1" if wide == 0 => Corpus::Random { seed, cases: cases as usize },
"random-v1" => return Err(refusal("certify", "argument", &"random-v1 has no wide circuits; pass wide = 0")),
"fabric-v1" => Corpus::Fabric { seed, cases: cases as usize, wide: wide as usize },
other => return Err(refusal("certify", "argument", &format!("corpus {other:?}: random-v1 or fabric-v1"))),
};
let cert = member.certify(gs, corpus).await.map_err(|e| member_refusal("certify", &e))?;
let json = serde_json::json!({
"kernel": cert.kernel,
"implementation": cert.implementation,
"member": cert.member,
"generator": cert.generator,
"parity": parity_label(&cert.parity),
"cases": cert.cases,
"corpus_hex": hex(&cert.corpus),
"certificate_hash_hex": hex(&cert.hash()),
});
Ok(JsValue::from_str(&json.to_string()))
})
}
/// Parse a `WQC1` or `WQC2` container and run it on the member. Resolves to
/// the statevector hash in hex, the hash the native simulator gives; rejects
/// with `<stage> (<class>): <reason>`: `WQC parse` with the class
/// `reject.json` lists for a container the reference reader refuses, or
/// `simulate` with the reference's class for a circuit it refuses
/// (`too-many-qubits`), or the member's (`too-wide` for more qubits than
/// its buffers hold, `out-of-range` where the reference's result leaves 32
/// bits, `device`, `memory`, ...).
pub fn statevector_hash_hex(&self, wqc: Vec<u8>) -> js_sys::Promise {
let member = Rc::clone(&self.member);
future_to_promise(async move {
let (circuit, _measure) = from_wqc(&wqc).map_err(|e| parse_refusal(&e))?;
drop(wqc);
let state = member.simulate(&circuit).await.map_err(|e| member_refusal("simulate", &e))?;
Ok(JsValue::from_str(&hex(&state.statevector_hash())))
})
}
}