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

wai-quantum 0.4.1: the fabric's WebGPU member — one WGSL kernel, certified bit-exact on every adapter #218

Merged
dcharlot merged 11 commits from wai-quantum-webgpu into main 2026-10-11 22:28:10 -04:00
Conversation 0 Commits 11 Files changed 26 +6410 -101
26 changed files with 6410 additions and 101 deletions
Download patch file Download diff file Expand all files Collapse all files

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

@ -490,7 +490,9 @@ jobs:
# 2. the library's unit tests in the modules that hold the circuit core, its
# container, its JSON reader and its pins (quantum, ijson, receipt_pins,
# quantum_fabric, quantum_source, quantum_channel, quantum_qudit, quantum_receipt),
# each required to have run;
# each required to have run, and each test of the fabric's groundwork for
# accelerator members required by name (the WebGPU member itself has no wasm32-
# wasip2 build; wai-quantum-webgpu.yml checks that the feature is inert there);
# 3. wai_quantum_vectors verify over all four trees of quantum-conformance/ (both
# capabilities, the reject corpus and the accept corpus), the same check the matrix
# above runs natively.
@ -558,6 +560,37 @@ jobs:
grep -q "^test $m::.* \.\.\. ok$" "$RUNNER_TEMP/wasip2-lib.log" \
|| { echo "::error::no test of $m ran"; exit 1; }
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.
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 \
quantum_fabric::tests::malformed_words_are_refused \
quantum_fabric::tests::z_on_the_most_negative_lane_is_flagged \
quantum_fabric::tests::simulate_checked_is_simulate \
quantum_fabric::tests::a_rounded_product_is_never_negated \
quantum_fabric::tests::the_guard_horizon_is_pinned \
quantum_fabric::tests::valid_circuits_stay_in_lanes \
quantum_fabric::tests::fabric_corpus_starts_with_random_v1 \
quantum_fabric::tests::the_fabric_v1_corpus_of_each_gateset_is_pinned \
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::wgsl::tests::implementation_id_is_pinned \
quantum_fabric::wgsl::tests::gate_record_layout \
quantum_fabric::wgsl::tests::grids_cover_every_item_once \
quantum_fabric::wgsl::tests::the_index_map_is_a_bijection \
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
grep -q "^test $t \.\.\. ok$" "$RUNNER_TEMP/wasip2-lib.log" \
|| { echo "::error::$t did not run"; exit 1; }
done
# The wai package is built as one codegen unit here. Its lib target is also a cdylib,
# and the wasm linker in Rust 1.98 (LLD 22) can crash with SIGSEGV linking that cdylib:

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

@ -0,0 +1,218 @@
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.
#
# Each member job declares the adapters its runner must have in
# WAI_QUANTUM_WEBGPU_EXPECT. The device test fails if the adapters found differ,
# and scripts/wai-quantum-webgpu-member.sh fails unless all tests ran, every
# declared adapter earned one bit-exact certificate per pinned digest, and every
# declared adapter killed every shader mutant.
#
# The GPU jobs then run the quantum conformance corpus through each member
# (`wai_quantum_vectors verify --webgpu`). wai-rs is built there with
# --no-default-features: its default feature, `codecs`, links system codec
# libraries that neither GPU host has, and the corpus needs none of them.
#
# 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.
concurrency:
group: ${{ github.workflow }}-${{ github.ref }}
cancel-in-progress: true
on:
push:
branches: [main]
paths:
- 'wai/wai-quantum/**'
- 'wai/quantum-conformance/**'
- 'wai/tools/quantum-wgsl-check/**'
- 'wai/wai-rs/src/bin/wai_quantum_vectors.rs'
# The GPU jobs build wai-rs with the forwarded feature, --locked.
- 'wai/wai-rs/Cargo.toml'
- 'wai/wai-rs/Cargo.lock'
- 'wai/wai-rs/src/lib.rs'
- 'rust-toolchain.toml'
- 'wai/wai-web/src/quantum_webgpu.rs'
- 'wai/wai-web/demo/quantum_webgpu*'
- 'scripts/wai-quantum-webgpu-member.sh'
- '.github/workflows/wai-quantum-webgpu.yml'
pull_request:
paths:
- 'wai/wai-quantum/**'
- 'wai/quantum-conformance/**'
- 'wai/tools/quantum-wgsl-check/**'
- 'wai/wai-rs/src/bin/wai_quantum_vectors.rs'
# The GPU jobs build wai-rs with the forwarded feature, --locked.
- 'wai/wai-rs/Cargo.toml'
- 'wai/wai-rs/Cargo.lock'
- 'wai/wai-rs/src/lib.rs'
- 'rust-toolchain.toml'
- 'wai/wai-web/src/quantum_webgpu.rs'
- 'wai/wai-web/demo/quantum_webgpu*'
- 'scripts/wai-quantum-webgpu-member.sh'
- '.github/workflows/wai-quantum-webgpu.yml'
jobs:
checks:
name: shader checks · browser build · wasip2 inertness · MSRV
runs-on: linux
timeout-minutes: 60
steps:
- uses: actions/checkout@v5
- name: Install Rust 1.98.0 (wasm32 targets) and 1.97.0
shell: bash
run: |
export PATH="$HOME/.cargo/bin:$PATH"
command -v rustup >/dev/null 2>&1 || \
curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --profile minimal --default-toolchain 1.98.0
rustup toolchain install 1.98.0 --profile minimal
rustup toolchain install 1.97.0 --profile minimal
rustup default 1.98.0
rustup target add --toolchain 1.98.0 wasm32-unknown-unknown wasm32-wasip2
echo "$HOME/.cargo/bin" >> "$GITHUB_PATH"
- name: SPIR-V validator
shell: bash
run: |
set -euo pipefail
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
shell: bash
working-directory: wai/tools/quantum-wgsl-check
run: |
set -euo pipefail
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"
- 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: On wasm32-wasip2 the feature is inert
shell: bash
working-directory: wai/wai-quantum
run: |
set -euo pipefail
cargo check --locked --lib --target wasm32-wasip2 --features full,quantum_fabric_webgpu
w="$RUNNER_TEMP/tree-wasip2.txt"; n="$RUNNER_TEMP/tree-native.txt"
cargo tree --locked --target wasm32-wasip2 --features full,quantum_fabric_webgpu -e normal --prefix none > "$w"
cargo tree --locked --target x86_64-unknown-linux-gnu --features quantum_fabric_webgpu -e normal --prefix none > "$n"
grep -q '^wai-quantum ' "$w" || { echo "::error::cargo tree printed no wasip2 tree"; exit 1; }
grep -q '^wgpu ' "$n" || { echo "::error::the native tree has no wgpu, so this check cannot see it"; exit 1; }
if grep -q '^wgpu' "$w"; then echo "::error::wgpu is in the wasm32-wasip2 tree"; exit 1; fi
- name: The package carries the shader
shell: bash
working-directory: wai/wai-quantum
run: |
set -euo pipefail
cargo package --locked --list --allow-dirty > "$RUNNER_TEMP/package.txt"
grep -qx 'src/quantum_fabric/lanes32.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
# The member's unit tests need no GPU (a failed native wait is fed to the
# guard as a panicking closure), and only the feature build compiles them.
- name: The member's unit tests, without a device
shell: bash
working-directory: wai/wai-quantum
run: |
set -euo pipefail
cargo test --locked --lib --features quantum_fabric_webgpu -- quantum_fabric_webgpu:: 2>&1 \
| tee "$RUNNER_TEMP/member-unit.log"
grep -q '^test quantum_fabric_webgpu::tests::a_failed_native_wait_is_a_typed_loss \.\.\. ok$' \
"$RUNNER_TEMP/member-unit.log" \
|| { echo "::error::the member's unit test did not run"; exit 1; }
member-software:
name: member · software Vulkan (linux)
runs-on: linux
timeout-minutes: 120
env:
WAI_QUANTUM_WEBGPU_EXPECT: swgpu-vulkan-10005
steps:
- uses: actions/checkout@v5
- name: Install Rust 1.98.0
shell: bash
run: |
export PATH="$HOME/.cargo/bin:$PATH"
command -v rustup >/dev/null 2>&1 || \
curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --profile minimal --default-toolchain 1.98.0
rustup toolchain install 1.98.0 --profile minimal
rustup default 1.98.0
echo "$HOME/.cargo/bin" >> "$GITHUB_PATH"
- name: The Vulkan loader and a software Vulkan driver
shell: bash
run: |
set -euo pipefail
SUDO=""; [ "$(id -u)" -ne 0 ] && SUDO="sudo"
$SUDO apt-get update -qq
$SUDO apt-get install -y -qq --no-install-recommends libvulkan1 mesa-vulkan-drivers
- name: The member's tests
shell: bash
run: scripts/wai-quantum-webgpu-member.sh
member-metal:
name: member · Metal (macos)
runs-on: macos
timeout-minutes: 90
env:
WAI_QUANTUM_WEBGPU_EXPECT: gpu-metal-0
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"
- name: The member's tests
shell: bash
run: nice -n 10 scripts/wai-quantum-webgpu-member.sh
- name: The conformance corpus through the member
shell: bash
working-directory: wai/wai-rs
run: nice -n 10 cargo run --locked --no-default-features --features quantum_fabric_webgpu --bin wai_quantum_vectors -- verify ../quantum-conformance --webgpu
member-hw:
name: member · discrete GPU, Vulkan (jetson-hub)
runs-on: linux-hw
timeout-minutes: 150
env:
# Set from the key list the first run prints; add any other ICD the host
# exposes (an integrated GPU, a software driver) as it reports them.
# The discrete GPU, and the software Vulkan driver the host also has. The wide
# sizes run on the GPU only.
WAI_QUANTUM_WEBGPU_EXPECT: gpu-vulkan-10de,swgpu-vulkan-10005
WAI_QUANTUM_WEBGPU_WIDE: '1'
CARGO_PROFILE_DEV_DEBUG: '0'
CARGO_BUILD_JOBS: '8'
steps:
- name: Fetch the source
shell: bash
run: |
set -euo pipefail
want="${{ github.event.pull_request.head.sha || github.sha }}"
rm -rf src && mkdir src && cd src
git init -q
# Fetch the commit itself, not the ref: a later commit on the ref that this
# workflow's paths do not match would otherwise be fetched instead.
git fetch -q --depth 1 http://100.64.224.0:3400/Transaction-Science/open-standards.git "$want"
git checkout -q FETCH_HEAD
got="$(git rev-parse HEAD)"
[ "$got" = "$want" ] || { echo "::error::fetched $got, not $want"; exit 1; }
- name: The member's tests
shell: bash
working-directory: src
run: CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum" nice -n 10 scripts/wai-quantum-webgpu-member.sh
- name: The conformance corpus through the member
shell: bash
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
# Windows, when its runner is back (DX12, translator naga+dxbc, and Vulkan, as separate members):
# member-windows:
# runs-on: windows
# env: { WAI_QUANTUM_WEBGPU_EXPECT: "<keys from the first run>" }
# steps: checkout; Set-ExecutionPolicy Bypass; run the script with `shell: bash`
# (Git Bash on PATH). Not a required status.

2
scripts/check-public-names.sh
View file

@ -76,7 +76,7 @@ else
files+=("$f")
done < <(git ls-files 'wai/*.md' 'wai/*.rs' 'wai/*.toml' 'wai/*.py' 'wai/*.sh' 'wai/*.mjs' 'wai/*.js' \
'wai/*.json' 'wai/*.html' 'wai/*.css' 'wai/*.ttl' 'wai/*.h' 'wai/*.c' 'wai/*.txt' \
'wai/*.gitignore')
'wai/*.gitignore' 'wai/*.wgsl')
fi
if [ ${#files[@]} -eq 0 ]; then
echo "::error::no file to check" >&2

49
scripts/wai-quantum-webgpu-member.sh Executable file
View file

@ -0,0 +1,49 @@
#!/usr/bin/env bash
#
# 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.
#
# WAI_QUANTUM_WEBGPU_EXPECT=<key[,key...]|none> [WAI_QUANTUM_WEBGPU_WIDE=1] \
# scripts/wai-quantum-webgpu-member.sh
set -euo pipefail
: "${WAI_QUANTUM_WEBGPU_EXPECT:?declare the adapter keys this runner must have, or none}"
cd "$(dirname "$0")/../wai/wai-quantum"
dir="${RUNNER_TEMP:-$(mktemp -d)}"
log="$dir/webgpu-member.log"
export WAI_QUANTUM_WEBGPU_REPORT="$dir/webgpu-member.report"
: > "$WAI_QUANTUM_WEBGPU_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 -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; }
[ "$WAI_QUANTUM_WEBGPU_EXPECT" = none ] && exit 0
pins=(
"wai.quantum.circuit/apply random-v1 d18454c639acfd9c7b66450f4ded1aac4ab46cc10a7bedbcf4958a96e204fb36"
"wai.quantum.circuit2/apply random-v1 3dbde276c1095bf827faec71d5980e050f8a55a5ba7d97cd3bbdaec3f43a13b2"
"wai.quantum.circuit/apply fabric-v1 41c726a0ac80f7defef92e15ecf893bdb7a3076b5e7f670fa0b7142e5c033df7"
"wai.quantum.circuit2/apply fabric-v1 9a84b55af68f1a8b07390b234cf562dd8d85cef3185a697ab85690c8790a7809"
)
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; }
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; }
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; }
fi
done

74
wai/CHANGELOG.md
View file

@ -6,6 +6,80 @@ wire format may change and such changes are recorded here.
## v0.4.0 — current
### wai-quantum 0.4.1: the fabric's WebGPU member
Additive. The default features, `full`, and every existing pin and digest are unchanged.
- **Certification one case at a time.** `quantum_fabric::Certification` (`new`, `add`,
`cases`, `finish`) folds a member's outputs into a certificate as they arrive, for a
member that produces them asynchronously. `certify` is now a loop over it and gives the
same certificates.
- **Corpora.** `Corpus::Random { seed, cases }` is the existing random corpus
(`random-v1`). `Corpus::Fabric { seed, cases, wide }` (`fabric-v1`) continues the same
stream after it with circuits of 11 to 17 qubits (`lanes::random_wide_circuit`), so that
every dispatch shape a 32-bit-lane member uses is exercised. `FABRIC_DEFAULT` is seed 7,
200 + 24 circuits; `GENERATOR_FABRIC`, `GENERATOR_FABRIC2` and `fabric_generator` name
it per gateset, and `Corpus::{generator, name, circuits}` describe any corpus. The
fabric-v1 lanes32 digests are pinned: `41c726a0…3df7` (circuit/1) and `9a84b55a…7809`
(circuit2). `lanes::certify_corpus` certifies the lane core on any corpus;
`certify_lanes` is it on the random corpus, with the same generator string.
- **The lane core, word by word.** `lanes::GateWords` and `gate_words` (a gate as the
eight 32-bit matrix words and the target and control masks a kernel reads),
`fxmul_wide` (the exact 64-bit `fx`), `PROBE_EDGES`, `apply_words` and `run_words` (the
general 2×2 path computed from exact 64-bit values, flagging a gate exactly when the
reference's result leaves 32 bits, and refusing malformed words), `simulate_checked`,
`output` and `reference_output`. `lanes::horizon(n, gs)` is the gate count below which a
valid circuit provably cannot leave 32 bits, computed in integers from the gateset's
table: 46,340 (circuit/1) and 46,339 (circuit2) at 26 qubits. The module documentation
and `fxmul` no longer claim that every result of a valid circuit fits in 32 bits; it
holds within the horizon.
- **`quantum_fabric::plan`**: `Program` (one step per op, the start reduced modulo 2^n),
`PlannedGate`, `Step`, `compile`, `LaneRun` and `execute`, the CPU executor a member's
output is held to; `PLAN_REVISION` is 1.
- **`quantum_fabric::wgsl`**: `SHADER`, the lane core transcribed into WGSL
(`src/quantum_fabric/lanes32.wgsl`: unsigned 32-bit values only, an exact 64-bit copy
of every result, literal shift amounts, no float, division, conversion, workgroup
memory or barrier); `implementation_id()` (`wgsl-lanes32/dc8e19f61e0e87ab`, a BLAKE3
over the shader text and `HOST_LAYOUT`, so any edit is a new implementation that earns
its certificates again); `gate_record`, `Grid`, `grid` and `gate_grid` (the 64-byte gate
record and the dispatch grid; `WORKGROUP` and `X_GROUPS_MAX` are 64); `AdapterIdentity`
with `member_id` (`local/gpu-…` or `local/swgpu-…`, sorting after every
`local/cpu-<arch>`), `key` and `class`; and `MUTANTS`, `mutate` and `Detector`, the
shader mutants every device must catch. The module documentation carries the argument
that the kernel is bit-exact. `wai/tools/quantum-wgsl-check` (unpublished) holds the
shader to that argument's rules on naga's IR and translates it to MSL, SPIR-V and HLSL.
- **The WebGPU member, feature `quantum_fabric_webgpu`** (opt-in, not in `full`). The
kernel through the platform's GPU interface natively (Metal, Vulkan, DX12) and through
a browser's WebGPU on `wasm32-unknown-unknown`; wgpu 30.0.1 enters only those two target
tables, and on every other target, `wasm32-wasip2` included, the module is compiled
out. `discover()` opens each adapter at the WebGPU downlevel limits with only the buffer
limits raised, pins the DX12 shader compiler (translator `naga+dxbc`), probes the
device's arithmetic against the lane core, and lists each adapter it cannot use with
the reason (`Discovery`, `Refusal`). `WebGpuMember` runs lanes, programs and circuits,
certifies a gateset on a corpus, and builds a fabric `Member`; a bit-exact member
reproduces the pinned random-v1 and fabric-v1 digests. A result that would leave 32
bits is refused (`WebGpuError::OutOfRange`), never returned. Every entry point is async
with no executor; natively `blocking` resolves each call in one poll, and a future still
pending is an error. Natively each submission is waited on before the next, so
`POLL_TIMEOUT` bounds one chunk of gates, not a run; a failed wait, which the GPU
interface reports by panicking, is caught and returned as `WebGpuError::Device`, and
the member treats its device as lost from then on (a `panic = "abort"` build still
aborts). With no energy figures, placement keeps the CPU reference, and a receipt over
a member's run is unmetered.
- **wai-rs.** The feature `quantum_fabric_webgpu` forwards to it and re-exports the
module. `wai_quantum_vectors verify <dir> --webgpu` runs the corpus check once per
WebGPU member, each state computed by the member; each accept file must also run on the
member as it does natively, so the 255-qubit files are refused as `too-many-qubits`.
Without the feature `--webgpu` is an error; with no member it exits 2.
- **CI.** `wai-quantum-webgpu.yml` checks the shader under naga and a SPIR-V validator,
the browser build, the feature's absence from the `wasm32-wasip2` dependency tree, the
shader file in the package, and the MSRV; and runs the member's tests
(`scripts/wai-quantum-webgpu-member.sh`) on software Vulkan, Metal, and a discrete GPU
through Vulkan. Each member job declares the adapters its runner must have, and fails
if the adapters found differ, if any test did not run, if a declared adapter lacks a
bit-exact certificate for each pinned digest, or if a shader mutant survives. The
`wasm32-wasip2` job requires the fabric's new unit tests by name.
### Packaging: refinement layers, stills, a piece's ladder, live windows and delivery
- **Refinement layers in packages** ([packaging](extensions/packaging.md) §5.4). For

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

@ -0,0 +1,645 @@
# This file is automatically @generated by Cargo.
# It is not intended for manual editing.
version = 4
[[package]]
name = "arrayvec"
version = "0.7.8"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d3fb67a6e08acf24fdeccbac2cb6ac4305825bd1f117462e0e6f2f193345ad56"
[[package]]
name = "autocfg"
version = "1.5.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53"
[[package]]
name = "base64ct"
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.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3ded4057c258ba199e2d26386d3af3780957ecaee6c4ef4041c6b4b8b97c0b06"
[[package]]
name = "blake3"
version = "1.8.7"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "6d9e454fc11f76977dc803893aff6304ed33d6a26efae8696573bea74baa27ae"
dependencies = [
"arrayvec",
"cc",
"cfg-if",
"constant_time_eq",
"cpufeatures 0.3.1",
]
[[package]]
name = "block-buffer"
version = "0.10.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3078c7629b62d3f0439517fa394996acacc5cbc91c5a20d8c658e77abd503a71"
dependencies = [
"generic-array",
]
[[package]]
name = "cc"
version = "1.4.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0ad534f4357a5264cce5019c989cf66a4f0dc4e0d1b1d15f8aacec0ff7360273"
dependencies = [
"find-msvc-tools",
"shlex",
]
[[package]]
name = "cfg-if"
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 = [
"unicode-width",
]
[[package]]
name = "const-oid"
version = "0.9.6"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c2459377285ad874054d797f3ccebf984978aa39129f6eafde5cdc8315b612f8"
[[package]]
name = "constant_time_eq"
version = "0.4.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3d52eff69cd5e647efe296129160853a42795992097e8af39800e1060caeea9b"
[[package]]
name = "cpufeatures"
version = "0.2.17"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "59ed5838eebb26a2bb2e58f6d5b5316989ae9d08bab10e0e6d103e656d1b0280"
dependencies = [
"libc",
]
[[package]]
name = "cpufeatures"
version = "0.3.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "5ca28b0ae3115b884660db4118d803791fd6756b6e88f39c0f3f7859060d7566"
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"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "78c8292055d1c1df0cce5d180393dc8cce0abec0a7102adb6c7b1eef6016d60a"
dependencies = [
"generic-array",
"typenum",
]
[[package]]
name = "curve25519-dalek"
version = "4.1.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "97fb8b7c4503de7d6ae7b42ab72a5a59857b4c937ec27a3d4539dba95b5ab2be"
dependencies = [
"cfg-if",
"cpufeatures 0.2.17",
"curve25519-dalek-derive",
"digest",
"fiat-crypto",
"rustc_version",
"subtle",
"zeroize",
]
[[package]]
name = "curve25519-dalek-derive"
version = "0.1.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f46882e17999c6cc590af592290432be3bce0428cb0d5f8b6715e4dc7b383eb3"
dependencies = [
"proc-macro2",
"quote",
"syn 2.0.119",
]
[[package]]
name = "der"
version = "0.7.10"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e7c1832837b905bbfb5101e07cc24c8deddf52f93225eee6ead5f4d63d53ddcb"
dependencies = [
"const-oid",
"zeroize",
]
[[package]]
name = "digest"
version = "0.10.7"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "9ed9a281f7bc9b7576e61468ba615a66a5c8cfdff42420a70aa82701a3b1e292"
dependencies = [
"block-buffer",
"crypto-common",
]
[[package]]
name = "ed25519"
version = "2.2.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "115531babc129696a58c64a4fef0a8bf9e9698629fb97e9e40767d235cfbcd53"
dependencies = [
"pkcs8",
"signature",
]
[[package]]
name = "ed25519-dalek"
version = "2.2.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "70e796c081cee67dc755e1a36a0a172b897fab85fc3f6bc48307991f64e4eca9"
dependencies = [
"curve25519-dalek",
"ed25519",
"serde",
"sha2",
"subtle",
"zeroize",
]
[[package]]
name = "equivalent"
version = "1.0.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f"
[[package]]
name = "fiat-crypto"
version = "0.2.9"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "28dea519a9695b9977216879a3ebfddf92f1c08c05d984f8996aecd6ecdc811d"
[[package]]
name = "find-msvc-tools"
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 = "generic-array"
version = "0.14.7"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "85649ca51fd72272d7821adaf274ad91c288277713d9c18820d8499a7ff69e9a"
dependencies = [
"typenum",
"version_check",
]
[[package]]
name = "getrandom"
version = "0.2.17"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ff2abc00be7fca6ebc474524697ae276ad847ad0a6b3faa4bcb027e9a4614ad0"
dependencies = [
"cfg-if",
"libc",
"wasi",
]
[[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.17.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ed5909b6e89a2db4456e54cd5f673791d7eca6732202bbf2a9cc504fe2f9b84a"
dependencies = [
"foldhash",
]
[[package]]
name = "indexmap"
version = "2.14.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "07aa2048142242915a31d35844fb311e0e53fcca590c3a0a40dcf1b841fa09eb"
dependencies = [
"equivalent",
"hashbrown",
]
[[package]]
name = "itoa"
version = "1.0.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "8f42a60cbdf9a97f5d2305f08a87dc4e09308d1276d28c869c684d7777685682"
[[package]]
name = "libc"
version = "0.2.189"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3eaf3ede3fee6db1a4c2ee091bf8a8b4dccdc6d17f656fb07896ee72867612f2"
[[package]]
name = "libm"
version = "0.2.16"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b6d2cec3eae94f9f509c767b45932f1ada8350c4bdb85af2fcab4a3c14807981"
[[package]]
name = "log"
version = "0.4.34"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f9f8bd3e56ce4dfc153cf470fffbfa98c7620958b312ca5c3a4b8d5181fd13c6"
[[package]]
name = "memchr"
version = "2.8.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98"
[[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",
"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",
"indexmap",
"rustc-hash",
"thiserror",
]
[[package]]
name = "num-traits"
version = "0.2.19"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841"
dependencies = [
"autocfg",
"libm",
]
[[package]]
name = "once_cell"
version = "1.21.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "9f7c3e4beb33f85d45ae3e3a1792185706c8e16d043238c593331cc7cd313b50"
[[package]]
name = "pkcs8"
version = "0.10.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f950b2377845cebe5cf8b5165cb3cc1a5e0fa5cfa3e1f7f55707d8fd82e0a7b7"
dependencies = [
"der",
"spki",
]
[[package]]
name = "proc-macro2"
version = "1.0.107"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9"
dependencies = [
"unicode-ident",
]
[[package]]
name = "quantum-wgsl-check"
version = "0.0.0"
dependencies = [
"naga",
"wai-quantum",
]
[[package]]
name = "quote"
version = "1.0.47"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001"
dependencies = [
"proc-macro2",
]
[[package]]
name = "rand_core"
version = "0.6.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ec0be4795e2f6a28069bec0b5ff3e2ac9bafc99e6a9a7dc3547996c5c816922c"
dependencies = [
"getrandom",
]
[[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"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "cfcb3a22ef46e85b45de6ee7e79d063319ebb6594faafcf1c225ea92ab6e9b92"
dependencies = [
"semver",
]
[[package]]
name = "semver"
version = "1.0.28"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "8a7852d02fc848982e0c167ef163aaff9cd91dc640ba85e263cb1ce46fae51cd"
[[package]]
name = "serde"
version = "1.0.229"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "4148590afebada386688f18773da617792bf2ef03ffc1e4cbd2b1d45b023e0ba"
dependencies = [
"serde_core",
]
[[package]]
name = "serde_core"
version = "1.0.229"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "67dca2c9c51e58a4791a4b1ed58308b39c64224d349a935ab5039aa360942a48"
dependencies = [
"serde_derive",
]
[[package]]
name = "serde_derive"
version = "1.0.229"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e7a5d71263a5a7d47b41f6b3f06ba276f10cc18b0931f1799f710578e2309348"
dependencies = [
"proc-macro2",
"quote",
"syn 3.0.4",
]
[[package]]
name = "serde_json"
version = "1.0.151"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c841b55ecdae098c80dcae9cf767f6f8a0c2cdb3416bbef72181df4d0fe73f14"
dependencies = [
"indexmap",
"itoa",
"memchr",
"serde",
"serde_core",
"zmij",
]
[[package]]
name = "sha2"
version = "0.10.9"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a7507d819769d01a365ab707794a4084392c824f54a7a6a7862f8c3d0892b283"
dependencies = [
"cfg-if",
"cpufeatures 0.2.17",
"digest",
]
[[package]]
name = "shlex"
version = "2.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "f8fadd59c855ef2080decdef8ff161eb6661b86933c9d82e5ba29dc602a55aba"
[[package]]
name = "signature"
version = "2.2.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "77549399552de45a898a580c1b41d445bf730df867cc44e6c0233bbc4b8329de"
dependencies = [
"rand_core",
]
[[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"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d91ed6c858b01f942cd56b37a94b3e0a1798290327d1236e4d9cf4eaca44d29d"
dependencies = [
"base64ct",
"der",
]
[[package]]
name = "subtle"
version = "2.6.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "13c2bddecc57b384dee18652358fb23172facb8a2c51ccc10d74c157bdea3292"
[[package]]
name = "syn"
version = "2.0.119"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297"
dependencies = [
"proc-macro2",
"quote",
"unicode-ident",
]
[[package]]
name = "syn"
version = "3.0.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e6275cddf4610d1775e6d1fe9469b2e77d0f39fd98fb7450901b821e0c53649f"
dependencies = [
"proc-macro2",
"quote",
"unicode-ident",
]
[[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 = "typenum"
version = "1.20.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b6f5e870be6c3b371b77fe0ee0bafb859fa4964b4404c27de1d380043c4dda20"
[[package]]
name = "unicode-ident"
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"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
[[package]]
name = "wai-quantum"
version = "0.4.1"
dependencies = [
"blake3",
"ed25519-dalek",
"serde_json",
]
[[package]]
name = "wasi"
version = "0.11.1+wasi-snapshot-preview1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ccf3ec651a847eb01de73ccad15eb7d99f80485de043efb2f370cd654f4ea44b"
[[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"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e13c156562582aa81c60cb29407084cdb54c4164760106ab78e6c5b0858cf64e"
[[package]]
name = "zmij"
version = "1.0.23"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "29666d0abbfad1e3dc4dcf6144730dd3a3ab225bbbdac83319345b1b44ccfc1b"

18
wai/tools/quantum-wgsl-check/Cargo.toml Normal file
View file

@ -0,0 +1,18 @@
# Checks the WGSL gate kernel of wai-quantum (quantum_fabric::wgsl::SHADER) under
# naga 30.0.1, the shader translator the native WebGPU member compiles it with:
# the rules the kernel's bit-exactness argument rests on, on naga's IR, and
# translation to MSL, SPIR-V and HLSL.
#
# Deliberately a SEPARATE crate with its own lockfile: naga is a dependency here
# only, and never compiles in wai-quantum's own builds or tests.
[package]
name = "quantum-wgsl-check"
version = "0.0.0"
edition = "2024"
rust-version = "1.97"
license = "Apache-2.0"
publish = false
[dependencies]
naga = { version = "30.0.1", default-features = false, features = ["wgsl-in", "spv-out", "msl-out", "hlsl-out"] }
wai-quantum = { path = "../../wai-quantum", default-features = false, features = ["quantum_fabric"] }

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

@ -0,0 +1,326 @@
//! `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
//! MSL, SPIR-V and HLSL.
//!
//! ```text
//! quantum-wgsl-check [--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`.
//!
//! The checks, in order; 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.
//! 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.
//! 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).
//!
//! The last line of a passing run is `quantum-wgsl-check: ok`.
use std::collections::BTreeMap;
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;
const NAME: &str = "quantum-wgsl-check";
type Check<T = ()> = Result<T, String>;
fn main() -> ExitCode {
let args: Vec<String> = std::env::args().skip(1).collect();
let (shader, out) = match parse_args(&args) {
Ok(a) => a,
Err(e) => {
eprintln!("{NAME}: {e}\nusage: {NAME} [--shader <file>] [<directory>]");
return ExitCode::from(2);
}
};
match run(shader.as_deref(), out.as_deref()) {
Ok(()) => {
println!("{NAME}: ok");
ExitCode::SUCCESS
}
Err(e) => {
eprintln!("{NAME}: FAILED: {e}");
ExitCode::FAILURE
}
}
}
fn parse_args(args: &[String]) -> Check<(Option<PathBuf>, Option<PathBuf>)> {
let (mut shader, mut out) = (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.starts_with("--") {
return Err(format!("unknown option {a}"));
} else if out.is_none() {
out = Some(PathBuf::from(a));
} else {
return Err(format!("unexpected argument {a}"));
}
}
Ok((shader, out))
}
fn run(shader: Option<&Path>, out: Option<&Path>) -> Check {
let text = match shader {
None => {
println!("{NAME}: checking wai_quantum::quantum_fabric::wgsl::SHADER");
wgsl::SHADER.to_string()
}
Some(p) => {
println!("{NAME}: checking {} in place of the compiled SHADER", 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");
for (name, f) in functions(&module) {
statements(&f.body, name)?;
}
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());
}
Ok(())
}
/// Step 1.
fn compile(text: &str) -> Check<(Module, ModuleInfo)> {
let module = naga::front::wgsl::parse_str(text).map_err(|e| format!("1. parse: {}", e.emit_to_string(text)))?;
let info = Validator::new(ValidationFlags::all(), Capabilities::empty()).validate(&module).map_err(|e| format!("1. validation: {}", e.emit_to_string(text)))?;
Ok((module, info))
}
/// Step 2.
fn entry_points(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\"]"));
}
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));
}
}
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)]);
let mut found = BTreeMap::new();
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::Storage { access } if access.contains(StorageAccess::LOAD | StorageAccess::STORE) => READ_WRITE,
AddressSpace::Uniform => UNIFORM,
other => return Err(format!("3. global `{name}` is in address space {other:?}")),
};
let Some(b) = v.binding else { return Err(format!("3. global `{name}` has no binding")) };
if b.group != 0 {
return Err(format!("3. global `{name}` is in group {}", b.group));
}
if found.insert(b.binding, space).is_some() {
return Err(format!("3. binding {} is declared twice", b.binding));
}
}
if found != want {
return Err(format!("3. the bindings are {found:?}, not {want:?}"));
}
Ok(())
}
fn scalar_ok(s: Scalar) -> bool {
s.kind == ScalarKind::Bool || (s.kind == ScalarKind::Uint && s.width == 4)
}
/// Whether a type holds only bool and 32-bit unsigned scalars. Pointers,
/// arrays and structs refer to other types of the arena, each checked on its
/// own.
fn type_ok(inner: &TypeInner) -> Check {
match *inner {
TypeInner::Scalar(s) | TypeInner::Vector { scalar: s, .. } | TypeInner::Atomic(s) | TypeInner::ValuePointer { scalar: s, .. } => {
if scalar_ok(s) { Ok(()) } else { Err(format!("{:?}, width {}", s.kind, s.width)) }
}
TypeInner::Pointer { .. } | TypeInner::Array { .. } | TypeInner::Struct { .. } => Ok(()),
ref other => Err(format!("{other:?}")),
}
}
/// The module's functions and entry points, by name.
fn functions(module: &Module) -> Vec<(&str, &naga::Function)> {
module
.functions
.iter()
.map(|(_, f)| (f.name.as_deref().unwrap_or("(unnamed)"), f))
.chain(module.entry_points.iter().map(|e| (e.name.as_str(), &e.function)))
.collect()
}
/// Step 4.
fn types(module: &Module, info: &ModuleInfo) -> Check {
let infos = module.functions.iter().map(|(h, _)| &info[h]).chain((0..module.entry_points.len()).map(|i| info.get_entry_point(i)));
for ((name, f), fi) in functions(module).into_iter().zip(infos) {
for (h, _) in f.expressions.iter() {
type_ok(fi[h].ty.inner_with(&module.types)).map_err(|e| {
let what = f.named_expressions.get(&h).map_or(format!("expression {h:?}"), |n| format!("`{n}`"));
format!("4. in `{name}`, {what} has a type that is not bool or u32: {e}")
})?;
}
}
for (h, _) in module.global_expressions.iter() {
type_ok(info[h].inner_with(&module.types)).map_err(|e| format!("4. at module scope, expression {h:?} has a type that is not bool or u32: {e}"))?;
}
for (_, t) in module.types.iter() {
type_ok(&t.inner).map_err(|e| format!("4. the type {} is not made of bool and u32: {e}", t.name.as_deref().unwrap_or("(unnamed)")))?;
}
Ok(())
}
/// A shift amount: a `u32` literal, or a constant initialised with one.
fn shift_amount(module: &Module, arena: &Arena<Expression>, h: Handle<Expression>) -> Option<u32> {
match arena[h] {
Expression::Literal(Literal::U32(v)) => Some(v),
Expression::Constant(c) => match module.global_expressions[module.constants[c].init] {
Expression::Literal(Literal::U32(v)) => Some(v),
_ => None,
},
_ => None,
}
}
/// Step 5, over every expression arena: module scope, functions, entry points.
fn expressions(module: &Module) -> Check {
let arenas = std::iter::once(("module scope", &module.global_expressions)).chain(functions(module).into_iter().map(|(n, f)| (n, &f.expressions)));
for (name, arena) in arenas {
for (h, e) in arena.iter() {
match *e {
Expression::As { kind, convert, .. } => return Err(format!("5. in `{name}`: a conversion or bitcast to {kind:?} ({convert:?}), expression {h:?}")),
Expression::Math { fun, .. } => return Err(format!("5. in `{name}`: the built-in function {fun:?}, expression {h:?}")),
Expression::Binary { op: op @ (BinaryOperator::Divide | BinaryOperator::Modulo), .. } => return Err(format!("5. in `{name}`: {op:?}, expression {h:?}")),
Expression::Binary { op: op @ (BinaryOperator::ShiftLeft | BinaryOperator::ShiftRight), right, .. } => match shift_amount(module, arena, right) {
Some(v) if (1..=31).contains(&v) => {}
Some(v) => return Err(format!("5. in `{name}`: {op:?} by {v}, outside 1..=31, expression {h:?}")),
None => return Err(format!("5. in `{name}`: {op:?} by an amount that is not a literal ({:?}), expression {h:?}", arena[right])),
},
_ => {}
}
}
}
Ok(())
}
/// Step 6, recursively through every block.
fn statements(block: &Block, name: &str) -> Check {
for s in block.iter() {
match s {
Statement::Atomic { fun, .. } => {
if !matches!(fun, AtomicFunction::InclusiveOr | AtomicFunction::Min) {
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::If { accept, reject, .. } => {
statements(accept, name)?;
statements(reject, name)?;
}
Statement::Switch { cases, .. } => {
for c in cases {
statements(&c.body, name)?;
}
}
Statement::Loop { body, continuing, .. } => {
statements(body, name)?;
statements(continuing, name)?;
}
_ => {}
}
}
Ok(())
}
/// Step 7.
fn translate(module: &Module, info: &ModuleInfo, out: Option<&Path>) -> Check {
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");
match out {
Some(dir) => {
let path = dir.join("lanes32.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());
}
None => println!("{NAME}: 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));
}
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));
}
compile(&mutated).map_err(|e| format!("8. mutant {}: {e}", m.name))?;
}
Ok(())
}

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

@ -698,7 +698,7 @@ dependencies = [
[[package]]
name = "wai-quantum"
version = "0.4.0"
version = "0.4.1"
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.0"
version = "0.4.1"
dependencies = [
"blake3",
"ed25519-dalek",

3
wai/wai-quantum/.gitattributes vendored Normal file
View file

@ -0,0 +1,3 @@
# The shader's bytes are hashed into its implementation id (quantum_fabric::wgsl):
# a checkout that converted line endings would change the id.
src/quantum_fabric/*.wgsl text eol=lf

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

File diff suppressed because it is too large Load diff

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

@ -1,10 +1,10 @@
[package]
name = "wai-quantum"
version = "0.4.0"
version = "0.4.1"
edition = "2024"
rust-version = "1.97"
license = "Apache-2.0"
description = "A deterministic quantum stack in pure Rust: byte-exact circuit simulation (statevector / stabilizer / tensor-network MPS / sparse-Pauli backends), sparse Pauli dynamics at utility scale (arbitrary angles, 1024 qubits), belief-propagation tensor networks on the hardware graph, error mitigation, qLDPC decoding, noise learning, circuit-equivalence proofs, a phasor interference-ML layer, information-theoretic limits, noisy channels and state tomography, and signed energy-accounted receipts. No QPU, no cloud, no system libraries — identical results native, in the browser, and as a WASI component at the edge."
description = "A deterministic quantum stack in pure Rust: byte-exact circuit simulation (statevector / stabilizer / tensor-network MPS / sparse-Pauli backends), sparse Pauli dynamics at utility scale (arbitrary angles, 1024 qubits), belief-propagation tensor networks on the hardware graph, error mitigation, qLDPC decoding, noise learning, circuit-equivalence proofs, a phasor interference-ML layer, information-theoretic limits, noisy channels and state tomography, and signed energy-accounted receipts. No QPU, no cloud, no system libraries in the default features — identical results native, in the browser, and as a WASI component at the edge; the opt-in WebGPU member reaches the platform's GPU interface."
repository = "https://git.transaction.science/Transaction-Science/open-standards"
homepage = "https://wai.transaction.science/quantum-ide"
documentation = "https://docs.rs/wai-quantum"
@ -26,6 +26,15 @@ serde_json = { version = "1", features = ["preserve_order"] }
# Receipt signatures. Optional — only the sealing features need it.
ed25519-dalek = { version = "2", optional = true }
# The WebGPU member of the fabric: the gate kernel in WGSL through the platform's
# GPU interface natively, and through a browser's WebGPU on wasm32-unknown-unknown.
# Absent on every other target; wasm32-wasip2 has no GPU interface.
[target.'cfg(not(target_arch = "wasm32"))'.dependencies]
wgpu = { version = "30.0.1", optional = true, default-features = false, features = ["std", "wgsl", "metal", "vulkan", "dx12"] }
[target.'cfg(all(target_arch = "wasm32", target_os = "unknown"))'.dependencies]
wgpu = { version = "30.0.1", optional = true, default-features = false, features = ["std", "wgsl", "webgpu"] }
[features]
default = ["quantum", "quantum_toolchain", "quantum_receipt"]
@ -124,6 +133,10 @@ quantum_qbom = ["quantum_ops"]
# placement by joules, and a signed receipt naming where a result ran. With the
# fixed-point core in 32-bit lanes, the kernel an integer-only fabric runs.
quantum_fabric = ["quantum", "quantum_ops"]
# The WebGPU member: one WGSL kernel on Metal, Vulkan, DX12 or a browser's
# WebGPU, certified bit-exact against the CPU reference on each adapter. Opt-in,
# and not in `full`: it is the one feature that reaches a platform interface.
quantum_fabric_webgpu = ["quantum_fabric", "dep:wgpu"]
# --- interference / ML layer ----------------------------------------------
quantum_vml = []
@ -142,6 +155,10 @@ full = [
"quantum_phasor_meter",
]
[[test]]
name = "webgpu_member"
required-features = ["quantum_fabric_webgpu"]
[[example]]
name = "qio_gset"
required-features = ["quantum_qio"]

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

@ -6,7 +6,9 @@ every stage.
No QPU, no cloud service, no vendor SDK, and **no system libraries** — so the
same source runs natively, in the browser via wasm, and as a WASI component at
the edge, producing **byte-identical results and receipts on all three**.
the edge, producing **byte-identical results and receipts on all three**. One
opt-in feature, `quantum_fabric_webgpu`, reaches the platform's GPU interface;
no default feature reaches a system library.
```toml
[dependencies]
@ -92,6 +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 |
| `full` | everything |
Default features are `quantum`, `quantum_toolchain`, `quantum_receipt`.

9
wai/wai-quantum/src/lib.rs
View file

@ -10,7 +10,9 @@
//! Rust, so the same source runs natively, in the browser via wasm, and as a
//! WASI component at the edge, producing **byte-identical** results and receipts
//! on all three. A fabric receipt is the exception by design: it names the
//! member that ran the work.
//! member that ran the work. One opt-in feature, `quantum_fabric_webgpu`,
//! reaches the platform's GPU interface; no default feature reaches a system
//! library.
//!
//! Everything is behind a feature; enable only what you use.
@ -70,6 +72,11 @@ pub mod quantum_compile;
#[cfg(feature = "quantum_fabric")]
pub mod quantum_fabric;
// The fabric's WebGPU member: natively, and in a browser on
// wasm32-unknown-unknown. No other wasm32 target has a GPU interface.
#[cfg(all(feature = "quantum_fabric_webgpu", any(not(target_arch = "wasm32"), target_os = "unknown")))]
pub mod quantum_fabric_webgpu;
#[cfg(feature = "quantum_control")]
pub mod quantum_control;

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

@ -18,7 +18,10 @@
//! same inputs and compare every output. The worst case over the corpus is
//! the parity. A [`Certificate`] names the kernel, the implementation, the
//! member and the corpus generator it was earned on, so whoever holds it can
//! re-run that corpus and check the digest.
//! re-run that corpus and check the digest. [`Certification`] does the same
//! one case at a time, for a member whose outputs arrive asynchronously, and
//! [`Corpus`] names the circuits: random-v1, or fabric-v1, which adds wide
//! circuits so that every dispatch shape a member uses is exercised.
//! - **Placement** ([`place`]): of the detected members whose certified
//! parity meets what the caller requires, the one that spends the fewest
//! joules per unit of work.
@ -28,11 +31,16 @@
//!
//! Numerics decide how much parity a fabric can earn. Most accelerators carry
//! no 64-bit float, and some no 64-bit integer. The gate kernel of the
//! fixed-point core needs neither: for the amplitudes and gate entries of a
//! valid circuit, each product fits in 64 bits and each result in 32. So a
//! fabric with only 32-bit integer lanes can reproduce it bit for bit.
//! [`lanes`] is that kernel written with only the operations such a fabric
//! has, and its certificate against the native simulator is bit-exact. The
//! fixed-point core needs neither: for a valid circuit of at most
//! [`lanes::horizon`] gates, every amplitude stays inside 32 bits, so a fabric
//! with only 32-bit integer lanes reproduces it bit for bit; past that horizon
//! 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
//! 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.
@ -52,13 +60,23 @@
//! for every input including the most negative 32-bit value.
//! - **The lane simulator is the native one.** Random circuits of every gate
//! the core supports, dyadic phases to `P(32)` and up to three controls, on
//! up to ten qubits, under each gateset: every final amplitude
//! bit-identical. The corpus digest of each gateset's certificate is pinned,
//! so a certificate issued under either one can be re-earned. The lane kernel
//! up to ten qubits (random-v1) and up to seventeen (fabric-v1), under each
//! gateset: every final amplitude bit-identical. The corpus digest of each
//! gateset's certificate is pinned for both corpora, so a certificate issued
//! under either one can be re-earned. The lane kernel
//! runs the general 2×2 path for every gate, where the native one takes
//! shortcuts; the shortcuts are exact by construction, and the certificate
//! confirms it. The lane simulator refuses every circuit the native one
//! refuses.
//! - **The range guard is exact.** [`lanes::apply_words`] computes every
//! output as an exact 64-bit value beside the stored 32-bit lane, and flags
//! a gate exactly when the reference's result leaves 32 bits: over random
//! gates and states across the whole 32-bit range, the stored lanes equal
//! [`lanes::apply`]'s and the flag equals a 128-bit recomputation. Valid
//! circuits, the QFT and a long run of phase gates never raise it; the
//! horizon below which it cannot rise is pinned per gateset. The executor of
//! [`plan`] equals the native simulator over both corpora, from more than
//! one start state, and refuses what it refuses.
//! - **Parity is compared, not asserted.** Bit-identical outputs certify
//! bit-exact; a one-unit difference certifies `Within` that bound; a NaN or
//! an infinity on either side, or a length or scale mismatch, certifies
@ -79,19 +97,22 @@
//!
//! # Honest boundaries
//!
//! - **Only detected members are placed, and the CPU is the only one
//! detected.** On x86-64 and aarch64, detection asks the running CPU for
//! its vector units. On other targets it reports the units the binary was
//! compiled for among those it knows (WebAssembly SIMD, 32-bit Arm NEON,
//! the RISC-V vector extension) and says so, with the standard library
//! alone. Other members can be listed as reachable or declared; they are
//! placed once a probe detects them and their own run of the corpus
//! certifies them.
//! - **Only detected members are placed.** [`detect_local`] detects the CPU:
//! on x86-64 and aarch64 it asks the running CPU for its vector units; on
//! other targets it reports the units the binary was compiled for among
//! those it knows (WebAssembly SIMD, 32-bit Arm NEON, the RISC-V vector
//! extension) and says so, with the standard library alone. With
//! `quantum_fabric_webgpu`, each WebGPU adapter that passes the arithmetic
//! probe in this process is detected too; it is placed once its own run of
//! the corpus certifies it. Other members can be listed as reachable or
//! declared; they are placed once a probe detects them and their own run of
//! the corpus certifies them.
//! - **A certificate names the member it ran on.** Run on another platform,
//! the same corpus gives the same parity and corpus digest, and a
//! certificate that differs in its member.
//! - **Lanes on the CPU are a stand-in.** They prove the algorithm a 32-bit
//! integer fabric would run, not that fabric's hardware.
//! integer fabric would run, not that fabric's hardware; the WebGPU member
//! (feature `quantum_fabric_webgpu`) runs them on real hardware.
//! - **Only the gate kernel is in lanes.** Sampling, probability weights and
//! fidelities in the core use 128-bit sums and are not ported.
@ -139,6 +160,23 @@ pub const fn random_generator(gs: Gateset) -> &'static str {
}
}
/// The corpus every fabric member is certified on, for `wai.quantum.circuit`:
/// the random-v1 circuits, then wide circuits of 11 to 17 qubits drawn from the
/// same stream, so every dispatch shape a member uses is exercised.
pub const GENERATOR_FABRIC: &str = "wai.quantum.circuit/fabric-v1";
/// The same corpus as [`GENERATOR_FABRIC`], in the circuit2 gateset.
pub const GENERATOR_FABRIC2: &str = "wai.quantum.circuit2/fabric-v1";
/// The name of the fabric corpus in `gs`: [`GENERATOR_FABRIC`] or
/// [`GENERATOR_FABRIC2`].
pub const fn fabric_generator(gs: Gateset) -> &'static str {
match gs {
Gateset::V1 => GENERATOR_FABRIC,
Gateset::V2 => GENERATOR_FABRIC2,
}
}
const DOMAIN_FABRIC: &[u8] = b"wai:quantum-fabric\x01";
const DOMAIN_FABRIC_ID: &[u8] = b"wai:quantum-fabric-id\x01";
const DOMAIN_CERT: &[u8] = b"wai:quantum-fabric-certificate\x01";
@ -539,8 +577,9 @@ pub struct Certificate {
/// The member id the corpus ran on.
pub member: String,
/// The corpus generator and its parameters, enough to draw the corpus
/// again: `"<generator> seed=<seed> cases=<n>"`, or `"definition"` for
/// the reference.
/// again: `"<generator> seed=<seed> cases=<n>"` (with ` wide=<w>` after
/// it for fabric-v1, [`Corpus::generator`]), or `"definition"` for the
/// reference.
pub generator: String,
pub parity: Parity,
/// Cases compared; zero for the reference itself.
@ -569,31 +608,120 @@ impl Certificate {
}
}
/// [`certify`], one case at a time, for members whose outputs arrive
/// asynchronously (a GPU, a browser): there is no need to hold the corpus.
/// The same cases in the same order give the same [`Certificate`].
#[derive(Clone, Debug)]
pub struct Certification {
hasher: blake3::Hasher,
parity: Option<Parity>,
cases: u64,
}
impl Certification {
/// Starts the corpus digest with the domain [`certify`] uses.
pub fn new() -> Certification {
let mut hasher = blake3::Hasher::new();
hasher.update(DOMAIN_CORPUS);
Certification { hasher, parity: None, cases: 0 }
}
/// Folds one case in exactly as [`certify`] does: the input, the
/// reference's bytes and the member's bytes, each prefixed with its length
/// (u64, little-endian); the parity is the worst so far.
pub fn add(&mut self, input: &[u8], reference: &Output, member: &Output) {
let (r, m) = (reference.bytes(), member.bytes());
for bytes in [input, r.as_slice(), m.as_slice()] {
self.hasher.update(&(bytes.len() as u64).to_le_bytes());
self.hasher.update(bytes);
}
let p = compare(reference, member);
self.parity = Some(self.parity.map_or(p, |q| q.worst(p)));
self.cases += 1;
}
/// The cases folded in so far.
pub fn cases(&self) -> u64 {
self.cases
}
/// The certificate the cases so far earn. With no case it certifies
/// nothing.
pub fn finish(self, kernel: &str, implementation: &str, member: &str, generator: &str) -> Certificate {
Certificate {
kernel: kernel.into(),
implementation: implementation.into(),
member: member.into(),
generator: generator.into(),
parity: self.parity.unwrap_or(Parity::Uncertified),
cases: self.cases,
corpus: *self.hasher.finalize().as_bytes(),
}
}
}
impl Default for Certification {
fn default() -> Self {
Self::new()
}
}
/// Certify `implementation` of `kernel` on `member`: each case is an input
/// and the reference's and the member's outputs on it. The parity is the
/// worst case's, and an empty corpus certifies nothing.
pub fn certify(kernel: &str, implementation: &str, member: &str, generator: &str, cases: impl IntoIterator<Item = (Vec<u8>, Output, Output)>) -> Certificate {
let mut h = blake3::Hasher::new();
h.update(DOMAIN_CORPUS);
let mut parity: Option<Parity> = None;
let mut n = 0u64;
for (input, reference, member) in cases {
for bytes in [input, reference.bytes(), member.bytes()] {
h.update(&(bytes.len() as u64).to_le_bytes());
h.update(&bytes);
}
let p = compare(&reference, &member);
parity = Some(parity.map_or(p, |q| q.worst(p)));
n += 1;
let mut c = Certification::new();
for (input, reference, out) in cases {
c.add(&input, &reference, &out);
}
Certificate {
kernel: kernel.into(),
implementation: implementation.into(),
member: member.into(),
generator: generator.into(),
parity: parity.unwrap_or(Parity::Uncertified),
cases: n,
corpus: *h.finalize().as_bytes(),
c.finish(kernel, implementation, member, generator)
}
/// A certification corpus of the gate kernel: which circuits, from which seed.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub enum Corpus {
/// `cases` circuits of [`lanes::random_circuit`] from `seed` (random-v1).
Random { seed: u64, cases: usize },
/// random-v1's `cases`, then `wide` circuits of
/// [`lanes::random_wide_circuit`] continuing the same stream (fabric-v1).
Fabric { seed: u64, cases: usize, wide: usize },
}
/// The corpus members are placed on: fabric-v1, seed 7, 200 random-v1
/// circuits and 24 wide ones.
pub const FABRIC_DEFAULT: Corpus = Corpus::Fabric { seed: 7, cases: 200, wide: 24 };
impl Corpus {
/// The certificate's generator string:
/// `"<generator> seed=<seed> cases=<cases>"` for random-v1
/// ([`random_generator`], the string [`certify_lanes`] has always
/// written), and `"<generator> seed=<seed> cases=<cases> wide=<wide>"` for
/// fabric-v1 ([`fabric_generator`]).
pub fn generator(&self, gs: Gateset) -> String {
match *self {
Corpus::Random { seed, cases } => format!("{} seed={seed} cases={cases}", random_generator(gs)),
Corpus::Fabric { seed, cases, wide } => format!("{} seed={seed} cases={cases} wide={wide}", fabric_generator(gs)),
}
}
/// The corpus's short name: `"random-v1"` or `"fabric-v1"`.
pub fn name(&self) -> &'static str {
match self {
Corpus::Random { .. } => "random-v1",
Corpus::Fabric { .. } => "fabric-v1",
}
}
/// The circuits, in order. Both gatesets draw the same ops.
pub fn circuits(&self, gs: Gateset) -> Vec<Circuit> {
let (seed, cases, wide) = match *self {
Corpus::Random { seed, cases } => (seed, cases, 0),
Corpus::Fabric { seed, cases, wide } => (seed, cases, wide),
};
let mut s = seed;
let mut v: Vec<Circuit> = (0..cases).map(|_| lanes::random_circuit(&mut s, gs)).collect();
v.extend((0..wide).map(|_| lanes::random_wide_circuit(&mut s, gs)));
v
}
}
@ -601,18 +729,10 @@ pub fn certify(kernel: &str, implementation: &str, member: &str, generator: &str
/// ([`apply_kernel`]), as implementation `"lanes32"` on `member`, against the
/// native simulator on `circuits` random circuits in `gs` drawn from `seed`
/// ([`random_generator`]). The two gatesets draw the same ops from the same
/// seed; their certificates differ in kernel, generator and corpus.
/// seed; their certificates differ in kernel, generator and corpus. The same
/// as [`lanes::certify_corpus`] on [`Corpus::Random`].
pub fn certify_lanes(gs: Gateset, member: &str, circuits: usize, seed: u64) -> Certificate {
let mut s = seed;
let cases = (0..circuits).map(move |_| {
let c = lanes::random_circuit(&mut s, gs);
let native = c.simulate_from(0).expect("random circuits are valid");
let lane = lanes::simulate(&c).expect("random circuits are valid");
let reference = Output::Fixed { values: native.amps.iter().flat_map(|a| [a.re, a.im]).collect(), frac: FRAC };
let port = Output::Fixed { values: lane.iter().flat_map(|a| [a[0] as i64, a[1] as i64]).collect(), frac: FRAC };
(lanes::circuit_bytes(&c), reference, port)
});
certify(apply_kernel(gs), "lanes32", member, &format!("{} seed={seed} cases={circuits}", random_generator(gs)), cases)
lanes::certify_corpus(gs, member, Corpus::Random { seed, cases: circuits })
}
/// Certify a sampling implementation of `kernel` on `member`: each case is an
@ -827,11 +947,14 @@ impl FabricReceipt {
/// The core's gate kernel in 32-bit lanes, using only what a fabric with
/// 32-bit integer lanes has: wrapping 32-bit add, subtract and multiply,
/// shifts and comparisons. A shader or a logic design that runs these steps
/// gets the native simulator's bits. The gate entries are prepared by the
/// host from the core's pinned table and handed to the lanes as 32-bit
/// values.
/// gets the native simulator's bits for as long as every result stays inside
/// 32 bits: always, for a valid circuit of at most [`horizon`](lanes::horizon)
/// gates, and [`apply_words`](lanes::apply_words) says exactly when it does
/// not. The gate entries are prepared by the host from the core's pinned table
/// and handed to the lanes as 32-bit values
/// ([`GateWords`](lanes::GateWords)).
pub mod lanes {
use super::{Circuit, MAX_QUBITS, QuantumError};
use super::{Certificate, Certification, Circuit, Corpus, FRAC, MAX_QUBITS, Output, QuantumError, apply_kernel, plan};
use crate::quantum::{BaseGate, Gate, Gateset, ONE};
/// The 64-bit two's-complement product of two 32-bit values, as
@ -862,9 +985,10 @@ pub mod lanes {
/// The core's fixed-point multiply, `(a·b + 2^29) >> 30`, in lanes.
/// Exact whenever the true result fits in 32 bits, which is exactly when
/// `−2^61 − 2^29 ≤ a·b < 2^61 − 2^29`. The amplitudes and gate entries of
/// a valid circuit are each at most `2^30` in size, so their products are
/// at most `2^60` and always inside.
/// `−2^61 − 2^29 ≤ a·b < 2^61 − 2^29`. Within [`horizon`] gates an
/// amplitude stays below `2^31` in magnitude and a gate entry is at most
/// `2^30`, so the product is inside that range. For every pair of 32-bit
/// inputs the result is the low 32 bits of the true one ([`fxmul_wide`]).
pub fn fxmul(a: i32, b: i32) -> i32 {
let (hi, lo) = mul_wide(a, b);
let lo2 = lo.wrapping_add(1 << 29);
@ -936,6 +1060,17 @@ pub mod lanes {
/// a seed are the same in both.
pub fn random_circuit(s: &mut u64, gs: Gateset) -> Circuit {
let n = 2 + (next(s) % 9) as u8;
random_ops(s, gs, n)
}
/// A random circuit like [`random_circuit`], on 11 to 17 qubits.
pub fn random_wide_circuit(s: &mut u64, gs: Gateset) -> Circuit {
let n = 11 + (next(s) % 7) as u8;
random_ops(s, gs, n)
}
/// The ops of a random circuit on `n` qubits: 10 to 59 of them.
fn random_ops(s: &mut u64, gs: Gateset, n: u8) -> Circuit {
let mut c = Circuit::with_gateset(n, gs);
let bases = [BaseGate::I, BaseGate::X, BaseGate::Y, BaseGate::Z, BaseGate::H, BaseGate::S, BaseGate::Sdg, BaseGate::T, BaseGate::Tdg, BaseGate::P];
for _ in 0..(10 + next(s) % 50) {
@ -967,8 +1102,218 @@ pub mod lanes {
}
o
}
/// A port's lanes as the harness compares them: [`Output::Fixed`] of the
/// lanes widened to 64 bits, `[re, im]` per amplitude, at [`FRAC`].
pub fn output(lanes: &[Lane]) -> Output {
Output::Fixed { values: lanes.iter().flat_map(|a| [i64::from(a[0]), i64::from(a[1])]).collect(), frac: FRAC }
}
/// The reference's output for a corpus circuit: `c.simulate_from(0)`, laid
/// out as [`output`] lays out lanes.
pub fn reference_output(c: &Circuit) -> Result<Output, QuantumError> {
let native = c.simulate_from(0)?;
Ok(Output::Fixed { values: native.amps.iter().flat_map(|a| [a.re, a.im]).collect(), frac: FRAC })
}
/// Certify the lane kernel on the CPU over `corpus`, against the native
/// simulator: implementation `"lanes32"` of [`apply_kernel`]`(gs)` on
/// `member`, with the generator string [`Corpus::generator`].
pub fn certify_corpus(gs: Gateset, member: &str, corpus: Corpus) -> Certificate {
let mut cert = Certification::new();
for c in corpus.circuits(gs) {
let reference = reference_output(&c).expect("corpus circuits are valid");
let port = output(&simulate(&c).expect("corpus circuits are valid"));
cert.add(&circuit_bytes(&c), &reference, &port);
}
cert.finish(apply_kernel(gs), "lanes32", member, &corpus.generator(gs))
}
// -----------------------------------------------------------------------
// Gate words and the exact range guard
// -----------------------------------------------------------------------
/// One gate as a 32-bit-lane member takes it: the matrix entries as 32-bit
/// two's-complement patterns (`m00.re, m00.im, m01.re, m01.im, m10.re,
/// m10.im, m11.re, m11.im`), the target's bit and the controls' bits.
#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
pub struct GateWords {
pub m: [u32; 8],
pub tbit: u32,
pub ctrl: u32,
}
/// The words of a validated gate of a circuit on at most
/// [`MAX_QUBITS`] qubits: the entries of the gate's matrix in `gs` as
/// 32-bit patterns, `tbit = 1 << target`, and `ctrl` the OR of
/// `1 << c` over the controls (a control a circuit/1 gate repeats folds
/// into one bit).
pub fn gate_words(g: &Gate, gs: Gateset) -> GateWords {
let m = g.base.matrix(g.param, gs);
let w = |a: crate::quantum::Amp| [a.re as i32 as u32, a.im as i32 as u32];
let ([a, b], [c, d], [e, f], [h, k]) = (w(m[0][0]), w(m[0][1]), w(m[1][0]), w(m[1][1]));
GateWords { m: [a, b, c, d, e, f, h, k], tbit: 1 << g.target, ctrl: g.controls.iter().fold(0, |acc, &q| acc | 1 << q) }
}
/// The core's `fx(a, b) = (a·b + 2^29) >> 30`, exactly: the product of two
/// 32-bit values is at most `2^62` in size, so the sum is inside 64 bits.
fn fx_exact(a: i32, b: i32) -> i64 {
(i64::from(a) * i64::from(b) + (1 << 29)) >> 30
}
/// The exact value of `fx(a, b) = (a·b + 2^29) >> 30` as a 64-bit
/// two's-complement `(high, low)` pair, for every pair of 32-bit inputs;
/// `low == fxmul(a, b) as u32`.
pub fn fxmul_wide(a: i32, b: i32) -> (u32, u32) {
let v = fx_exact(a, b);
((v >> 32) as u32, v as u32)
}
/// The operands every multiply probe runs pairwise: zero, `±1`, `±2`,
/// `±2^29`, `±(2^30 − 1)`, `±2^30`, `±46341`, `±3·2^29`, `±65535`,
/// `±65536`, and the largest, smallest and next-to-smallest 32-bit values.
pub const PROBE_EDGES: [i32; 22] = [
0,
1,
-1,
2,
-2,
1 << 29,
-(1 << 29),
(1 << 30) - 1,
-((1 << 30) - 1),
1 << 30,
-(1 << 30),
46341,
-46341,
3 << 29,
-(3 << 29),
65535,
-65535,
65536,
-65536,
i32::MAX,
i32::MIN,
i32::MIN + 1,
];
/// One gate on a state of `2^n` lanes (`n ≥ 1`): the general 2×2 path on
/// every pair whose control bits are all set, from exact 64-bit values.
/// The stored lanes are the low 32 bits of those values, and equal
/// [`apply`]'s.
///
/// Returns `Some(true)` if every exact output fits in 32 bits, and
/// `Some(false)` if any does not: then the reference's result has left
/// 32 bits and no lane can hold it. Returns `None`, with the state
/// untouched, if the length is not a power of two of at least 2 or the
/// words are malformed for it: `tbit` zero, not a power of two, or above
/// half the length; `ctrl` sharing `tbit`; `ctrl` with a bit at or above
/// the length.
pub fn apply_words(amps: &mut [Lane], w: &GateWords) -> Option<bool> {
let len = amps.len();
if len < 2 || !len.is_power_of_two() {
return None;
}
let (tbit, ctrl) = (w.tbit as usize, w.ctrl as usize);
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));
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];
amps[j] = [o1[0] as i32, o1[1] as i32];
}
}
Some(fits)
}
/// 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
/// ([`apply_words`]).
pub fn run_words(state: &[Lane], gates: &[GateWords]) -> Option<plan::LaneRun> {
let mut amps = state.to_vec();
let mut out_of_range = None;
for (k, w) in gates.iter().enumerate() {
if !apply_words(&mut amps, w)? && out_of_range.is_none() {
out_of_range = Some(k as u32);
}
}
Some(plan::LaneRun { amps, out_of_range })
}
/// [`Circuit::simulate_from`] in lanes, through [`plan::compile`] and
/// [`plan::execute`], with the first op whose exact result left 32 bits.
/// Refuses what the native simulator refuses, with the same error; the
/// state starts at index `start % 2^n`, as the native one does.
pub fn simulate_checked(c: &Circuit, start: usize) -> Result<plan::LaneRun, QuantumError> {
Ok(plan::execute(&plan::compile(c, start)?))
}
/// `⌈√x⌉`.
fn ceil_sqrt(x: u128) -> u128 {
let r = x.isqrt();
if r * r < x { r + 1 } else { r }
}
/// `2^60·δ` for [`horizon`], rounded up: the largest of
/// `⌈√(|e|²·2^60)⌉ − 2^60` over every `P(k)` entry `e` of the table, and
/// `⌈√(2h²·2^60)⌉ − 2^60` for `H`, whose entries are `±h`; zero if none
/// exceeds one. Every other gate's matrix has one entry of size one in
/// each row and column.
pub(super) fn table_excess(gs: Gateset) -> u128 {
const S: u128 = 1 << 60;
let excess = |sq: u128| ceil_sqrt(sq * S).saturating_sub(S);
let sq = |re: i64, im: i64| (re.unsigned_abs() as u128).pow(2) + (im.unsigned_abs() as u128).pow(2);
let h = gs.inv_sqrt2();
gs.table()[1..].iter().map(|p| excess(sq(p.re, p.im))).chain([excess(2 * sq(h, 0))]).max().unwrap_or(0)
}
/// The number of gates of a valid circuit in `gs` on `n_qubits`, from a
/// basis state, that provably cannot take an amplitude out of 32 bits:
/// every such circuit runs in lanes with the native simulator's bits.
///
/// Each rounded product is off the exact one by at most half a unit, so a
/// gate moves the state by at most `ε = 2√2·2^(n/2)·2^−30` (in units of
/// one) beyond the fixed-point matrix's own action, and that matrix
/// stretches the state by at most `1 + δ`, with `δ` the most by which the
/// size of a `P(k)` entry of the table, or `√2·h` for `H`'s entries `±h`,
/// exceeds one (zero if none does; every other gate's matrix has one entry
/// of size one in each row and column). After `k` gates the state's norm
/// is at most `(1 + δ)^k·(1 + kε)`, and a component reaches `2^31` units
/// only once the norm reaches two. With `k(ε + 2δ) < 1`,
/// `1 + kε < 2(1 − kδ)` and `(1 + δ)^k ≤ 1/(1 − kδ)`, so the norm stays
/// below two. In integers at scale `2^60`:
/// `(2^60 − 1) / (⌈√(2^(63+n))⌉ + 2·⌈2^60·δ⌉)`.
pub fn horizon(n_qubits: u8, gs: Gateset) -> u64 {
const S: u128 = 1 << 60;
// From 58 qubits on, ε alone is more than one: no gate is safe. The
// cut keeps `2^(63+n)` inside 128 bits.
if n_qubits >= 58 {
return 0;
}
let e = ceil_sqrt(1u128 << (63 + u32::from(n_qubits)));
((S - 1) / (e + 2 * table_excess(gs))) as u64
}
}
pub mod plan;
pub mod wgsl;
/// The fixed-point one is a 32-bit lane value.
const _: () = assert!(ONE <= i32::MAX as i64);
@ -1090,8 +1435,8 @@ mod tests {
}
}
#[test]
fn the_lane_simulator_refuses_what_the_native_one_refuses() {
/// Five circuits the native simulator refuses, in each gateset.
pub(super) fn refused_circuits() -> Vec<Circuit> {
let gate = |base, controls: Vec<u8>, target, param| Gate { base, controls, target, param };
let bad = [
(1, vec![gate(BaseGate::H, vec![], 0, 0), gate(BaseGate::P, vec![], 0, 0)]),
@ -1100,14 +1445,18 @@ mod tests {
(2, vec![gate(BaseGate::X, vec![0], 0, 0)]),
(2, vec![gate(BaseGate::X, vec![5], 1, 0)]),
];
for gateset in Gateset::ALL {
for (n, ops) in bad.clone() {
let c = Circuit { n_qubits: n, ops, gateset };
let native = c.simulate_from(0).unwrap_err();
assert_eq!(lanes::simulate(&c).unwrap_err(), native);
}
Gateset::ALL.into_iter().flat_map(|gateset| bad.clone().into_iter().map(move |(n_qubits, ops)| Circuit { n_qubits, ops, gateset })).collect()
}
#[test]
fn the_lane_simulator_refuses_what_the_native_one_refuses() {
for c in refused_circuits() {
let native = c.simulate_from(0).unwrap_err();
assert_eq!(lanes::simulate(&c).unwrap_err(), native);
assert_eq!(lanes::simulate_checked(&c, 0).unwrap_err(), native);
}
assert!(matches!(lanes::simulate(&Circuit::new(MAX_QUBITS + 1)), Err(QuantumError::TooManyQubits(_))));
assert!(matches!(lanes::simulate_checked(&Circuit::new(MAX_QUBITS + 1), 0), Err(QuantumError::TooManyQubits(_))));
}
#[test]
@ -1342,4 +1691,339 @@ mod tests {
assert_eq!((k.kernel.as_str(), k.implementation.as_str()), (apply_kernel(gs), REFERENCE));
}
}
fn widened(lanes: &[lanes::Lane]) -> Vec<[i64; 2]> {
lanes.iter().map(|a| [i64::from(a[0]), i64::from(a[1])]).collect()
}
fn native(c: &Circuit, start: usize) -> Vec<[i64; 2]> {
c.simulate_from(start).unwrap().amps.iter().map(|a| [a.re, a.im]).collect()
}
#[test]
fn certification_is_certify() {
let mut s = 7u64;
let cases: Vec<(Vec<u8>, Output, Output)> = (0..50)
.map(|_| {
let c = lanes::random_circuit(&mut s, Gateset::V1);
(lanes::circuit_bytes(&c), lanes::reference_output(&c).unwrap(), lanes::output(&lanes::simulate(&c).unwrap()))
})
.collect();
// The same cases with one member output a unit off, so that the worst
// case is not the first.
let mut off = cases.clone();
let Output::Fixed { values, .. } = &mut off[17].2 else { panic!() };
values[0] += 1;
for (corpus, parity) in [(cases, Parity::BitExact), (off, Parity::Within { max_ulp: 1, max_abs: 1.0 / (1u64 << 30) as f64 }), (Vec::new(), Parity::Uncertified)] {
let whole = certify("k", "port", "m", "g", corpus.clone());
let mut one = Certification::new();
for (input, reference, member) in &corpus {
one.add(input, reference, member);
}
assert_eq!(one.cases(), corpus.len() as u64);
let one = one.finish("k", "port", "m", "g");
assert_eq!(one, whole);
assert_eq!((one.parity, one.cases), (parity, corpus.len() as u64));
}
let empty = Certification::default().finish("k", "port", "m", "g");
assert_eq!((empty.parity, empty.cases), (Parity::Uncertified, 0));
assert_eq!(empty, certify("k", "port", "m", "g", Vec::new()));
}
#[test]
fn fxmul_wide_is_the_exact_product() {
let mut pairs: Vec<(i32, i32)> = lanes::PROBE_EDGES.iter().flat_map(|&a| lanes::PROBE_EDGES.iter().map(move |&b| (a, b))).collect();
let mut s = 11u64;
for _ in 0..400_000 {
pairs.push((next(&mut s) as i32, next(&mut s) as i32));
}
for (a, b) in pairs {
let (hi, lo) = lanes::fxmul_wide(a, b);
let exact = (i128::from(a) * i128::from(b) + (1 << 29)) >> 30;
assert_eq!(i128::from(((u64::from(hi) << 32) | u64::from(lo)) as i64), exact, "{a} × {b}");
assert_eq!(lo, lanes::fxmul(a, b) as u32, "{a} × {b}");
}
let mut edges = lanes::PROBE_EDGES.to_vec();
edges.sort();
edges.dedup();
assert_eq!(edges.len(), lanes::PROBE_EDGES.len());
}
#[test]
fn apply_words_is_apply() {
let fx = |a: i32, b: i32| (i128::from(a) * i128::from(b) + (1 << 29)) >> 30;
let mut s = 13u64;
let (mut inside, mut outside) = (0, 0);
for case in 0..2000 {
let n = 1 + (next(&mut s) % 10) as u8;
let len = 1usize << n;
// A third of the cases across the whole 32-bit range, some of the
// values edges; a third with every output inside 32 bits; a third
// at the boundary, where a sum of four products may leave it.
let mode = case % 3;
let value = |s: &mut u64, shift: u32| -> i32 {
if mode == 0 && next(s).is_multiple_of(4) { lanes::PROBE_EDGES[(next(s) % 22) as usize] } else { (next(s) as i32) >> shift }
};
let (entry_shift, state_shift) = [(0, 0), (1, 3), (1, 1)][mode];
let entries: [i32; 8] = std::array::from_fn(|_| value(&mut s, entry_shift));
let state: Vec<lanes::Lane> = (0..len).map(|_| [value(&mut s, state_shift), value(&mut s, state_shift)]).collect();
let target = (next(&mut s) % u64::from(n)) as u8;
let tbit = 1usize << target;
let ctrl = if next(&mut s).is_multiple_of(2) { 0 } else { (next(&mut s) as usize) & (len - 1) & !tbit };
let m = [[[entries[0], entries[1]], [entries[2], entries[3]]], [[entries[4], entries[5]], [entries[6], entries[7]]]];
let w = lanes::GateWords { m: entries.map(|e| e as u32), tbit: tbit as u32, ctrl: ctrl as u32 };
let (mut a, mut b) = (state.clone(), state.clone());
let fits = lanes::apply_words(&mut a, &w).unwrap();
lanes::apply(&mut b, target, ctrl, m);
assert_eq!(a, b, "case {case}");
let mut want = true;
for i in (0..len).filter(|i| i & tbit == 0 && i & ctrl == ctrl) {
let (a0, a1) = (state[i], state[i | tbit]);
for [p, q] in m {
let re = fx(p[0], a0[0]) - fx(p[1], a0[1]) + fx(q[0], a1[0]) - fx(q[1], a1[1]);
let im = fx(p[0], a0[1]) + fx(p[1], a0[0]) + fx(q[0], a1[1]) + fx(q[1], a1[0]);
want &= i32::try_from(re).is_ok() && i32::try_from(im).is_ok();
}
}
assert_eq!(fits, want, "case {case}");
if fits { inside += 1 } else { outside += 1 }
}
assert!(inside >= 100 && outside >= 100, "{inside} inside, {outside} outside");
}
#[test]
fn malformed_words_are_refused() {
let mut s = 5u64;
let state: Vec<lanes::Lane> = (0..8).map(|_| [next(&mut s) as i32, next(&mut s) as i32]).collect();
let good = lanes::GateWords { m: [ONE as u32, 0, 0, 0, 0, 0, ONE as u32, 0], tbit: 2, ctrl: 1 };
let mut a = state.clone();
assert_eq!(lanes::apply_words(&mut a, &good), Some(true));
assert_eq!(a, state);
// tbit 0, 3 and 8; ctrl sharing tbit; ctrl at or above the length.
for (tbit, ctrl) in [(0, 0), (3, 0), (8, 0), (2, 2), (2, 3), (1, 8), (4, 9)] {
let mut a = state.clone();
assert_eq!(lanes::apply_words(&mut a, &lanes::GateWords { tbit, ctrl, ..good }), None, "tbit {tbit} ctrl {ctrl}");
assert_eq!(a, state);
}
// A state that is not a power of two of at least two lanes.
for len in [0, 1, 3, 6] {
let mut a = vec![[1, 2]; len];
assert_eq!(lanes::apply_words(&mut a, &lanes::GateWords { tbit: 1, ctrl: 0, ..good }), None);
assert_eq!(a, vec![[1, 2]; len]);
}
assert_eq!(lanes::run_words(&state, &[good, lanes::GateWords { tbit: 8, ..good }]), None);
}
#[test]
fn z_on_the_most_negative_lane_is_flagged() {
let gate = |base| Gate { base, controls: vec![], target: 0, param: 0 };
for gs in Gateset::ALL {
let (z, i) = (lanes::gate_words(&gate(BaseGate::Z), gs), lanes::gate_words(&gate(BaseGate::I), gs));
let state = [[0, 0], [i32::MIN, 0]];
let mut amps = state.to_vec();
assert_eq!(lanes::apply_words(&mut amps, &z), Some(false));
assert_eq!(amps, state);
// The exact value is 2^31, one past the largest lane.
assert_eq!(lanes::fxmul_wide(-(ONE as i32), i32::MIN), (0, 1 << 31));
// A run names the first gate that left 32 bits.
assert_eq!(lanes::run_words(&state, &[i, z, i]).unwrap().out_of_range, Some(1));
assert_eq!(lanes::run_words(&state, &[z, i, z]).unwrap().out_of_range, Some(0));
assert_eq!(lanes::run_words(&state, &[i, i]).unwrap(), plan::LaneRun { amps: state.to_vec(), out_of_range: None });
}
}
#[test]
fn simulate_checked_is_simulate() {
for gs in Gateset::ALL {
for c in (Corpus::Random { seed: 7, cases: 300 }).circuits(gs) {
let dim = 1usize << c.n_qubits;
for start in [0, 1, dim + 3] {
let run = lanes::simulate_checked(&c, start).unwrap();
assert_eq!(run.out_of_range, None);
assert_eq!(widened(&run.amps), native(&c, start));
}
assert_eq!(lanes::simulate_checked(&c, dim + 3), lanes::simulate_checked(&c, 3));
// The same gates as words, from the same state.
let mut basis = vec![[0, 0]; dim];
basis[0] = [ONE as i32, 0];
let words: Vec<lanes::GateWords> = c.ops.iter().map(|g| lanes::gate_words(g, gs)).collect();
assert_eq!(lanes::run_words(&basis, &words), Some(lanes::simulate_checked(&c, 0).unwrap()));
}
}
}
#[test]
fn a_rounded_product_is_never_negated() {
let v2 = [[759250126, 0], [1, 0], [759250126, 0], [1, 0]];
let v1 = [[759250122, 0], [0, 0], [759250122, 0], [0, 0]];
for (gs, want) in [(Gateset::V2, v2), (Gateset::V1, v1)] {
let mut c = Circuit::with_gateset(2, gs);
c.h(0).h(1).h(0);
assert_eq!(native(&c, 0), widened(&want));
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));
// 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;
let mut a = [[ONE as i32, 0], [0, 0], [0, 0], [0, 0]];
for t in [0, 1, 0] {
let bit = 1usize << t;
for i in (0..4).filter(|i| i & bit == 0) {
let j = i | bit;
for k in 0..2 {
let (p, q) = (lanes::fxmul(h, a[i][k]), lanes::fxmul(h, a[j][k]));
(a[i][k], a[j][k]) = (p + q, p - q);
}
}
}
match gs {
// A tie: fx(h, h) = 2^29 exactly, and fx(h, 2^29) and
// fx(−h, 2^29) differ by more than a sign.
Gateset::V2 => {
assert_eq!((a[1][0], a[3][0]), (0, 0));
assert_ne!(a, want);
}
// An even h: no tie, the same bits either way.
Gateset::V1 => assert_eq!(a, want),
}
}
}
#[test]
fn the_guard_horizon_is_pinned() {
assert_eq!((lanes::table_excess(Gateset::V1), lanes::table_excess(Gateset::V2)), (0, 386_868_465));
for (n, v1, v2) in [(10, 11_863_283, 11_769_578), (20, 370_727, 370_635), (26, 46_340, 46_339)] {
assert_eq!((lanes::horizon(n, Gateset::V1), lanes::horizon(n, Gateset::V2)), (v1, v2), "n = {n}");
}
for n in 1..=26 {
assert!(lanes::horizon(n, Gateset::V1) >= lanes::horizon(n, Gateset::V2), "n = {n}");
for gs in Gateset::ALL {
assert!(lanes::horizon(n, gs) <= lanes::horizon(n - 1, gs), "{gs:?} n = {n}");
}
}
// Rounding alone can carry a state of 57 qubits or more out of 32
// bits in one gate.
for gs in Gateset::ALL {
assert_eq!((lanes::horizon(56, gs), lanes::horizon(57, gs), lanes::horizon(58, gs), lanes::horizon(u8::MAX, gs)), (1, 0, 0, 0));
}
}
#[test]
fn valid_circuits_stay_in_lanes() {
for gs in Gateset::ALL {
for c in (Corpus::Random { seed: 7, cases: 1000 }).circuits(gs) {
assert_eq!(lanes::simulate_checked(&c, 0).unwrap().out_of_range, None);
}
for n in 1..=16 {
let mut c = Circuit::qft(n);
c.gateset = gs;
assert_eq!(lanes::simulate_checked(&c, 0).unwrap().out_of_range, None, "qft({n}) {gs:?}");
}
// A long run of H, T and P(3..=32) on four qubits.
let mut c = Circuit::with_gateset(4, gs);
let mut s = 21u64;
for _ in 0..10_000 {
let q = (next(&mut s) % 4) as u8;
match next(&mut s) % 3 {
0 => c.h(q),
1 => c.t(q),
_ => c.p(3 + (next(&mut s) % 30) as u16, q),
};
}
assert!(c.ops.len() as u64 <= lanes::horizon(4, gs));
let run = lanes::simulate_checked(&c, 0).unwrap();
assert_eq!(run.out_of_range, None, "{gs:?}");
assert_eq!(widened(&run.amps), native(&c, 0), "{gs:?}");
}
}
#[test]
fn fabric_corpus_starts_with_random_v1() {
let random_v1 = Corpus::Random { seed: 7, cases: 200 };
assert_eq!((FABRIC_DEFAULT.name(), random_v1.name()), ("fabric-v1", "random-v1"));
for (gs, pin) in [(Gateset::V1, "d18454c639acfd9c7b66450f4ded1aac4ab46cc10a7bedbcf4958a96e204fb36"), (Gateset::V2, "3dbde276c1095bf827faec71d5980e050f8a55a5ba7d97cd3bbdaec3f43a13b2")] {
let fabric = FABRIC_DEFAULT.circuits(gs);
assert_eq!(fabric.len(), 224);
assert_eq!(fabric[..200], random_v1.circuits(gs)[..]);
// One stream: random-v1's circuits, then the wide ones after them.
let mut s = 7u64;
for c in &fabric[..200] {
assert_eq!(*c, lanes::random_circuit(&mut s, gs));
}
for c in &fabric[200..] {
assert_eq!(*c, lanes::random_wide_circuit(&mut s, gs));
assert!((11..=17).contains(&c.n_qubits), "{}", c.n_qubits);
}
let cert = lanes::certify_corpus(gs, "local/test", random_v1);
assert_eq!(hex(&cert.corpus), pin);
assert_eq!(cert.generator, format!("{} seed=7 cases=200", random_generator(gs)));
assert_eq!(cert, certify_lanes(gs, "local/test", 200, 7));
}
}
/// The fabric-v1 corpus of each gateset, seed 7, 200 + 24 cases. The
/// digests were computed by this test's first run, when the corpus was
/// introduced; they must never be edited to make this test pass.
#[test]
fn the_fabric_v1_corpus_of_each_gateset_is_pinned() {
for (gs, pin) in [(Gateset::V1, "41c726a0ac80f7defef92e15ecf893bdb7a3076b5e7f670fa0b7142e5c033df7"), (Gateset::V2, "9a84b55af68f1a8b07390b234cf562dd8d85cef3185a697ab85690c8790a7809")] {
let cert = lanes::certify_corpus(gs, "local/test", FABRIC_DEFAULT);
assert_eq!(cert.parity, Parity::BitExact, "{cert:?}");
assert_eq!(cert.cases, 224);
assert_eq!(cert.generator, format!("{} seed=7 cases=200 wide=24", fabric_generator(gs)));
assert_eq!(hex(&cert.corpus), pin, "{gs:?}");
}
}
/// With 64 invocations to a workgroup, one per amplitude pair, and at most
/// 64 workgroups to a row: up to 6 qubits fill part of one workgroup, 7 to
/// 13 one row, and 14 or more a grid (`wgsl::tests::grids_cover_every_item_once`
/// checks these bounds against `wgsl::gate_grid`). If a count is ever zero
/// for the fixed seed, the corpus gains wide circuits; the assertion is
/// never relaxed.
#[test]
fn fabric_v1_covers_every_dispatch_shape() {
for gs in Gateset::ALL {
let cs = FABRIC_DEFAULT.circuits(gs);
let circuits = |f: &dyn Fn(&Circuit) -> bool| cs.iter().filter(|c| f(c)).count();
let ops = |f: &dyn Fn(&Gate) -> bool| cs.iter().flat_map(|c| &c.ops).filter(|g| f(g)).count();
let shapes = [
("circuit on at most 6 qubits", circuits(&|c| c.n_qubits <= 6)),
("circuit on 7 to 13 qubits", circuits(&|c| (7..=13).contains(&c.n_qubits))),
("circuit on 14 qubits or more", circuits(&|c| c.n_qubits >= 14)),
("op with a control", ops(&|g| !g.controls.is_empty())),
("op with two controls or more", ops(&|g| g.controls.len() >= 2)),
("op on a target of 10 or more", ops(&|g| g.target >= 10)),
("P(k) with k of 19 or more", ops(&|g| g.base == BaseGate::P && g.param >= 19)),
];
for (what, count) in shapes {
assert!(count > 0, "{gs:?}: no {what}");
}
}
}
#[test]
fn gate_words_are_the_table() {
use BaseGate::*;
let gates = [I, X, Y, Z, H, S, Sdg, T, Tdg].map(|b| (b, 0u16)).into_iter().chain((1..=32).map(|k| (P, k)));
for gs in Gateset::ALL {
for (base, param) in gates.clone() {
for (target, controls) in [(0u8, vec![]), (3, vec![0, 5]), (25, vec![24, 1, 7])] {
let g = Gate { base, controls: controls.clone(), target, param };
let w = lanes::gate_words(&g, gs);
let m = base.matrix(param, gs);
let table: Vec<i64> = [m[0][0], m[0][1], m[1][0], m[1][1]].iter().flat_map(|a| [a.re, a.im]).collect();
assert_eq!(w.m.map(|x| i64::from(x as i32)).to_vec(), table, "{base:?}({param}) {gs:?}");
assert_eq!(w.m.to_vec(), table.iter().map(|&e| e as i32 as u32).collect::<Vec<_>>());
assert_eq!(w.tbit, 1 << target);
assert_eq!(w.ctrl, controls.iter().fold(0u32, |acc, &c| acc | 1 << c));
}
}
}
// A circuit/1 gate may repeat a control: one bit.
let c = Circuit { n_qubits: 2, ops: vec![Gate { base: X, controls: vec![0, 0], target: 1, param: 0 }], gateset: Gateset::V1 };
c.validate().unwrap();
assert_eq!(lanes::gate_words(&c.ops[0], Gateset::V1).ctrl, 1);
}
}

152
wai/wai-quantum/src/quantum_fabric/lanes32.wgsl Normal file
View file

@ -0,0 +1,152 @@
// The gate kernel of wai.quantum.circuit and wai.quantum.circuit2 in 32-bit
// lanes: a line-for-line transcription of quantum_fabric::lanes (mul_wide,
// fxmul, cmul, cadd, apply) that also carries every result as an exact 64-bit
// value and reports any result that leaves the 32-bit range.
//
// Every value is a u32. A signed amplitude or matrix entry is the bit pattern
// of its i32 value. This file holds no signed type, no float, no division, no
// conversion or bitcast, no workgroup memory, no barrier, and no shift whose
// amount is not a literal; no constant expression overflows. Any edit to this
// file changes the implementation id, and every certificate must be earned again.
struct Gate {
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
pairs: u32, // 2^(n-1)
row_stride: u32, // invocations per grid row: 64 x workgroups along x
index: u32, // the gate's position in the op-log
pad0: u32,
pad1: u32,
pad2: u32,
}
@group(0) @binding(0) var<storage, read_write> re: array<u32>;
@group(0) @binding(1) var<storage, read_write> im: array<u32>;
// report[0]: bit 1, a result left the 32-bit range; bit 2, an index outside the
// state; bit 4, a malformed record. report[1]: the lowest op-log index that set
// bit 1. The host writes (0, 0xffffffff) before a run.
@group(0) @binding(2) var<storage, read_write> report: array<atomic<u32>, 2>;
@group(0) @binding(3) var<uniform> g: Gate;
const SIGN: u32 = 0x80000000u;
const HALF: u32 = 0x20000000u; // 2^29, the rounding constant of fx
const LOW16: u32 = 0x0000ffffu;
const EXT30: u32 = 0xfffffffcu; // the bits an arithmetic shift by 30 copies from the sign
const ONES: u32 = 0xffffffffu;
fn neg(x: u32) -> bool { return (x & SIGN) != 0u; }
fn bit(c: bool) -> u32 { return select(0u, 1u, c); }
// lanes::mul_wide: the 64-bit two's-complement product of two i32 patterns, as (hi, lo).
fn mul_wide(a: u32, b: u32) -> vec2<u32> {
let negative = neg(a) != neg(b);
let x = select(a, 0u - a, neg(a)); // i32::unsigned_abs
let y = select(b, 0u - b, neg(b));
let x1 = x >> 16u;
let x0 = x & LOW16;
let y1 = y >> 16u;
let y0 = y & LOW16;
let p00 = x0 * y0; // no partial product reaches 2^32
let p01 = x0 * y1;
let p10 = x1 * y0;
let p11 = x1 * y1;
let mid = p01 + p10;
let mid_carry = bit(mid < p01);
let lo = p00 + (mid << 16u);
let lo_carry = bit(lo < p00);
let hi = p11 + (mid >> 16u) + (mid_carry << 16u) + lo_carry;
if (negative) {
let lo_n = ~lo + 1u;
let hi_n = ~hi + bit(lo_n == 0u);
return vec2<u32>(hi_n, lo_n);
}
return vec2<u32>(hi, lo);
}
// The exact 64-bit value of fx(a, b) = (a*b + 2^29) >> 30, as (hi, lo).
// lo is lanes::fxmul(a, b).
fn fx64(a: u32, b: u32) -> vec2<u32> {
let w = mul_wide(a, b);
let lo2 = w.y + HALF;
let hi2 = w.x + bit(lo2 < w.y);
let lo = (lo2 >> 30u) | (hi2 << 2u);
let hi = (hi2 >> 30u) | select(0u, EXT30, neg(hi2));
return vec2<u32>(hi, lo);
}
fn add64(x: vec2<u32>, y: vec2<u32>) -> vec2<u32> {
let lo = x.y + y.y;
return vec2<u32>(x.x + y.x + bit(lo < x.y), lo);
}
fn sub64(x: vec2<u32>, y: vec2<u32>) -> vec2<u32> {
return vec2<u32>(x.x - y.x - bit(x.y < y.y), x.y - y.y);
}
// Whether a 64-bit value is an i32: its high word repeats its low word's sign.
fn fits(v: vec2<u32>) -> bool { return v.x == select(0u, ONES, neg(v.y)); }
// One output amplitude m_r0*a0 + m_r1*a1, m = (m_r0.re, m_r0.im, m_r1.re, m_r1.im),
// each part exact; its low words are lanes::cadd(cmul(m_r0, a0), cmul(m_r1, a1)).
// Returns (re.hi, re.lo, im.hi, im.lo).
fn row(m: vec4<u32>, a0: vec2<u32>, a1: vec2<u32>) -> vec4<u32> {
let r = add64(sub64(fx64(m.x, a0.x), fx64(m.y, a0.y)),
sub64(fx64(m.z, a1.x), fx64(m.w, a1.y)));
let i = add64(add64(fx64(m.x, a0.y), fx64(m.y, a0.x)),
add64(fx64(m.z, a1.y), fx64(m.w, a1.x)));
return vec4<u32>(r, i);
}
@compute @workgroup_size(64)
fn apply(@builtin(global_invocation_id) id: vec3<u32>) {
let tbit = g.tbit;
let span = g.pairs << 1u; // 2^n
// A malformed record is reported and never run.
if (tbit == 0u || (tbit & (tbit - 1u)) != 0u || tbit > g.pairs
|| (g.ctrl & tbit) != 0u || (g.ctrl & ~(span - 1u)) != 0u) {
atomicOr(&report[0], 4u);
return;
}
let p = id.y * g.row_stride + id.x;
if (p >= g.pairs) { return; }
let low = p & (tbit - 1u);
let i = ((p ^ low) << 1u) | low; // p with a 0 inserted at the target bit
let j = i | tbit;
if ((i & g.ctrl) != g.ctrl) { return; }
if (j >= arrayLength(&re) || j >= arrayLength(&im)) {
atomicOr(&report[0], 2u);
return;
}
let a0 = vec2<u32>(re[i], im[i]);
let a1 = vec2<u32>(re[j], im[j]);
let o0 = row(g.row0, a0, a1);
let o1 = row(g.row1, a0, a1);
re[i] = o0.y;
im[i] = o0.w;
re[j] = o1.y;
im[j] = o1.w;
if (!(fits(o0.xy) && fits(o0.zw) && fits(o1.xy) && fits(o1.zw))) {
atomicOr(&report[0], 1u);
atomicMin(&report[1], g.index);
}
}
// The arithmetic probe: mul_wide and fx64 of each case, for detection and for
// the full-range check. Its bindings are distinct from apply's.
struct Probe { a: u32, b: u32, hi: u32, lo: u32, fx_hi: u32, fx_lo: u32, pad0: u32, pad1: u32, }
@group(0) @binding(4) var<storage, read_write> cases: array<Probe>;
@group(0) @binding(5) var<uniform> probe_grid: vec4<u32>; // x: case count, y: invocations per grid row
@compute @workgroup_size(64)
fn probe(@builtin(global_invocation_id) id: vec3<u32>) {
let k = id.y * probe_grid.y + id.x;
if (k >= probe_grid.x || k >= arrayLength(&cases)) { return; }
let w = mul_wide(cases[k].a, cases[k].b);
let f = fx64(cases[k].a, cases[k].b);
cases[k].hi = w.x;
cases[k].lo = w.y;
cases[k].fx_hi = f.x;
cases[k].fx_lo = f.y;
}

162
wai/wai-quantum/src/quantum_fabric/plan.rs Normal file
View file

@ -0,0 +1,162 @@
//! The program a fabric member runs for a circuit, and the CPU executor of
//! exactly that program. A member is correct when its runs equal [`execute`]'s;
//! `execute` is certified against the native simulator for every
//! 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.
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;
/// One gate of a program.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub struct PlannedGate {
pub words: GateWords,
/// The gate's position in the op-log (what a range report names).
pub index: u32,
/// `m01 = m10 = (0, 0)`: the gate acts on each amplitude alone.
pub diag: bool,
}
/// 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 },
}
/// What a member runs for one circuit from one basis state.
#[derive(Clone, Debug, PartialEq, Eq)]
pub struct Program {
pub n_qubits: u8,
/// The basis state the run starts from, already reduced modulo `2^n`.
pub start: usize,
pub gates: Vec<PlannedGate>,
pub steps: Vec<Step>,
}
/// 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.
pub fn compile(c: &Circuit, start: usize) -> 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 })
}
/// A run in lanes.
#[derive(Clone, Debug, PartialEq, Eq)]
pub struct LaneRun {
pub amps: Vec<Lane>,
/// The lowest op-log index whose exact result left 32 bits; `None` if
/// none did. When it is set, `amps` holds only the low 32 bits of the
/// reference's values from that op on.
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.
pub fn execute(p: &Program) -> LaneRun {
let mut amps = vec![[0i32; 2]; 1usize << p.n_qubits];
amps[p.start] = [ONE as i32, 0];
let mut out_of_range: Option<u32> = None;
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)));
}
}
}
}
LaneRun { amps, out_of_range }
}
#[cfg(test)]
mod tests {
use super::*;
use crate::quantum::{BaseGate, Gate, Gateset};
use crate::quantum_fabric::{Corpus, FABRIC_DEFAULT};
fn widened(lanes: &[Lane]) -> Vec<[i64; 2]> {
lanes.iter().map(|a| [i64::from(a[0]), i64::from(a[1])]).collect()
}
fn native(c: &Circuit, start: usize) -> Vec<[i64; 2]> {
c.simulate_from(start).unwrap().amps.iter().map(|a| [a.re, a.im]).collect()
}
#[test]
fn execute_is_the_native_simulator() {
for gs in Gateset::ALL {
let random = Corpus::Random { seed: 7, cases: 300 }.circuits(gs);
let fabric = FABRIC_DEFAULT.circuits(gs);
for (name, circuits) in [("random-v1", random), ("fabric-v1", fabric)] {
for c in &circuits {
for start in [0, 5] {
let p = compile(c, start).unwrap();
assert_eq!(p.start, start % (1 << c.n_qubits));
assert_eq!((p.gates.len(), p.steps.len()), (c.ops.len(), c.ops.len()));
let run = execute(&p);
assert_eq!(run.out_of_range, None, "{name} {gs:?}");
assert_eq!(widened(&run.amps), native(c, start), "{name} {gs:?} start {start}");
}
}
}
}
// The program is the op-log: one step per op, in order, each gate
// naming its position; `diag` exactly when the off-diagonal is zero.
let mut c = Circuit::with_gateset(3, Gateset::V1);
c.h(0).t(1).cx(0, 2).p(19, 2).y(1).z(0);
let p = compile(&c, 11).unwrap();
assert_eq!(p.start, 3);
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]);
// 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:?}");
}
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(&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);
both.ops.push(Gate { base: BaseGate::P, controls: vec![], target: 0, param: 0 });
assert_eq!(compile(&both, 0).unwrap_err(), both.simulate_from(0).unwrap_err());
assert!(matches!(compile(&both, 0), Err(QuantumError::Invalid(_))));
}
}

645
wai/wai-quantum/src/quantum_fabric/wgsl.rs Normal file
View file

@ -0,0 +1,645 @@
//! The gate kernel in WGSL, and everything a WebGPU member needs that is not
//! the device. Pure Rust: compiled and tested on every target.
//!
//! [`SHADER`] is the kernel: [`lanes`] transcribed line for line into WGSL,
//! with every value a `u32` holding a 32-bit two's-complement pattern, and an
//! exact 64-bit copy of every result beside it so that the kernel reports any
//! result that leaves 32 bits. Its entry point `apply` runs one gate over the
//! whole state, one invocation per amplitude pair, in place; `probe` runs the
//! wide multiply and the exact `fx` on listed operands. Around it, this module
//! holds what a host needs: the implementation id ([`implementation_id`]), the
//! 64-byte gate record ([`gate_record`]), the dispatch grid ([`grid`],
//! [`gate_grid`]), the id of a member built on an adapter
//! ([`AdapterIdentity`]), and the textual edits of the kernel a device must be
//! caught running ([`MUTANTS`]).
//!
//! The state is two planes, `re` and `im`, each `2^n` `u32` words. The report
//! is two atomic words: `report[0]` collects bit 1 (a result left 32 bits),
//! bit 2 (an index outside the state) and bit 4 (a malformed record), and
//! `report[1]` the lowest op-log index that set bit 1. The host writes
//! `(0, 0xffffffff)` before a run.
//!
//! # Why the kernel is bit-exact
//!
//! **Lemma 1 (the same function).** Line for line:
//!
//! | [`lanes`] (Rust) | WGSL | Why equal |
//! |---|---|---|
//! | `a < 0` | `neg(a)` | the two's-complement sign bit |
//! | `a.unsigned_abs()` | `select(a, 0u - a, neg(a))` | for `a < 0`, \|a\| = 2^32 − pattern(a) = `0u − a` modulo 2^32, the most negative value (→ 2^31) included |
//! | `x >> 16`, `x & 0xffff` | the same | logical shift and mask |
//! | `wrapping_mul` of halves | `*` | every partial product is below 2^32: exact in both |
//! | `wrapping_add`, `u32::from(mid < p01)` | `+`, `bit(mid < p01)` | `u32` addition modulo 2^32; `u32` comparison |
//! | `p00.wrapping_add(mid << 16)` | `p00 + (mid << 16u)` | `<<` discards the high bits in both |
//! | `(!lo).wrapping_add(1)`, … | `~lo + 1u`, … | complement, then addition modulo 2^32 |
//! | `fxmul`: `lo2`, `hi2`, `(lo2 >> 30) \| (hi2 << 2)` | `fx64(..).y` | identical expressions |
//! | `cmul`, `cadd` with wrapping 32-bit operations | the low words of `add64`, `sub64` | the low word of a 64-bit sum is the 32-bit wrapping sum of the low words |
//! | `(a0, a1)` read before both stores | `a0`, `a1` loaded before any store | the same data dependence |
//!
//! No partial product reaches 2^32 (`x0·y1 ≤ 0xffff·0x8000`), so no `u32`
//! multiply wraps; `mid_carry` is always zero and is kept so that the
//! transcription stays literal. The only intended wraps are the carry
//! additions, `mid << 16u`, the two's-complement negation, and the low words of
//! `add64` and `sub64`. So for every state and every record, the stored words
//! equal those of [`lanes::apply_words`], and report bit 1 is raised exactly
//! when it returns `Some(false)`. For a record whose `pairs` is half the
//! state's length, report bit 4 is raised exactly when it returns `None`: the
//! same predicate on `tbit` and `ctrl`.
//!
//! **Lemma 2 (`fx64` is exact).** `mul_wide` is exact for every pair of 32-bit
//! values; `|a·b| ≤ 2^62`, so adding 2^29 cannot overflow 64 bits; and an
//! arithmetic shift of `(hi2:lo2)` by 30 has as its low word bits 30 to 61,
//! `(lo2 >> 30) | (hi2 << 2)`, and as its high word `hi2 >> 30` with the sign
//! copied into bits 2 to 31, which is `EXT30`. Each output part is a sum of four
//! such values, each at most about 2^32 in size, so `add64` and `sub64` hold
//! it exactly, and `fits` (the high word repeats the low word's sign) is true
//! exactly when the value is a 32-bit one.
//!
//! **Lemma 3 (no dynamic error).** The map from `p` to `i`, a zero inserted at
//! bit `t`, is strictly increasing on `[0, 2^(n−1))` with inverse "delete bit
//! `t`": a bijection onto the indices with bit `t` clear, the set
//! [`lanes::apply`] visits. `p = id.y·row_stride + id.x` with
//! `id.x < row_stride` is injective, and the host's grid ([`gate_grid`])
//! covers every `p` below `pairs`. Every array access is guarded by
//! `arrayLength` and `i < j`, and a malformed record never reaches an access.
//! The pairs partition the indices, so no two invocations touch one location;
//! `report` is touched only by atomics; the record and `arrayLength` are only
//! read.
//!
//! **Lemma 4 (order does not matter).** Each invocation reads and writes only
//! its own pair, so a dispatch's result does not depend on the order in which
//! invocations run, and equals the sequential loop of [`lanes::apply`], whose
//! iterations also touch disjoint pairs. Dispatches run in queue order, each its
//! own usage scope. The final values of `atomicOr` and `atomicMin` do not depend
//! on order.
//!
//! **Theorem.** On an implementation that conforms for the constructs the
//! kernel uses (`u32` `+ − *` modulo 2^32; `~ & | ^`; shifts by literal amounts
//! below 32; `== != <`; `select`; relaxed `atomicOr` and `atomicMin`;
//! little-endian buffer layout; dispatches in queue order), dispatching `apply`
//! once per gate in op-log order over [`gate_grid`] gives the lanes and
//! `out_of_range` of [`plan::execute`]. With the exact range guard, if no gate
//! is reported, the state equals that of
//! [`Circuit::simulate_from`](crate::quantum::Circuit::simulate_from); a valid
//! circuit of at most [`lanes::horizon`] gates cannot have a gate reported.
//!
//! The kernel holds no signed type, no float, no division, no conversion or
//! bitcast, no workgroup memory and no barrier, and every shift amount is a
//! literal between 1 and 31: so the kernel brings no signed overflow, no shift
//! by 32 or more and no signed right shift into any language it is translated
//! to. The rules are checked on the translator's own reading of the text by
//! the unpublished crate `wai/tools/quantum-wgsl-check`.
//!
//! # 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.
//!
//! [`lanes`]: super::lanes
//! [`lanes::apply`]: super::lanes::apply
//! [`lanes::apply_words`]: super::lanes::apply_words
//! [`lanes::horizon`]: super::lanes::horizon
//! [`plan::execute`]: super::plan::execute
use super::lanes::GateWords;
use super::{FabricClass, put};
use crate::quantum::MAX_QUBITS;
/// The kernel, `lanes32.wgsl`, exactly as it is compiled.
pub const SHADER: &str = include_str!("lanes32.wgsl");
/// The implementation name the id starts with.
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.
pub const HOST_LAYOUT: u32 = 1;
/// Invocations per workgroup, as the kernel declares.
pub const WORKGROUP: u32 = 64;
/// Workgroups along x, at most. Production and certification use the same
/// value, so two-dimensional grids (14 qubits or more) occur inside the
/// certification corpus.
pub const X_GROUPS_MAX: u32 = 64;
/// Bytes of one gate record, the kernel's `Gate` uniform.
pub const GATE_RECORD_BYTES: usize = 64;
const DOMAIN_IMPLEMENTATION: &[u8] = b"wai:quantum-fabric-wgsl\x01";
const DOMAIN_MEMBER: &[u8] = b"wai:quantum-fabric-webgpu-member\x01";
fn hex(bytes: &[u8]) -> String {
bytes.iter().map(|b| format!("{b:02x}")).collect()
}
/// `"wgsl-lanes32/<16 hex>"`: the first 8 bytes of
/// BLAKE3(`"wai:quantum-fabric-wgsl\x01"` ‖ put([`SHADER`]) ‖ [`HOST_LAYOUT`]
/// as a big-endian `u32`), in hex. Any edit to the kernel's text or to the host
/// layout changes it.
pub fn implementation_id() -> String {
let mut o = DOMAIN_IMPLEMENTATION.to_vec();
put(&mut o, SHADER.as_bytes());
o.extend_from_slice(&HOST_LAYOUT.to_be_bytes());
format!("{IMPLEMENTATION}/{}", hex(&blake3::hash(&o).as_bytes()[..8]))
}
/// The 64-byte little-endian uniform record of one gate: `m[0..8]`, `tbit`,
/// `ctrl`, `pairs`, `row_stride`, `index`, then three zero words.
pub fn gate_record(w: &GateWords, pairs: u32, row_stride: u32, index: u32) -> [u8; GATE_RECORD_BYTES] {
let mut r = [0u8; GATE_RECORD_BYTES];
let words = w.m.into_iter().chain([w.tbit, w.ctrl, pairs, row_stride, index, 0, 0, 0]);
for (k, v) in words.enumerate() {
r[4 * k..4 * k + 4].copy_from_slice(&v.to_le_bytes());
}
r
}
/// A dispatch: `x` by `y` workgroups of [`WORKGROUP`] invocations, where the
/// invocation at `(gx, gy)` handles item `gy·row_stride + gx` if it is below
/// `items`.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub struct Grid {
pub x: u32,
pub y: u32,
pub row_stride: u32,
pub items: u32,
}
/// Workgroups covering `items` invocations: `groups = ⌈items / 64⌉`,
/// `x = min(groups, max_x)`, `y = ⌈groups / x⌉`, `row_stride = 64·x`.
///
/// Panics if `items` or `max_x` is zero, or if `64·x` does not fit in 32 bits.
pub fn grid(items: u32, max_x: u32) -> Grid {
assert!(items >= 1 && max_x >= 1, "a grid covers at least one item, with at least one workgroup along x");
let groups = items.div_ceil(WORKGROUP);
let x = groups.min(max_x);
let row_stride = WORKGROUP.checked_mul(x).expect("the row stride fits in 32 bits");
Grid { x, y: groups.div_ceil(x), row_stride, items }
}
/// The grid of one gate on `n_qubits`: `grid(2^(n−1), X_GROUPS_MAX)`, one
/// invocation per amplitude pair.
///
/// Panics unless `1 ≤ n_qubits ≤ MAX_QUBITS`.
pub fn gate_grid(n_qubits: u8) -> Grid {
assert!((1..=MAX_QUBITS).contains(&n_qubits), "a gate acts on 1 to {MAX_QUBITS} qubits, not {n_qubits}");
grid(1 << (n_qubits - 1), X_GROUPS_MAX)
}
/// 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.
#[derive(Clone, Debug, PartialEq, Eq)]
pub struct AdapterIdentity {
/// `"metal"`, `"vulkan"`, `"dx12"` or `"webgpu"`.
pub backend: String,
pub name: String,
pub vendor: u32,
pub device: u32,
/// `"discrete"`, `"integrated"`, `"virtual"`, `"cpu"` or `"other"`.
pub device_type: String,
pub pci_bus: String,
pub driver: String,
pub driver_info: String,
/// What turns WGSL into the platform's shader language, as the device
/// module names it: `"naga"` natively, followed where a bytecode compiler
/// is pinned by the bytecode format it emits (`"naga+dxbc"` on DX12), and
/// `"user-agent"` for a browser's own compiler.
pub translator: String,
pub max_storage_binding: u64,
pub max_buffer: u64,
}
impl AdapterIdentity {
fn software(&self) -> bool {
self.device_type == "cpu"
}
/// [`FabricClass::Cpu`] for a software adapter (device type `"cpu"`),
/// [`FabricClass::Gpu`] otherwise.
pub fn class(&self) -> FabricClass {
if self.software() { FabricClass::Cpu } else { FabricClass::Gpu }
}
/// `"local/gpu-<backend>-<d12>"`, or `"local/swgpu-<backend>-<d12>"` for a
/// software adapter. `d12` is the first 12 hex digits of
/// BLAKE3(`"wai:quantum-fabric-webgpu-member\x01"` ‖ put(backend) ‖
/// put(name) ‖ vendor ‖ device ‖ put(device_type) ‖ put(pci_bus) ‖
/// put(driver) ‖ put(driver_info) ‖ put(translator) ‖ max_storage_binding ‖
/// max_buffer), the numbers big-endian.
///
/// Every such id sorts after every `local/cpu-<arch>`, so with no energy
/// figures [`place`](super::place) keeps the CPU reference: a software
/// adapter is class `Cpu`, but its id does not begin `local/cpu-`.
pub fn member_id(&self) -> String {
let mut o = DOMAIN_MEMBER.to_vec();
put(&mut o, self.backend.as_bytes());
put(&mut o, self.name.as_bytes());
o.extend_from_slice(&self.vendor.to_be_bytes());
o.extend_from_slice(&self.device.to_be_bytes());
for s in [&self.device_type, &self.pci_bus, &self.driver, &self.driver_info, &self.translator] {
put(&mut o, s.as_bytes());
}
o.extend_from_slice(&self.max_storage_binding.to_be_bytes());
o.extend_from_slice(&self.max_buffer.to_be_bytes());
let d12 = hex(&blake3::hash(&o).as_bytes()[..6]);
format!("local/{}-{}-{d12}", self.prefix(), self.backend)
}
/// The key a runner declares in CI: `"<gpu|swgpu>-<backend>-<vendor in
/// lowercase hex>"`, for instance `"gpu-metal-0"` or `"swgpu-vulkan-10005"`.
pub fn key(&self) -> String {
format!("{}-{}-{:x}", self.prefix(), self.backend, self.vendor)
}
fn prefix(&self) -> &'static str {
if self.software() { "swgpu" } else { "gpu" }
}
}
/// How a device that runs a mutant must be caught.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub enum Detector {
/// The arithmetic probe's outputs differ from the CPU's on the probe edges.
Probe,
/// The corpus digest differs from the pinned one.
Certificate,
/// A result that left 32 bits is not reported.
Range,
}
/// A textual edit of [`SHADER`] that a conforming device must be caught
/// running.
#[derive(Clone, Copy, Debug)]
pub struct Mutant {
pub name: &'static str,
pub find: &'static str,
pub replace: &'static str,
pub detector: Detector,
}
/// The mutants. Each `find` occurs exactly once in [`SHADER`]. `mid_carry` is
/// always zero, so an edit of it would change nothing and is not among them.
pub const MUTANTS: [Mutant; 7] = [
// fx64 takes the wrong 32 bits of the rounded product.
Mutant { name: "M1", find: "let lo = (lo2 >> 30u) | (hi2 << 2u);", replace: "let lo = (lo2 >> 29u) | (hi2 << 3u);", detector: Detector::Probe },
// mul_wide never negates.
Mutant { name: "M2", find: "let negative = neg(a) != neg(b);", replace: "let negative = false;", detector: Detector::Probe },
// The rounding constant is one short.
Mutant { name: "M3", find: "const HALF: u32 = 0x20000000u;", replace: "const HALF: u32 = 0x1fffffffu;", detector: Detector::Probe },
// The carry out of the low word is dropped.
Mutant { name: "M4", find: "let lo_carry = bit(lo < p00);", replace: "let lo_carry = 0u;", detector: Detector::Probe },
// The control test is inverted.
Mutant { name: "M5", find: "if ((i & g.ctrl) != g.ctrl) { return; }", replace: "if ((i & g.ctrl) == g.ctrl) { return; }", detector: Detector::Certificate },
// The range guard never fires.
Mutant { name: "M6", find: "fn fits(v: vec2<u32>) -> bool { return v.x == select(0u, ONES, neg(v.y)); }", replace: "fn fits(v: vec2<u32>) -> bool { return true; }", detector: Detector::Range },
// The arithmetic shift by 30 loses the sign.
Mutant { name: "M7", find: "const EXT30: u32 = 0xfffffffcu;", replace: "const EXT30: u32 = 0x00000000u;", detector: Detector::Probe },
];
/// [`SHADER`] with `m` applied. Panics unless `m.find` occurs exactly once.
pub fn mutate(m: &Mutant) -> String {
let n = SHADER.matches(m.find).count();
assert_eq!(n, 1, "mutant {}: `{}` occurs {n} times in the kernel", m.name, m.find);
SHADER.replacen(m.find, m.replace, 1)
}
#[cfg(test)]
mod tests {
use super::*;
use crate::quantum::{BaseGate, Gate, Gateset};
use crate::quantum_fabric::lanes::{self, Lane};
use crate::quantum_fabric::{Certificate, Corpus, FABRIC_DEFAULT, KernelEntry, Member, Parity, Presence, apply_kernel, detect_local, place};
/// Changes only with a reviewed edit to the kernel and to the argument in
/// this module's documentation (or to the host layout, with
/// [`HOST_LAYOUT`]): every certificate of the old id must then be earned
/// again.
const PINNED_IMPLEMENTATION: &str = "wgsl-lanes32/dc8e19f61e0e87ab";
#[test]
fn implementation_id_is_pinned() {
let id = implementation_id();
assert_eq!(id, PINNED_IMPLEMENTATION);
let digest = id.strip_prefix("wgsl-lanes32/").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_eq!(implementation_id(), id);
// The kernel is stored with LF line endings.
assert!(!SHADER.contains('\r'));
}
#[test]
fn gate_record_layout() {
let w = GateWords { m: [1, 2, 3, 4, 5, 6, 7, 8], tbit: 4, ctrl: 3 };
let r = gate_record(&w, 16, 64, 9);
let word = |k: usize| u32::from_le_bytes(r[4 * k..4 * k + 4].try_into().unwrap());
for k in 0..8 {
assert_eq!(r[4 * k..4 * k + 4], (k as u32 + 1).to_le_bytes(), "m[{k}]");
}
assert_eq!((word(8), word(9), word(10), word(11), word(12)), (4, 3, 16, 64, 9));
assert_eq!(r[52..64], [0; 12]);
// Little-endian: the low byte first.
let r = gate_record(&GateWords { m: [0x8000_0001, 0, 0, 0, 0, 0, 0, 0], tbit: 1, ctrl: 0 }, 1, 64, 0x0102_0304);
assert_eq!(r[0..4], [0x01, 0x00, 0x00, 0x80]);
assert_eq!(r[48..52], [0x04, 0x03, 0x02, 0x01]);
}
#[test]
fn grids_cover_every_item_once() {
for n in 1..=26u32 {
let items = 1u32 << (n - 1);
for max_x in [1, 3, X_GROUPS_MAX] {
let g = grid(items, max_x);
assert_eq!((g.items, g.row_stride), (items, WORKGROUP * g.x));
assert!(g.x >= 1 && g.x <= max_x && g.y >= 1, "{g:?}");
let per_row = u64::from(g.x) * u64::from(WORKGROUP);
assert!(per_row * u64::from(g.y) >= u64::from(items), "n = {n}, {g:?}");
assert!(per_row * u64::from(g.y - 1) < u64::from(items), "n = {n}, {g:?}");
if max_x == X_GROUPS_MAX {
assert!(g.y <= 65535, "n = {n}, {g:?}");
assert_eq!(gate_grid(n as u8), g);
}
if n <= 12 {
let mut hits = vec![0u32; items as usize];
for gy in 0..g.y {
for gx in 0..g.row_stride {
let p = gy * g.row_stride + gx;
if p < items {
hits[p as usize] += 1;
}
}
}
assert!(hits.iter().all(|&h| h == 1), "n = {n}, {g:?}");
}
}
// The dispatch shapes fabric-v1 is checked to cover: up to 6
// qubits part of one workgroup, 7 to 13 one row, 14 or more a grid.
let g = gate_grid(n as u8);
assert_eq!(n <= 6, items < WORKGROUP, "n = {n}");
assert_eq!((7..=13).contains(&n), items >= WORKGROUP && g.y == 1, "n = {n}");
assert_eq!(n >= 14, g.y > 1, "n = {n}");
}
}
/// The kernel's index map: `p` with a zero inserted at the target bit.
fn insert0(p: u32, tbit: u32) -> u32 {
let low = p & (tbit - 1);
((p ^ low) << 1) | low
}
#[test]
fn the_index_map_is_a_bijection() {
for n in 1..=12u32 {
for t in 0..n {
let tbit = 1u32 << t;
let mapped: Vec<u32> = (0..1u32 << (n - 1)).map(|p| insert0(p, tbit)).collect();
let clear: Vec<u32> = (0..1u32 << n).filter(|i| i & tbit == 0).collect();
// Strictly increasing and onto the indices with bit t clear.
assert_eq!(mapped, clear, "n = {n}, t = {t}");
}
}
}
fn hardware() -> AdapterIdentity {
AdapterIdentity {
backend: "metal".into(),
name: "Test GPU".into(),
vendor: 0,
device: 0,
device_type: "integrated".into(),
pci_bus: String::new(),
driver: String::new(),
driver_info: String::new(),
translator: "naga".into(),
max_storage_binding: 1 << 28,
max_buffer: 1 << 28,
}
}
fn software() -> AdapterIdentity {
AdapterIdentity {
backend: "vulkan".into(),
name: "Test software rasterizer".into(),
vendor: 0x10005,
device: 0,
device_type: "cpu".into(),
pci_bus: String::new(),
driver: "test-driver".into(),
driver_info: "1.0".into(),
translator: "naga".into(),
max_storage_binding: 1 << 27,
max_buffer: 1 << 28,
}
}
/// The ids of [`hardware`] and [`software`]. They change only with a
/// reviewed change to what a member id hashes; every certificate earned
/// under an old id must then be earned again.
const PINNED_HARDWARE: &str = "local/gpu-metal-a47619e299a8";
const PINNED_SOFTWARE: &str = "local/swgpu-vulkan-917c9bb1d948";
#[test]
fn member_ids_are_pinned() {
let (hw, sw) = (hardware(), software());
assert_eq!((hw.member_id(), sw.member_id()), (PINNED_HARDWARE.to_string(), PINNED_SOFTWARE.to_string()));
assert_eq!((hw.class(), sw.class()), (FabricClass::Gpu, FabricClass::Cpu));
assert_eq!((hw.key(), sw.key()), ("gpu-metal-0".to_string(), "swgpu-vulkan-10005".to_string()));
// Every field is in the id.
let edits: [fn(&mut AdapterIdentity); 11] = [
|a| a.backend.push('x'),
|a| a.name.push('x'),
|a| a.vendor ^= 1,
|a| a.device ^= 1,
|a| a.device_type = if a.device_type == "cpu" { "virtual".into() } else { "discrete".into() },
|a| a.pci_bus.push('0'),
|a| a.driver.push('x'),
|a| a.driver_info.push('x'),
|a| a.translator.push_str("+bytecode"),
|a| a.max_storage_binding += 1,
|a| a.max_buffer += 1,
];
for base in [hw, sw] {
for (k, edit) in edits.iter().enumerate() {
let mut a = base.clone();
edit(&mut a);
assert_ne!(a.member_id(), base.member_id(), "field {k} of {base:?}");
}
}
// Fields do not run into each other.
let mut a = hardware();
let mut b = hardware();
(a.driver, a.driver_info) = ("ab".into(), String::new());
(b.driver, b.driver_info) = ("a".into(), "b".into());
assert_ne!(a.member_id(), b.member_id());
}
#[test]
fn the_cpu_reference_is_placed_first_when_nothing_is_metered() {
for arch in ["aarch64", "x86_64", "x86", "wasm32", "riscv64", "arm", "s390x", "powerpc64", "loongarch64"] {
let cpu = format!("local/cpu-{arch}");
assert!(cpu.as_str() < PINNED_HARDWARE && cpu.as_str() < PINNED_SOFTWARE, "{cpu}");
}
let local = detect_local();
for gs in Gateset::ALL {
let kernel = apply_kernel(gs);
let member = |a: &AdapterIdentity| {
let id = a.member_id();
let certificate = Certificate {
kernel: kernel.into(),
implementation: implementation_id(),
member: id.clone(),
generator: FABRIC_DEFAULT.generator(gs),
parity: Parity::BitExact,
cases: 224,
corpus: [7; 32],
};
let entry = KernelEntry { kernel: kernel.into(), implementation: implementation_id(), certificate, energy: None };
Member { id, class: a.class(), presence: Presence::Detected, kernels: vec![entry], detail: String::new() }
};
let (gpu, swgpu) = (member(&hardware()), member(&software()));
// Each member is placeable on its own.
for m in [&gpu, &swgpu] {
let (placed, k) = place(std::slice::from_ref(m), kernel, &Parity::BitExact).unwrap();
assert_eq!((placed.id.as_str(), k.implementation.as_str()), (m.id.as_str(), implementation_id().as_str()));
}
// Beside the CPU, with no energy figure anywhere, the reference stays.
for members in [vec![local.clone(), gpu.clone()], vec![gpu.clone(), swgpu.clone(), local.clone()]] {
let (m, k) = place(&members, kernel, &Parity::BitExact).unwrap();
assert_eq!((m.id.as_str(), k.implementation.as_str()), (local.id.as_str(), crate::quantum_fabric::REFERENCE), "{gs:?}");
}
}
}
#[test]
fn mutants_apply_once() {
let names: Vec<&str> = MUTANTS.iter().map(|m| m.name).collect();
assert_eq!(names, ["M1", "M2", "M3", "M4", "M5", "M6", "M7"]);
for m in &MUTANTS {
let text = mutate(m);
assert_ne!(text, SHADER, "{}", m.name);
assert_eq!(text.matches(m.replace).count(), 1, "{}", m.name);
assert_eq!(text.len() as isize - SHADER.len() as isize, m.replace.len() as isize - m.find.len() as isize);
}
let detectors: Vec<Detector> = MUTANTS.iter().map(|m| m.detector).collect();
use Detector::*;
assert_eq!(detectors, [Probe, Probe, Probe, Probe, Certificate, Range, Probe]);
}
/// The kernel's `mul_wide` and `fx64` in Rust, line for line, with one of
/// the probe-detected mutants applied (or none): `[hi, lo, fx_hi, fx_lo]`,
/// what the probe entry point writes.
fn probe_mirror(a: u32, b: u32, mutant: &str) -> [u32; 4] {
let neg = |x: u32| x & 0x8000_0000 != 0;
let bit = |c: bool| u32::from(c);
let half: u32 = if mutant == "M3" { 0x1fff_ffff } else { 0x2000_0000 };
let ext30: u32 = if mutant == "M7" { 0 } else { 0xffff_fffc };
// mul_wide
let negative = if mutant == "M2" { false } else { neg(a) != neg(b) };
let x = if neg(a) { 0u32.wrapping_sub(a) } else { a };
let y = if neg(b) { 0u32.wrapping_sub(b) } else { b };
let (x1, x0, y1, y0) = (x >> 16, x & 0xffff, y >> 16, y & 0xffff);
let (p00, p01, p10, p11) = (x0 * y0, x0 * y1, x1 * y0, x1 * y1);
let mid = p01.wrapping_add(p10);
let mid_carry = bit(mid < p01);
let lo = p00.wrapping_add(mid << 16);
let lo_carry = if mutant == "M4" { 0 } else { bit(lo < p00) };
let hi = p11.wrapping_add(mid >> 16).wrapping_add(mid_carry << 16).wrapping_add(lo_carry);
let (hi, lo) = if negative {
let lo_n = (!lo).wrapping_add(1);
((!hi).wrapping_add(bit(lo_n == 0)), lo_n)
} else {
(hi, lo)
};
// fx64
let lo2 = lo.wrapping_add(half);
let hi2 = hi.wrapping_add(bit(lo2 < lo));
let fx_lo = if mutant == "M1" { (lo2 >> 29) | (hi2 << 3) } else { (lo2 >> 30) | (hi2 << 2) };
let fx_hi = (hi2 >> 30) | if neg(hi2) { ext30 } else { 0 };
[hi, lo, fx_hi, fx_lo]
}
fn probe_reference(a: i32, b: i32) -> [u32; 4] {
let ((hi, lo), (fx_hi, fx_lo)) = (lanes::mul_wide(a, b), lanes::fxmul_wide(a, b));
[hi, lo, fx_hi, fx_lo]
}
/// `apply_words` with the control test inverted (M5): a pair is updated
/// when its control bits are not all set.
fn apply_inverted_control(amps: &mut [Lane], w: &GateWords) {
let fx = |a: i32, b: i32| {
let (hi, lo) = lanes::fxmul_wide(a, b);
((u64::from(hi) << 32) | u64::from(lo)) as i64
};
let e = |k: usize| w.m[k] as i32;
let (tbit, ctrl) = (w.tbit as usize, w.ctrl as usize);
for i in 0..amps.len() {
if i & tbit == 0 && i & ctrl != ctrl {
let j = i | tbit;
let (a0, a1) = (amps[i], amps[j]);
let row = |r: usize| -> Lane {
let (mr, mi, nr, ni) = (e(4 * r), e(4 * r + 1), e(4 * r + 2), e(4 * r + 3));
[
(fx(mr, a0[0]) - fx(mi, a0[1]) + fx(nr, a1[0]) - fx(ni, a1[1])) as i32,
(fx(mr, a0[1]) + fx(mi, a0[0]) + fx(nr, a1[1]) + fx(ni, a1[0])) as i32,
]
};
(amps[i], amps[j]) = (row(0), row(1));
}
}
}
#[test]
fn mutant_witnesses_exist() {
let edges = lanes::PROBE_EDGES;
let pairs: Vec<(i32, i32)> = edges.iter().flat_map(|&a| edges.iter().map(move |&b| (a, b))).collect();
// The mirror without a mutant is the CPU's arithmetic, on the edges
// and across the whole 32-bit range.
let mut s = 17u64;
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 random: Vec<(i32, i32)> = (0..20_000).map(|_| (next() as i32, next() as i32)).collect();
for &(a, b) in pairs.iter().chain(&random) {
assert_eq!(probe_mirror(a as u32, b as u32, ""), probe_reference(a, b), "{a} × {b}");
}
// Each probe-detected mutant changes the probe's output on some pair
// of edges. The counts are the ones computed, independently of this
// crate, when the mutants were chosen.
for (name, differing) in [("M1", 436), ("M2", 220), ("M3", 80), ("M4", 48), ("M7", 184)] {
let m = MUTANTS.iter().find(|m| m.name == name).unwrap();
assert_eq!(m.detector, Detector::Probe);
let count = pairs.iter().filter(|&&(a, b)| probe_mirror(a as u32, b as u32, name) != probe_reference(a, b)).count();
assert!(count >= 1, "{name}");
assert_eq!(count, differing, "{name}");
}
// M5 changes the output of a random-v1 circuit, so the certificate's
// digest moves.
let mut changed = 0;
for c in (Corpus::Random { seed: 7, cases: 200 }).circuits(Gateset::V1) {
let mut amps = vec![[0i32; 2]; 1 << c.n_qubits];
amps[0] = [crate::quantum::ONE as i32, 0];
for g in &c.ops {
apply_inverted_control(&mut amps, &lanes::gate_words(g, c.gateset));
}
if lanes::output(&amps) != lanes::reference_output(&c).unwrap() {
changed += 1;
}
}
assert!(changed >= 1);
// M6 is witnessed by `z_on_the_most_negative_lane_is_flagged`: a
// correct guard raises the flag there, and M6's never does.
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));
}
}

1047
wai/wai-quantum/src/quantum_fabric_webgpu.rs Normal file
View file

File diff suppressed because it is too large Load diff

543
wai/wai-quantum/tests/webgpu_member.rs Normal file
View file

@ -0,0 +1,543 @@
//! The fabric's WebGPU member on the adapters of the machine that runs this.
//!
//! The machine declares the adapters it must have in
//! `WAI_QUANTUM_WEBGPU_EXPECT`: their keys, comma-separated and sorted, or
//! `none`. Every test that needs a member fails unless the adapters found are
//! exactly the declared ones, so a machine whose GPU went missing, or a run
//! that declared nothing, fails instead of passing over no member. Each test
//! prints what it established and, when `WAI_QUANTUM_WEBGPU_REPORT` names a
//! file, appends the same lines to it, which a CI step can count per adapter.
//!
//! ```text
//! WAI_QUANTUM_WEBGPU_EXPECT=gpu-metal-0 \
//! cargo test --features quantum_fabric_webgpu --test webgpu_member -- --test-threads=1
//! ```
//!
//! `WAI_QUANTUM_WEBGPU_WIDE=1` adds the widest states a GPU member holds.
#![cfg(not(target_arch = "wasm32"))]
use std::future::Future;
use std::io::Write;
use std::path::Path;
use std::sync::OnceLock;
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::{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};
/// The CPU's corpus digests: random-v1 (seed 7, 200 cases) and fabric-v1
/// (seed 7, 200 + 24 cases), per gateset. A bit-exact member reproduces them.
const PINS: [(Gateset, &str, &str); 2] = [
(Gateset::V1, "d18454c639acfd9c7b66450f4ded1aac4ab46cc10a7bedbcf4958a96e204fb36", "41c726a0ac80f7defef92e15ecf893bdb7a3076b5e7f670fa0b7142e5c033df7"),
(Gateset::V2, "3dbde276c1095bf827faec71d5980e050f8a55a5ba7d97cd3bbdaec3f43a13b2", "9a84b55af68f1a8b07390b234cf562dd8d85cef3185a697ab85690c8790a7809"),
];
const RANDOM_V1: Corpus = Corpus::Random { seed: 7, cases: 200 };
static DISCOVERY: OnceLock<Discovery> = OnceLock::new();
fn discovery() -> &'static Discovery {
DISCOVERY.get_or_init(|| blocking::discover().expect("discovery resolves in one poll"))
}
/// Prints `line` and appends it to the file `WAI_QUANTUM_WEBGPU_REPORT` names.
fn report(line: &str) {
println!("{line}");
if let Some(path) = std::env::var_os("WAI_QUANTUM_WEBGPU_REPORT") {
let mut f = std::fs::OpenOptions::new().create(true).append(true).open(&path).expect("the report file opens");
writeln!(f, "{line}").expect("the report line is written");
}
}
/// The declared adapter set, as written, if it is exactly what was found;
/// otherwise why not, with what was found and refused and how to declare it.
fn declared(d: &Discovery) -> Result<String, String> {
let mut found: Vec<&str> = d.found.iter().map(|m| m.key()).collect();
found.sort_unstable();
let found_list = if found.is_empty() { "none".to_string() } else { found.join(",") };
let refused: Vec<String> = d.refused.iter().map(|r| format!("{} ({:?}): {}", r.key, r.adapter, r.error)).collect();
let refused_list = if refused.is_empty() { "none".to_string() } else { refused.join("; ") };
let how = format!("found: {found_list}; refused: {refused_list}. Declare the adapters this machine must have: WAI_QUANTUM_WEBGPU_EXPECT={found_list}");
let expect = match std::env::var("WAI_QUANTUM_WEBGPU_EXPECT") {
Ok(v) if !v.is_empty() => v,
_ => return Err(format!("WAI_QUANTUM_WEBGPU_EXPECT is not set. {how}")),
};
let keys: Vec<&str> = if expect == "none" { Vec::new() } else { expect.split(',').collect() };
let well_formed = keys.windows(2).all(|w| w[0] <= w[1]) && keys.iter().all(|k| !k.is_empty() && *k != "none" && k.trim() == *k);
if !well_formed {
return Err(format!("WAI_QUANTUM_WEBGPU_EXPECT={expect:?} is not a sorted, comma-separated list of keys, nor none. {how}"));
}
for k in &keys {
let (want, have) = (keys.iter().filter(|x| *x == k).count(), found.iter().filter(|x| *x == k).count());
if have < want
&& let Some(r) = d.refused.iter().find(|r| r.key == *k)
{
return Err(format!("the declared adapter {k} was refused: {}. {how}", r.error));
}
}
if found != keys {
return Err(format!("the adapters found are not the declared ones ({expect}). {how}"));
}
Ok(expect)
}
/// The members, once the declared adapters are exactly the ones found.
fn members() -> &'static [WebGpuMember] {
let d = discovery();
if let Err(why) = declared(d) {
panic!("{why}");
}
&d.found
}
/// A member call's result; every call resolves in one poll.
fn attempt<T>(f: impl Future<Output = Result<T, WebGpuError>>) -> Result<T, WebGpuError> {
blocking::resolve(f).expect("a native call resolves in one poll")
}
/// A member call that must succeed.
fn run<T>(f: impl Future<Output = Result<T, WebGpuError>>) -> T {
attempt(f).unwrap_or_else(|e| panic!("the member refused: {e}"))
}
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)
}
/// A value in `[-bound, bound]`.
fn within(s: &mut u64, bound: i64) -> i32 {
((next(s) % (2 * bound as u64 + 1)) as i64 - bound) as i32
}
/// A random circuit of `ops` gates on `n` qubits with the wide corpus's gate
/// mix: 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
}
/// Five circuits the native simulator refuses, in each gateset.
fn refused_circuits() -> Vec<Circuit> {
let gate = |base, controls: Vec<u8>, target, param| Gate { base, controls, target, param };
let bad = [
(1, vec![gate(BaseGate::H, vec![], 0, 0), gate(BaseGate::P, vec![], 0, 0)]),
(1, vec![gate(BaseGate::P, vec![], 0, 33)]),
(1, vec![gate(BaseGate::X, vec![], 3, 0)]),
(2, vec![gate(BaseGate::X, vec![0], 0, 0)]),
(2, vec![gate(BaseGate::X, vec![5], 1, 0)]),
];
let mut out = Vec::new();
for gs in Gateset::ALL {
for (n, ops) in bad.clone() {
let mut c = Circuit::with_gateset(n, gs);
c.ops = ops;
out.push(c);
}
}
out
}
/// `Z` on qubit 0: on `[(0, 0), (i32::MIN, 0)]` its exact result is `2^31`,
/// which leaves 32 bits.
fn z_words() -> GateWords {
lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 0, param: 0 }, Gateset::V2)
}
const Z_STATE: [Lane; 2] = [[0, 0], [i32::MIN, 0]];
#[test]
fn g0_adapters_are_declared() {
let d = discovery();
for r in &d.refused {
println!("webgpu refused: key={} adapter={:?}: {}", r.key, r.adapter, r.error);
}
let expect = declared(d).unwrap_or_else(|why| panic!("{why}"));
for m in &d.found {
assert_eq!(m.implementation(), wgsl::implementation_id());
assert_eq!((m.identity().member_id(), m.identity().key()), (m.id().to_string(), m.key().to_string()));
assert_eq!(m.class(), m.identity().class());
let prefix = if m.class() == FabricClass::Cpu { "local/swgpu-" } else { "local/gpu-" };
assert!(m.id().starts_with(prefix) && m.id() > "local/cpu-zzzz", "{}", m.id());
assert!((25..=MAX_QUBITS).contains(&m.max_qubits()), "{}: {}", m.key(), m.max_qubits());
assert!(m.detail().ends_with("probe passed in this process"), "{}", m.detail());
}
report(&format!("webgpu adapters: as declared ({expect})"));
for m in &d.found {
report(&format!(
"webgpu member: key={} id={} class={} backend={} max_qubits={} implementation={}",
m.key(),
m.id(),
m.class().label(),
m.backend(),
m.max_qubits(),
m.implementation()
));
}
}
#[test]
fn g1_probe_matches_lanes() {
let pairs = probe_pairs(1 << 20, 11);
let expected: Vec<((u32, u32), (u32, u32))> = pairs.iter().map(|&(a, b)| (lanes::mul_wide(a, b), lanes::fxmul_wide(a, b))).collect();
for m in members() {
// A slot the device never wrote keeps its sentinel, which no correct
// result can equal under both.
for sentinel in [0xA5A5_A5A5u32, 0x5A5A_5A5A] {
let got = run(m.probe_with_sentinel(&pairs, sentinel));
assert_eq!(got.len(), pairs.len(), "{}", m.key());
if let Some(k) = (0..pairs.len()).find(|&k| got[k] != expected[k]) {
panic!("{}, sentinel {sentinel:#x}: {} × {} gives {:x?}, the CPU {:x?}", m.key(), pairs[k].0, pairs[k].1, got[k], expected[k]);
}
}
report(&format!("webgpu probe: key={} ok", m.key()));
}
}
#[test]
fn g2_kernel_matches_run_words() {
for m in members() {
let mut s = 0x6b65_726e_656c_u64;
let (mut flagged, mut clean) = (0, 0);
for case in 0..2000 {
let n = 1 + (next(&mut s) % 12) as u32;
// Three regimes: the whole 32-bit range (nearly always out of
// range), small states and entries (never), and states and
// entries up to 2^30 (either).
let mode = case % 3;
let value = |s: &mut u64| -> i32 {
match mode {
0 => next(s) as i32,
1 => within(s, 1 << 22),
_ => within(s, 1 << 30),
}
};
let entry = |s: &mut u64| -> u32 {
match mode {
0 => next(s) as u32,
1 => within(s, 1 << 29) as u32,
_ => within(s, 1 << 30) 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) % 8)
.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));
GateWords { m: w, tbit: 1 << t, ctrl }
})
.collect();
let want = lanes::run_words(&state, &gates).expect("valid shapes");
let got = run(m.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());
if want.out_of_range.is_some() {
flagged += 1;
} else {
clean += 1;
}
}
assert!(flagged >= 100 && clean >= 100, "{}: {flagged} flagged, {clean} clean", m.key());
report(&format!("webgpu kernel: key={} ok", m.key()));
}
}
#[test]
fn g3_certificates_reproduce_the_pins() {
let ms = members();
for (gs, random_pin, fabric_pin) in PINS {
let cpu = lanes::certify_corpus(gs, "local/test", RANDOM_V1);
assert_eq!(hex(&cpu.corpus), random_pin, "{gs:?}");
let other = Corpus::Random { seed: 8, cases: 300 };
let cpu_other = lanes::certify_corpus(gs, "local/test", other);
for m in ms {
for (corpus, cases, pin) in [(RANDOM_V1, 200, random_pin), (FABRIC_DEFAULT, 224, fabric_pin)] {
let c = blocking::certify(m, gs, corpus).unwrap_or_else(|e| panic!("{} {gs:?} {}: {e}", m.key(), corpus.name()));
assert_eq!(c.parity, Parity::BitExact, "{} {gs:?} {}", m.key(), corpus.name());
assert_eq!(c.cases, cases);
assert_eq!(hex(&c.corpus), pin, "{} {gs:?} {}", m.key(), corpus.name());
assert_eq!((c.kernel.as_str(), c.implementation.as_str(), c.member.as_str()), (apply_kernel(gs), m.implementation(), m.id()));
assert_eq!(c.generator, corpus.generator(gs));
if corpus == RANDOM_V1 {
assert_eq!(c.corpus, cpu.corpus);
}
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)
));
}
// A corpus with no pin: the member and the CPU agree on it.
let c = blocking::certify(m, gs, other).unwrap_or_else(|e| panic!("{} {gs:?} seed 8: {e}", m.key()));
assert_eq!((c.parity, c.cases), (Parity::BitExact, 300), "{} {gs:?}", m.key());
assert_eq!(c.corpus, cpu_other.corpus, "{} {gs:?}", m.key());
}
}
}
#[test]
fn g4_forced_geometry_gives_the_same_digests() {
for m in members() {
// One workgroup to a row (many rows), and three (a ragged last row).
for max_x in [1, 3] {
let forced = m.with_max_x(max_x);
for (gs, pin, _) in PINS {
let c = run(forced.certify(gs, RANDOM_V1));
assert_eq!(c.parity, Parity::BitExact, "{} max_x {max_x} {gs:?}", m.key());
assert_eq!(hex(&c.corpus), pin, "{} max_x {max_x} {gs:?}", m.key());
}
}
// Submissions and readback in small pieces. A ring of three records,
// rewritten between submissions of at most three gates, and state
// planes read back twelve bytes at a time: the smallest plane (two
// qubits) is sixteen bytes, so every plane takes more than one piece,
// the last one short.
let small = m.with_ring_records(3).with_readback_chunk(12);
for (gs, pin, _) in PINS {
let c = run(small.certify(gs, RANDOM_V1));
assert_eq!(c.parity, Parity::BitExact, "{} small pieces {gs:?}", m.key());
assert_eq!(hex(&c.corpus), pin, "{} small pieces {gs:?}", m.key());
}
// fabric-v1 through a ring of two, read back in pieces of 12 KiB and
// four bytes, so that the wide circuits' planes come back in many
// pieces, the last one short.
let pieces = m.with_ring_records(2).with_readback_chunk(12 * 1024 + 4);
for (gs, _, pin) in PINS {
let c = run(pieces.certify(gs, FABRIC_DEFAULT));
assert_eq!(c.parity, Parity::BitExact, "{} fabric-v1 in pieces {gs:?}", m.key());
assert_eq!(hex(&c.corpus), pin, "{} fabric-v1 in pieces {gs:?}", m.key());
}
report(&format!("webgpu geometry: key={} ok", m.key()));
}
}
#[test]
fn g5_conformance_vectors() {
let root = Path::new(env!("CARGO_MANIFEST_DIR")).join("../quantum-conformance");
for m in members() {
let mut total = 0;
for (dir, manifest) in [("vectors", "expected_hashes.json"), ("v2/vectors", "v2/expected_hashes.json")] {
let text = std::fs::read_to_string(root.join(manifest)).expect("the manifest reads");
let manifest: serde_json::Value = serde_json::from_str(&text).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 entry = manifest.get(&name).unwrap_or_else(|| panic!("{dir}/{name} is not in its manifest"));
let (c, _) = from_wqc(&std::fs::read(&path).unwrap()).unwrap_or_else(|e| panic!("{dir}/{name}: {e}"));
let state = run(m.simulate(&c));
assert_eq!(Some(hex(&state.statevector_hash()).as_str()), entry["statevector_hash"].as_str(), "{} {dir}/{name}", m.key());
if let Some(measure) = entry.get("measure") {
let (seed, shots) = (measure["seed"].as_u64().unwrap(), measure["shots"].as_u64().unwrap());
assert_eq!(Some(hex(&state.histogram_hash(seed, shots)).as_str()), measure["histogram_hash"].as_str(), "{} {dir}/{name} histogram", m.key());
}
total += 1;
}
}
report(&format!("webgpu vectors: key={} {total} ok", m.key()));
}
}
#[test]
fn g6_wide_circuits_match_native() {
let wide = std::env::var("WAI_QUANTUM_WEBGPU_WIDE").is_ok_and(|v| v == "1");
for m in members() {
let mut sizes = vec![12u8, 16, 20];
if wide {
if m.class() == FabricClass::Gpu {
sizes.push(24);
if m.max_qubits() > 24 {
sizes.push(m.max_qubits());
}
} else {
report(&format!("webgpu wide: key={} software adapter: wide sizes not run", m.key()));
}
}
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, start) in [(&qft, 0usize), (&random, 0), (&random, (1usize << n) + 5)] {
let native = c.simulate_from(start).unwrap().statevector_hash();
let member = blocking::simulate_from(m, c, start).unwrap_or_else(|e| panic!("{} n = {n}: {e}", m.key()));
assert_eq!(hex(&member.statevector_hash()), hex(&native), "{} n = {n} start {start}", m.key());
}
report(&format!("webgpu wide: key={} n={n} matches native", m.key()));
}
// Long circuits on few qubits: more gates than one submission holds,
// so the ring is rewritten between submissions; and the same through a
// ring of three records.
for gs in Gateset::ALL {
let long = random_circuit(4, 5000, 0x6c6f_6e67, gs);
let native = hex(&long.simulate().unwrap().statevector_hash());
for member in [m.clone(), m.with_ring_records(3)] {
let got = blocking::simulate_from(&member, &long, 0).unwrap_or_else(|e| panic!("{} long {gs:?}: {e}", m.key()));
assert_eq!(hex(&got.statevector_hash()), native, "{} long {gs:?} {member:?}", 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, not in its submission.
let identity = lanes::gate_words(&Gate { base: BaseGate::I, controls: vec![], target: 0, param: 0 }, Gateset::V2);
for (member, first, second) in [(m.clone(), 4100usize, 4500usize), (m.with_ring_records(3), 4, 7)] {
let gates: Vec<GateWords> = (0..5000).map(|k| if k == first || k == second { z_words() } else { identity }).collect();
let want = lanes::run_words(&Z_STATE, &gates).expect("valid shapes");
assert_eq!(want.out_of_range, Some(first as u32));
assert_eq!(run(member.run_lanes(&Z_STATE, &gates)), want, "{} first out of range at {first}", m.key());
}
report(&format!("webgpu long: key={} matches native", m.key()));
}
}
#[test]
fn g7_refusals_are_typed() {
for m in members() {
for c in refused_circuits() {
let native = c.simulate_from(0).unwrap_err();
assert_eq!(attempt(m.simulate_from(&c, 0)).unwrap_err(), WebGpuError::Circuit(native), "{c:?}");
}
assert_eq!(attempt(m.simulate(&Circuit::new(MAX_QUBITS + 1))).unwrap_err(), WebGpuError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1)));
if m.max_qubits() < MAX_QUBITS {
let n = m.max_qubits() + 1;
assert_eq!(attempt(m.simulate(&Circuit::new(n))).unwrap_err(), WebGpuError::TooWide { n_qubits: n, max: m.max_qubits() });
}
// Past the member's width, whatever the core allows.
let wide = Program { n_qubits: m.max_qubits() + 1, start: 0, gates: Vec::new(), steps: Vec::new() };
assert_eq!(attempt(m.run_program(&wide)).unwrap_err(), WebGpuError::TooWide { n_qubits: m.max_qubits() + 1, max: m.max_qubits() });
// A result that leaves 32 bits is reported at its op, and the lane
// holds its low 32 bits.
let r = run(m.run_lanes(&Z_STATE, &[z_words()]));
assert_eq!(r.out_of_range, Some(0), "{}", m.key());
assert_eq!(r.amps[1][0] as u32, 0x8000_0000);
assert_eq!(attempt(m.simulate_from(&Circuit::new(1), 0)).map(|s| s.amps.len()), Ok(2));
// A malformed record, and a state the host cannot lay out, are host
// defects, never results.
let bad = GateWords { tbit: 3, ..z_words() };
assert_eq!(attempt(m.run_lanes(&[[0, 0]; 4], &[bad])).unwrap_err(), WebGpuError::Internal(4));
assert_eq!(attempt(m.run_lanes(&[[0, 0]; 3], &[z_words()])).unwrap_err(), WebGpuError::Internal(2));
report(&format!("webgpu refusals: key={} typed", m.key()));
}
}
#[test]
fn g8_mutants_are_killed() {
let edges = probe_pairs(0, 0);
assert_eq!(edges.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);
killed += 1;
}
assert_eq!(killed, MUTANTS.len());
report(&format!("webgpu mutants: key={} killed {killed} of {}", m.key(), MUTANTS.len()));
}
}
#[test]
fn g9_receipt_and_placement() {
for m in members() {
let member = blocking::member(m, FABRIC_DEFAULT).unwrap_or_else(|e| panic!("{}: {e}", m.key()));
assert_eq!((member.id.as_str(), member.class, member.presence), (m.id(), m.class(), Presence::Detected));
assert_eq!(member.kernels.len(), 2);
for (entry, gs) in member.kernels.iter().zip(Gateset::ALL) {
assert_eq!((entry.kernel.as_str(), entry.implementation.as_str()), (apply_kernel(gs), m.implementation()));
assert!(entry.energy.is_none());
assert_eq!(entry.certificate.parity, Parity::BitExact, "{gs:?}");
assert_eq!(entry.certificate.member, m.id());
}
// A receipt over a member's run.
let qft = Circuit::qft(5);
let state = blocking::simulate_from(m, &qft, 0).unwrap();
assert_eq!(state, qft.simulate().unwrap());
let entry = member.kernels.iter().find(|k| k.kernel == apply_kernel(qft.gateset)).unwrap();
let signer = SigningKey::from_bytes(&[7; 32]);
let receipt = FabricReceipt::seal(&signer, "local/test", state.statevector_hash(), &member, entry, 0, None);
assert!(receipt.verify());
assert!(receipt.certificate_matches(&entry.certificate));
let mut moved = receipt.clone();
moved.member_id.push('x');
assert!(!moved.verify());
// Placement: beside the CPU reference, with no energy figure, the
// reference; alone, the member.
let local = detect_local();
for gs in Gateset::ALL {
let both = [local.clone(), member.clone()];
let (placed, k) = place(&both, apply_kernel(gs), &Parity::BitExact).unwrap();
assert_eq!((placed.id.as_str(), k.implementation.as_str()), (local.id.as_str(), REFERENCE));
let (placed, k) = place(std::slice::from_ref(&member), apply_kernel(gs), &Parity::BitExact).unwrap();
assert_eq!((placed.id.as_str(), k.implementation.as_str()), (m.id(), m.implementation()));
}
report(&format!("webgpu receipt: key={} ok", m.key()));
}
}
#[test]
fn g10_repeated_runs_agree() {
let c = random_circuit(18, 200, 0x7265_7065_6174, Gateset::V1);
let native = hex(&c.simulate().unwrap().statevector_hash());
for m in members() {
let hashes: Vec<String> = (0..5).map(|_| hex(&blocking::simulate_from(m, &c, 0).unwrap().statevector_hash())).collect();
assert!(hashes.iter().all(|h| *h == native), "{}: {hashes:?}, native {native}", m.key());
report(&format!("webgpu repeated: key={} 5 runs, one hash", m.key()));
}
}
#[test]
fn g11_native_futures_resolve_in_one_poll() {
// A future that is still pending after one poll is an error, never a spin.
assert_eq!(blocking::resolve(std::future::pending::<()>()), Err(WebGpuError::Pending));
assert_eq!(blocking::resolve(std::future::ready(7)), Ok(7));
// Every other test reached its results through `blocking::resolve`.
report("webgpu futures: one poll each");
}

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

@ -69,6 +69,12 @@ dependencies = [
"alloc-no-stdlib",
]
[[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"
@ -116,6 +122,15 @@ dependencies = [
"stable_deref_trait",
]
[[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 = "asn1-rs"
version = "0.7.2"
@ -290,6 +305,21 @@ dependencies = [
"smallvec",
]
[[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 = "1.3.2"
@ -370,6 +400,15 @@ dependencies = [
"hybrid-array",
]
[[package]]
name = "block2"
version = "0.6.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "cdeb9d870516001442e364c5220d3574d2da8dc765554b4a617230d33fa58ef5"
dependencies = [
"objc2",
]
[[package]]
name = "bon"
version = "3.10.0"
@ -476,6 +515,20 @@ 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 = "byteorder"
@ -666,6 +719,12 @@ 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 = "chrono"
version = "0.4.45"
@ -722,6 +781,17 @@ dependencies = [
"cc",
]
[[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 = "combine"
version = "4.6.8"
@ -1110,6 +1180,16 @@ dependencies = [
"crypto-common 0.2.2",
]
[[package]]
name = "dispatch2"
version = "0.3.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "1e0e367e4e7da84520dedcac1901e4da967309406d1e51017ae1abfb97adbd38"
dependencies = [
"bitflags 2.13.1",
"objc2",
]
[[package]]
name = "displaydoc"
version = "0.2.7"
@ -1121,6 +1201,15 @@ dependencies = [
"syn 3.0.4",
]
[[package]]
name = "document-features"
version = "0.2.12"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d4b8a88685455ed29a21542a33abd9cb6510b6b129abadabdcef0f4c55bc8f61"
dependencies = [
"litrs",
]
[[package]]
name = "dyn-clone"
version = "1.0.20"
@ -1299,6 +1388,12 @@ dependencies = [
"miniz_oxide 0.8.9",
]
[[package]]
name = "foldhash"
version = "0.2.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "77ce24cb58228fbb8aa041425bb1050850ac19177686ea6e0f41a70416f56fdb"
[[package]]
name = "form_urlencoded"
version = "1.2.2"
@ -1393,6 +1488,20 @@ version = "0.3.4"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e4eba85ea1d0a966a983acd07deee566e67395d2d96b6fb39e62b5a833f1eb0b"
[[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 2.0.20",
"windows",
]
[[package]]
name = "group"
version = "0.13.0"
@ -1412,6 +1521,7 @@ checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b"
dependencies = [
"cfg-if",
"crunchy",
"num-traits",
"zerocopy",
]
@ -1421,11 +1531,25 @@ version = "0.12.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "8a9ee70c43aaf417c914396645a0fa852624801b24ebb7ae78fe8272889ac888"
[[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"
@ -2005,6 +2129,16 @@ dependencies = [
"cc",
]
[[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"
@ -2023,6 +2157,21 @@ version = "0.8.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "47d9d19d1d6efa0109d2f65ff4c85cddd50bd572e5a00127ab10987290bcefae"
[[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"
@ -2117,6 +2266,44 @@ dependencies = [
"pxfm",
]
[[package]]
name = "naga"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "a616d2fb8c89516ac2723a581f69d6c18576046bed761bd6b305e5618e6ae130"
dependencies = [
"arrayvec",
"bit-set",
"bitflags 2.13.1",
"cfg-if",
"cfg_aliases",
"codespan-reporting",
"half",
"hashbrown 0.17.1",
"indexmap 2.14.0",
"libm",
"log",
"naga-types",
"num-traits",
"once_cell",
"rustc-hash",
"spirv",
"thiserror 2.0.20",
"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 2.14.0",
"rustc-hash",
"thiserror 2.0.20",
]
[[package]]
name = "nasm-rs"
version = "0.2.5"
@ -2293,6 +2480,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 2.13.1",
"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 2.13.1",
"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 2.13.1",
"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 2.13.1",
"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 2.13.1",
"objc2",
"objc2-core-foundation",
"objc2-core-graphics",
"objc2-foundation",
"objc2-metal",
]
[[package]]
name = "ogg"
version = "0.9.2"
@ -2332,6 +2605,15 @@ dependencies = [
"audiopus_sys",
]
[[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 = "ort"
version = "2.0.0-rc.13"
@ -2394,6 +2676,29 @@ dependencies = [
"sha2 0.10.9",
]
[[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 = "parsenic"
version = "0.2.1"
@ -2554,6 +2859,12 @@ dependencies = [
"zerocopy",
]
[[package]]
name = "presser"
version = "0.3.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e8cf8e6a8aa66ce33f63993ffc4ea4271eb5b0530a9002db8455ea6050c77bfa"
[[package]]
name = "prettyplease"
version = "0.3.0"
@ -2756,6 +3067,12 @@ dependencies = [
"rand_core 0.9.5",
]
[[package]]
name = "range-alloc"
version = "0.1.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ca45419789ae5a7899559e9512e58ca889e41f04f1f2445e9f4b290ceccd1d08"
[[package]]
name = "range-set"
version = "0.1.1"
@ -2935,6 +3252,24 @@ dependencies = [
"rgb",
]
[[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 = "rawpointer"
version = "0.2.1"
@ -2961,6 +3296,15 @@ dependencies = [
"crossbeam-utils",
]
[[package]]
name = "redox_syscall"
version = "0.5.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "ed2bf2547551a7053d6fdfafda3f938979645c44812fbfcda098faae3f1a362d"
dependencies = [
"bitflags 2.13.1",
]
[[package]]
name = "ref-cast"
version = "1.0.27"
@ -3010,6 +3354,12 @@ version = "0.8.11"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4"
[[package]]
name = "renderdoc-sys"
version = "1.1.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "19b30a45b0cd0bcca8037f3d0dc3421eaf95327a17cad11964fb8179b4fc4832"
[[package]]
name = "rfc6979"
version = "0.4.0"
@ -3067,6 +3417,12 @@ dependencies = [
"zeroize",
]
[[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"
@ -3226,6 +3582,12 @@ dependencies = [
"serde_json",
]
[[package]]
name = "scopeguard"
version = "1.2.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "94143f37725109f92c262ed2cf5e59bce7498c01bcc1502d7b9afe439a4e9f49"
[[package]]
name = "sec1"
version = "0.7.3"
@ -3524,6 +3886,15 @@ version = "0.9.9"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "3763264f6b73151db08c50ff20d7d8a0b8796e021cdea7ceedad07b80155fa0e"
[[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 2.13.1",
]
[[package]]
name = "spki"
version = "0.7.3"
@ -3690,6 +4061,15 @@ dependencies = [
"windows-sys 0.61.2",
]
[[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 = "1.0.69"
@ -3936,6 +4316,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 = "untrusted"
version = "0.9.0"
@ -4073,11 +4459,12 @@ dependencies = [
[[package]]
name = "wai-quantum"
version = "0.4.0"
version = "0.4.1"
dependencies = [
"blake3",
"ed25519-dalek",
"serde_json",
"wgpu",
]
[[package]]
@ -4198,6 +4585,159 @@ dependencies = [
"rustls-pki-types",
]
[[package]]
name = "wgpu"
version = "30.0.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "527ccdf43dd5b2e8676eed9984ce00e2bbb0a1b85b70c1969dcb6cd2eb55ab9e"
dependencies = [
"arrayvec",
"bitflags 2.13.1",
"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 2.13.1",
"bytemuck",
"cfg_aliases",
"document-features",
"hashbrown 0.17.1",
"indexmap 2.14.0",
"log",
"naga",
"naga-types",
"once_cell",
"parking_lot",
"portable-atomic",
"profiling",
"raw-window-handle",
"rustc-hash",
"smallvec",
"thiserror 2.0.20",
"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 2.13.1",
"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 2.0.20",
"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 2.13.1",
"bytemuck",
"js-sys",
"log",
"naga-types",
"raw-window-handle",
"static_assertions",
"web-sys",
]
[[package]]
name = "winapi"
version = "0.3.9"
@ -4229,6 +4769,27 @@ version = "0.4.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f"
[[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"
@ -4242,6 +4803,17 @@ dependencies = [
"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"
@ -4270,6 +4842,16 @@ 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"
@ -4322,6 +4904,15 @@ dependencies = [
"windows_x86_64_msvc",
]
[[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 = "windows_aarch64_gnullvm"
version = "0.52.6"

7
wai/wai-rs/Cargo.toml
View file

@ -100,7 +100,7 @@ curve25519-dalek = { version = "4", optional = true }
# ---- quantum (optional, behind the `quantum*` features) ----
# The quantum stack now lives in its own dependency-light crate so it can be used
# without the media/codec chain. `wai` re-exports it unchanged.
wai-quantum = { version = "0.4.0", path = "../wai-quantum", optional = true, default-features = false }
wai-quantum = { version = "0.4.1", path = "../wai-quantum", optional = true, default-features = false }
[features]
default = ["codecs"]
@ -408,6 +408,11 @@ quantum_qir = ["quantum_toolchain", "wai-quantum/quantum_qir"]
# The fabric: members, parity against the CPU reference, certification,
# placement by joules, and the core in 32-bit lanes.
quantum_fabric = ["quantum", "quantum_ops", "wai-quantum/quantum_fabric"]
# The fabric's WebGPU member: the 32-bit-lane kernel in WGSL through the
# platform's GPU interface, certified bit-exact per adapter, and
# `wai_quantum_vectors verify --webgpu`. Opt-in: it is the one feature that
# reaches a platform interface.
quantum_fabric_webgpu = ["quantum_fabric", "wai-quantum/quantum_fabric_webgpu"]
# Variational quantum machine learning (extensions/quantum-sim § VML,
# `wai.quantum.vml`): the interference-based ML the thermodynamic-computing pivot
# left behind — a Quantum Circuit Born Machine (P_θ(x)=|ψ_θ(x)|²) trained by the

180
wai/wai-rs/src/bin/wai_quantum_vectors.rs
View file

@ -24,15 +24,23 @@
//! names and reproduces its hashes, each reject file fails with its class, and
//! each accept file reads as listed.
//!
//! `verify <dir> --webgpu` (with the `quantum_fabric_webgpu` feature) runs the
//! same check once per WebGPU member this machine has, each vector's state
//! computed by the member instead of the native simulator. Each accept file is
//! also run on the member and must give what the native simulator gives: the
//! same state, or a refusal of the same class (the 255-qubit files are refused
//! as `too-many-qubits`). With no member it exits 2, listing the adapters
//! refused and why.
//!
//! `appendix` writes the phase tables of the extension's Appendix B, rendered
//! from the reference constants, into the generated part of the file it is
//! given (`extensions/quantum-sim.md`), so the tables are never typed by hand.
//!
//! wai_quantum_vectors gen <dir> | verify <dir> | appendix <quantum-sim.md>
//! wai_quantum_vectors gen <dir> | verify <dir> [--webgpu] | appendix <quantum-sim.md>
use std::collections::BTreeSet;
use std::path::Path;
use wai::quantum::{accept_corpus, circuit_hash, from_wqc, reject_corpus, to_wqc, BaseGate, Circuit, Gate, Gateset, Measure, DYADIC_MAX};
use wai::quantum::{accept_corpus, circuit_hash, from_wqc, reject_corpus, to_wqc, BaseGate, Circuit, Gate, Gateset, Measure, StateVector, DYADIC_MAX};
/// The circuit2 tree, under the corpus directory.
const V2_DIR: &str = "v2";
@ -49,20 +57,76 @@ fn main() {
let a: Vec<String> = std::env::args().collect();
match (a.get(1).map(String::as_str), a.get(2)) {
(Some("gen"), Some(d)) => generate(Path::new(d)),
(Some("verify"), Some(d)) => {
let fails = verify(Path::new(d));
if fails > 0 {
eprintln!("{fails} failed — NOT conformant");
std::process::exit(1);
(Some("verify"), Some(d)) => match a.get(3).map(String::as_str) {
None => {
if !summary(verify(Path::new(d))) {
std::process::exit(1);
}
}
println!("all vectors pass: statevector-equivalence holds, every reject file is refused as listed, and every accept file reads as listed");
}
Some("--webgpu") => verify_webgpu(Path::new(d)),
Some(_) => usage(),
},
(Some("appendix"), Some(f)) => write_appendix(Path::new(f)),
_ => {
eprintln!("usage: wai_quantum_vectors <gen|verify> <corpus-dir> | wai_quantum_vectors appendix <quantum-sim.md>");
std::process::exit(2);
_ => usage(),
}
}
fn usage() -> ! {
eprintln!("usage: wai_quantum_vectors gen <corpus-dir> | wai_quantum_vectors verify <corpus-dir> [--webgpu] | wai_quantum_vectors appendix <quantum-sim.md>");
std::process::exit(2);
}
/// Print the outcome of a `verify`; whether it passed.
fn summary(fails: usize) -> bool {
if fails > 0 {
eprintln!("{fails} failed — NOT conformant");
return false;
}
println!("all vectors pass: statevector-equivalence holds, every reject file is refused as listed, and every accept file reads as listed");
true
}
/// `verify <dir> --webgpu`: the whole check once per WebGPU member, each state
/// computed by the member. Exits 2 when there is no member, 1 when any member
/// fails.
#[cfg(all(feature = "quantum_fabric_webgpu", not(target_arch = "wasm32")))]
fn verify_webgpu(dir: &Path) {
use wai::quantum_fabric_webgpu::{blocking, WebGpuError};
let d = blocking::discover().unwrap_or_else(|e| {
eprintln!("verify --webgpu: adapter discovery did not resolve: {e}");
std::process::exit(2);
});
for r in &d.refused {
eprintln!("verify --webgpu: adapter {} ({}) refused: {}", r.key, r.adapter, r.error);
}
if d.found.is_empty() {
eprintln!("verify --webgpu: no WebGPU member on this machine");
std::process::exit(2);
}
let mut failed = 0;
for m in &d.found {
println!("verify --webgpu: member {}", m.id());
let run = |c: &Circuit| {
blocking::simulate_from(m, c, 0).map_err(|e| match e {
WebGpuError::Circuit(q) => (q.kind(), q.to_string()),
other => ("member", other.to_string()),
})
};
if !summary(verify_with(dir, Some(&run))) {
failed += 1;
}
}
if failed > 0 {
eprintln!("verify --webgpu: {failed} of {} members failed", d.found.len());
std::process::exit(1);
}
}
/// Without the WebGPU member in this build, `--webgpu` is an error.
#[cfg(not(all(feature = "quantum_fabric_webgpu", not(target_arch = "wasm32"))))]
fn verify_webgpu(_dir: &Path) {
eprintln!("verify --webgpu: this build has no WebGPU member; build wai_quantum_vectors natively with --features quantum_fabric_webgpu");
std::process::exit(2);
}
fn hex(b: &[u8]) -> String {
@ -280,12 +344,25 @@ fn write_appendix(path: &Path) {
// ---- Verification ----------------------------------------------------------------
/// A member's run of a circuit: its state, or a refusal as (class, message).
/// The class of a refusal the circuit itself earns is its `QuantumError`'s
/// kind; a refusal of the member's own (device, width, range) is `"member"`.
type Simulate<'a> = dyn Fn(&Circuit) -> Result<StateVector, (&'static str, String)> + 'a;
/// Verify all four trees under `dir`; the number of failures.
fn verify(dir: &Path) -> usize {
verify_tree(dir, "")
+ verify_tree(&dir.join(V2_DIR), &format!("{V2_DIR}/"))
verify_with(dir, None)
}
/// Verify all four trees under `dir`, each vector's state computed by `member`
/// when one is given and by the native simulator otherwise; with a member,
/// each accept file must also run on it as it does natively. The number of
/// failures.
fn verify_with(dir: &Path, member: Option<&Simulate>) -> usize {
verify_tree(dir, "", member)
+ verify_tree(&dir.join(V2_DIR), &format!("{V2_DIR}/"), member)
+ verify_rejects(&dir.join(REJECT_DIR), &format!("{REJECT_DIR}/"))
+ verify_accepts(&dir.join(ACCEPT_DIR), &format!("{ACCEPT_DIR}/"))
+ verify_accepts(&dir.join(ACCEPT_DIR), &format!("{ACCEPT_DIR}/"), member)
}
/// The `.wqc` file names in `dir`.
@ -337,7 +414,7 @@ fn same_files(listed: BTreeSet<String>, present: Result<BTreeSet<String>, String
}
/// A vector tree: `expected_hashes.json` and `vectors/`.
fn verify_tree(dir: &Path, label: &str) -> usize {
fn verify_tree(dir: &Path, label: &str, member: Option<&Simulate>) -> usize {
let m = match read_object(&dir.join("expected_hashes.json")) {
Ok(m) => m,
Err(why) => return report(&format!("{label}expected_hashes.json"), Err(why)),
@ -345,13 +422,14 @@ fn verify_tree(dir: &Path, label: &str) -> usize {
let v = dir.join("vectors");
let mut fails = report(&format!("{label}vectors/ holds the listed files"), same_files(m.keys().cloned().collect(), wqc_names(&v)));
for (file, spec) in &m {
fails += report(&format!("{label}{file}"), check_vector(&v.join(file), spec));
fails += report(&format!("{label}{file}"), check_vector(&v.join(file), spec, member));
}
fails
}
/// One vector against its manifest entry.
fn check_vector(path: &Path, spec: &serde_json::Value) -> Result<String, String> {
/// One vector against its manifest entry, its state computed by `member` when
/// one is given.
fn check_vector(path: &Path, spec: &serde_json::Value, member: Option<&Simulate>) -> Result<String, String> {
let field = |k: &str| spec.get(k).ok_or(format!("the entry has no {k}"));
let text = |k: &str| field(k).and_then(|v| v.as_str().ok_or(format!("{k} is not a string")));
let number = |v: &serde_json::Value, k: &str| v.get(k).and_then(|x| x.as_u64()).ok_or(format!("{k} is not an unsigned integer"));
@ -369,7 +447,10 @@ fn check_vector(path: &Path, spec: &serde_json::Value) -> Result<String, String>
if u64::from(circuit.n_qubits) != number(spec, "n_qubits")? || circuit.ops.len() as u64 != number(spec, "n_ops")? {
return Err("n_qubits or n_ops differ from the entry".into());
}
let sv = circuit.simulate().map_err(|e| e.to_string())?;
let sv = match member {
None => circuit.simulate().map_err(|e| e.to_string())?,
Some(run) => run(&circuit).map_err(|(_, why)| format!("the member refused it: {why}"))?,
};
let mut bad = Vec::new();
if hex(&circuit_hash(&circuit).map_err(|e| e.to_string())?) != text("circuit_hash")? {
bad.push("circuit_hash");
@ -415,8 +496,9 @@ fn verify_rejects(dir: &Path, label: &str) -> usize {
fails
}
/// The accept tree: every listed file is accepted and reads as listed.
fn verify_accepts(dir: &Path, label: &str) -> usize {
/// The accept tree: every listed file is accepted and reads as listed; with a
/// member, each also runs on it as it does natively.
fn verify_accepts(dir: &Path, label: &str, member: Option<&Simulate>) -> usize {
let m = match read_object(&dir.join(ACCEPT_LISTING)) {
Ok(m) => m,
Err(why) => return report(&format!("{label}{ACCEPT_LISTING}"), Err(why)),
@ -427,17 +509,37 @@ fn verify_accepts(dir: &Path, label: &str) -> usize {
let bytes = std::fs::read(dir.join(file)).map_err(|e| e.to_string())?;
let (c, measure) = from_wqc(&bytes).map_err(|e| format!("refused as {}: {e}", e.kind()))?;
let got = reading(&c, measure);
if &got == want { Ok(" (accepted, read as listed)".into()) } else { Err(format!("read as {got}, the listing says {want}")) }
if &got != want {
return Err(format!("read as {got}, the listing says {want}"));
}
match member {
None => Ok(" (accepted, read as listed)".into()),
Some(run) => runs_as_native(&c, run).map(|how| format!(" (accepted, read as listed; {how})")),
}
})();
fails += report(&format!("{label}{file}"), r);
}
fails
}
/// `c` on the member gives what it gives natively: the same state, or a
/// refusal of the same class.
fn runs_as_native(c: &Circuit, run: &Simulate) -> Result<String, String> {
match (c.simulate(), run(c)) {
(Ok(native), Ok(on)) if native == on => Ok("the member's state is the native one".into()),
(Ok(_), Ok(_)) => Err("the member's state differs from the native one".into()),
(Err(e), Err((class, _))) if e.kind() == class => Ok(format!("refused as {class} natively and by the member")),
(Err(e), Err((class, why))) => Err(format!("refused natively as {}, by the member as {class}: {why}", e.kind())),
(Ok(_), Err((_, why))) => Err(format!("the member refused it: {why}")),
(Err(e), Ok(_)) => Err(format!("the member ran it; natively it is refused: {e}")),
}
}
#[cfg(test)]
mod tests {
use super::*;
use std::path::PathBuf;
use wai::quantum::MAX_QUBITS;
fn corpus() -> PathBuf {
Path::new(env!("CARGO_MANIFEST_DIR")).join("../quantum-conformance")
@ -551,4 +653,36 @@ mod tests {
std::fs::remove_dir_all(s.0.join("v2")).unwrap();
assert_eq!(verify(&s.0), 3);
}
/// With a member, `verify` passes when the member gives the native states
/// and refusals, and fails each vector and accept file where it does not:
/// the 255-qubit accept files must be refused as too-many-qubits.
#[test]
fn verify_with_a_member_holds_it_to_the_native_simulator() {
let dir = corpus();
let native = |c: &Circuit| c.simulate().map_err(|e| (e.kind(), e.to_string()));
assert_eq!(verify_with(&dir, Some(&native)), 0);
let accepts = accept_corpus();
let wide = accepts.iter().filter(|(_, _, c, _)| c.simulate().is_err()).count();
assert!(wide > 0, "an accept file is refused at simulate");
assert!(
accepts.iter().all(|(_, _, c, _)| c.simulate().is_err() == (c.n_qubits > MAX_QUBITS)),
"the accept files refused at simulate are the ones too wide to simulate"
);
let n_vectors = 2 * vectors(Gateset::V1).len();
// A member that changes one amplitude: every vector, and every accept
// file but the ones both refuse.
let off_by_one = |c: &Circuit| {
native(c).map(|mut s| {
s.amps[0].re += 1;
s
})
};
assert_eq!(verify_with(&dir, Some(&off_by_one)), n_vectors + accepts.len() - wide);
// A member that refuses everything as its own: every vector and every
// accept file, the 255-qubit ones included (its refusal is not the
// circuit's class).
let refuses = |_: &Circuit| -> Result<StateVector, (&'static str, String)> { Err(("member", "no device".to_string())) };
assert_eq!(verify_with(&dir, Some(&refuses)), n_vectors + accepts.len());
}
}

5
wai/wai-rs/src/lib.rs
View file

@ -737,6 +737,11 @@ pub use wai_quantum::quantum_ldpc;
#[cfg(feature = "quantum_fabric")]
pub use wai_quantum::quantum_fabric;
// The fabric's WebGPU member: natively, and in a browser on
// wasm32-unknown-unknown. No other wasm32 target has a GPU interface.
#[cfg(all(feature = "quantum_fabric_webgpu", any(not(target_arch = "wasm32"), target_os = "unknown")))]
pub use wai_quantum::quantum_fabric_webgpu;
#[cfg(feature = "quantum_qio")]
pub use wai_quantum::quantum_qio;

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

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