From 7443ce1992991dbd8e5b078e2f43e0442c7cf300 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Wed, 14 Oct 2026 01:14:01 -0400 Subject: [PATCH] wai-quantum 0.4.4: the fabric's CUDA member MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A second accelerator member beside the WebGPU one: the gate kernel in CUDA C, compiled at run time by the CUDA runtime compiler for each device's own architecture, certified bit-exact against the CPU reference's pinned corpus digests. Feature `quantum_fabric_cuda`, opt-in, native targets only. The kernel (src/quantum_fabric/lanes64.cu) is a transcription of lanes::exact_row and apply_words in 64-bit integers, which every CUDA device has: each fx(a, b) = (a·b + 2^29) >> 30 is exact in 64 bits, each row's sum too, the lane keeps the low 32 bits, and a sum outside 32 bits records the gate's op-log index, the lowest winning. One thread per amplitude pair, one launch per gate, so no thread reads what another writes within a launch. The driver and the runtime compiler are loaded when first used, so a build with the feature runs on a host with no CUDA and finds no member, with the reason (Discovery::unavailable). The member mirrors the WebGPU member's surface: discovery with typed refusals, the arithmetic probe (shared now as lanes::probe_pairs and first_mismatch), run_lanes / run_program, simulate_from and simulate_metered, certify, member and member_metered, a Runner for the Fabric, with_max_qubits and with_source for tests. Its id binds the device, UUID, PCI location, driver, compiler and memory: local/gpu-cuda-<12 hex>, keyed in CI as gpu-cuda-sm. Device tests (tests/cuda_member.rs), on the RTX 4050 (sm_89): the probe agrees on 65,536 pairs; the kernel matches run_words on 2,000 random states and gate words, out-of-range runs included; all four pinned certificates reproduced; random circuits up to 22 qubits and QFTs equal to native; typed refusals; all six kernel mutants caught; a receipt, a placement and a Fabric run; repeated runs identical. On a host with no CUDA the same tests pass over no member. CI: wai-quantum-cuda.yml, on the GPU host (requires the four certificates and every mutant caught) and on a runner with no CUDA (requires the member to say so). Every lockfile that includes wai-quantum moves to 0.4.4 by its version line alone. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-cuda.yml | 105 +++ wai/tools/quantum-meter/Cargo.lock | 2 +- wai/tools/quantum-wgsl-check/Cargo.lock | 2 +- wai/wai-pack/Cargo.lock | 2 +- wai/wai-quantum-serve/Cargo.lock | 2 +- wai/wai-quantum/Cargo.lock | 26 +- wai/wai-quantum/Cargo.toml | 12 +- wai/wai-quantum/README.md | 10 +- wai/wai-quantum/src/lib.rs | 9 +- wai/wai-quantum/src/quantum_fabric.rs | 23 + wai/wai-quantum/src/quantum_fabric/cuda.rs | 247 +++++++ wai/wai-quantum/src/quantum_fabric/lanes64.cu | 95 +++ wai/wai-quantum/src/quantum_fabric_cuda.rs | 645 ++++++++++++++++++ wai/wai-quantum/src/quantum_fabric_webgpu.rs | 24 +- wai/wai-quantum/tests/cuda_member.rs | 429 ++++++++++++ wai/wai-rs/Cargo.lock | 2 +- wai/wai-web/Cargo.lock | 2 +- 17 files changed, 1600 insertions(+), 37 deletions(-) create mode 100644 .github/workflows/wai-quantum-cuda.yml create mode 100644 wai/wai-quantum/src/quantum_fabric/cuda.rs create mode 100644 wai/wai-quantum/src/quantum_fabric/lanes64.cu create mode 100644 wai/wai-quantum/src/quantum_fabric_cuda.rs create mode 100644 wai/wai-quantum/tests/cuda_member.rs diff --git a/.github/workflows/wai-quantum-cuda.yml b/.github/workflows/wai-quantum-cuda.yml new file mode 100644 index 000000000..9dd4a2739 --- /dev/null +++ b/.github/workflows/wai-quantum-cuda.yml @@ -0,0 +1,105 @@ +name: wai-quantum CUDA member + +# The fabric's CUDA member: the gate kernel in CUDA C, compiled at run time for +# each device's own architecture and certified bit-exact against the CPU +# reference's pinned corpus digests. +# +# - cuda-hw (jetson-hub, the RTX 4050): the device tests must find exactly the +# declared device, agree with the CPU on the arithmetic probe and on random +# states (out-of-range runs included), reproduce all four pinned certificates +# (two gatesets, random-v1 and fabric-v1), and catch every kernel mutant. +# - cuda-absent (a container with no GPU and no CUDA libraries): the same build +# finds no member and says why, rather than failing; the device tests run over +# no member with the declaration `none`. + +concurrency: + group: ${{ github.workflow }}-${{ github.ref }} + cancel-in-progress: true + +on: + push: + branches: [main] + paths: + - 'wai/wai-quantum/**' + - 'rust-toolchain.toml' + - '.github/workflows/wai-quantum-cuda.yml' + pull_request: + paths: + - 'wai/wai-quantum/**' + - 'rust-toolchain.toml' + - '.github/workflows/wai-quantum-cuda.yml' + +jobs: + cuda-hw: + name: member · CUDA (jetson-hub) + runs-on: linux-hw + timeout-minutes: 90 + env: + WAI_QUANTUM_CUDA_EXPECT: gpu-cuda-sm89 + 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 unit tests + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + export PATH="$HOME/.cargo/bin:$PATH" + CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-cuda" nice -n 10 cargo test --locked --features quantum_fabric_cuda --lib -- quantum_fabric_cuda:: quantum_fabric::cuda:: 2>&1 | tee "$RUNNER_TEMP/unit.log" + grep -qE '^test result: ok\. [1-9][0-9]* passed; 0 failed' "$RUNNER_TEMP/unit.log" || { echo "::error::no unit test ran"; exit 1; } + - name: The member's device tests + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + export PATH="$HOME/.cargo/bin:$PATH" + report="$RUNNER_TEMP/cuda-report.txt"; : > "$report" + WAI_QUANTUM_CUDA_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-cuda" nice -n 10 \ + cargo test --locked --release --features quantum_fabric_cuda --test cuda_member -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/device.log" + grep -qE '^test result: ok\. 9 passed; 0 failed' "$RUNNER_TEMP/device.log" || { echo "::error::not every device test ran and passed"; exit 1; } + cat "$report" + key="$WAI_QUANTUM_CUDA_EXPECT" + grep -qx "cuda devices: as declared ($key)" "$report" || { echo "::error::the devices found are not the declared ones"; exit 1; } + n="$(grep -c "^cuda certified: key=$key .* parity=bit-exact " "$report")" + [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; } + grep -qx "cuda mutants: key=$key killed 6 of 6" "$report" || { echo "::error::a kernel mutant survived"; exit 1; } + + cuda-absent: + name: member · CUDA absent (linux) + runs-on: linux + timeout-minutes: 60 + env: + WAI_QUANTUM_CUDA_EXPECT: none + 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: No CUDA here, and the member says so + shell: bash + working-directory: wai/wai-quantum + run: | + set -euo pipefail + report="$RUNNER_TEMP/cuda-report.txt"; : > "$report" + WAI_QUANTUM_CUDA_REPORT="$report" cargo test --locked --features quantum_fabric_cuda --test cuda_member -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/device.log" + grep -qE '^test result: ok\. 9 passed; 0 failed' "$RUNNER_TEMP/device.log" || { echo "::error::not every device test ran and passed"; exit 1; } + grep -qx 'cuda devices: as declared (none)' "$report" || { echo "::error::the absent case was not established"; exit 1; } diff --git a/wai/tools/quantum-meter/Cargo.lock b/wai/tools/quantum-meter/Cargo.lock index f4645f51d..402a34400 100644 --- a/wai/tools/quantum-meter/Cargo.lock +++ b/wai/tools/quantum-meter/Cargo.lock @@ -1017,7 +1017,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" [[package]] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/tools/quantum-wgsl-check/Cargo.lock b/wai/tools/quantum-wgsl-check/Cargo.lock index 92290ce3b..1c74ca644 100644 --- a/wai/tools/quantum-wgsl-check/Cargo.lock +++ b/wai/tools/quantum-wgsl-check/Cargo.lock @@ -599,7 +599,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" [[package]] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-pack/Cargo.lock b/wai/wai-pack/Cargo.lock index 7e2ff5f8e..89652ca29 100644 --- a/wai/wai-pack/Cargo.lock +++ b/wai/wai-pack/Cargo.lock @@ -698,7 +698,7 @@ dependencies = [ [[package]] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-quantum-serve/Cargo.lock b/wai/wai-quantum-serve/Cargo.lock index 6982538a9..923efee37 100644 --- a/wai/wai-quantum-serve/Cargo.lock +++ b/wai/wai-quantum-serve/Cargo.lock @@ -589,7 +589,7 @@ dependencies = [ [[package]] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-quantum/Cargo.lock b/wai/wai-quantum/Cargo.lock index 893115472..a1ce41a3c 100644 --- a/wai/wai-quantum/Cargo.lock +++ b/wai/wai-quantum/Cargo.lock @@ -29,7 +29,7 @@ version = "0.38.0+1.3.281" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f" dependencies = [ - "libloading", + "libloading 0.8.9", ] [[package]] @@ -201,6 +201,15 @@ dependencies = [ "typenum", ] +[[package]] +name = "cudarc" +version = "0.19.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7359bf1de037ddada72729c3acec9b040f3468400c6e9e740b01ba0cd2e89562" +dependencies = [ + "libloading 0.9.0", +] + [[package]] name = "curve25519-dalek" version = "4.1.3" @@ -449,6 +458,16 @@ dependencies = [ "windows-link", ] +[[package]] +name = "libloading" +version = "0.9.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "754ca22de805bb5744484a5b151a9e1a8e837d5dc232c2d7d8c2e3492edc8b60" +dependencies = [ + "cfg-if", + "windows-link", +] + [[package]] name = "libm" version = "0.2.16" @@ -995,9 +1014,10 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" [[package]] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" dependencies = [ "blake3", + "cudarc", "ed25519-dalek", "serde_json", "wgpu", @@ -1173,7 +1193,7 @@ dependencies = [ "gpu-allocator", "hashbrown 0.17.1", "libc", - "libloading", + "libloading 0.8.9", "log", "naga", "naga-types", diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml index aacf5cfdf..9a4a1866a 100644 --- a/wai/wai-quantum/Cargo.toml +++ b/wai/wai-quantum/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" edition = "2024" rust-version = "1.97" license = "Apache-2.0" @@ -34,6 +34,10 @@ ed25519-dalek = { version = "2", optional = true } # 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"] } +# The CUDA member: the driver and the runtime compiler are loaded when first used, +# so a build with the feature needs no CUDA installed and finds no member where +# there is none. +cudarc = { version = "0.19.10", optional = true, default-features = false, features = ["std", "driver", "nvrtc", "dynamic-loading", "cuda-12080"] } [target.'cfg(all(target_arch = "wasm32", target_os = "unknown"))'.dependencies] wgpu = { version = "30.0.1", optional = true, default-features = false, features = ["std", "wgsl", "webgpu"] } @@ -140,6 +144,8 @@ quantum_fabric = ["quantum", "quantum_ops"] # 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"] +# The CUDA member of the fabric (an NVIDIA GPU, native targets only). +quantum_fabric_cuda = ["quantum_fabric", "dep:cudarc"] # --- interference / ML layer ---------------------------------------------- quantum_vml = [] @@ -162,6 +168,10 @@ full = [ name = "webgpu_member" required-features = ["quantum_fabric_webgpu"] +[[test]] +name = "cuda_member" +required-features = ["quantum_fabric_cuda"] + [[test]] name = "fabric_expected" required-features = ["quantum_fabric"] diff --git a/wai/wai-quantum/README.md b/wai/wai-quantum/README.md index b4a936433..2e671289f 100644 --- a/wai/wai-quantum/README.md +++ b/wai/wai-quantum/README.md @@ -6,9 +6,10 @@ 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**. One -opt-in feature, `quantum_fabric_webgpu`, reaches the platform's GPU interface; -no default feature reaches a system library. +the edge, producing **byte-identical results and receipts on all three**. Two +opt-in features reach a platform interface: `quantum_fabric_webgpu` (the +platform's GPU interface) and `quantum_fabric_cuda` (the CUDA driver and runtime +compiler, loaded when first used); no default feature reaches a system library. ```toml [dependencies] @@ -94,7 +95,8 @@ 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`; a second implementation applies runs of gates tile by tile, one dispatch per run, and is certified on its own. Each adapter is probed against the CPU's arithmetic in the process that uses it and certified bit-exact by reproducing the CPU reference's pinned corpus digests; a result that would leave 32 bits is refused, never returned. The one feature that reaches a platform interface | +| `quantum_fabric_webgpu` | **the fabric's WebGPU member**, opt-in and not in `full`: the 32-bit-lane gate kernel in WGSL, one dispatch per gate, through Metal, Vulkan or DX12 natively and through a browser's WebGPU on `wasm32-unknown-unknown`; a second implementation applies runs of gates tile by tile, one dispatch per run, and is certified on its own. Each adapter is probed against the CPU's arithmetic in the process that uses it and certified bit-exact by reproducing the CPU reference's pinned corpus digests; a result that would leave 32 bits is refused, never returned. One of the two features that reach a platform interface | +| `quantum_fabric_cuda` | **the fabric's CUDA member**, opt-in and not in `full`, native targets only: the gate kernel in CUDA C with exact 64-bit integers, compiled at run time for each device's own architecture, one launch per gate. The CUDA driver and runtime compiler are loaded when first used, so a build with the feature runs on a host without CUDA and finds no member there, saying why. Each device 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 | | `full` | everything | Default features are `quantum`, `quantum_toolchain`, `quantum_receipt`. diff --git a/wai/wai-quantum/src/lib.rs b/wai/wai-quantum/src/lib.rs index 4df5cf199..d0b9f8607 100644 --- a/wai/wai-quantum/src/lib.rs +++ b/wai/wai-quantum/src/lib.rs @@ -10,9 +10,10 @@ //! 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. One opt-in feature, `quantum_fabric_webgpu`, -//! reaches the platform's GPU interface; no default feature reaches a system -//! library. +//! member that ran the work. Two opt-in features reach a platform interface: +//! `quantum_fabric_webgpu` (the platform's GPU interface) and +//! `quantum_fabric_cuda` (the CUDA driver and runtime compiler, loaded when +//! first used); no default feature reaches a system library. //! //! Everything is behind a feature; enable only what you use. @@ -76,6 +77,8 @@ pub mod quantum_fabric; // 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(all(feature = "quantum_fabric_cuda", not(target_arch = "wasm32")))] +pub mod quantum_fabric_cuda; #[cfg(feature = "quantum_control")] pub mod quantum_control; diff --git a/wai/wai-quantum/src/quantum_fabric.rs b/wai/wai-quantum/src/quantum_fabric.rs index df9e6a5db..241cc9140 100644 --- a/wai/wai-quantum/src/quantum_fabric.rs +++ b/wai/wai-quantum/src/quantum_fabric.rs @@ -1815,6 +1815,28 @@ pub mod lanes { i32::MIN + 1, ]; + /// The probe's operands, as every GPU member runs them: every pair of + /// [`PROBE_EDGES`], then `random` pairs of splitmix64 draws from `seed`. + #[doc(hidden)] + pub fn probe_pairs(random: usize, seed: u64) -> Vec<(i32, i32)> { + let edges = PROBE_EDGES; + let mut pairs: Vec<(i32, i32)> = edges.iter().flat_map(|&a| edges.iter().map(move |&b| (a, b))).collect(); + let mut s = seed; + pairs.extend((0..random).map(|_| (next(&mut s) as i32, next(&mut s) as i32))); + pairs + } + + /// The first pair whose device result differs from the CPU's + /// ([`mul_wide`], [`fxmul_wide`]); a result list of the wrong length + /// fails at its first pair. + #[doc(hidden)] + pub fn first_mismatch(pairs: &[(i32, i32)], out: &[((u32, u32), (u32, u32))]) -> Option<(i32, i32)> { + if pairs.len() != out.len() { + return pairs.first().copied().or(Some((0, 0))); + } + pairs.iter().zip(out).find(|((a, b), o)| **o != (mul_wide(*a, *b), fxmul_wide(*a, *b))).map(|(p, _)| *p) + } + /// 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 @@ -1933,6 +1955,7 @@ pub mod lanes { } } +pub mod cuda; pub mod plan; pub mod wgsl; diff --git a/wai/wai-quantum/src/quantum_fabric/cuda.rs b/wai/wai-quantum/src/quantum_fabric/cuda.rs new file mode 100644 index 000000000..ffb0dd41f --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/cuda.rs @@ -0,0 +1,247 @@ +//! The gate kernel for CUDA devices, and what names a CUDA member: the +//! kernel's text (`lanes64.cu`), its implementation id, a device's identity +//! and the mutants a conforming device must be caught running. No device is +//! touched here, so all of it is testable on any host; the CUDA member itself +//! is `quantum_fabric_cuda`. +//! +//! **Why the kernel is bit-exact.** Every CUDA device has 64-bit integer +//! arithmetic, so the kernel is a transcription of [`lanes::exact_row`] and +//! [`lanes::apply_words`] rather than of the 32-bit [`lanes`] steps a WGSL +//! kernel needs: +//! +//! - a gate entry and an amplitude are 32-bit values, so their product is at +//! most `2^62` in size, and `fx(a, b) = (a·b + 2^29) >> 30` is exact in 64 +//! bits, with `>>` arithmetic as Rust's is on `i64`; +//! - a row's sum of four such terms is at most `2^34` in size, exact too; +//! - the lane keeps the sum's low 32 bits, as [`lanes::apply_words`] does, +//! and a sum outside 32 bits records the gate's op-log index, the lowest +//! winning ([`lanes::run_words`]). +//! +//! Each thread owns one amplitude pair and a launch is one gate, so no thread +//! reads what another writes within a launch, and the order threads run in +//! cannot change a bit. The member's probe checks the device's products and +//! shifts against [`lanes::mul_wide`] and [`lanes::fxmul_wide`] before the +//! device is a member, and its certificates check whole runs against the CPU's +//! pinned corpus digests. +//! +//! [`lanes::exact_row`]: super::lanes +//! [`lanes::apply_words`]: super::lanes::apply_words +//! [`lanes::run_words`]: super::lanes::run_words +//! [`lanes::mul_wide`]: super::lanes::mul_wide +//! [`lanes::fxmul_wide`]: super::lanes::fxmul_wide +//! [`lanes`]: super::lanes + +use super::FabricClass; +use super::wgsl::{Detector, Mutant}; + +/// The kernel, `lanes64.cu`, exactly as it is compiled. +pub const SOURCE: &str = include_str!("lanes64.cu"); + +/// The implementation name the id starts with. +pub const IMPLEMENTATION: &str = "cuda-lanes64"; + +/// Version of the host-side layout: interleaved `int2` lanes, one launch per +/// gate with the gate's words as kernel parameters, the report word (the +/// lowest out-of-range op index, `u32::MAX` for none), the probe's `int2` +/// cases and `uint4` outputs, and the launch shape ([`BLOCK`], +/// [`GRID_MAX`]). Bumped on any change to them; it is part of the +/// implementation id. +pub const HOST_LAYOUT: u32 = 1; + +/// Threads per block. +pub const BLOCK: u32 = 256; + +/// Blocks per launch, at most; the kernel strides over the rest. +pub const GRID_MAX: u32 = 4096; + +/// The options the runtime compiler is given, besides the target +/// architecture, which is the device's own: no fused multiply-add, though the +/// kernel has no floating point for it to touch. Fast math is off unless +/// asked for, and is not. +pub const COMPILE_OPTIONS: [&str; 2] = ["--std=c++17", "--fmad=false"]; + +const DOMAIN_IMPLEMENTATION: &[u8] = b"wai:quantum-fabric-cuda\x01"; +const DOMAIN_MEMBER: &[u8] = b"wai:quantum-fabric-cuda-member\x01"; + +fn hex(bytes: &[u8]) -> String { + bytes.iter().map(|b| format!("{b:02x}")).collect() +} + +/// A length-prefixed byte string, so that no field reads as part of another. +fn put(o: &mut Vec, bytes: &[u8]) { + o.extend_from_slice(&(bytes.len() as u64).to_be_bytes()); + o.extend_from_slice(bytes); +} + +/// `"cuda-lanes64/<16 hex>"`: the first 8 bytes of +/// BLAKE3(`"wai:quantum-fabric-cuda\x01"` ‖ put([`SOURCE`]) ‖ put of each of +/// [`COMPILE_OPTIONS`] ‖ [`HOST_LAYOUT`] as a big-endian `u32`), in hex. Any +/// edit to the kernel's text, its compile options or the host layout changes +/// it. The target architecture is the device's and is part of the member id +/// instead. +pub fn implementation_id() -> String { + implementation_id_of(SOURCE) +} + +/// [`implementation_id`] for another kernel text, such as a mutant's. +pub fn implementation_id_of(source: &str) -> String { + let mut o = DOMAIN_IMPLEMENTATION.to_vec(); + put(&mut o, source.as_bytes()); + for option in COMPILE_OPTIONS { + put(&mut o, option.as_bytes()); + } + o.extend_from_slice(&HOST_LAYOUT.to_be_bytes()); + format!("{IMPLEMENTATION}/{}", hex(&blake3::hash(&o).as_bytes()[..8])) +} + +/// What a CUDA member is, as its id hashes it. The member fills it from the +/// driver; the fields are plain values here so that the id is testable +/// without a device. +#[derive(Clone, Debug, PartialEq, Eq)] +pub struct DeviceIdentity { + pub name: String, + /// The device's UUID, in hex. + pub uuid: String, + /// The compute capability, `(major, minor)`: the architecture the kernel + /// is compiled for. + pub compute: (u32, u32), + /// `"::"`, in hex. + pub pci: String, + /// The driver's CUDA version, as `cuDriverGetVersion` reports it + /// (`major·1000 + minor·10`). + pub driver: u32, + /// The runtime compiler's version, `(major, minor)`. + pub compiler: (u32, u32), + /// The device's memory, bytes. + pub memory: u64, +} + +impl DeviceIdentity { + /// Always [`FabricClass::Gpu`]. + pub fn class(&self) -> FabricClass { + FabricClass::Gpu + } + + /// The target architecture the kernel is compiled for: + /// `"compute_"`. + pub fn arch(&self) -> String { + format!("compute_{}{}", self.compute.0, self.compute.1) + } + + /// `"local/gpu-cuda-"`. `d12` is the first 12 hex digits of + /// BLAKE3(`"wai:quantum-fabric-cuda-member\x01"` ‖ put(name) ‖ put(uuid) ‖ + /// compute major ‖ compute minor ‖ put(pci) ‖ driver ‖ compiler major ‖ + /// compiler minor ‖ memory), the numbers big-endian. It sorts after every + /// `local/cpu-`, so with no energy figures + /// [`place`](super::place) keeps the CPU reference. + pub fn member_id(&self) -> String { + let mut o = DOMAIN_MEMBER.to_vec(); + put(&mut o, self.name.as_bytes()); + put(&mut o, self.uuid.as_bytes()); + o.extend_from_slice(&self.compute.0.to_be_bytes()); + o.extend_from_slice(&self.compute.1.to_be_bytes()); + put(&mut o, self.pci.as_bytes()); + o.extend_from_slice(&self.driver.to_be_bytes()); + o.extend_from_slice(&self.compiler.0.to_be_bytes()); + o.extend_from_slice(&self.compiler.1.to_be_bytes()); + o.extend_from_slice(&self.memory.to_be_bytes()); + format!("local/gpu-cuda-{}", hex(&blake3::hash(&o).as_bytes()[..6])) + } + + /// The key a runner declares in CI: `"gpu-cuda-sm"`, for + /// instance `"gpu-cuda-sm89"`. + pub fn key(&self) -> String { + format!("gpu-cuda-sm{}{}", self.compute.0, self.compute.1) + } +} + +/// The mutants: textual edits of [`SOURCE`] that a conforming device must be +/// caught running, each by the detector named. Each `find` occurs exactly +/// once in [`SOURCE`]. +pub const MUTANTS: [Mutant; 6] = [ + // The rounding constant is half what it should be. + Mutant { name: "C1", find: "((i64)a * (i64)b + ((i64)1 << 29)) >> 30;", replace: "((i64)a * (i64)b + ((i64)1 << 28)) >> 30;", detector: Detector::Probe }, + // The product is shifted by one bit too few. + Mutant { name: "C2", find: "((i64)a * (i64)b + ((i64)1 << 29)) >> 30;", replace: "((i64)a * (i64)b + ((i64)1 << 29)) >> 29;", detector: Detector::Probe }, + // The probe's product is taken in 32 bits. + Mutant { name: "C3", find: "u64 prod = (u64)((i64)c.x * (i64)c.y);", replace: "u64 prod = (u64)(unsigned int)(c.x * c.y);", detector: Detector::Probe }, + // The control test is inverted. + Mutant { name: "C4", find: "if ((i & (u64)ctrl) != (u64)ctrl) {", replace: "if ((i & (u64)ctrl) == (u64)ctrl) {", detector: Detector::Certificate }, + // A sign of row 0's real part is wrong, on the imaginary part of the + // top-right entry, which `Y` has. (The top-left entry is real for every + // gate of the set, so the same edit on its imaginary part changes no run.) + Mutant { name: "C5", find: "fx(m2, a1.x) - fx(m3, a1.y);", replace: "fx(m2, a1.x) + fx(m3, a1.y);", detector: Detector::Certificate }, + // The range guard never fires. + Mutant { name: "C6", find: "if (!(fits32(o0r) && fits32(o0i) && fits32(o1r) && fits32(o1i))) {", replace: "if (false) {", detector: Detector::Range }, +]; + +/// [`SOURCE`] with `m` applied. Panics unless `m.find` occurs exactly once. +pub fn mutate(m: &Mutant) -> String { + let n = SOURCE.matches(m.find).count(); + assert_eq!(n, 1, "mutant {}: `{}` occurs {n} times in the kernel", m.name, m.find); + SOURCE.replacen(m.find, m.replace, 1) +} + +#[cfg(test)] +mod tests { + use super::*; + + fn identity() -> DeviceIdentity { + DeviceIdentity { + name: "a device".into(), + uuid: "00112233445566778899aabbccddeeff".into(), + compute: (8, 9), + pci: "0000:01:00".into(), + driver: 13020, + compiler: (12, 8), + memory: 6 << 30, + } + } + + #[test] + fn every_field_of_the_identity_is_in_the_member_id() { + let base = identity(); + let id = base.member_id(); + assert!(id.starts_with("local/gpu-cuda-") && id.len() == "local/gpu-cuda-".len() + 12, "{id}"); + assert!(id.as_str() > "local/cpu-x86_64" && id.as_str() > "local/cpu-aarch64", "sorts after every CPU reference"); + let variants = [ + DeviceIdentity { name: "another".into(), ..identity() }, + DeviceIdentity { uuid: "ff".into(), ..identity() }, + DeviceIdentity { compute: (9, 0), ..identity() }, + DeviceIdentity { compute: (8, 6), ..identity() }, + DeviceIdentity { pci: "0000:02:00".into(), ..identity() }, + DeviceIdentity { driver: 12080, ..identity() }, + DeviceIdentity { compiler: (13, 0), ..identity() }, + DeviceIdentity { compiler: (12, 9), ..identity() }, + DeviceIdentity { memory: 8 << 30, ..identity() }, + ]; + for v in variants { + assert_ne!(v.member_id(), id, "{v:?}"); + } + assert_eq!(base.key(), "gpu-cuda-sm89"); + assert_eq!(base.arch(), "compute_89"); + assert_eq!(base.class(), FabricClass::Gpu); + } + + #[test] + fn the_implementation_id_follows_the_kernel_text() { + let id = implementation_id(); + assert!(id.starts_with("cuda-lanes64/") && id.len() == "cuda-lanes64/".len() + 16, "{id}"); + for m in &MUTANTS { + assert_ne!(implementation_id_of(&mutate(m)), id, "mutant {}", m.name); + } + } + + #[test] + fn every_mutant_edits_the_kernel_once_and_names_distinct_edits() { + let mut seen = std::collections::BTreeSet::new(); + for m in &MUTANTS { + let edited = mutate(m); + assert_ne!(edited, SOURCE, "mutant {} changes nothing", m.name); + assert!(seen.insert(edited), "mutant {} duplicates another", m.name); + } + for d in [Detector::Probe, Detector::Certificate, Detector::Range] { + assert!(MUTANTS.iter().any(|m| m.detector == d), "no mutant for {d:?}"); + } + } +} diff --git a/wai/wai-quantum/src/quantum_fabric/lanes64.cu b/wai/wai-quantum/src/quantum_fabric/lanes64.cu new file mode 100644 index 000000000..278e5eb8b --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/lanes64.cu @@ -0,0 +1,95 @@ +// wai-quantum fabric: the gate kernel for CUDA devices. +// +// A transcription of `lanes::exact_row` and `lanes::apply_words` in 64-bit +// integers, which every CUDA device has natively. Amplitudes are pairs of +// 32-bit fixed-point values (30 fractional bits). Each product +// `fx(a, b) = (a·b + 2^29) >> 30` is exact in 64 bits (a 32×32 product is at +// most 2^62 in size), and so is each row's sum of four products. The lane +// stores the low 32 bits of each sum; a sum outside 32 bits records the +// gate's op-log index in `report[0]`, the lowest such index winning, so the +// host refuses the run with the gate the CPU reference names. +// +// Nothing here is floating point, depends on the order threads run in, or +// reads an amplitude another thread of the same launch writes: each thread +// owns one pair (i, j), and a launch is one gate. +// +// `>>` on a negative signed value is arithmetic in CUDA C++, as Rust's is on +// i64; the host's probe checks this, among the rest of the arithmetic, before +// a device is a member. + +typedef long long i64; +typedef unsigned long long u64; + +// The exact fixed-point product. +__device__ __forceinline__ i64 fx(int a, int b) { + return ((i64)a * (i64)b + ((i64)1 << 29)) >> 30; +} + +// The low 32 bits of `v`, as the lane keeps them. +__device__ __forceinline__ int low32(i64 v) { + return (int)(unsigned int)(u64)v; +} + +// Whether `v` fits a 32-bit lane. +__device__ __forceinline__ bool fits32(i64 v) { + return v >= -(i64)2147483648LL && v <= (i64)2147483647LL; +} + +// One gate on a state of `2 * pairs` lanes, one thread per pair. +// +// `m0..m7` are the gate's words: row 0 is (m0 + i·m1)·a0 + (m2 + i·m3)·a1, +// row 1 is (m4 + i·m5)·a0 + (m6 + i·m7)·a1. `tbit` is 1 << target; `ctrl` +// the OR of the controls' bits. `index` is the gate's position in the op +// log. The host checks the words before a launch (`apply_words` refuses the +// same malformed words), so the kernel trusts them. +extern "C" __global__ void apply_gate( + int2* __restrict__ amps, + u64 pairs, + unsigned int tbit, + unsigned int ctrl, + int m0, int m1, int m2, int m3, int m4, int m5, int m6, int m7, + unsigned int index, + unsigned int* __restrict__ report) +{ + u64 p = (u64)blockIdx.x * (u64)blockDim.x + (u64)threadIdx.x; + u64 stride = (u64)gridDim.x * (u64)blockDim.x; + for (; p < pairs; p += stride) { + // The p-th index whose target bit is clear: p with a zero inserted at + // the target's position. + u64 low = p & ((u64)tbit - 1); + u64 i = ((p - low) << 1) | low; + if ((i & (u64)ctrl) != (u64)ctrl) { + continue; + } + u64 j = i | (u64)tbit; + int2 a0 = amps[i]; + int2 a1 = amps[j]; + i64 o0r = fx(m0, a0.x) - fx(m1, a0.y) + fx(m2, a1.x) - fx(m3, a1.y); + i64 o0i = fx(m0, a0.y) + fx(m1, a0.x) + fx(m2, a1.y) + fx(m3, a1.x); + i64 o1r = fx(m4, a0.x) - fx(m5, a0.y) + fx(m6, a1.x) - fx(m7, a1.y); + i64 o1i = fx(m4, a0.y) + fx(m5, a0.x) + fx(m6, a1.y) + fx(m7, a1.x); + if (!(fits32(o0r) && fits32(o0i) && fits32(o1r) && fits32(o1i))) { + atomicMin(&report[0], index); + } + amps[i] = make_int2(low32(o0r), low32(o0i)); + amps[j] = make_int2(low32(o1r), low32(o1i)); + } +} + +// The arithmetic probe: for each case (a, b), the 64-bit two's-complement +// product a·b and the exact fx(a, b), each as (high, low) words, the outputs +// of `lanes::mul_wide` and `lanes::fxmul_wide`. +extern "C" __global__ void probe( + const int2* __restrict__ cases, + uint4* __restrict__ out, + u64 count) +{ + u64 k = (u64)blockIdx.x * (u64)blockDim.x + (u64)threadIdx.x; + u64 stride = (u64)gridDim.x * (u64)blockDim.x; + for (; k < count; k += stride) { + int2 c = cases[k]; + u64 prod = (u64)((i64)c.x * (i64)c.y); + u64 f = (u64)fx(c.x, c.y); + out[k] = make_uint4((unsigned int)(prod >> 32), (unsigned int)prod, (unsigned int)(f >> 32), (unsigned int)f); + } +} diff --git a/wai/wai-quantum/src/quantum_fabric_cuda.rs b/wai/wai-quantum/src/quantum_fabric_cuda.rs new file mode 100644 index 000000000..21d8a375a --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric_cuda.rs @@ -0,0 +1,645 @@ +//! The fabric's CUDA member: the gate kernel of +//! [`quantum_fabric::cuda`](crate::quantum_fabric::cuda) on an NVIDIA GPU, +//! compiled at run time for the device's own architecture. See that module +//! for the kernel and the argument that it is bit-exact. +//! +//! - **Discovery** ([`discover`]): the driver and the runtime compiler are +//! loaded when the process first asks for them, so a build with this feature +//! runs on a host with no CUDA at all and finds no member there +//! ([`Discovery::unavailable`] says why). Every device the driver lists is +//! opened, the kernel is compiled for it, and its arithmetic is probed +//! against [`lanes`] in this process. A device that cannot host the kernel, +//! or whose arithmetic disagrees, is listed with the reason ([`Refusal`]); +//! one that passes is a [`CudaMember`], present as [`Presence::Detected`]. +//! - **Running** ([`CudaMember::run_program`]): the state stays on the device; +//! one launch per gate, in op-log order, on one stream; then the report word +//! and the state come back. The kernel reports exactly when the reference's +//! result leaves 32 bits; the member then returns +//! [`CudaError::OutOfRange`], never a state. +//! - **Certification** ([`CudaMember::certify`], [`CudaMember::member`]): the +//! member's own run of a corpus, folded case by case into a +//! [`Certification`]. The corpus digest covers inputs and outputs only, so a +//! bit-exact member reproduces the CPU's pinned digests; a device that runs +//! nothing, skips a case or gets one bit wrong cannot. +//! - **Placement, receipts and metering** are as for the WebGPU member: +//! [`place`](crate::quantum_fabric::place), +//! [`FabricReceipt`](crate::quantum_fabric::FabricReceipt) and a caller's +//! [`RunMeter`], whose window runs from the first upload to the last +//! readback. +//! +//! Every call is synchronous; the device is waited on before a call returns. +//! No type of the CUDA bindings appears in this module's public API. +//! +//! # Honest boundaries +//! +//! - **A member is what this process measured.** The probe and the +//! certificates are earned per process; the member id binds the device, its +//! UUID, the driver, the runtime compiler and the architecture. +//! - **Plans are per gate.** [`Config::PER_GATE`]: a tile kernel in shared +//! memory, as the WebGPU member has, is not here yet. +//! - **Unmetered by default**, as for every member. + +use crate::quantum::{Amp, Circuit, Gateset, MAX_QUBITS, ONE, QuantumError, StateVector}; +use crate::quantum_energy::EnergyClass; +use crate::quantum_fabric::cuda::{self, DeviceIdentity}; +use crate::quantum_fabric::lanes::{self, GateWords, Lane}; +use crate::quantum_fabric::plan::{self, Config, LaneRun, Program, Step}; +use crate::quantum_fabric::{ + Certificate, Certification, Corpus, EnergyCalibration, FabricClass, KernelEntry, Member, MemberError, Presence, RunFuture, RunMeter, Runner, Unmetered, Window, + apply_kernel, pair_updates, +}; +use cudarc::driver::{CudaContext, CudaFunction, CudaSlice, CudaStream, DriverError, LaunchConfig, PushKernelArg}; +use cudarc::nvrtc::{CompileOptions, compile_ptx_with_opts}; +use std::sync::{Arc, Mutex, OnceLock}; + +/// One probe case's device results: [`lanes::mul_wide`] and +/// [`lanes::fxmul_wide`], each as `(high, low)` words. +pub type ProbeOutput = ((u32, u32), (u32, u32)); + +/// A run's lanes and report, with the meter's accepted reading (zero and +/// `None` when unmetered). +type MeteredRun = (LaneRun, (u64, Option)); + +/// Why a device is not a member, or a run has no state. +#[derive(Debug, PartialEq, Eq)] +#[non_exhaustive] +pub enum CudaError { + /// The driver or the runtime compiler could not be used, and why. + Unavailable(String), + /// The runtime compiler refused the kernel for this device. + Compile(String), + /// The driver reported an error, or the device was lost. + Device(String), + /// Out of memory on the device or the host. + Memory(String), + /// Exactly what the CPU reference refuses. + Circuit(QuantumError), + /// More qubits than this member's memory holds. + TooWide { n_qubits: u8, max: u8 }, + /// The reference's result leaves 32 bits at this op; the member returns + /// no state. + OutOfRange { gate: u32 }, + /// A program the host cannot lay out (malformed words, a start outside the + /// state, a step this member does not run): a defect of the host, never a + /// result. + Internal(&'static str), + /// The device's arithmetic disagreed with the CPU's on this pair. + Probe { a: i32, b: i32 }, +} + +impl std::fmt::Display for CudaError { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + match self { + CudaError::Unavailable(why) => write!(f, "CUDA is unavailable: {why}"), + CudaError::Compile(e) => write!(f, "the kernel did not compile: {e}"), + CudaError::Device(e) => write!(f, "device: {e}"), + CudaError::Memory(e) => write!(f, "out of memory: {e}"), + CudaError::Circuit(e) => write!(f, "{e}"), + CudaError::TooWide { n_qubits, max } => write!(f, "{n_qubits} qubits; this member holds at most {max}"), + CudaError::OutOfRange { gate } => write!(f, "the result of op {gate} leaves 32 bits"), + CudaError::Internal(why) => write!(f, "host defect: {why}"), + CudaError::Probe { a, b } => write!(f, "the device's arithmetic differs from the CPU's at ({a}, {b})"), + } + } +} + +impl std::error::Error for CudaError {} + +impl From for MemberError { + fn from(e: CudaError) -> MemberError { + match e { + CudaError::TooWide { n_qubits, max } => MemberError::TooWide { n_qubits, max }, + CudaError::OutOfRange { gate } => MemberError::OutOfRange { gate }, + CudaError::Memory(why) => MemberError::Memory(why), + CudaError::Device(why) => MemberError::Device(why), + CudaError::Circuit(e) => MemberError::Circuit(e), + other => MemberError::Other(other.to_string()), + } + } +} + +/// A driver error as this module reports it: out of memory as +/// [`CudaError::Memory`], anything else as [`CudaError::Device`]. +fn driver(e: DriverError) -> CudaError { + if e.0 == cudarc::driver::sys::cudaError_enum::CUDA_ERROR_OUT_OF_MEMORY { + CudaError::Memory(e.to_string()) + } else { + CudaError::Device(e.to_string()) + } +} + +/// A device that is not a member, and why. +#[derive(Debug)] +pub struct Refusal { + /// The device's name, or its ordinal when the driver gave none. + pub device: String, + /// Its CI key ([`DeviceIdentity::key`]), when the device could be read. + pub key: String, + pub error: CudaError, +} + +/// What [`discover`] found. +#[derive(Debug)] +pub struct Discovery { + /// Devices that compiled the kernel and passed the probe, by ordinal. + pub found: Vec, + /// Devices that did not, with why. + pub refused: Vec, + /// Why no device was even listed: no driver, no runtime compiler, or the + /// driver failed to start. `None` when the driver listed its devices. + pub unavailable: Option, +} + +/// Whether the driver and the runtime compiler can be loaded in this process, +/// or why not. +pub fn available() -> Result<(), String> { + // SAFETY: each only tries to open a shared library by its known names. + if !unsafe { cudarc::driver::sys::is_culib_present() } { + return Err("no CUDA driver library on this host".into()); + } + // SAFETY: as above. + if !unsafe { cudarc::nvrtc::sys::is_culib_present() } { + return Err("no CUDA runtime compiler (NVRTC) library on this host".into()); + } + Ok(()) +} + +/// Every device the driver lists, each opened, compiled for and probed in this +/// process before it is listed. On a host without the driver or the runtime +/// compiler, nothing is found and [`Discovery::unavailable`] says why. +pub fn discover() -> Discovery { + let mut d = Discovery { found: Vec::new(), refused: Vec::new(), unavailable: None }; + if let Err(why) = available() { + d.unavailable = Some(why); + return d; + } + let count = match CudaContext::device_count() { + Ok(n) => usize::try_from(n).unwrap_or(0), + Err(e) => { + d.unavailable = Some(format!("the driver did not start: {e}")); + return d; + } + }; + for ordinal in 0..count { + match open(ordinal) { + Ok(m) => d.found.push(m), + Err(r) => d.refused.push(r), + } + } + d +} + +/// The target architecture as the runtime compiler's options take it: a +/// `&'static str`, one per distinct architecture for the life of the process. +fn arch_option(arch: String) -> &'static str { + static ARCHES: OnceLock>> = OnceLock::new(); + let mut arches = ARCHES.get_or_init(|| Mutex::new(Vec::new())).lock().unwrap_or_else(|p| p.into_inner()); + if let Some(a) = arches.iter().find(|a| **a == arch) { + return a; + } + let leaked: &'static str = Box::leak(arch.into_boxed_str()); + arches.push(leaked); + leaked +} + +/// The driver's CUDA version (`major·1000 + minor·10`), or 0. +fn driver_version() -> u32 { + let mut v: std::ffi::c_int = 0; + // SAFETY: writes one integer; the driver library is loaded. + let ok = unsafe { cudarc::driver::sys::cuDriverGetVersion(&mut v) } == cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS; + if ok { u32::try_from(v).unwrap_or(0) } else { 0 } +} + +/// The runtime compiler's version, or `(0, 0)`. +fn compiler_version() -> (u32, u32) { + let (mut major, mut minor): (std::ffi::c_int, std::ffi::c_int) = (0, 0); + // SAFETY: writes two integers; the compiler library is loaded. + let ok = unsafe { cudarc::nvrtc::sys::nvrtcVersion(&mut major, &mut minor) } == cudarc::nvrtc::sys::nvrtcResult::NVRTC_SUCCESS; + if ok { (u32::try_from(major).unwrap_or(0), u32::try_from(minor).unwrap_or(0)) } else { (0, 0) } +} + +/// The widest state of `2^n` interleaved 8-byte lanes that half of `free` +/// bytes holds, at most [`MAX_QUBITS`]. +fn qubits_for(free: u64) -> u8 { + let lanes = free / 2 / 8; + if lanes < 2 { 0 } else { (lanes.ilog2() as u8).min(MAX_QUBITS) } +} + +/// Reads the device at `ordinal`, compiles the kernel for it and probes it. +fn open(ordinal: usize) -> Result { + let unread = |error: CudaError| Refusal { device: format!("device {ordinal}"), key: String::new(), error }; + let ctx = CudaContext::new(ordinal).map_err(|e| unread(driver(e)))?; + let name = ctx.name().map_err(|e| unread(driver(e)))?; + let (major, minor) = ctx.compute_capability().map_err(|e| unread(driver(e)))?; + let uuid = ctx.uuid().map_err(|e| unread(driver(e)))?; + let attr = |a| ctx.attribute(a).map(|v| u32::try_from(v).unwrap_or(0)); + use cudarc::driver::sys::CUdevice_attribute::*; + let pci = match (attr(CU_DEVICE_ATTRIBUTE_PCI_DOMAIN_ID), attr(CU_DEVICE_ATTRIBUTE_PCI_BUS_ID), attr(CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID)) { + (Ok(d), Ok(b), Ok(v)) => format!("{d:04x}:{b:02x}:{v:02x}"), + _ => String::new(), + }; + let memory = ctx.total_mem().map_err(|e| unread(driver(e)))? as u64; + let identity = DeviceIdentity { + name: name.clone(), + uuid: uuid.bytes.iter().map(|b| format!("{:02x}", *b as u8)).collect(), + compute: (u32::try_from(major).unwrap_or(0), u32::try_from(minor).unwrap_or(0)), + pci, + driver: driver_version(), + compiler: compiler_version(), + memory, + }; + let key = identity.key(); + let refuse = |error: CudaError| Refusal { device: name.clone(), key: key.clone(), error }; + let kernel = Kernel::compile(&ctx, &identity, cuda::SOURCE).map_err(refuse)?; + let (free, _) = ctx.mem_get_info().map_err(|e| refuse(driver(e)))?; + let max_qubits = qubits_for(free as u64); + let detail = format!( + "cuda; device {name:?}; sm_{major}{minor}; uuid {}; pci {}; driver {}; runtime compiler {}.{}; memory {memory} bytes, {free} free at discovery; up to {max_qubits} qubits; a metered run's window runs from the first upload to the last readback", + identity.uuid, identity.pci, identity.driver, identity.compiler.0, identity.compiler.1, + ); + let member = CudaMember { + stream: ctx.default_stream(), + ctx, + kernel, + id: identity.member_id(), + key: key.clone(), + implementation: cuda::implementation_id(), + detail, + max_qubits, + identity, + }; + let pairs = lanes::probe_pairs(4096, 0x5eed); + match member.probe(&pairs) { + Ok(out) => match lanes::first_mismatch(&pairs, &out) { + None => Ok(member), + Some((a, b)) => Err(refuse(CudaError::Probe { a, b })), + }, + Err(e) => Err(refuse(e)), + } +} + +/// The kernel compiled for one device: its two entry points. +#[derive(Clone)] +struct Kernel { + apply: CudaFunction, + probe: CudaFunction, +} + +impl Kernel { + fn compile(ctx: &Arc, identity: &DeviceIdentity, source: &str) -> Result { + let opts = CompileOptions { + arch: Some(arch_option(identity.arch())), + options: cuda::COMPILE_OPTIONS.iter().map(|s| s.to_string()).collect(), + ..Default::default() + }; + let ptx = compile_ptx_with_opts(source, opts).map_err(|e| CudaError::Compile(e.to_string()))?; + let module = ctx.load_module(ptx).map_err(driver)?; + Ok(Kernel { apply: module.load_function("apply_gate").map_err(driver)?, probe: module.load_function("probe").map_err(driver)? }) + } +} + +/// The launch shape for `work` items: [`cuda::BLOCK`] threads per block and +/// at most [`cuda::GRID_MAX`] blocks, the kernel striding over the rest. +fn launch(work: u64) -> LaunchConfig { + let blocks = work.div_ceil(u64::from(cuda::BLOCK)).clamp(1, u64::from(cuda::GRID_MAX)) as u32; + LaunchConfig { grid_dim: (blocks, 1, 1), block_dim: (cuda::BLOCK, 1, 1), shared_mem_bytes: 0 } +} + +/// One CUDA device as a fabric member: a context on it, the kernel compiled +/// for it, and the arithmetic probe passed in this process. +#[derive(Clone)] +pub struct CudaMember { + ctx: Arc, + stream: Arc, + kernel: Kernel, + identity: DeviceIdentity, + id: String, + key: String, + implementation: String, + detail: String, + max_qubits: u8, +} + +impl std::fmt::Debug for CudaMember { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("CudaMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).field("max_qubits", &self.max_qubits).finish() + } +} + +impl CudaMember { + /// The member id ([`DeviceIdentity::member_id`]). + pub fn id(&self) -> &str { + &self.id + } + + /// The key a runner declares for it ([`DeviceIdentity::key`]). + pub fn key(&self) -> &str { + &self.key + } + + /// Always [`FabricClass::Gpu`]. + pub fn class(&self) -> FabricClass { + self.identity.class() + } + + /// What discovery found: device, architecture, driver, compiler, memory, + /// and what a metered run's window covers. Not signed. + pub fn detail(&self) -> &str { + &self.detail + } + + /// The most qubits a run holds on this member. + pub fn max_qubits(&self) -> u8 { + self.max_qubits + } + + /// The implementation id ([`cuda::implementation_id`], or a mutant's). + pub fn implementation(&self) -> &str { + &self.implementation + } + + /// What the member id hashes. + pub fn identity(&self) -> &DeviceIdentity { + &self.identity + } + + /// The device's [`lanes::mul_wide`] and [`lanes::fxmul_wide`] of each + /// pair, in order. + pub fn probe(&self, pairs: &[(i32, i32)]) -> Result, CudaError> { + if pairs.is_empty() { + return Ok(Vec::new()); + } + let flat: Vec = pairs.iter().flat_map(|&(a, b)| [a, b]).collect(); + let cases = self.stream.clone_htod(&flat).map_err(driver)?; + let mut out: CudaSlice = self.stream.alloc_zeros(pairs.len() * 4).map_err(driver)?; + let count = pairs.len() as u64; + let mut args = self.stream.launch_builder(&self.kernel.probe); + args.arg(&cases).arg(&mut out).arg(&count); + // SAFETY: `cases` holds `count` int2 cases and `out` `count` uint4 + // outputs, the kernel's parameters in order. + unsafe { args.launch(launch(count)) }.map_err(driver)?; + let words = self.stream.clone_dtoh(&out).map_err(driver)?; + Ok(words.as_chunks::<4>().0.iter().map(|w| ((w[0], w[1]), (w[2], w[3]))).collect()) + } + + /// The planning this member runs: [`Config::PER_GATE`]. + pub fn config(&self) -> Config { + Config::PER_GATE + } + + /// Runs `p` from its basis state: the lanes, and the first op whose exact + /// result left 32 bits. A program another config planned (one with tile + /// steps) is [`CudaError::Internal`]: this member runs gate steps only. + pub fn run_program(&self, p: &Program) -> Result { + Ok(self.run_program_metered(p, &mut Unmetered)?.0) + } + + fn run_program_metered(&self, p: &Program, meter: &mut M) -> Result { + let n = p.n_qubits; + if n == 0 { + return Ok((LaneRun { amps: vec![[ONE as i32, 0]], out_of_range: None }, (0, None))); + } + if n > self.max_qubits { + return Err(CudaError::TooWide { n_qubits: n, max: self.max_qubits }); + } + if n > MAX_QUBITS || p.start >= 1usize << n { + return Err(CudaError::Internal("the start is outside the state")); + } + let len = 1usize << n; + let mut gates = Vec::with_capacity(p.steps.len()); + for step in &p.steps { + let Step::Gate { gate } = *step else { return Err(CudaError::Internal("a tile step; this member runs gate steps only")) }; + let g = p.gates.get(gate as usize).ok_or(CudaError::Internal("a step names a gate the program does not hold"))?; + if !g.is_well_formed() { + return Err(CudaError::Internal("malformed gate words for this state")); + } + gates.push((&g.words, g.index)); + } + let mut init: Vec = Vec::new(); + init.try_reserve_exact(2 * len).map_err(|e| CudaError::Memory(e.to_string()))?; + init.resize(2 * len, 0); + init[2 * p.start] = ONE as i32; + self.execute(init, &gates, meter) + } + + /// Gates in order from `state`, as [`lanes::run_words`] runs them: the + /// lanes, and the lowest position in `gates` of a gate whose exact result + /// left 32 bits. A state that is not `2^n` lanes (`n ≥ 1`) or words + /// [`lanes::apply_words`] refuses for it are [`CudaError::Internal`]. + pub fn run_lanes(&self, state: &[Lane], gates: &[GateWords]) -> Result { + let len = state.len(); + if len < 2 || !len.is_power_of_two() { + return Err(CudaError::Internal("a state that is not 2^n lanes, n at least 1")); + } + let n = len.trailing_zeros() as u8; + if n > self.max_qubits { + return Err(CudaError::TooWide { n_qubits: n, max: self.max_qubits }); + } + let planned: Vec<(&GateWords, u32)> = gates.iter().zip(0u32..).collect(); + let mut init: Vec = Vec::new(); + init.try_reserve_exact(2 * len).map_err(|e| CudaError::Memory(e.to_string()))?; + init.extend(state.iter().flat_map(|l| [l[0], l[1]])); + Ok(self.execute(init, &planned, &mut Unmetered)?.0) + } + + /// Uploads `init` (interleaved lanes, `2·2^n` words), runs `gates` in order + /// inside a window of `meter`, and reads the report and the state back. + /// Every gate's words are checked against the state first: the words + /// [`lanes::apply_words`] refuses, the host refuses. + fn execute(&self, init: Vec, gates: &[(&GateWords, u32)], meter: &mut M) -> Result { + let len = init.len() / 2; + for (w, _) in gates { + 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 Err(CudaError::Internal("malformed gate words for this state")); + } + } + let pairs = (len / 2) as u64; + let mut window = Window::new(meter); + window.open(); + let mut amps = self.stream.clone_htod(&init).map_err(driver)?; + drop(init); + let mut report = self.stream.clone_htod(&[u32::MAX]).map_err(driver)?; + for &(w, index) in gates { + let m: [i32; 8] = w.m.map(|x| x as i32); + let mut args = self.stream.launch_builder(&self.kernel.apply); + args.arg(&mut amps).arg(&pairs).arg(&w.tbit).arg(&w.ctrl); + for v in &m { + args.arg(v); + } + args.arg(&index).arg(&mut report); + // SAFETY: `amps` holds `2·pairs` int2 lanes and `report` one word, + // and the words were checked against the state above. + unsafe { args.launch(launch(pairs)) }.map_err(driver)?; + } + let bad = self.stream.clone_dtoh(&report).map_err(driver)?; + let words = self.stream.clone_dtoh(&s).map_err(driver)?; + let reading = window.close(); + let out_of_range = bad.first().copied().filter(|&k| k != u32::MAX); + let mut lanes_out: Vec = Vec::new(); + lanes_out.try_reserve_exact(len).map_err(|e| CudaError::Memory(e.to_string()))?; + lanes_out.extend(words.as_chunks::<2>().0.iter().copied()); + Ok((LaneRun { amps: lanes_out, out_of_range }, reading)) + } + + /// Equal to `c.simulate_from(start)`, or a typed refusal: what the CPU + /// reference refuses ([`CudaError::Circuit`]), more qubits than the member + /// holds, or a result that leaves 32 bits. + pub fn simulate_from(&self, c: &Circuit, start: usize) -> Result { + Ok(self.simulate_with(c, start, &mut Unmetered)?.0) + } + + /// [`Self::simulate_from`] inside a window of `meter`: the state, and the + /// microjoules and class the meter read, which are zero and `None` unless + /// the crate's label check accepts them. The window opens just before the + /// first upload and closes after the last readback; planning and the + /// conversion to a state vector are outside it. + pub fn simulate_metered(&self, c: &Circuit, start: usize, meter: &mut dyn RunMeter) -> Result<(StateVector, u64, Option), CudaError> { + self.simulate_with(c, start, meter) + } + + fn simulate_with(&self, c: &Circuit, start: usize, meter: &mut M) -> Result<(StateVector, u64, Option), CudaError> { + let p = plan::compile_with(c, start, &self.config()).map_err(CudaError::Circuit)?; + if p.n_qubits > self.max_qubits { + return Err(CudaError::TooWide { n_qubits: p.n_qubits, max: self.max_qubits }); + } + let (run, (joules_micro, class)) = self.run_program_metered(&p, meter)?; + if let Some(gate) = run.out_of_range { + return Err(CudaError::OutOfRange { gate }); + } + let mut amps = Vec::new(); + amps.try_reserve_exact(run.amps.len()).map_err(|e| CudaError::Memory(e.to_string()))?; + amps.extend(run.amps.iter().map(|l| Amp { re: i64::from(l[0]), im: i64::from(l[1]) })); + Ok((StateVector { n_qubits: p.n_qubits, amps }, joules_micro, class)) + } + + /// [`Self::simulate_from`] the all-zeros state. + pub fn simulate(&self, c: &Circuit) -> Result { + self.simulate_from(c, 0) + } + + /// The certificate of [`apply_kernel`]`(gs)` as [`Self::implementation`] + /// on this member, over `corpus`: each circuit run from `|0…0⟩`, against + /// the native simulator. A range report inside a corpus is an error: its + /// circuits are valid and within the horizon. + pub fn certify(&self, gs: Gateset, corpus: Corpus) -> Result { + Ok(self.certify_metered(gs, corpus, &mut Unmetered)?.0) + } + + fn certify_metered(&self, gs: Gateset, corpus: Corpus, meter: &mut M) -> Result<(Certificate, Option<(u64, EnergyClass)>), CudaError> { + let mut cert = Certification::new(); + let mut calibration = EnergyCalibration::new(); + for c in corpus.circuits(gs) { + let reference = lanes::reference_output(&c).map_err(CudaError::Circuit)?; + let (run, (joules_micro, class)) = self.run_program_metered(&plan::compile_with(&c, 0, &self.config()).map_err(CudaError::Circuit)?, &mut *meter)?; + if let Some(gate) = run.out_of_range { + return Err(CudaError::OutOfRange { gate }); + } + cert.add(&lanes::circuit_bytes(&c), &reference, &lanes::output(&run.amps)); + calibration.add(pair_updates(&c), joules_micro, class.as_ref()); + } + Ok((cert.finish(apply_kernel(gs), &self.implementation, &self.id, &corpus.generator(gs)), calibration.finish())) + } + + /// A fabric [`Member`]: present as [`Presence::Detected`], with one entry + /// per gateset, each certified in this call on `corpus`, and no energy + /// figure. + pub fn member(&self, corpus: Corpus) -> Result { + self.member_with(corpus, &mut Unmetered) + } + + /// [`Self::member`], each corpus run inside its own window of `meter`: + /// every entry also carries the energy figure its runs earn + /// ([`EnergyCalibration`]), or none if any run went unmetered. + pub fn member_metered(&self, corpus: Corpus, meter: &mut dyn RunMeter) -> Result { + self.member_with(corpus, meter) + } + + fn member_with(&self, corpus: Corpus, meter: &mut M) -> Result { + let mut kernels = Vec::new(); + for gs in Gateset::ALL { + let (certificate, energy) = self.certify_metered(gs, corpus, &mut *meter)?; + kernels.push(KernelEntry { kernel: apply_kernel(gs).into(), implementation: self.implementation.clone(), certificate, energy }); + } + Ok(Member { id: self.id.clone(), class: self.class(), presence: Presence::Detected, kernels, detail: self.detail.clone() }) + } + + /// The same member holding at most `max_qubits` qubits (no more than it + /// holds), so that tests can make it refuse a circuit as too wide. + #[doc(hidden)] + pub fn with_max_qubits(&self, max_qubits: u8) -> CudaMember { + CudaMember { max_qubits: max_qubits.min(self.max_qubits), ..self.clone() } + } + + /// The same device running `source` in place of the kernel, under the + /// implementation id that text earns: for the mutation tests, which + /// require every mutant to be caught. + #[doc(hidden)] + pub fn with_source(&self, source: &str) -> Result { + let kernel = Kernel::compile(&self.ctx, &self.identity, source)?; + Ok(CudaMember { kernel, implementation: cuda::implementation_id_of(source), ..self.clone() }) + } +} + +/// The member on the fabric ([`Fabric`](crate::quantum_fabric::Fabric)): its +/// own implementation, through [`CudaMember::simulate_metered`]. Pair it with +/// a [`Member`] whose entries name that implementation, such as +/// [`CudaMember::member`]'s; another implementation is +/// [`MemberError::Other`]. +impl Runner for CudaMember { + fn member_id(&self) -> String { + self.id.clone() + } + + fn run<'r>(&'r self, implementation: &'r str, c: &'r Circuit, start: usize, meter: &'r mut dyn RunMeter) -> RunFuture<'r> { + Box::pin(async move { + if implementation != self.implementation { + return Err(MemberError::Other(format!("member {} runs implementation {}, not {implementation:?}", self.id, self.implementation))); + } + Ok(self.simulate_with(c, start, meter)?) + }) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn a_driver_out_of_memory_is_a_memory_error_and_every_other_a_device_error() { + use cudarc::driver::sys::cudaError_enum::*; + assert!(matches!(driver(DriverError(CUDA_ERROR_OUT_OF_MEMORY)), CudaError::Memory(_))); + assert!(matches!(driver(DriverError(CUDA_ERROR_LAUNCH_FAILED)), CudaError::Device(_))); + assert!(MemberError::from(CudaError::Memory("x".into())).falls_back()); + assert!(MemberError::from(CudaError::OutOfRange { gate: 3 }).falls_back()); + assert!(!MemberError::from(CudaError::Probe { a: 1, b: 2 }).falls_back()); + assert!(!MemberError::from(CudaError::Internal("x")).falls_back()); + } + + #[test] + fn the_state_a_member_holds_is_half_its_free_memory_at_most() { + assert_eq!(qubits_for(0), 0); + assert_eq!(qubits_for(31), 0); + assert_eq!(qubits_for(32), 1); + // 6 GiB free: 3 GiB of 8-byte lanes is 2^28.58 lanes, so 28 qubits, capped. + assert_eq!(qubits_for(6 << 30), MAX_QUBITS.min(28)); + assert_eq!(qubits_for(u64::MAX), MAX_QUBITS); + } + + #[test] + fn the_launch_shape_covers_the_work_and_caps_the_blocks() { + assert_eq!(launch(0).grid_dim.0, 1); + assert_eq!(launch(1).grid_dim.0, 1); + assert_eq!(launch(u64::from(cuda::BLOCK) + 1).grid_dim.0, 2); + assert_eq!(launch(1 << 40).grid_dim.0, cuda::GRID_MAX); + assert_eq!(launch(7).block_dim, (cuda::BLOCK, 1, 1)); + } + + #[test] + fn an_architecture_is_leaked_once_per_process() { + let a = arch_option("compute_89".into()); + let b = arch_option("compute_89".into()); + assert!(std::ptr::eq(a, b)); + assert_eq!(arch_option("compute_90".into()), "compute_90"); + } +} diff --git a/wai/wai-quantum/src/quantum_fabric_webgpu.rs b/wai/wai-quantum/src/quantum_fabric_webgpu.rs index e1246b334..508c6a824 100644 --- a/wai/wai-quantum/src/quantum_fabric_webgpu.rs +++ b/wai/wai-quantum/src/quantum_fabric_webgpu.rs @@ -423,33 +423,17 @@ fn lock(sink: &Mutex) -> std::sync::MutexGuard<'_, Sink> { sink.lock().unwrap_or_else(|p| p.into_inner()) } -/// The probe's operands: every pair of [`lanes::PROBE_EDGES`], then `random` -/// pairs of splitmix64 draws from `seed`. +/// The probe's operands ([`lanes::probe_pairs`]). #[doc(hidden)] pub fn probe_pairs(random: usize, seed: u64) -> Vec<(i32, i32)> { - let edges = lanes::PROBE_EDGES; - let mut pairs: Vec<(i32, i32)> = edges.iter().flat_map(|&a| edges.iter().map(move |&b| (a, b))).collect(); - let mut s = seed; - pairs.extend((0..random).map(|_| (splitmix(&mut s) as i32, splitmix(&mut s) as i32))); - pairs -} - -fn splitmix(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) + lanes::probe_pairs(random, seed) } /// The first pair whose device result differs from the CPU's -/// ([`lanes::mul_wide`], [`lanes::fxmul_wide`]). +/// ([`lanes::first_mismatch`]). #[doc(hidden)] pub fn first_mismatch(pairs: &[(i32, i32)], out: &[((u32, u32), (u32, u32))]) -> Option<(i32, i32)> { - if pairs.len() != out.len() { - return pairs.first().copied().or(Some((0, 0))); - } - pairs.iter().zip(out).find(|((a, b), o)| **o != (lanes::mul_wide(*a, *b), lanes::fxmul_wide(*a, *b))).map(|(p, _)| *p) + lanes::first_mismatch(pairs, out) } /// The three error scopes a step of work runs inside. The fields drop in diff --git a/wai/wai-quantum/tests/cuda_member.rs b/wai/wai-quantum/tests/cuda_member.rs new file mode 100644 index 000000000..1a75216cc --- /dev/null +++ b/wai/wai-quantum/tests/cuda_member.rs @@ -0,0 +1,429 @@ +//! The fabric's CUDA member on the devices of the machine that runs this. +//! +//! The machine declares the devices it must have in `WAI_QUANTUM_CUDA_EXPECT`: +//! their keys, comma-separated and sorted, or `none`. Every test that needs a +//! member fails unless the devices found are exactly the declared ones, so a +//! machine whose GPU or driver 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_CUDA_REPORT` names a file, appends the +//! same lines to it, which a CI step can count per device. +//! +//! ```text +//! WAI_QUANTUM_CUDA_EXPECT=gpu-cuda-sm89 \ +//! cargo test --features quantum_fabric_cuda --test cuda_member -- --test-threads=1 +//! ``` +#![cfg(not(target_arch = "wasm32"))] + +use std::io::Write; +use std::sync::OnceLock; +use wai_quantum::ed25519_dalek::SigningKey; +use wai_quantum::quantum::{BaseGate, Circuit, Gate, Gateset, MAX_QUBITS, QuantumError}; +use wai_quantum::quantum_fabric::cuda::{self, MUTANTS}; +use wai_quantum::quantum_fabric::lanes::{self, GateWords, Lane}; +use wai_quantum::quantum_fabric::plan::{Program, Step}; +use wai_quantum::quantum_fabric::wgsl::{Detector, Mutant}; +use wai_quantum::quantum_fabric::{ + Corpus, CpuReference, FABRIC_DEFAULT, Fabric, FabricClass, FabricReceipt, MemberError, Parity, Presence, REFERENCE, apply_kernel, detect_local, place, ready, +}; +use wai_quantum::quantum_fabric_cuda::{CudaError, CudaMember, Discovery, discover}; + +/// 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 = OnceLock::new(); + +fn discovery() -> &'static Discovery { + DISCOVERY.get_or_init(discover) +} + +/// Prints `line` and appends it to the file `WAI_QUANTUM_CUDA_REPORT` names. +fn report(line: &str) { + println!("{line}"); + if let Some(path) = std::env::var_os("WAI_QUANTUM_CUDA_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 device 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 { + 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 = d.refused.iter().map(|r| format!("{} ({:?}): {}", r.key, r.device, r.error)).collect(); + let refused_list = if refused.is_empty() { "none".to_string() } else { refused.join("; ") }; + let unavailable = d.unavailable.as_deref().unwrap_or("no"); + let how = format!("found: {found_list}; refused: {refused_list}; unavailable: {unavailable}. Declare the devices this machine must have: WAI_QUANTUM_CUDA_EXPECT={found_list}"); + let expect = match std::env::var("WAI_QUANTUM_CUDA_EXPECT") { + Ok(v) if !v.is_empty() => v, + _ => return Err(format!("WAI_QUANTUM_CUDA_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_CUDA_EXPECT={expect:?} is not a sorted, comma-separated list of keys, nor none. {how}")); + } + for k in &keys { + if !found.contains(k) + && let Some(r) = d.refused.iter().find(|r| r.key == *k) + { + return Err(format!("the declared device {k} was refused: {}. {how}", r.error)); + } + } + if found != keys { + return Err(format!("the devices found are not the declared ones ({expect}). {how}")); + } + Ok(expect) +} + +/// The members, once the declared devices are exactly the ones found. +fn members() -> &'static [CudaMember] { + let d = discovery(); + if let Err(why) = declared(d) { + panic!("{why}"); + } + &d.found +} + +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: 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 { + let gate = |base, controls: Vec, 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 c0_devices_are_declared() { + let d = discovery(); + match declared(d) { + Ok(expect) => report(&format!("cuda devices: as declared ({expect})")), + Err(why) => panic!("{why}"), + } + for m in &d.found { + assert_eq!(m.class(), FabricClass::Gpu); + assert!(m.id().starts_with("local/gpu-cuda-"), "{}", m.id()); + assert_eq!(m.implementation(), cuda::implementation_id()); + assert_eq!(m.identity().member_id(), m.id()); + assert!(m.max_qubits() >= 20, "{}: holds only {} qubits", m.key(), m.max_qubits()); + report(&format!("cuda member: key={} id={} {}", m.key(), m.id(), m.detail())); + } +} + +#[test] +fn c1_probe_matches_lanes() { + for m in members() { + let pairs = lanes::probe_pairs(1 << 16, 0x0cda); + let out = m.probe(&pairs).unwrap_or_else(|e| panic!("{}: {e}", m.key())); + assert_eq!(lanes::first_mismatch(&pairs, &out), None, "{}", m.key()); + report(&format!("cuda probe: key={} {} pairs agree", m.key(), pairs.len())); + } +} + +#[test] +fn c2_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 = (0..1usize << n).map(|_| [value(&mut s), value(&mut s)]).collect(); + let gates: Vec = (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).is_multiple_of(4)).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 = m.run_lanes(&state, &gates).unwrap_or_else(|e| panic!("{}: case {case}: {e}", m.key())); + 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!("cuda kernel: key={} ok", m.key())); + } +} + +#[test] +fn c3_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 = m.certify(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)); + report(&format!( + "cuda 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 = m.certify(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 c4_wide_circuits_match_native() { + for m in members() { + let widest = m.max_qubits().min(22); + let mut s = 0x0077_6964_6500_u64; + for (k, n) in [12u8, 16, 18, 20, widest].into_iter().enumerate() { + for gs in Gateset::ALL { + let c = random_circuit(n, 40, 0x1000 + k as u64, gs); + let start = (next(&mut s) as usize) % (1usize << n); + let want = c.simulate_from(start).expect("the native run"); + let got = m.simulate_from(&c, start).unwrap_or_else(|e| panic!("{}: n = {n} {gs:?}: {e}", m.key())); + assert!(got == want, "{}: n = {n} {gs:?}: the states differ", m.key()); + } + } + // The transforms the conformance corpus carries. + for n in [3u8, 5, 8, 12] { + let qft = Circuit::qft(n); + assert!(m.simulate(&qft).unwrap() == qft.simulate().unwrap(), "{}: QFT-{n}", m.key()); + } + report(&format!("cuda wide: key={} up to {widest} qubits match native", m.key())); + } +} + +#[test] +fn c5_refusals_are_typed() { + for m in members() { + for c in refused_circuits() { + let native = c.simulate_from(0).unwrap_err(); + assert_eq!(m.simulate_from(&c, 0).unwrap_err(), CudaError::Circuit(native), "{c:?}"); + } + assert_eq!(m.simulate(&Circuit::new(MAX_QUBITS + 1)).unwrap_err(), CudaError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1))); + let small = m.with_max_qubits(4); + assert_eq!(small.simulate(&Circuit::new(5)).unwrap_err(), CudaError::TooWide { n_qubits: 5, max: 4 }); + let wide = Program { n_qubits: 5, start: 0, gates: Vec::new(), steps: Vec::new() }; + assert_eq!(small.run_program(&wide).unwrap_err(), CudaError::TooWide { n_qubits: 5, max: 4 }); + // A result that leaves 32 bits is reported at its op, and the lane + // holds its low 32 bits. + let r = m.run_lanes(&Z_STATE, &[z_words()]).unwrap(); + assert_eq!(r.out_of_range, Some(0), "{}", m.key()); + assert_eq!(r.amps[1][0] as u32, 0x8000_0000); + // Malformed words, a state the host cannot lay out and a tile step are + // host defects, never results. + let bad = GateWords { tbit: 3, ..z_words() }; + assert!(matches!(m.run_lanes(&[[0, 0]; 4], &[bad]), Err(CudaError::Internal(_)))); + assert!(matches!(m.run_lanes(&[[0, 0]; 3], &[z_words()]), Err(CudaError::Internal(_)))); + let tiled = Program { n_qubits: 2, start: 0, gates: Vec::new(), steps: vec![Step::Tile { first: 0, count: 1, size: 2 }] }; + assert!(matches!(m.run_program(&tiled), Err(CudaError::Internal(_)))); + // The refusals that may still fall back to another member, do. + assert!(MemberError::from(CudaError::TooWide { n_qubits: 5, max: 4 }).falls_back()); + report(&format!("cuda refusals: key={} typed", m.key())); + } +} + +/// Whether a member running `mutant`'s text is caught by the mutant's +/// detector. +fn caught(mm: &CudaMember, mutant: &Mutant, key: &str) -> bool { + match mutant.detector { + Detector::Probe => { + let edges = lanes::probe_pairs(0, 0); + let out = mm.probe(&edges).expect("the mutant's probe runs"); + let differing = edges.iter().zip(&out).filter(|&(&(a, b), o)| *o != (lanes::mul_wide(a, b), lanes::fxmul_wide(a, b))).count(); + println!("cuda mutant {} on {key}: {differing} of 484 edge pairs differ", mutant.name); + differing >= 1 + } + Detector::Certificate => Gateset::ALL.into_iter().all(|gs| match mm.certify(gs, RANDOM_V1) { + Err(e) => { + println!("cuda mutant {} on {key}: {gs:?} certification refused: {e}", mutant.name); + true + } + Ok(c) => { + let cpu = lanes::certify_corpus(gs, "local/test", RANDOM_V1); + c.parity != Parity::BitExact || c.corpus != cpu.corpus + } + }), + Detector::Range => mm.run_lanes(&Z_STATE, &[z_words()]).map(|r| r.out_of_range) != Ok(Some(0)), + other => panic!("no CUDA mutant uses the {other:?} detector"), + } +} + +#[test] +fn c6_mutants_are_killed() { + assert_eq!(lanes::probe_pairs(0, 0).len(), 484); + for m in members() { + let mut killed = 0; + for mutant in &MUTANTS { + let mm = m.with_source(&cuda::mutate(mutant)).unwrap_or_else(|e| panic!("{}: mutant {} does not compile: {e}", m.key(), mutant.name)); + assert_ne!(mm.implementation(), m.implementation()); + assert!(caught(&mm, mutant, m.key()), "{}: mutant {} survived", m.key(), mutant.name); + killed += 1; + } + assert_eq!(killed, MUTANTS.len()); + report(&format!("cuda mutants: key={} killed {killed} of {}", m.key(), MUTANTS.len())); + } +} + +#[test] +fn c7_receipt_placement_and_the_fabric() { + for m in members() { + let member = m.member(FABRIC_DEFAULT).unwrap_or_else(|e| panic!("{}: {e}", m.key())); + assert_eq!((member.id.as_str(), member.class, member.presence), (m.id(), FabricClass::Gpu, 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:?}"); + } + // A receipt over a member's run. + let qft = Circuit::qft(5); + let state = m.simulate(&qft).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)); + // Placement: beside the CPU reference with no energy figures, 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())); + } + // On the fabric: the member alone runs the circuit; a pairing with + // another member's entries is refused. + let mut fabric = Fabric::new(); + fabric.add(member.clone(), m); + let run = ready(fabric.run(&qft, 0, &Parity::BitExact, &signer, "local/test")).expect("resolves").expect("runs"); + assert_eq!(run.state, state); + assert_eq!(run.receipt.member_id, m.id()); + assert!(run.receipt.verify()); + let mut mismatched = Fabric::new(); + mismatched.add(local.clone(), m); + assert!(ready(mismatched.run(&qft, 0, &Parity::BitExact, &signer, "local/test")).expect("resolves").is_err()); + let mut both = Fabric::new(); + both.add(local.clone(), CpuReference).add(member.clone(), m); + let run = ready(both.run(&qft, 0, &Parity::BitExact, &signer, "local/test")).expect("resolves").expect("runs"); + assert_eq!(run.state, state); + report(&format!("cuda receipt: key={} ok", m.key())); + } +} + +#[test] +fn c8_repeated_runs_agree() { + for m in members() { + let c = random_circuit(16, 60, 0xabc, Gateset::V2); + let first = m.simulate(&c).unwrap(); + for _ in 0..5 { + assert!(m.simulate(&c).unwrap() == first, "{}", m.key()); + } + report(&format!("cuda repeat: key={} ok", m.key())); + } +} diff --git a/wai/wai-rs/Cargo.lock b/wai/wai-rs/Cargo.lock index 721e47a05..261d05f7e 100644 --- a/wai/wai-rs/Cargo.lock +++ b/wai/wai-rs/Cargo.lock @@ -4459,7 +4459,7 @@ dependencies = [ [[package]] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-web/Cargo.lock b/wai/wai-web/Cargo.lock index 9862ac0f3..65ed3a854 100644 --- a/wai/wai-web/Cargo.lock +++ b/wai/wai-web/Cargo.lock @@ -1226,7 +1226,7 @@ dependencies = [ [[package]] name = "wai-quantum" -version = "0.4.3" +version = "0.4.4" dependencies = [ "blake3", "ed25519-dalek", -- 2.54.0 (Apple Git-157)