From 540c1670b7bd4b6d23836e825ed2fce06830094e Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 01:49:20 -0400 Subject: [PATCH 01/19] wai-quantum 0.5.2: the dataflow device's vector engine as a member On the device generation whose vector engine multiplies 23-bit numbers (PCI 1e52:b140), a second dataflow member, dataflow-1e52-b140-vector, runs the gate's arithmetic on that engine, 32 lanes at a time, on 32 compute tiles. - Each tile's first core writes the unpacker's and the packer's configuration, then hands the tile its instructions inline (a word with low bits other than 0b11 in its own instruction stream): the unpacker copies the pairs' columns into the destination registers as 32-bit integers, the vector engine forms each pair's four rows, the packer copies them back, and the core signals when the packer is done. - Exact arithmetic: each operand split into 23 and 9 bits; four partial products give the unsigned high word, corrected for sign to P_hi, and the bits of the low word the rounding reaches; fx = P_hi * 2^(32-s) + t, held as 30 low bits and the rest; each row adds or subtracts its terms with a carry after each. The three mutants are the other members'. - quantum_fabric::tensix: the instructions used, transcribed from the device simulator's own table; CI checks the transcription against it. quantum_fabric::vector_engine builds the programs (rows and probe); the assembler check covers them (14 programs, identical words). - A pass is sized to its share of the pairs in steps of 64 lanes: one iteration covers the even or the odd columns of four destination rows, so two cover the next 64 pairs in order. Measured on the simulator, built from source: 66020 probe pairs agree, 2000 kernel cases match run_words, all four pinned certificates, circuits up to 20 qubits match native, 3 of 3 mutants killed; the instruction table matches for all 20 instructions; the other generation's cores member is unchanged (9 of 9). Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-dataflow.yml | 15 +- wai/CHANGELOG.md | 22 ++ wai/tools/quantum-meter/Cargo.lock | 2 +- wai/tools/quantum-wgsl-check/Cargo.lock | 2 +- wai/wai-pack/Cargo.lock | 2 +- wai/wai-quantum-qpu/Cargo.lock | 2 +- wai/wai-quantum-serve/Cargo.lock | 2 +- wai/wai-quantum/Cargo.lock | 2 +- wai/wai-quantum/Cargo.toml | 2 +- wai/wai-quantum/README.md | 2 +- wai/wai-quantum/examples/rv32_listing.rs | 11 + wai/wai-quantum/src/quantum_fabric.rs | 4 + wai/wai-quantum/src/quantum_fabric/rv32.rs | 15 + wai/wai-quantum/src/quantum_fabric/tensix.rs | 147 ++++++++ .../src/quantum_fabric/vector_engine.rs | 315 ++++++++++++++++++ .../src/quantum_fabric_dataflow.rs | 228 +++++++++++-- wai/wai-quantum/tests/dataflow_member.rs | 6 +- wai/wai-rs/Cargo.lock | 2 +- wai/wai-web/Cargo.lock | 2 +- 19 files changed, 733 insertions(+), 50 deletions(-) create mode 100644 wai/wai-quantum/src/quantum_fabric/tensix.rs create mode 100644 wai/wai-quantum/src/quantum_fabric/vector_engine.rs diff --git a/.github/workflows/wai-quantum-dataflow.yml b/.github/workflows/wai-quantum-dataflow.yml index c0f66523b..0873f83c1 100644 --- a/.github/workflows/wai-quantum-dataflow.yml +++ b/.github/workflows/wai-quantum-dataflow.yml @@ -65,7 +65,7 @@ jobs: include: - pci: '1e52:b140' build: bh - key: dataflow-1e52-b140 + key: dataflow-1e52-b140,dataflow-1e52-b140-vector - pci: '1e52:401e' build: wh key: dataflow-1e52-401e @@ -115,7 +115,10 @@ jobs: git -C "$b" checkout -q FETCH_HEAD [ "$(git -C "$b" rev-parse HEAD)" = "$SIM_COMMIT" ] || { echo "::error::the simulator is not at the pinned commit"; exit 1; } (cd "$b" && ./make.py "src/_out/release_${{ matrix.build }}/libttsim.so" > build.log 2>&1) || { tail -30 "$b/build.log"; exit 1; } - mkdir -p "$(dirname "$lib")" && cp "$b/src/_out/release_${{ matrix.build }}/libttsim.so" "$lib.tmp" && mv "$lib.tmp" "$lib" + mkdir -p "$(dirname "$lib")" + # The instruction table too: the vector programs' encoder is checked against it. + cp "$b/data/${{ matrix.build }}/tensix_isa.json" "$(dirname "$lib")/tensix_isa.json" + cp "$b/src/_out/release_${{ matrix.build }}/libttsim.so" "$lib.tmp" && mv "$lib.tmp" "$lib" rm -rf "$b" fi export RUSTUP_HOME="$cache/rustup" CARGO_HOME="$cache/cargo" @@ -145,7 +148,7 @@ jobs: cmp "$p.bin" "$p.as.bin" || { echo "::error::$(basename "$p"): the encoder's words differ from the assembler's"; exit 1; } n=$((n + 1)) done - [ "$n" = 7 ] || { echo "::error::$n of 7 programs were compared"; exit 1; } + [ "$n" = 14 ] || { echo "::error::$n of 14 programs were compared"; exit 1; } echo "all $n programs: the encoder's words are the assembler's" - name: The member's tests shell: bash @@ -158,8 +161,12 @@ jobs: export RUSTUP_HOME="$cache/rustup" CARGO_HOME="$cache/cargo" CARGO_TARGET_DIR="$cache/target" export PATH="$CARGO_HOME/bin:$PATH" export WAI_QUANTUM_DATAFLOW_LIB="$cache/sim-$SIM_COMMIT/${{ matrix.build }}/libttsim.so" - cargo test --locked --release -j 6 --features quantum_fabric_dataflow --lib -- quantum_fabric_dataflow:: quantum_fabric::rv32:: 2>&1 | tee "$RUNNER_TEMP/unit.log" + export WAI_QUANTUM_TENSIX_ISA="$cache/sim-$SIM_COMMIT/${{ matrix.build }}/tensix_isa.json" + cargo test --locked --release -j 6 --features quantum_fabric_dataflow --lib -- quantum_fabric_dataflow:: quantum_fabric::rv32:: quantum_fabric::tensix:: quantum_fabric::vector_engine:: --nocapture 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; } + # The vector engine's encoder against the simulator's own table (the + # test returns early without the table, so require its line). + grep -qE '^tensix layouts: [0-9]+ instructions match the table' "$RUNNER_TEMP/unit.log" || { echo "::error::the instruction table was not checked"; exit 1; } report="$RUNNER_TEMP/dataflow-report.txt"; : > "$report" WAI_QUANTUM_DATAFLOW_REPORT="$report" cargo test --locked --release -j 6 --features quantum_fabric_dataflow --test dataflow_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; } diff --git a/wai/CHANGELOG.md b/wai/CHANGELOG.md index 24a33e6d3..23ecd492f 100644 --- a/wai/CHANGELOG.md +++ b/wai/CHANGELOG.md @@ -6,6 +6,28 @@ wire format may change and such changes are recorded here. ## v0.4.0 — current +### wai-quantum 0.5.2: the dataflow device's vector engine + +Additive. The default features, `full`, and every pinned digest are unchanged. + +- **The vector engine as a member** (`quantum_fabric_dataflow`, key `dataflow-1e52-b140-vector`): + on the device generation whose vector engine multiplies 23-bit numbers, the gate's + arithmetic runs on that engine, 32 lanes at a time, on 32 compute tiles. Each tile's + first core configures the unpacker and the packer and hands the tile its instructions + inline; the unpacker copies the pairs' columns into the destination registers as + 32-bit integers, the vector engine forms the rows, the packer copies them back. Each + 32×32 product comes from four 23-bit partial products, corrected for sign; `fx` is + held as 30 low bits and the rest, and each row's terms are added or subtracted with a + carry after each. A pass is sized to its share of the pairs, in steps of 64. +- **`quantum_fabric::tensix`**: the compute tile's instructions the programs use, their + opcodes and fields as the device simulator's own table gives them; CI checks the + transcription against that table. **`quantum_fabric::vector_engine`** builds the + programs, the probe included; the assembler check covers them too. +- **Measured** on the device simulator, built from source: 66020 probe pairs agree, all + four pinned certificates, circuits up to 20 qubits match native, 3 of 3 kernel mutants + (the same three the other members carry) killed. Every certificate is labelled + simulator. + ### wai-quantum 0.5.1: neural-processor, matrix-extension, vector-unit, signal-processor and dataflow members Additive. The default features, `full`, and every pinned digest are unchanged. diff --git a/wai/tools/quantum-meter/Cargo.lock b/wai/tools/quantum-meter/Cargo.lock index 2e1ea15ea..cd24b9ed9 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.5.1" +version = "0.5.2" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/tools/quantum-wgsl-check/Cargo.lock b/wai/tools/quantum-wgsl-check/Cargo.lock index f99c82d2c..60e84a45d 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.5.1" +version = "0.5.2" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-pack/Cargo.lock b/wai/wai-pack/Cargo.lock index ba68033a7..55ab937cf 100644 --- a/wai/wai-pack/Cargo.lock +++ b/wai/wai-pack/Cargo.lock @@ -703,7 +703,7 @@ dependencies = [ [[package]] name = "wai-quantum" -version = "0.5.1" +version = "0.5.2" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-quantum-qpu/Cargo.lock b/wai/wai-quantum-qpu/Cargo.lock index 716a80636..1b3261977 100644 --- a/wai/wai-quantum-qpu/Cargo.lock +++ b/wai/wai-quantum-qpu/Cargo.lock @@ -422,7 +422,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" [[package]] name = "wai-quantum" -version = "0.5.1" +version = "0.5.2" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-quantum-serve/Cargo.lock b/wai/wai-quantum-serve/Cargo.lock index cb7ca6cb4..15d5d9cbe 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.5.1" +version = "0.5.2" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-quantum/Cargo.lock b/wai/wai-quantum/Cargo.lock index 78db5d69b..c4a208df5 100644 --- a/wai/wai-quantum/Cargo.lock +++ b/wai/wai-quantum/Cargo.lock @@ -1079,7 +1079,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" [[package]] name = "wai-quantum" -version = "0.5.1" +version = "0.5.2" dependencies = [ "blake3", "block2", diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml index 5771600a7..4e721a632 100644 --- a/wai/wai-quantum/Cargo.toml +++ b/wai/wai-quantum/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "wai-quantum" -version = "0.5.1" +version = "0.5.2" edition = "2024" rust-version = "1.97" license = "Apache-2.0" diff --git a/wai/wai-quantum/README.md b/wai/wai-quantum/README.md index e9d6a06b8..2c8c0c998 100644 --- a/wai/wai-quantum/README.md +++ b/wai/wai-quantum/README.md @@ -112,7 +112,7 @@ never creates information. | `quantum_fabric_sme` | **the fabric's matrix-extension member**, opt-in and not in `full`, aarch64 only: the gate's fixed-point arithmetic in the processor's streaming vector mode (SME), 64-bit lanes, the host gathering each gate's pairs and checking each row's range. Every instruction it uses exists only in that mode, so a kernel that ran, ran there. Detected by the processor's feature flags (on Linux, the kernel's `AT_HWCAP2`, which a user-mode emulator supplies); under an emulator the runner says so (`WAI_QUANTUM_SME_EMULATOR`) and every certificate carries it. Three kernel mutants are compiled in and must be caught | | `quantum_fabric_dsp` | **the fabric's DSP vector-unit member**, opt-in and not in `full`, in a build for the signal processor with its vector unit in 128-byte mode: the gate's fixed-point arithmetic in 32-bit lanes, each 64-bit value a pair of words, every product the vector unit's exact 32×32 multiply and every carry and borrow kept; the host gathers each gate's pairs and checks each row's range. Its instructions exist only on that unit, so a kernel that ran, ran there. Under an emulator the runner says so (`WAI_QUANTUM_DSP_EMULATOR`) and every certificate carries it. Three kernel mutants are compiled in and must be caught. Its class is `dsp` | | `quantum_fabric_vector` | **the fabric's CPU vector-unit members**, opt-in and not in `full`: the gate's fixed-point arithmetic in the processor's own vector extension, one kernel per instruction set — `avx2` (x86-64, 256-bit), `sve` (aarch64, the scalable vector extension, any length) and `rvv` (riscv64, the vector extension 1.0, any length) — each operand sign-extended to 64-bit lanes, the host gathering each gate's pairs and checking each row's range. The scalable kernels set their length per pass; the member reads the length from the unit and names it in its id and detail. Detected by the processor's feature flags (on Linux, `AT_HWCAP` for the vector extension, which a user-mode emulator supplies); under an emulator the runner says so (`WAI_QUANTUM_VECTOR_EMULATOR`). Its id extends the CPU reference's, so placement with no energy figures keeps the reference. Three kernel mutants per instruction set are compiled in and must be caught | -| `quantum_fabric_dataflow` | **the fabric's dataflow member**, opt-in and not in `full`, Linux: the gate's fixed-point arithmetic on the first RISC-V core of each of 32 compute tiles of a dataflow accelerator, each over a contiguous slice of a gate's pairs. The program is built in this crate (`quantum_fabric::rv32`, an RV32IM encoder whose listings an independent assembler reproduces word for word): every 64-bit value a pair of words, `mul`/`mulh` for products, `sltu` for carries and borrows. The device's library (`WAI_QUANTUM_DATAFLOW_LIB`, loaded at run time) presents the device as a PCI endpoint; the host opens two memory windows on each tile, holds its cores in reset, writes the program at the reset address and that tile's parameters and slice beside it, releases the first cores and clocks the device until every tile reports done. The device is named by its PCI ids (`dataflow-1e52-`); the library is a simulator, and every certificate says so (`WAI_QUANTUM_DATAFLOW_SIMULATOR`). Three kernel mutants are compiled in | +| `quantum_fabric_dataflow` | **the fabric's dataflow member**, opt-in and not in `full`, Linux: the gate's fixed-point arithmetic on the first RISC-V core of each of 32 compute tiles of a dataflow accelerator, each over a contiguous slice of a gate's pairs. The program is built in this crate (`quantum_fabric::rv32`, an RV32IM encoder whose listings an independent assembler reproduces word for word): every 64-bit value a pair of words, `mul`/`mulh` for products, `sltu` for carries and borrows. The device's library (`WAI_QUANTUM_DATAFLOW_LIB`, loaded at run time) presents the device as a PCI endpoint; the host opens two memory windows on each tile, holds its cores in reset, writes the program at the reset address and that tile's parameters and slice beside it, releases the first cores and clocks the device until every tile reports done. The device is named by its PCI ids (`dataflow-1e52-`); the library is a simulator, and every certificate says so (`WAI_QUANTUM_DATAFLOW_SIMULATOR`). Three kernel mutants are compiled in. On the device generation whose vector engine multiplies 23-bit numbers, a second member (`-vector`) runs the gate on that engine, 32 lanes at a time: the first core configures the tile's unpacker and packer and hands the tile its instructions inline (`quantum_fabric::tensix`, transcribed from the simulator's own table and checked against it; `quantum_fabric::vector_engine`), each 32×32 product from four 23-bit partial products | | `full` | everything | Default features are `quantum`, `quantum_toolchain`, `quantum_receipt`. diff --git a/wai/wai-quantum/examples/rv32_listing.rs b/wai/wai-quantum/examples/rv32_listing.rs index cd5eb2256..f988aa541 100644 --- a/wai/wai-quantum/examples/rv32_listing.rs +++ b/wai/wai-quantum/examples/rv32_listing.rs @@ -9,6 +9,7 @@ //! //! In ``, for each program ``: `.s` (the listing, with a //! `.text` header) and `.bin` (the words, little-endian). +use wai_quantum::quantum_fabric::vector_engine::{self, Constants}; use wai_quantum::quantum_fabric_dataflow::{Variant, programs}; fn main() { @@ -24,6 +25,16 @@ fn main() { all.push((format!("probe{tag}"), probe)); } } + // The vector engine's, over a full pass, with a sample gate's entries. + let m = [0x2000_0000, -0x2000_0000, 0x1234_5678, -0x0765_4321, 1, -1, i32::MAX, i32::MIN]; + for (tag, v) in [("", Variant::Conforming), ("_m1", Variant::M1), ("_m2", Variant::M2), ("_m3", Variant::M3)] { + let (rnd, shift, last_add) = v.constants(); + let c = Constants { round_bit: rnd.trailing_zeros(), shift, last_add }; + all.push((format!("vector_rows{tag}"), vector_engine::rows_program(&m, c, 16))); + if v != Variant::M3 { + all.push((format!("vector_probe{tag}"), vector_engine::probe_program(c, 16))); + } + } for (name, (words, text)) in all { std::fs::write(dir.join(format!("{name}.s")), format!(".text\n.option norvc\n{text}")).expect("the listing"); std::fs::write(dir.join(format!("{name}.bin")), words.iter().flat_map(|w| w.to_le_bytes()).collect::>()).expect("the words"); diff --git a/wai/wai-quantum/src/quantum_fabric.rs b/wai/wai-quantum/src/quantum_fabric.rs index 9ec9d1bb3..89ae10b3d 100644 --- a/wai/wai-quantum/src/quantum_fabric.rs +++ b/wai/wai-quantum/src/quantum_fabric.rs @@ -2054,6 +2054,10 @@ pub mod opencl; pub mod plan; #[cfg(feature = "quantum_fabric_dataflow")] pub mod rv32; +#[cfg(feature = "quantum_fabric_dataflow")] +pub mod tensix; +#[cfg(feature = "quantum_fabric_dataflow")] +pub mod vector_engine; pub mod sampling; #[cfg(feature = "quantum_sv")] pub mod sv64; diff --git a/wai/wai-quantum/src/quantum_fabric/rv32.rs b/wai/wai-quantum/src/quantum_fabric/rv32.rs index ff5eda95d..bbe67b617 100644 --- a/wai/wai-quantum/src/quantum_fabric/rv32.rs +++ b/wai/wai-quantum/src/quantum_fabric/rv32.rs @@ -41,6 +41,7 @@ pub const A5: R = R(15); pub const A6: R = R(16); pub const A7: R = R(17); pub const S2: R = R(18); +pub const S3: R = R(19); pub const T3: R = R(28); pub const T4: R = R(29); pub const T5: R = R(30); @@ -162,6 +163,20 @@ impl Asm { self.branch(0, "beq", rs1, rs2, label) } + pub fn bne(&mut self, rs1: R, rs2: R, label: &'static str) -> &mut Self { + self.branch(1, "bne", rs1, rs2, label) + } + + /// A word for the tile, carried inline: the core passes on any word in its + /// instruction stream whose low two bits are not `0b11`, rotated right by + /// two, so the stream holds the word rotated left by two. `what` names it + /// in the listing. + pub fn tt(&mut self, word: u32, what: &str) -> &mut Self { + let inline = word.rotate_left(2); + assert_ne!(inline & 3, 3, "{word:#010x} would be a RISC-V instruction"); + self.emit(inline, format!(".word {inline:#010x} # {what}")) + } + /// `jal zero, label`. pub fn j(&mut self, label: &'static str) -> &mut Self { self.fixups.push(Fixup { at: self.words.len(), label, f3: u32::MAX, rs1: ZERO, rs2: ZERO }); diff --git a/wai/wai-quantum/src/quantum_fabric/tensix.rs b/wai/wai-quantum/src/quantum_fabric/tensix.rs new file mode 100644 index 000000000..f7be541d0 --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/tensix.rs @@ -0,0 +1,147 @@ +//! The instructions of a dataflow accelerator's compute tile that the vector +//! member's programs use: their opcodes and fields, transcribed from the device +//! simulator's own instruction table, and checked against that table in CI. +//! +//! A word is `opcode << 24 | fields`. The tile's first RISC-V core passes a +//! word on to the tile when it meets one in its own instruction stream whose +//! low two bits are not `0b11`, rotated right by two: so a program carries +//! [`inline`] of the word. + +/// One instruction: its name, opcode, and fields as (name, high bit, low bit). +pub struct Layout { + pub name: &'static str, + pub opcode: u8, + pub fields: &'static [(&'static str, u8, u8)], +} + +/// The instructions used, as the simulator's table gives them. +pub const LAYOUTS: &[Layout] = &[ + Layout { name: "INCRWC", opcode: 0x38, fields: &[("rwc_a", 9, 6), ("rwc_b", 13, 10), ("rwc_cr", 20, 18), ("rwc_d", 17, 14)] }, + Layout { + name: "PACR", + opcode: 0x41, + fields: &[ + ("addr_cnt_context", 14, 13), + ("addr_mode", 16, 15), + ("cfg_context", 22, 21), + ("concat", 4, 4), + ("ctxt_ctrl", 3, 2), + ("dst_access_mode", 17, 17), + ("flush", 1, 1), + ("last", 0, 0), + ("ovrd_thread_id", 7, 7), + ("read_intf_sel", 11, 8), + ("row_pad_zero", 20, 18), + ("zero_write", 12, 12), + ], + }, + Layout { name: "SETADCXX", opcode: 0x5e, fields: &[("cnt_set_mask", 23, 21), ("x_end2", 19, 10), ("x_start", 9, 0)] }, + Layout { name: "SETADCXY", opcode: 0x51, fields: &[("bit_mask", 3, 0), ("ch0_x", 8, 6), ("ch0_y", 11, 9), ("ch1_x", 14, 12), ("ch1_y", 17, 15), ("cnt_set_mask", 23, 21)] }, + Layout { name: "SETADCZW", opcode: 0x54, fields: &[("bit_mask", 3, 0), ("ch0_x", 8, 6), ("ch0_y", 11, 9), ("ch1_x", 14, 12), ("ch1_y", 17, 15), ("cnt_set_mask", 23, 21)] }, + Layout { name: "SETC16", opcode: 0xb2, fields: &[("setc16_reg", 23, 16), ("setc16_value", 15, 0)] }, + Layout { name: "SETGPR", opcode: 0x45, fields: &[("payload_sig_sel", 21, 8), ("payload_sig_sel_size", 23, 22), ("reg_index_16b", 6, 0), ("set_signals_mode", 7, 7)] }, + Layout { name: "SETRWC", opcode: 0x37, fields: &[("bit_mask", 3, 0), ("clear_ab_vld", 23, 22), ("rwc_a", 9, 6), ("rwc_b", 13, 10), ("rwc_cr", 21, 18), ("rwc_d", 17, 14)] }, + Layout { name: "SFPAND", opcode: 0x7e, fields: &[("imm12_math", 15, 12), ("instr_mod1", 3, 0), ("lreg_c", 11, 8), ("lreg_dest", 7, 4)] }, + Layout { name: "SFPCONFIG", opcode: 0x91, fields: &[("config_dest", 7, 4), ("imm16_math", 23, 8), ("instr_mod1", 3, 0)] }, + Layout { name: "SFPIADD", opcode: 0x79, fields: &[("imm12_math", 23, 12), ("instr_mod1", 3, 0), ("lreg_c", 11, 8), ("lreg_dest", 7, 4)] }, + Layout { name: "SFPLOAD", opcode: 0x70, fields: &[("dest_reg_addr", 9, 0), ("instr_mod0", 19, 16), ("lreg_ind", 23, 20), ("sfpu_addr_mode", 15, 13)] }, + Layout { name: "SFPLOADI", opcode: 0x71, fields: &[("imm16", 15, 0), ("instr_mod0", 19, 16), ("lreg_ind", 23, 20)] }, + Layout { + name: "SFPMUL24", + opcode: 0x98, + fields: &[("instr_mod1", 3, 0), ("lreg_dest", 7, 4), ("lreg_src_a", 19, 16), ("lreg_src_b", 15, 12), ("lreg_src_c", 11, 8)], + }, + Layout { name: "SFPNOP", opcode: 0x8f, fields: &[] }, + Layout { name: "SFPOR", opcode: 0x7f, fields: &[("imm12_math", 15, 12), ("instr_mod1", 3, 0), ("lreg_c", 11, 8), ("lreg_dest", 7, 4)] }, + Layout { name: "SFPSHFT", opcode: 0x7a, fields: &[("imm12_math", 23, 12), ("instr_mod1", 3, 0), ("lreg_c", 11, 8), ("lreg_dest", 7, 4)] }, + Layout { name: "SFPSTORE", opcode: 0x72, fields: &[("dest_reg_addr", 9, 0), ("instr_mod0", 19, 16), ("lreg_ind", 23, 20), ("sfpu_addr_mode", 15, 13)] }, + Layout { name: "STALLWAIT", opcode: 0xa2, fields: &[("stall_res", 23, 15), ("wait_res", 12, 0)] }, + Layout { + name: "UNPACR", + opcode: 0x42, + fields: &[ + ("addr_cnt_context_id", 9, 8), + ("addr_mode", 22, 15), + ("auto_inc_context_id", 3, 3), + ("cfg_context_cnt_inc", 13, 13), + ("cfg_context_id", 12, 10), + ("last", 0, 0), + ("ovrd_thread_id", 7, 7), + ("row_search", 2, 2), + ("search_cache_flush", 1, 1), + ("set_dat_valid", 6, 6), + ("srcb_bcast", 5, 5), + ("unpack_block_selection", 23, 23), + ("zero_write2", 4, 4), + ], + }, +]; + +/// The word for `name` with `fields` set (every other field zero). +pub fn word(name: &str, fields: &[(&str, u32)]) -> u32 { + let l = LAYOUTS.iter().find(|l| l.name == name).unwrap_or_else(|| panic!("no instruction {name}")); + let mut w = u32::from(l.opcode) << 24; + for &(f, v) in fields { + let &(_, hi, lo) = l.fields.iter().find(|x| x.0 == f).unwrap_or_else(|| panic!("{name} has no field {f}")); + let width = u32::from(hi - lo + 1); + assert!(u64::from(v) < (1u64 << width), "{name}.{f} = {v} does not fit {width} bits"); + w |= v << lo; + } + w +} + +/// The word as the core's instruction stream carries it. +pub fn inline(word: u32) -> u32 { + word.rotate_left(2) +} + +#[cfg(test)] +mod tests { + use super::*; + + /// Words worked out independently from the same table. + #[test] + fn words_match_an_independent_reading_of_the_table() { + let checks: [(u32, u32); 12] = [ + (0x5420_000F, word("SETADCZW", &[("bit_mask", 0xF), ("cnt_set_mask", 1)])), + (0x5E23_FC00, word("SETADCXX", &[("x_end2", 255), ("cnt_set_mask", 1)])), + (0x4208_8081, word("UNPACR", &[("last", 1), ("ovrd_thread_id", 1), ("addr_mode", 0b0001_0001)])), + (0xA280_0002, word("STALLWAIT", &[("stall_res", 1 << 8), ("wait_res", 1 << 1)])), + (0x7004_0000, word("SFPLOAD", &[("instr_mod0", 4)])), + (0x7264_0000, word("SFPSTORE", &[("lreg_ind", 6), ("instr_mod0", 4)])), + (0x3800_8000, word("INCRWC", &[("rwc_d", 2)])), + (0x3700_0004, word("SETRWC", &[("bit_mask", 4)])), + (0x4100_0001, word("PACR", &[("last", 1)])), + (0x4500_017E, word("SETGPR", &[("reg_index_16b", 126), ("payload_sig_sel", 1)])), + (0x9100_00C0, word("SFPCONFIG", &[("config_dest", 12)])), + (0x9802_3940, word("SFPMUL24", &[("lreg_src_a", 2), ("lreg_src_b", 3), ("lreg_src_c", 9), ("lreg_dest", 4)])), + ]; + for (want, got) in checks { + assert_eq!(got, want, "{got:#010x} vs {want:#010x}"); + } + // An inline word is never a 32-bit RISC-V instruction: its low bits are not 0b11. + for l in LAYOUTS { + assert!(l.opcode < 0xc0); + assert_ne!(inline(u32::from(l.opcode) << 24) & 3, 3, "{}", l.name); + } + assert_eq!(inline(0xA280_0002).rotate_right(2), 0xA280_0002); + } + + /// With `WAI_QUANTUM_TENSIX_ISA` naming the simulator's table + /// (`data/bh/tensix_isa.json`), every transcribed layout must be the table's. + #[test] + fn layouts_are_the_simulators_table() { + let Some(path) = std::env::var_os("WAI_QUANTUM_TENSIX_ISA") else { return }; + let table: serde_json::Value = serde_json::from_slice(&std::fs::read(&path).expect("the table reads")).expect("the table is JSON"); + for l in LAYOUTS { + let t = &table[l.name]; + assert_eq!(t["opcode"].as_u64(), Some(u64::from(l.opcode)), "{} opcode", l.name); + let args = t["args"].as_object().unwrap_or_else(|| panic!("{} has no args", l.name)); + assert_eq!(args.len(), l.fields.len(), "{} field count", l.name); + for &(f, hi, lo) in l.fields { + assert_eq!(args.get(f).and_then(|v| v.as_str()), Some(format!("{hi}:{lo}").as_str()), "{}.{f}", l.name); + } + } + println!("tensix layouts: {} instructions match the table", LAYOUTS.len()); + } +} diff --git a/wai/wai-quantum/src/quantum_fabric/vector_engine.rs b/wai/wai-quantum/src/quantum_fabric/vector_engine.rs new file mode 100644 index 000000000..b86e697ba --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/vector_engine.rs @@ -0,0 +1,315 @@ +//! The gate on a compute tile's vector engine: the programs the tile's first +//! RISC-V core runs to drive it, built with [`rv32`] and [`tensix`]. +//! +//! **One pass.** The host writes up to [`PAIRS`] pairs' four columns to the +//! tile's memory at [`IN`]. The core writes the unpacker's and the packer's +//! configuration, then hands the tile its instructions inline: the unpacker +//! copies the columns, as 32-bit integers, into the destination registers; the +//! vector engine forms each pair's four rows, 32 lanes at a time; the packer +//! copies the rows back to [`OUT`]; and the core writes [`DONE`] once the +//! packer has finished. +//! +//! **Exact arithmetic.** The vector engine multiplies 23-bit numbers +//! (`SFPMUL24`: the low or the high 23 bits of the 46-bit product). Each +//! operand is split into 23 and 9 bits; the four partial products give the +//! unsigned product's high word, corrected for sign to the signed product's +//! `P_hi`, and the bits of the low word that the rounding reaches. With the +//! rounding bit `r` and the shift `s`, `fx = (P + 2^r) >> s` is +//! `P_hi · 2^(32−s) + t`, `t = ((P_lo >> r) + 1) >> (s − r)`, held as 30 low +//! bits and the rest; a row adds or subtracts its four terms in that form, +//! carrying after each, and is stored as its 64-bit value's two words. + +use crate::quantum_fabric::rv32::{A0, Asm, S1, T0, T1, T2, T3, ZERO}; +use crate::quantum_fabric::tensix::word; + +/// Pairs per pass, per tile: half a 32 × 32 tile of each operand. +pub const PAIRS: usize = 512; +/// Where the columns go in the tile's memory: four of [`PAIRS`] 32-bit values. +pub const IN: u32 = 0x4_0000; +/// Where the rows come back: for each row, [`PAIRS`] low words then [`PAIRS`] +/// high words. +pub const OUT: u32 = 0x5_0000; +/// The word the core sets to 1 when the pass is done. +pub const DONE: u32 = 0x3_F000; + +/// The kernel's constants: the rounding bit `r` (rounding constant `2^r`), +/// the shift, and whether row 0's last real term is added. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct Constants { + pub round_bit: u32, + pub shift: u32, + pub last_add: bool, +} + +const INT32: u32 = 4; +/// The vector register holding `2^30 − 1`. +const M30: u32 = 12; +/// Where the operands sit in the destination registers (rows). +const INPUTS: u32 = 256; + +fn ld(l: u32, addr: u32) -> u32 { + word("SFPLOAD", &[("lreg_ind", l), ("instr_mod0", INT32), ("dest_reg_addr", addr)]) +} +fn st(l: u32, addr: u32) -> u32 { + word("SFPSTORE", &[("lreg_ind", l), ("instr_mod0", INT32), ("dest_reg_addr", addr)]) +} +fn imm12(v: i32) -> u32 { + (v as u32) & 0xfff +} +/// `d = c << n` (n ≥ 0) or `d = c >>u −n`. +fn shift(d: u32, c: u32, n: i32) -> u32 { + word("SFPSHFT", &[("lreg_dest", d), ("lreg_c", c), ("imm12_math", imm12(n)), ("instr_mod1", 5)]) +} +/// `d = c >>a n`. +fn shra(d: u32, c: u32, n: i32) -> u32 { + word("SFPSHFT", &[("lreg_dest", d), ("lreg_c", c), ("imm12_math", imm12(-n)), ("instr_mod1", 7)]) +} +/// `d = c + d` (the condition codes untouched). +fn add(d: u32, c: u32) -> u32 { + word("SFPIADD", &[("lreg_dest", d), ("lreg_c", c), ("instr_mod1", 4)]) +} +/// `d = c − d`. +fn rsub(d: u32, c: u32) -> u32 { + word("SFPIADD", &[("lreg_dest", d), ("lreg_c", c), ("instr_mod1", 6)]) +} +/// `d = c + v`. +fn addi(d: u32, c: u32, v: i32) -> u32 { + word("SFPIADD", &[("lreg_dest", d), ("lreg_c", c), ("imm12_math", imm12(v)), ("instr_mod1", 5)]) +} +/// `d = L[l] & c`. +fn and_with(d: u32, l: u32, c: u32) -> u32 { + word("SFPAND", &[("lreg_dest", d), ("lreg_c", c), ("imm12_math", l), ("instr_mod1", 1)]) +} +/// `d = d & c`. +fn and(d: u32, c: u32) -> u32 { + word("SFPAND", &[("lreg_dest", d), ("lreg_c", c)]) +} +/// `d = d | c`. +fn or(d: u32, c: u32) -> u32 { + word("SFPOR", &[("lreg_dest", d), ("lreg_c", c)]) +} +/// `d = lo23(a · b)` or `hi23` (`high`), then a gap: silicon does not wait for +/// the multiplier's result before the next instruction reads it. +fn mul(w: &mut Vec<(u32, String)>, d: u32, a: u32, b: u32, high: bool) { + let m = word("SFPMUL24", &[("lreg_src_a", a), ("lreg_src_b", b), ("lreg_src_c", 9), ("lreg_dest", d), ("instr_mod1", u32::from(high))]); + w.push((m, format!("SFPMUL24 L{d} = {}(L{a} * L{b})", if high { "hi23" } else { "lo23" }))); + w.push((word("SFPNOP", &[]), "SFPNOP".into())); +} +fn push(w: &mut Vec<(u32, String)>, word: u32, what: &str) { + w.push((word, what.into())); +} + +/// The product of the operands in L0 and L1: `P_hi` into L3, the 9 bits of +/// `P_lo` above bit 22 (`d1 & 511`, at the bottom of L2's word), with L0 and +/// L1 kept. +fn product(w: &mut Vec<(u32, String)>) { + push(w, shift(2, 0, -23), "L2 = a1 = a >>u 23"); + push(w, shift(3, 1, -23), "L3 = e1 = e >>u 23"); + mul(w, 4, 2, 3, false); // a1·e1, exact + mul(w, 5, 2, 1, true); + push(w, add(4, 5), "d2 += hi(a1, e0)"); + mul(w, 5, 0, 3, true); + push(w, add(4, 5), "d2 += hi(a0, e1)"); + mul(w, 2, 2, 1, false); + mul(w, 3, 0, 3, false); + push(w, add(2, 3), "d1 = lo(a1, e0) + lo(a0, e1)"); + mul(w, 3, 0, 1, true); + push(w, add(2, 3), "d1 += hi(a0, e0)"); + push(w, shra(3, 0, 31), "L3 = a < 0 ? -1 : 0"); + push(w, and(3, 1), "L3 = a < 0 ? e : 0"); + push(w, shra(5, 1, 31), "L5 = e < 0 ? -1 : 0"); + push(w, and(5, 0), "L5 = e < 0 ? a : 0"); + push(w, add(3, 5), "the sign correction"); + push(w, shift(4, 4, 14), "d2 << 14"); + push(w, shift(5, 2, -9), "d1 >>u 9"); + push(w, add(4, 5), "U_hi"); + push(w, rsub(3, 4), "L3 = P_hi = U_hi - correction"); +} + +/// From [`product`]: `fx` as 30 low bits in L4 and the rest in L3. +fn fx_from_product(w: &mut Vec<(u32, String)>, c: Constants) { + let k = 32 - c.shift as i32; + push(w, shift(2, 2, 23), "the bits of P_lo the rounding reaches, at the top"); + push(w, shift(2, 2, -(c.round_bit as i32)), "P_lo >>u r"); + push(w, addi(2, 2, 1), "+ 1"); + if c.shift > c.round_bit { + push(w, shift(2, 2, -((c.shift - c.round_bit) as i32)), "t = ((P_lo >> r) + 1) >> (s - r)"); + } + push(w, shift(4, 3, k), "P_hi << (32 - s)"); + push(w, add(4, 2), "+ t"); + push(w, and_with(4, M30, 4), "fx low 30 bits"); + push(w, shift(2, 2, -k), "t >> (32 - s)"); + push(w, add(3, 2), "P_hi + (t >> (32 - s))"); + push(w, shra(3, 3, 30 - k), "fx >> 30"); +} + +/// The row accumulator (L6 low 30 bits, L7 the rest) plus or minus `fx`, carried. +fn accumulate(w: &mut Vec<(u32, String)>, sign: i32, first: bool) { + if first { + push(w, addi(6, 4, 0), "R_lo = fx_lo"); + push(w, addi(7, 3, 0), "R_hi = fx_hi"); + return; + } + if sign > 0 { + push(w, add(6, 4), "R_lo += fx_lo"); + push(w, add(7, 3), "R_hi += fx_hi"); + } else { + push(w, rsub(4, 6), "R_lo - fx_lo"); + push(w, addi(6, 4, 0), "R_lo -= fx_lo"); + push(w, rsub(3, 7), "R_hi - fx_hi"); + push(w, addi(7, 3, 0), "R_hi -= fx_hi"); + } + push(w, shra(5, 6, 30), "the carry"); + push(w, add(7, 5), "R_hi += carry"); + push(w, and_with(6, M30, 6), "R_lo &= 2^30 - 1"); +} + +/// One iteration: the four rows of 32 lanes. `m` the gate's entries. +fn rows_iteration(m: &[i32; 8], c: Constants) -> Vec<(u32, String)> { + let last = if c.last_add { 1 } else { -1 }; + let rows: [[(usize, u32, i32); 4]; 4] = [ + [(0, 0, 1), (1, 1, -1), (2, 2, 1), (3, 3, last)], + [(0, 1, 1), (1, 0, 1), (2, 3, 1), (3, 2, 1)], + [(4, 0, 1), (5, 1, -1), (6, 2, 1), (7, 3, -1)], + [(4, 1, 1), (5, 0, 1), (6, 3, 1), (7, 2, 1)], + ]; + let mut w = Vec::new(); + for (r, terms) in rows.iter().enumerate() { + for (i, &(e, op, sign)) in terms.iter().enumerate() { + push(&mut w, ld(0, INPUTS + 32 * op), &format!("L0 = column {op}")); + let v = m[e] as u32; + push(&mut w, word("SFPLOADI", &[("lreg_ind", 1), ("instr_mod0", 10), ("imm16", v & 0xffff)]), &format!("L1 = m{e} (low half)")); + push(&mut w, word("SFPLOADI", &[("lreg_ind", 1), ("instr_mod0", 8), ("imm16", v >> 16)]), &format!("L1 = m{e} (high half)")); + product(&mut w); + fx_from_product(&mut w, c); + accumulate(&mut w, sign, i == 0); + } + push(&mut w, shift(4, 7, 30), "R_hi << 30"); + push(&mut w, or(4, 6), "the row's low word"); + push(&mut w, shra(7, 7, 2), "the row's high word"); + push(&mut w, st(4, 64 * r as u32), &format!("row {r} low")); + push(&mut w, st(7, 64 * r as u32 + 32), &format!("row {r} high")); + } + push(&mut w, word("INCRWC", &[("rwc_d", 2)]), "the next 32 lanes"); + w +} + +/// One iteration of the probe: each pair's product and `fx`, both 64-bit. +fn probe_iteration(c: Constants) -> Vec<(u32, String)> { + let mut w = Vec::new(); + push(&mut w, ld(0, INPUTS), "L0 = a"); + push(&mut w, ld(1, INPUTS + 32), "L1 = b"); + product(&mut w); + mul(&mut w, 6, 0, 1, false); + push(&mut w, shift(5, 2, 23), "(d1 & 511) << 23"); + push(&mut w, or(6, 5), "P_lo"); + push(&mut w, st(6, 0), "the product's low word"); + push(&mut w, st(3, 32), "the product's high word"); + fx_from_product(&mut w, c); + push(&mut w, shift(5, 3, 30), "fx_hi << 30"); + push(&mut w, or(5, 4), "fx low word"); + push(&mut w, shra(3, 3, 2), "fx high word"); + push(&mut w, st(5, 64), "fx low"); + push(&mut w, st(3, 96), "fx high"); + push(&mut w, word("INCRWC", &[("rwc_d", 2)]), "the next 32 lanes"); + w +} + +/// The configuration word `idx` of the tile, through the core's view of it. +fn cfg(p: &mut Asm, idx: u32, value: u32) { + p.li(T1, value).sw(T1, (4 * idx) as i32, T0); +} + +/// Lanes one iteration covers. +pub const LANES: usize = 32; + +/// The whole program around one iteration's `body`, run `iterations` times +/// (1 to 16, each [`LANES`] more pairs): configure, unpack `faces` faces of +/// columns, iterate, pack `rows` destination rows, signal. +fn program(body: &[(u32, String)], iterations: u32, faces: u32, rows: u32) -> (Vec, String) { + assert!((1..=(PAIRS / LANES) as u32).contains(&iterations)); + let mut p = Asm::default(); + p.li(T0, 0xffef_0000); + // Unpacker 0: 32-bit integers into the destination registers from row 256. + for (idx, v) in [(64, 0x18), (65, 0x0020_0001), (72, 0x808), (73, 0x11), (76, IN / 16 - 1), (84, 0x1040_1040), (86, 0x0100_0100), (57, 0x400)] { + cfg(&mut p, idx, v); + } + // Packer 0: 32-bit integers from the destination registers from row 0. + for (idx, v) in [(1, 0x1000_0000), (12, 0x0040_0000), (13, 0x1000_0400), (18, 1), (24, 0xffff), (28, 0x1000), (69, OUT / 16 - 1), (70, 0x881)] { + cfg(&mut p, idx, v); + } + // Read a configuration word back before the tile acts on them. + p.lw(T1, 4 * 76, T0).add(T1, T1, ZERO); + p.li(T2, 0xffe0_0000).sw(ZERO, 0xfc, T2).lw(T1, 0xfc, T2); + p.tt(word("SETC16", &[("setc16_reg", 37), ("setc16_value", 4)]), "SETC16: the packer steps 4 rows"); + p.tt(word("SETADCZW", &[("bit_mask", 0xf), ("cnt_set_mask", 1)]), "SETADCZW unpacker 0"); + p.tt(word("SETADCXY", &[("bit_mask", 0xf), ("cnt_set_mask", 1)]), "SETADCXY unpacker 0"); + p.tt(word("SETADCXX", &[("x_end2", 255), ("cnt_set_mask", 1)]), "SETADCXX unpacker 0: a face"); + for f in 0..faces { + p.tt(word("UNPACR", &[("last", 1), ("ovrd_thread_id", 1), ("addr_mode", 0b0001_0001)]), &format!("UNPACR face {f}")); + } + p.tt(word("STALLWAIT", &[("stall_res", 1 << 8), ("wait_res", 1 << 1)]), "STALLWAIT: the vector engine waits for the unpacker"); + p.tt(word("SFPLOADI", &[("lreg_ind", 0), ("instr_mod0", 10), ("imm16", 0xffff)]), "L0 = 2^30 - 1 (low half)"); + p.tt(word("SFPLOADI", &[("lreg_ind", 0), ("instr_mod0", 8), ("imm16", 0x3fff)]), "L0 = 2^30 - 1 (high half)"); + p.tt(word("SFPCONFIG", &[("config_dest", M30)]), "L12 = L0"); + p.tt(word("SETRWC", &[("bit_mask", 4)]), "SETRWC: lanes from row 0"); + p.li(S1, iterations); + p.label("iteration"); + for (w, what) in body { + p.tt(*w, what); + } + p.addi(S1, S1, -1).bne(S1, ZERO, "iteration"); + p.tt(word("STALLWAIT", &[("stall_res", 1 << 2), ("wait_res", 1 << 11)]), "STALLWAIT: the packer waits for the vector engine"); + p.tt(word("SETADCXY", &[("bit_mask", 0xb), ("cnt_set_mask", 4)]), "SETADCXY packer"); + p.tt(word("SETADCZW", &[("bit_mask", 0xf), ("cnt_set_mask", 4)]), "SETADCZW packer"); + p.tt(word("SETADCXX", &[("x_end2", 15), ("cnt_set_mask", 4)]), "SETADCXX packer: a row"); + for k in 0..rows / 4 { + p.tt(word("PACR", &[("last", u32::from(k == rows / 4 - 1))]), &format!("PACR rows {}..{}", 4 * k, 4 * k + 3)); + } + p.tt(word("STALLWAIT", &[("stall_res", 1 << 5), ("wait_res", 1 << 3)]), "STALLWAIT: until the packer is done"); + p.tt(word("SETGPR", &[("reg_index_16b", 126), ("payload_sig_sel", 1)]), "SETGPR: register 63 = 1"); + p.label("wait").lw(T1, 0xfc, T2).beq(T1, ZERO, "wait"); + p.li(T3, DONE).addi(A0, ZERO, 1).sw(A0, 0, T3); + p.label("idle").j("idle"); + p.finish() +} + +/// The rows program for one gate over `iterations` × [`LANES`] pairs a tile: +/// the gate's entries are in the program. +pub fn rows_program(m: &[i32; 8], c: Constants, iterations: u32) -> (Vec, String) { + program(&rows_iteration(m, c), iterations, 8, 256) +} + +/// The probe program: each pair's `a·b` and `fx(a, b)` (a in column 0, b in +/// column 1); the results are four blocks of [`PAIRS`] words (the product's +/// low and high words, then `fx`'s). +pub fn probe_program(c: Constants, iterations: u32) -> (Vec, String) { + program(&probe_iteration(c), iterations, 4, 128) +} + +/// Clocks a pass of `iterations` takes, about: under nine hundred instructions +/// an iteration, and the unpacking and packing around them. +pub fn pass_clocks(iterations: u32) -> u64 { + 2_000 + 900 * u64::from(iterations) +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn the_programs_fit_and_their_mutants_differ() { + let c = Constants { round_bit: 29, shift: 30, last_add: false }; + let m = [1, -2, 3, -4, 5, -6, 7, -8]; + let (w, text) = rows_program(&m, c, 16); + assert!(4 * w.len() < DONE as usize && text.contains("SFPMUL24")); + for other in [Constants { round_bit: 28, ..c }, Constants { shift: 29, ..c }, Constants { last_add: true, ..c }] { + assert_ne!(rows_program(&m, other, 16).1, text); + } + assert_ne!(rows_program(&[0; 8], c, 16).0, w, "the entries are in the program"); + assert_ne!(rows_program(&m, c, 1).0, w, "so is the pass's length"); + let (pw, _) = probe_program(c, 16); + assert!(4 * pw.len() < DONE as usize); + assert!(IN + 4 * 4 * PAIRS as u32 <= OUT && OUT + 8 * 4 * PAIRS as u32 <= 0x8_0000); + } +} diff --git a/wai/wai-quantum/src/quantum_fabric_dataflow.rs b/wai/wai-quantum/src/quantum_fabric_dataflow.rs index c93fb30ca..1f0fd5ebc 100644 --- a/wai/wai-quantum/src/quantum_fabric_dataflow.rs +++ b/wai/wai-quantum/src/quantum_fabric_dataflow.rs @@ -25,6 +25,7 @@ use crate::quantum::{Amp, Circuit, Gateset, MAX_QUBITS, ONE, QuantumError, State use crate::quantum_fabric::lanes::{self, GateWords, Lane}; use crate::quantum_fabric::plan::{self, Config, LaneRun, Program, Step}; use crate::quantum_fabric::rv32::{self, Layout}; +use crate::quantum_fabric::vector_engine::{self, Constants}; use crate::quantum_fabric::{Certificate, Certification, Corpus, FabricClass, KernelEntry, Member, MemberError, Presence, Provenance, RunFuture, RunMeter, Runner, apply_kernel}; use std::sync::{Mutex, OnceLock}; @@ -113,8 +114,33 @@ pub fn programs(v: Variant) -> ((Vec, String), (Vec, String)) { (rv32::rows_program(LAYOUT, rnd, sh, last_add), rv32::probe_program(LAYOUT, rnd, sh)) } +/// What runs the gate on a tile: its first RISC-V core, or the tile's vector +/// engine driven by that core. +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub enum Engine { + Cores, + Vector, +} + +impl Engine { + pub fn label(self) -> &'static str { + match self { + Engine::Cores => "cores", + Engine::Vector => "vector", + } + } +} + +/// The vector engine's kernel constants for a variant. +fn vector_constants(v: Variant) -> Constants { + let (rnd, sh, last_add) = v.constants(); + Constants { round_bit: rnd.trailing_zeros(), shift: sh, last_add } +} + /// The implementation name the id starts with. pub const IMPLEMENTATION: &str = "rv32-lanes64"; +/// The vector engine's. +pub const VECTOR_IMPLEMENTATION: &str = "sfpu-lanes64"; /// Version of the host layout: the gather of a gate's pairs into four columns, /// the batch of [`LAYOUT`], the rows as word pairs, and the host's range check. @@ -124,18 +150,40 @@ pub const HOST_LAYOUT: u32 = 1; /// `"rv32-lanes64/<16 hex>"` over a variant's two listings and the layout. pub fn implementation_id_of(v: Variant) -> String { let ((_, rows), (_, probe)) = programs(v); - let mut o = b"wai:quantum-fabric-dataflow\x01".to_vec(); - for text in [&rows, &probe] { + id_over(b"wai:quantum-fabric-dataflow\x01", IMPLEMENTATION, &[&rows, &probe]) +} + +pub fn implementation_id() -> String { + implementation_id_of(Variant::Conforming) +} + +/// `"sfpu-lanes64/<16 hex>"` over the vector engine's two listings (the rows +/// program with every entry zero; a gate's entries are its program's +/// immediates) and the layout. +pub fn vector_implementation_id_of(v: Variant) -> String { + let c = vector_constants(v); + let (_, rows) = vector_engine::rows_program(&[0; 8], c, 16); + let (_, probe) = vector_engine::probe_program(c, 16); + id_over(b"wai:quantum-fabric-dataflow-vector\x01", VECTOR_IMPLEMENTATION, &[&rows, &probe]) +} + +/// The implementation id of `engine` running `v`. +pub fn implementation_id_for(engine: Engine, v: Variant) -> String { + match engine { + Engine::Cores => implementation_id_of(v), + Engine::Vector => vector_implementation_id_of(v), + } +} + +fn id_over(domain: &[u8], name: &str, texts: &[&String]) -> String { + let mut o = domain.to_vec(); + for text in texts { o.extend_from_slice(&(text.len() as u64).to_be_bytes()); o.extend_from_slice(text.as_bytes()); } o.extend_from_slice(&HOST_LAYOUT.to_be_bytes()); let h: String = blake3::hash(&o).as_bytes()[..8].iter().map(|b| format!("{b:02x}")).collect(); - format!("{IMPLEMENTATION}/{h}") -} - -pub fn implementation_id() -> String { - implementation_id_of(Variant::Conforming) + format!("{name}/{h}") } /// The compute tiles a run spreads its pairs over, the first core of each. @@ -152,6 +200,9 @@ pub struct Chip { /// A window's size, and the stride of its configuration registers. window: u64, cfg_stride: u64, + /// Whether its vector engine has the 23-bit multiplier the vector kernel + /// is built on. + pub vector: bool, } /// The PCI vendor id of the devices this member knows. @@ -183,9 +234,9 @@ const TILES_B140: [(u32, u32); TILES] = { pub const CHIPS: [Chip; 2] = [ // One 64-bit configuration register per 1 MiB window. - Chip { device: 0x401e, tiles: &TILES_401E, window: 1 << 20, cfg_stride: 8 }, + Chip { device: 0x401e, tiles: &TILES_401E, window: 1 << 20, cfg_stride: 8, vector: false }, // Three 32-bit configuration registers per 2 MiB window. - Chip { device: 0xb140, tiles: &TILES_B140, window: 1 << 21, cfg_stride: 12 }, + Chip { device: 0xb140, tiles: &TILES_B140, window: 1 << 21, cfg_stride: 12, vector: true }, ]; /// The tile's soft-reset register: bit 11 holds the first core, 12–14 and 18 @@ -333,10 +384,10 @@ impl Device { bytes.as_chunks::<4>().0.iter().map(|&c| u32::from_le_bytes(c)).collect() } - /// Runs `program` on the first `counts.len()` tiles, tile `j` over - /// `counts[j]` pairs already written to it, until each writes its done - /// word; at most `budget` clocks. - fn run(&mut self, program: &[u32], counts: &[u32], budget: u64) -> Result<(), DataflowError> { + /// Runs `program` on the first `tiles` tiles, each with its inputs + /// already written, until each writes 1 at `done`; a first stretch of + /// `expected` clocks, then short ones, at most `budget` in all. + fn run(&mut self, program: &[u32], tiles: usize, done: u32, expected: u64, budget: u64) -> Result<(), DataflowError> { for j in 0..TILES { self.reg_write(j, HOLD_ALL); } @@ -346,23 +397,21 @@ impl Device { } self.loaded = Some(program.to_vec()); } - for (j, &count) in counts.iter().enumerate() { - self.l1_write_words(j, LAYOUT.base + Layout::COUNT, &[count]); - self.l1_write_words(j, LAYOUT.base + Layout::DONE, &[0]); + for j in 0..tiles { + self.l1_write_words(j, done, &[0]); } - for j in 0..counts.len() { + for j in 0..tiles { self.reg_write(j, RUN_FIRST); } let mut spent = 0u64; - // A first stretch near what the program needs, then short ones. - let mut step = budget.div_ceil(2).clamp(1, u64::from(u32::MAX)); - let mut waiting: Vec = (0..counts.len()).collect(); + let mut step = expected.clamp(1, budget.min(u64::from(u32::MAX))); + let mut waiting: Vec = (0..tiles).collect(); loop { // SAFETY: advances the device's clock; nothing else is touched. unsafe { (self.lib.clock)(step as u32) }; spent += step; self.clocks += step; - waiting.retain(|&j| self.l1_read_words(j, LAYOUT.base + Layout::DONE, 1)[0] != 1); + waiting.retain(|&j| self.l1_read_words(j, done, 1)[0] != 1); if waiting.is_empty() { break; } @@ -374,11 +423,20 @@ impl Device { } step = 4096.min(budget - spent).max(1); } - for j in 0..counts.len() { + for j in 0..tiles { self.reg_write(j, HOLD_ALL); } Ok(()) } + + /// The cores' programs: each tile's pair count, then the run. + fn run_counts(&mut self, program: &[u32], counts: &[u32], per_pair: u64) -> Result<(), DataflowError> { + for (j, &count) in counts.iter().enumerate() { + self.l1_write_words(j, LAYOUT.base + Layout::COUNT, &[count]); + } + let b = budget(counts[0], per_pair); + self.run(program, counts.len(), LAYOUT.base + Layout::DONE, b.div_ceil(2), b) + } } fn device() -> Result<&'static Mutex, String> { @@ -403,21 +461,42 @@ pub fn discover() -> Discovery { return d; } }; - match DataflowMember::open(chip, Variant::Conforming) { - Ok(m) => d.found.push(m), - Err(e) => d.refused.push((key_of(chip), e)), + let engines: &[Engine] = if chip.vector { &[Engine::Cores, Engine::Vector] } else { &[Engine::Cores] }; + for &engine in engines { + match DataflowMember::open(chip, engine, Variant::Conforming) { + Ok(m) => d.found.push(m), + Err(e) => d.refused.push((key_for(chip, engine), e)), + } } d } -/// The key a runner declares: `dataflow--`. +/// The key a runner declares for the cores: `dataflow--`. pub fn key_of(chip: Chip) -> String { format!("dataflow-{VENDOR:04x}-{:04x}", chip.device) } +/// The key for `engine`: the vector engine's adds `-vector`. +pub fn key_for(chip: Chip, engine: Engine) -> String { + match engine { + Engine::Cores => key_of(chip), + Engine::Vector => format!("{}-vector", key_of(chip)), + } +} + /// The most qubits a run holds. pub const MAX_DATAFLOW_QUBITS: u8 = if 20 < MAX_QUBITS { 20 } else { MAX_QUBITS }; +/// A round of the vector engine: `len` pairs spread over the tiles, at most a +/// pass each, in whole pairs of iterations. An iteration covers the even or +/// the odd columns of four destination rows, so two cover the next 64 pairs +/// in order. The pairs each tile takes, and the iterations its pass runs. +fn vector_share(len: usize) -> (usize, u32) { + let unit = 2 * vector_engine::LANES; + let share = (len.div_ceil(TILES).div_ceil(unit).max(1) * unit).min(vector_engine::PAIRS); + (share, (share / vector_engine::LANES) as u32) +} + /// Clocks a program needs for `count` pairs, with room to spare: the rows /// program executes about 220 instructions a pair, the probe about 25. fn budget(count: u32, per_pair: u64) -> u64 { @@ -428,6 +507,7 @@ fn budget(count: u32, per_pair: u64) -> u64 { #[derive(Clone, Debug)] pub struct DataflowMember { chip: Chip, + engine: Engine, variant: Variant, key: String, id: String, @@ -438,10 +518,10 @@ pub struct DataflowMember { } impl DataflowMember { - fn open(chip: Chip, variant: Variant) -> Result { + fn open(chip: Chip, engine: Engine, variant: Variant) -> Result { let label = std::env::var("WAI_QUANTUM_DATAFLOW_SIMULATOR").ok().filter(|s| !s.trim().is_empty()).unwrap_or_else(|| "the device library".into()); let provenance = Provenance::Simulator(label.trim().to_string()); - let key = key_of(chip); + let key = key_for(chip, engine); let mut o = b"wai:quantum-fabric-dataflow-member\x01".to_vec(); o.extend_from_slice(key.as_bytes()); for &(x, y) in chip.tiles { @@ -450,17 +530,22 @@ impl DataflowMember { } o.extend_from_slice(provenance.label().as_bytes()); let id = format!("local/dataflow-{}", blake3::hash(&o).as_bytes()[..6].iter().map(|b| format!("{b:02x}")).collect::()); + let what = match engine { + Engine::Cores => "the first RISC-V core (RV32IM)", + Engine::Vector => "the vector engine, driven by the first RISC-V core,", + }; let member = DataflowMember { chip, + engine, variant, detail: format!( - "dataflow; pci {VENDOR:04x}:{:04x}, the first RISC-V core (RV32IM) of {TILES} compute tiles; {}; up to {MAX_DATAFLOW_QUBITS} qubits; unmetered", + "dataflow; pci {VENDOR:04x}:{:04x}, {what} of {TILES} compute tiles; {}; up to {MAX_DATAFLOW_QUBITS} qubits; unmetered", chip.device, provenance.label() ), key, id, - implementation: implementation_id_of(variant), + implementation: implementation_id_for(engine, variant), provenance, max_qubits: MAX_DATAFLOW_QUBITS, }; @@ -477,7 +562,12 @@ impl DataflowMember { /// The same device running a mutant: for the mutation tests. #[doc(hidden)] pub fn with_variant(&self, v: Variant) -> DataflowMember { - DataflowMember { variant: v, implementation: implementation_id_of(v), ..self.clone() } + DataflowMember { variant: v, implementation: implementation_id_for(self.engine, v), ..self.clone() } + } + + /// What runs the gate: the tiles' cores or their vector engines. + pub fn engine(&self) -> Engine { + self.engine } #[doc(hidden)] @@ -537,6 +627,9 @@ impl DataflowMember { /// The kernel's [`lanes::mul_wide`] and [`lanes::fxmul_wide`] of each pair. pub fn probe(&self, pairs: &[(i32, i32)]) -> Result, DataflowError> { + if self.engine == Engine::Vector { + return self.vector_probe(pairs); + } let (_, (program, _)) = programs(self.variant); let mut dev = self.device()?; let mut out = Vec::with_capacity(pairs.len()); @@ -547,7 +640,7 @@ impl DataflowMember { dev.l1_write_words(j, LAYOUT.column(1), &slice.iter().map(|p| p.1 as u32).collect::>()); } let counts: Vec = slices.iter().map(|s| s.len() as u32).collect(); - dev.run(&program, &counts, budget(counts[0], 25))?; + dev.run_counts(&program, &counts, 25)?; for (j, slice) in slices.iter().enumerate() { let w = dev.l1_read_words(j, LAYOUT.results(), 4 * slice.len()); out.extend(w.as_chunks::<4>().0.iter().map(|c| ((c[1], c[0]), (c[3], c[2])))); @@ -558,6 +651,9 @@ impl DataflowMember { /// The four rows of each pair, from the kernel. fn rows(&self, m: &[i32; 8], cols: &[Vec; 4]) -> Result<[Vec; 4], DataflowError> { + if self.engine == Engine::Vector { + return self.vector_rows(m, cols); + } let ((program, _), _) = programs(self.variant); let mut dev = self.device()?; let n = cols[0].len(); @@ -574,7 +670,7 @@ impl DataflowMember { } } let counts: Vec = slices.iter().map(|&(s, e)| (e - s) as u32).collect(); - dev.run(&program, &counts, budget(counts[0], 220))?; + dev.run_counts(&program, &counts, 220)?; for (j, &(s, e)) in slices.iter().enumerate() { let w = dev.l1_read_words(j, LAYOUT.results(), 8 * (e - s)); for pair in w.as_chunks::<8>().0 { @@ -587,6 +683,60 @@ impl DataflowMember { Ok(o) } + /// The vector engine's probe: up to [`vector_engine::PAIRS`] pairs a tile. + fn vector_probe(&self, pairs: &[(i32, i32)]) -> Result, DataflowError> { + let n = vector_engine::PAIRS; + let mut dev = self.device()?; + let mut out = Vec::with_capacity(pairs.len()); + for round in pairs.chunks(TILES * n) { + let (share, iterations) = vector_share(round.len()); + let (program, _) = vector_engine::probe_program(vector_constants(self.variant), iterations); + let slices: Vec<&[(i32, i32)]> = round.chunks(share).collect(); + for (j, slice) in slices.iter().enumerate() { + dev.l1_write_words(j, vector_engine::IN, &slice.iter().map(|p| p.0 as u32).collect::>()); + dev.l1_write_words(j, vector_engine::IN + 4 * n as u32, &slice.iter().map(|p| p.1 as u32).collect::>()); + } + let clocks = vector_engine::pass_clocks(iterations); + dev.run(&program, slices.len(), vector_engine::DONE, clocks, 20 * clocks)?; + for (j, slice) in slices.iter().enumerate() { + let w = dev.l1_read_words(j, vector_engine::OUT, 4 * n); + out.extend((0..slice.len()).map(|i| ((w[n + i], w[i]), (w[3 * n + i], w[2 * n + i])))); + } + } + Ok(out) + } + + /// The vector engine's rows: up to [`vector_engine::PAIRS`] pairs a tile, + /// the gate's entries in the program. + fn vector_rows(&self, m: &[i32; 8], cols: &[Vec; 4]) -> Result<[Vec; 4], DataflowError> { + let n = vector_engine::PAIRS; + let mut dev = self.device()?; + let len = cols[0].len(); + let mut o: [Vec; 4] = Default::default(); + for start in (0..len).step_by(TILES * n) { + let end = (start + TILES * n).min(len); + let (share, iterations) = vector_share(end - start); + let (program, _) = vector_engine::rows_program(m, vector_constants(self.variant), iterations); + let slices: Vec<(usize, usize)> = (start..end).step_by(share).map(|s| (s, (s + share).min(end))).collect(); + for (j, &(s, e)) in slices.iter().enumerate() { + for (k, col) in cols.iter().enumerate() { + dev.l1_write_words(j, vector_engine::IN + (4 * n * k) as u32, &col[s..e].iter().map(|&v| v as u32).collect::>()); + } + } + let clocks = vector_engine::pass_clocks(iterations); + dev.run(&program, slices.len(), vector_engine::DONE, clocks, 20 * clocks)?; + for (j, &(s, e)) in slices.iter().enumerate() { + let w = dev.l1_read_words(j, vector_engine::OUT, 8 * n); + for i in 0..e - s { + for (r, row) in o.iter_mut().enumerate() { + row.push((i64::from(w[(2 * r + 1) * n + i] as i32) << 32) | i64::from(w[2 * r * n + i])); + } + } + } + } + Ok(o) + } + /// Gates in order from `state`, as [`lanes::run_words`] runs them. pub fn run_lanes(&self, state: &[Lane], gates: &[GateWords]) -> Result { let len = state.len(); @@ -759,6 +909,18 @@ mod tests { assert!(TILES_401E.iter().all(|&(x, y)| x != 5 && x != 0 && (1..=5).contains(&y))); } + #[test] + fn a_vector_round_covers_its_pairs_in_whole_pairs_of_iterations() { + for len in [1usize, 31, 32, 33, 1000, 1024, 1025, 16_383, 16_384] { + let (share, it) = vector_share(len); + assert_eq!(share, it as usize * vector_engine::LANES); + assert!(it % 2 == 0, "{len}: an odd number of iterations leaves pairs out of order"); + assert!(share <= vector_engine::PAIRS && share * TILES >= len, "{len}: {share}"); + } + assert_eq!(vector_share(1), (64, 2)); + assert_eq!(vector_share(TILES * vector_engine::PAIRS), (vector_engine::PAIRS, 16)); + } + #[test] fn the_keys_name_the_pci_ids() { assert_eq!(key_of(CHIPS[0]), "dataflow-1e52-401e"); diff --git a/wai/wai-quantum/tests/dataflow_member.rs b/wai/wai-quantum/tests/dataflow_member.rs index e357be9c6..9945fc2ed 100644 --- a/wai/wai-quantum/tests/dataflow_member.rs +++ b/wai/wai-quantum/tests/dataflow_member.rs @@ -9,7 +9,7 @@ //! same lines to it, which a CI step can count per device. //! //! ```text -//! WAI_QUANTUM_DATAFLOW_LIB= WAI_QUANTUM_DATAFLOW_EXPECT=dataflow-1e52-b140 \ +//! WAI_QUANTUM_DATAFLOW_LIB= WAI_QUANTUM_DATAFLOW_EXPECT=dataflow-1e52-b140,dataflow-1e52-b140-vector \ //! cargo test --release --features quantum_fabric_dataflow --test dataflow_member -- --test-threads=1 //! ``` #![cfg(not(target_arch = "wasm32"))] @@ -169,9 +169,9 @@ fn c0_devices_are_declared() { for m in &d.found { assert_eq!(m.class(), FabricClass::Dataflow); assert!(m.id().starts_with("local/dataflow-"), "{}", m.id()); - assert_eq!(m.implementation(), dataflow::implementation_id()); + assert_eq!(m.implementation(), dataflow::implementation_id_for(m.engine(), Variant::Conforming)); assert!(m.max_qubits() >= 17, "{}: holds only {} qubits", m.key(), m.max_qubits()); - report(&format!("dataflow member: key={} id={} {}", m.key(), m.id(), m.detail())); + report(&format!("dataflow member: key={} id={} engine={} {}", m.key(), m.id(), m.engine().label(), m.detail())); } } diff --git a/wai/wai-rs/Cargo.lock b/wai/wai-rs/Cargo.lock index 35e758dfe..5206f40a0 100644 --- a/wai/wai-rs/Cargo.lock +++ b/wai/wai-rs/Cargo.lock @@ -4300,7 +4300,7 @@ version = "0.1.0" [[package]] name = "wai-quantum" -version = "0.5.1" +version = "0.5.2" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-web/Cargo.lock b/wai/wai-web/Cargo.lock index 33a654a8c..0c3b53fde 100644 --- a/wai/wai-web/Cargo.lock +++ b/wai/wai-web/Cargo.lock @@ -1231,7 +1231,7 @@ version = "0.1.0" [[package]] name = "wai-quantum" -version = "0.5.1" +version = "0.5.2" dependencies = [ "blake3", "ed25519-dalek", -- 2.54.0 (Apple Git-157) From 38a699aaa0910353915822158afd34085238e5e8 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 04:06:59 -0400 Subject: [PATCH 02/19] wai-quantum: the gate as a portable tensor program MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit quantum_fabric::stablehlo writes the gate's rows and the arithmetic probe as a StableHLO module: every value a 64-bit integer that never overflows (|e·a| < 2^62; a row of four fx below 2^35), so any tensor compiler must give the exact rows. The stablehlo_gate example writes the module, its three mutants, a batch of 4096 pairs reaching the 32-bit edges, and the exact outputs (raw little-endian i64, compared under the module's own signless type). The wai-quantum-tensor-program workflow builds it with a pinned open-source tensor compiler (IREE 3.12.0) and has its runtime compare: the CPU build and the three mutants in a container, plus builds for three GPU generations the fleet lacks (compiled only); every Vulkan device of the hardware host and CUDA; Metal on the macOS runner. Each GPU leg also runs one mutant, so its comparison is shown to catch a difference. Co-Authored-By: Claude Opus 5.5 --- .../workflows/wai-quantum-tensor-program.yml | 233 ++++++++++++++++++ wai/CHANGELOG.md | 16 +- wai/wai-quantum/Cargo.toml | 4 + wai/wai-quantum/README.md | 2 +- wai/wai-quantum/examples/stablehlo_gate.rs | 69 ++++++ wai/wai-quantum/src/quantum_fabric.rs | 1 + .../src/quantum_fabric/stablehlo.rs | 190 ++++++++++++++ 7 files changed, 513 insertions(+), 2 deletions(-) create mode 100644 .github/workflows/wai-quantum-tensor-program.yml create mode 100644 wai/wai-quantum/examples/stablehlo_gate.rs create mode 100644 wai/wai-quantum/src/quantum_fabric/stablehlo.rs diff --git a/.github/workflows/wai-quantum-tensor-program.yml b/.github/workflows/wai-quantum-tensor-program.yml new file mode 100644 index 000000000..ce2038b83 --- /dev/null +++ b/.github/workflows/wai-quantum-tensor-program.yml @@ -0,0 +1,233 @@ +name: wai-quantum portable tensor program + +# The gate as a portable tensor program (quantum_fabric::stablehlo, StableHLO +# as MLIR text): every value a 64-bit integer that never overflows, so a +# compiler for any device must give the exact rows. An open-source tensor +# compiler and runtime (IREE 3.12.0, Apache-2.0, pinned) builds it for each +# device here, and its runtime compares the outputs with the exact ones the +# crate writes: +# +# - cpu (the container runner): the x86-64 build; the three mutants, each of +# which must fail where it should; builds for three GPU generations the +# fleet has none of (compiled, not run); +# - gpu (the hardware host): every Vulkan device the host has (its discrete +# and integrated GPUs and the software one), and the discrete GPU through +# CUDA; +# - metal (the macOS runner): the platform's GPU through Metal. +# Each GPU leg also runs one mutant, so the comparison on that device is +# shown to catch a difference. + +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-tensor-program.yml' + pull_request: + paths: + - 'wai/wai-quantum/**' + - 'rust-toolchain.toml' + - '.github/workflows/wai-quantum-tensor-program.yml' + +env: + IREE_VERSION: 3.12.0 + +jobs: + # Does this pull request's own diff touch the paths below? This forge matches a + # workflow's `paths:` against the push, so a merge of main into a branch fires the + # workflow whatever the pull request changes; every job here waits on this answer. + # A push to main, a schedule, or anything the script cannot determine answers + # `true` (scripts/pr-touches-paths.py). + relevant: + name: does this change touch it + runs-on: linux-arm64 + outputs: + run: ${{ steps.gate.outputs.run }} + steps: + - uses: actions/checkout@v5 + with: + fetch-depth: 0 + - id: gate + run: | + set -euo pipefail + r=$(python3 scripts/pr-touches-paths.py .github/workflows/wai-quantum-tensor-program.yml \ + "${{ github.event.pull_request.base.sha }}" \ + "${{ github.event.pull_request.head.sha }}") + echo "this change touches wai-quantum-tensor-program: $r" + echo "run=$r" >> "$GITHUB_OUTPUT" + + cpu: + name: tensor program · CPU, mutants, GPU builds + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux + timeout-minutes: 60 + container: + image: debian@sha256:913f6706df59a68922d1dd08f78c2476560a8d367897200a6005b00e5f67c2d5 + options: --memory=6g --memory-swap=6g + steps: + - name: Tools + shell: bash + run: | + set -euo pipefail + apt-get update -qq + apt-get install -y -qq --no-install-recommends ca-certificates curl git build-essential python3 python3-venv >/dev/null + # No Node for a checkout action: the forge's archive of the exact commit, + # its embedded commit id checked. + - name: Fetch the source + shell: bash + run: | + set -euo pipefail + want="${{ github.event.pull_request.head.sha || github.sha }}" + rm -rf src src.tar.gz src.tar && mkdir src + curl -sSf -o src.tar.gz "http://100.64.224.0:3400/Transaction-Science/open-standards/archive/$want.tar.gz" + gzip -dc src.tar.gz > src.tar + got="$(git get-tar-commit-id < src.tar)" + [ "$got" = "$want" ] || { echo "::error::the archive is of $got, not $want"; exit 1; } + tar -xf src.tar -C src --strip-components=1 + rm src.tar.gz src.tar + - name: The program, built and run on the CPU; its mutants; the GPU builds + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + cache=/ci-cache/wai-quantum-tensor-program + mkdir -p "$cache" + exec 9>"$cache/.lock" + flock 9 + venv="$cache/iree-$IREE_VERSION" + [ -x "$venv/bin/iree-compile" ] || { rm -rf "$venv"; python3 -m venv "$venv"; "$venv/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } + export RUSTUP_HOME="$cache/rustup" CARGO_HOME="$cache/cargo" CARGO_TARGET_DIR="$cache/target" + export PATH="$venv/bin:$CARGO_HOME/bin:$PATH" + if ! command -v rustup >/dev/null; then + curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --no-modify-path --profile minimal --default-toolchain none + fi + rustup toolchain install + cargo test --locked --features quantum_fabric --lib -- quantum_fabric::stablehlo:: 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; } + d="$RUNNER_TEMP/program" + cargo run --locked --release -q --example stablehlo_gate --features quantum_fabric -- "$d" + cd "$d" + n=4096 + rows="--function=rows --input=@e0.npy --input=@e1.npy --input=@e2.npy --input=@e3.npy --input=@e4.npy --input=@e5.npy --input=@e6.npy --input=@e7.npy --input=@a0r.npy --input=@a0i.npy --input=@a1r.npy --input=@a1i.npy --expected_output=${n}xi64=@r0.bin --expected_output=${n}xi64=@r1.bin --expected_output=${n}xi64=@r2.bin --expected_output=${n}xi64=@r3.bin" + probe="--function=probe --input=@pa.npy --input=@pb.npy --expected_output=${n}xi64=@prod.bin --expected_output=${n}xi64=@fx.bin" + cpu="--iree-hal-target-device=local --iree-hal-local-target-device-backends=llvm-cpu" + # gate: rows and probe exact; m1, m2 (the rounding, the shift): both wrong; + # m3 (row 0's sign): the rows wrong, the probe exact. + for g in gate gate_m1 gate_m2 gate_m3; do + iree-compile $g.mlir $cpu -o $g.vmfb + r=0; iree-run-module --device=local-task --module=$g.vmfb $rows > $g.rows.log 2>&1 || r=1 + p=0; iree-run-module --device=local-task --module=$g.vmfb $probe > $g.probe.log 2>&1 || p=1 + echo "$g on the CPU: rows $([ $r = 0 ] && echo exact || echo differ), probe $([ $p = 0 ] && echo exact || echo differ)" + case "$g/$r$p" in + gate/00|gate_m1/11|gate_m2/11|gate_m3/10) ;; + *) echo "::error::$g gave rows=$r probe=$p"; head -20 $g.rows.log $g.probe.log; exit 1 ;; + esac + done + grep -q "all function outputs matched" gate.rows.log || { echo "::error::the runtime did not compare the rows"; exit 1; } + for t in gfx942 gfx1100 gfx1201; do + iree-compile gate.mlir --iree-hal-target-backends=rocm --iree-rocm-target=$t -o gate.$t.vmfb + echo "built for $t: $(stat -c %s gate.$t.vmfb) bytes" + done + + gpu: + name: tensor program · GPUs (jetson-hub) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux-hw + timeout-minutes: 60 + env: + # The Vulkan devices this host has: its discrete and integrated GPUs and + # the software one. Every one must give the exact rows. + VULKAN_DEVICES: '3' + 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 program on every Vulkan device and through CUDA + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + export PATH="$HOME/.cargo/bin:$PATH" + venv="$HOME/iree-$IREE_VERSION" + [ -x "$venv/bin/iree-compile" ] || { rm -rf "$venv"; python3 -m venv "$venv"; "$venv/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } + export PATH="$venv/bin:$PATH" + d="$RUNNER_TEMP/program" + CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-tensor-program" nice -n 10 cargo run --locked --release -q --example stablehlo_gate --features quantum_fabric -- "$d" + cd "$d" + n=4096 + rows="--function=rows --input=@e0.npy --input=@e1.npy --input=@e2.npy --input=@e3.npy --input=@e4.npy --input=@e5.npy --input=@e6.npy --input=@e7.npy --input=@a0r.npy --input=@a0i.npy --input=@a1r.npy --input=@a1i.npy --expected_output=${n}xi64=@r0.bin --expected_output=${n}xi64=@r1.bin --expected_output=${n}xi64=@r2.bin --expected_output=${n}xi64=@r3.bin" + probe="--function=probe --input=@pa.npy --input=@pb.npy --expected_output=${n}xi64=@prod.bin --expected_output=${n}xi64=@fx.bin" + # A capability profile with 64-bit integers; the default has none. + for g in gate gate_m3; do + iree-compile $g.mlir --iree-hal-target-device=vulkan --iree-vulkan-target=ampere -o $g.vulkan.vmfb + iree-compile $g.mlir --iree-hal-target-device=cuda --iree-cuda-target=sm_89 -o $g.cuda.vmfb + done + devices="$(iree-run-module --list_devices 2>/dev/null | grep -oE 'vulkan://[^ ]+' | sort -u)" + count="$(grep -c . <<<"$devices")" + [ "$count" = "$VULKAN_DEVICES" ] || { echo "::error::$count Vulkan devices, not $VULKAN_DEVICES"; iree-run-module --list_devices; exit 1; } + # On each device: the program exact, and a mutant's rows caught (its + # probe is the program's, so exact too). + for dev in $devices cuda; do + api=vulkan; [ "$dev" = cuda ] && api=cuda + iree-run-module --device="$dev" --module=gate.$api.vmfb $rows > rows.log 2>&1 || { echo "::error::$dev: the rows differ"; head -10 rows.log; exit 1; } + iree-run-module --device="$dev" --module=gate.$api.vmfb $probe > probe.log 2>&1 || { echo "::error::$dev: the probe differs"; head -10 probe.log; exit 1; } + if iree-run-module --device="$dev" --module=gate_m3.$api.vmfb $rows > m3.log 2>&1; then echo "::error::$dev: the mutant's rows passed"; exit 1; fi + iree-run-module --device="$dev" --module=gate_m3.$api.vmfb $probe > m3p.log 2>&1 || { echo "::error::$dev: the mutant's probe differs"; head -10 m3p.log; exit 1; } + echo "$dev: rows and probe exact; the mutant's rows caught" + done + + metal: + name: tensor program · Metal (macos) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: macos + timeout-minutes: 60 + steps: + - uses: actions/checkout@v5 + # launchd's PATH omits cargo on the macOS host runner. + - name: Put cargo on PATH + shell: bash + run: echo "$HOME/.cargo/bin" >> "$GITHUB_PATH" + - name: The program on the platform's GPU + shell: bash + working-directory: wai/wai-quantum + run: | + set -euo pipefail + # The system Python is older than the compiler's packages support. + py="$(command -v /opt/homebrew/bin/python3.12 python3.12 python3 | head -1)" + venv="$HOME/iree-$IREE_VERSION" + [ -x "$venv/bin/iree-compile" ] || { rm -rf "$venv"; "$py" -m venv "$venv"; "$venv/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } + export PATH="$venv/bin:$PATH" + d="$RUNNER_TEMP/program" + nice -n 10 cargo run --locked --release -q --example stablehlo_gate --features quantum_fabric -- "$d" + cd "$d" + n=4096 + rows="--function=rows --input=@e0.npy --input=@e1.npy --input=@e2.npy --input=@e3.npy --input=@e4.npy --input=@e5.npy --input=@e6.npy --input=@e7.npy --input=@a0r.npy --input=@a0i.npy --input=@a1r.npy --input=@a1i.npy --expected_output=${n}xi64=@r0.bin --expected_output=${n}xi64=@r1.bin --expected_output=${n}xi64=@r2.bin --expected_output=${n}xi64=@r3.bin" + probe="--function=probe --input=@pa.npy --input=@pb.npy --expected_output=${n}xi64=@prod.bin --expected_output=${n}xi64=@fx.bin" + # Shader source in the module, compiled by the system at load: no + # developer toolchain needed. + for g in gate gate_m3; do + iree-compile $g.mlir --iree-hal-target-device=metal --iree-metal-compile-to-metallib=false -o $g.metal.vmfb + done + iree-run-module --device=metal --module=gate.metal.vmfb $rows > rows.log 2>&1 || { echo "::error::the rows differ"; head -10 rows.log; exit 1; } + iree-run-module --device=metal --module=gate.metal.vmfb $probe > probe.log 2>&1 || { echo "::error::the probe differs"; head -10 probe.log; exit 1; } + if iree-run-module --device=metal --module=gate_m3.metal.vmfb $rows > m3.log 2>&1; then echo "::error::the mutant's rows passed"; exit 1; fi + iree-run-module --device=metal --module=gate_m3.metal.vmfb $probe > m3p.log 2>&1 || { echo "::error::the mutant's probe differs"; head -10 m3p.log; exit 1; } + echo "metal: rows and probe exact; the mutant's rows caught" diff --git a/wai/CHANGELOG.md b/wai/CHANGELOG.md index 23ecd492f..93f958108 100644 --- a/wai/CHANGELOG.md +++ b/wai/CHANGELOG.md @@ -6,7 +6,7 @@ wire format may change and such changes are recorded here. ## v0.4.0 — current -### wai-quantum 0.5.2: the dataflow device's vector engine +### wai-quantum 0.5.2: the dataflow device's vector engine, and the gate as a portable tensor program Additive. The default features, `full`, and every pinned digest are unchanged. @@ -27,6 +27,20 @@ Additive. The default features, `full`, and every pinned digest are unchanged. four pinned certificates, circuits up to 20 qubits match native, 3 of 3 kernel mutants (the same three the other members carry) killed. Every certificate is labelled simulator. +- **The gate as a portable tensor program** (`quantum_fabric::stablehlo`): the gate's rows + and the arithmetic probe as a StableHLO module, the input of compilers for tensor + processors, GPUs and CPUs. Every value is a 64-bit integer that never overflows + (`|e·a| < 2^62`; a row of four `fx` stays below `2^35`), so a compiler must give the + exact rows whatever it does with them. It is not a member: the program leaves the + crate, and the compiler's own runtime compares its outputs with the exact ones. The + `stablehlo_gate` example writes the module, its three mutants (the same three), a + batch of 4096 pairs whose gate and states reach the 32-bit edges, and the outputs a + build must give. +- **Measured** with an open-source tensor compiler (IREE 3.12.0): exact on the CPU; on + every Vulkan device of the hardware host (a discrete GPU, an integrated GPU, and a + software one) and through CUDA; through Metal on two generations of the macOS + platform's GPU. Each mutant fails where it should. Builds for three GPU generations + the fleet lacks compile but are not run. CI: `wai-quantum-tensor-program`. ### wai-quantum 0.5.1: neural-processor, matrix-extension, vector-unit, signal-processor and dataflow members diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml index 4e721a632..11d067100 100644 --- a/wai/wai-quantum/Cargo.toml +++ b/wai/wai-quantum/Cargo.toml @@ -241,6 +241,10 @@ required-features = ["quantum_fabric_dataflow"] name = "fabric_expected" required-features = ["quantum_fabric"] +[[example]] +name = "stablehlo_gate" +required-features = ["quantum_fabric"] + [[example]] name = "rv32_listing" required-features = ["quantum_fabric_dataflow"] diff --git a/wai/wai-quantum/README.md b/wai/wai-quantum/README.md index 2c8c0c998..3eddc770d 100644 --- a/wai/wai-quantum/README.md +++ b/wai/wai-quantum/README.md @@ -103,7 +103,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` | **the compute fabric**: members (the CPU and accelerators, quantum processors among them), each kernel implementation certified against the CPU reference in the process that uses it (bit-exact, within a stated tolerance, or statistical by classical fidelity over stated shots), with the certificate saying where it was earned (silicon, an emulator, or a simulator); placement by parity, provenance and metered energy; a signed receipt naming the member and implementation behind a result. Two kernels: apply (a circuit's state) and sample (its measurement counts), the latter with the CPU reference sampler, a corpus shaped for quantum processors, and a virtual quantum processor whose stated noise model lets every step of a processor's path run and be tested without one; its certificates say they were earned in a simulator | +| `quantum_fabric` | **the compute fabric**: members (the CPU and accelerators, quantum processors among them), each kernel implementation certified against the CPU reference in the process that uses it (bit-exact, within a stated tolerance, or statistical by classical fidelity over stated shots), with the certificate saying where it was earned (silicon, an emulator, or a simulator); placement by parity, provenance and metered energy; a signed receipt naming the member and implementation behind a result. Two kernels: apply (a circuit's state) and sample (its measurement counts), the latter with the CPU reference sampler, a corpus shaped for quantum processors, and a virtual quantum processor whose stated noise model lets every step of a processor's path run and be tested without one; its certificates say they were earned in a simulator. `quantum_fabric::stablehlo` writes the gate as a portable tensor program (StableHLO), every value a 64-bit integer that never overflows, for any tensor compiler to build for its own devices; the `stablehlo_gate` example writes it with a batch of inputs and the exact outputs a build must give | | `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 features that reach outside the crate | | `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. With `quantum_sv` too, the **dense f64 state vector** runs on the device bit for bit as on the CPU: the CPU's own launches (diagonal, one-qubit and grouped formulas), each an IEEE binary64 operation with no fused multiply-add, certified against a pinned corpus digest | | `quantum_fabric_opencl` | **the fabric's OpenCL member**, opt-in and not in `full`, native targets only: the same exact 64-bit kernel in OpenCL C, built at run time by each device's own compiler, one launch per gate, on any OpenCL GPU or CPU (an OpenCL accelerator is refused: no fabric class names it). The ICD loader is opened when first used, so a build with the feature runs on a host without OpenCL and finds no member there, saying why. Each device is probed and certified as the CUDA member's are. With `quantum_sv` too, the dense f64 state vector runs bit for bit as on the CPU on every device with binary64 (contraction off); a device without it refuses, saying so | diff --git a/wai/wai-quantum/examples/stablehlo_gate.rs b/wai/wai-quantum/examples/stablehlo_gate.rs new file mode 100644 index 000000000..fb5bbfd6d --- /dev/null +++ b/wai/wai-quantum/examples/stablehlo_gate.rs @@ -0,0 +1,69 @@ +//! Writes the gate as a portable tensor program, a batch of inputs, and the +//! exact outputs a compiler's build of it must give: +//! +//! ```text +//! cargo run --example stablehlo_gate --features quantum_fabric -- +//! ``` +//! +//! In ``: `gate.mlir` (conforming) and `gate_m1.mlir` … `gate_m3.mlir` +//! (mutants, each of which must fail); for `rows`, the inputs `e0.npy` … +//! `e7.npy`, `a0r.npy` … `a1i.npy` and the expected `r0.bin` … `r3.bin`; for +//! `probe`, `pa.npy`, `pb.npy` and the expected `prod.bin`, `fx.bin`. Each +//! expected file is the raw little-endian 64-bit values, `4096xi64`, so a +//! runtime compares them under the module's own (signless) element type. +use wai_quantum::quantum_fabric::lanes; +use wai_quantum::quantum_fabric::stablehlo::{CONFORMING, Constants, module, npy_i32, npy_scalar_i64, reference_rows}; + +/// Pairs per batch. +const N: usize = 4096; + +fn raw(v: &[i64]) -> Vec { + v.iter().flat_map(|x| x.to_le_bytes()).collect() +} + +fn write(dir: &std::path::Path, name: &str, bytes: &[u8]) { + std::fs::write(dir.join(name), bytes).unwrap_or_else(|e| panic!("{name}: {e}")); +} + +fn main() { + let dir = std::path::PathBuf::from(std::env::args().nth(1).expect("an output directory")); + std::fs::create_dir_all(&dir).expect("the directory"); + write(&dir, "gate.mlir", module(N, CONFORMING).as_bytes()); + for (name, c) in [ + ("gate_m1.mlir", Constants { round: 1 << 28, ..CONFORMING }), + ("gate_m2.mlir", Constants { shift: 29, ..CONFORMING }), + ("gate_m3.mlir", Constants { last_add: true, ..CONFORMING }), + ] { + write(&dir, name, module(N, c).as_bytes()); + } + // A gate whose entries reach the 32-bit edges, and states that do too. + let m: [i64; 8] = [i64::from(i32::MAX), i64::from(i32::MIN), 1 << 30, -(1 << 30), 0x1234_5678, -0x0765_4321, 1, -1]; + let mut s = 0x5eed_u64; + let mut next = move || { + 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)) as u32 as i32 + }; + let edges = [0, 1, -1, i32::MAX, i32::MIN, 1 << 29, -(1 << 29) - 1]; + let cols: [Vec; 4] = std::array::from_fn(|k| (0..N).map(|i| if (i + k) % 9 == 0 { edges[(i / 9 + k) % edges.len()] } else { next() }).collect()); + for (k, e) in m.iter().enumerate() { + write(&dir, &format!("e{k}.npy"), &npy_scalar_i64(*e)); + } + for (name, c) in ["a0r", "a0i", "a1r", "a1i"].iter().zip(&cols) { + write(&dir, &format!("{name}.npy"), &npy_i32(c)); + } + for (r, v) in reference_rows(&m, &cols, CONFORMING).iter().enumerate() { + write(&dir, &format!("r{r}.bin"), &raw(v)); + } + // The probe: every pair of the edge set, then random pairs, to N. + let mut pairs = lanes::probe_pairs(N, 0x0cda); + pairs.truncate(N); + let (pa, pb): (Vec, Vec) = pairs.iter().copied().unzip(); + write(&dir, "pa.npy", &npy_i32(&pa)); + write(&dir, "pb.npy", &npy_i32(&pb)); + write(&dir, "prod.bin", &raw(&pairs.iter().map(|&(a, b)| i64::from(a) * i64::from(b)).collect::>())); + write(&dir, "fx.bin", &raw(&pairs.iter().map(|&(a, b)| (i64::from(a) * i64::from(b) + (1 << 29)) >> 30).collect::>())); + println!("wrote the module, its mutants, and a batch of {N} pairs to {}", dir.display()); +} diff --git a/wai/wai-quantum/src/quantum_fabric.rs b/wai/wai-quantum/src/quantum_fabric.rs index 89ae10b3d..955b716c9 100644 --- a/wai/wai-quantum/src/quantum_fabric.rs +++ b/wai/wai-quantum/src/quantum_fabric.rs @@ -2059,6 +2059,7 @@ pub mod tensix; #[cfg(feature = "quantum_fabric_dataflow")] pub mod vector_engine; pub mod sampling; +pub mod stablehlo; #[cfg(feature = "quantum_sv")] pub mod sv64; pub mod vqpu; diff --git a/wai/wai-quantum/src/quantum_fabric/stablehlo.rs b/wai/wai-quantum/src/quantum_fabric/stablehlo.rs new file mode 100644 index 000000000..2d679efcb --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/stablehlo.rs @@ -0,0 +1,190 @@ +//! The gate as a portable tensor program (StableHLO, as MLIR text): the input +//! of compilers for tensor processors, GPUs and CPUs. Every value is a 64-bit +//! integer that never overflows: `|e·a| < 2^62` and a row of four +//! `fx = (e·a + 2^29) >> 30` stays below 2^35, so each compiler must give the +//! exact rows whatever it does with them. +//! +//! - `rows(e0, …, e7, a0r, a0i, a1r, a1i)`: the gate's entries as scalars, a +//! batch of `n` pairs' columns as 32-bit integers; the four rows, 64-bit. +//! - `probe(a, b)`: each pair's `a·b` and `fx(a, b)`, 64-bit. +//! +//! [`npy`] writes inputs and expected outputs in the array format such tools +//! read. + +use std::fmt::Write as _; + +/// The kernel's constants: the rounding constant, the shift, and whether row +/// 0's last real term is added. +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub struct Constants { + pub round: i64, + pub shift: i64, + pub last_add: bool, +} + +pub const CONFORMING: Constants = Constants { round: 1 << 29, shift: 30, last_add: false }; + +struct Body { + text: String, + next: usize, + n: usize, +} + +impl Body { + fn ty(&self) -> String { + format!("tensor<{}xi64>", self.n) + } + fn op(&mut self, rhs: String) -> String { + let v = format!("%v{}", self.next); + self.next += 1; + let _ = writeln!(self.text, " {v} = {rhs}"); + v + } + fn constant(&mut self, c: i64) -> String { + let t = self.ty(); + self.op(format!("stablehlo.constant dense<{c}> : {t}")) + } + fn binary(&mut self, name: &str, a: &str, b: &str) -> String { + let t = self.ty(); + self.op(format!("stablehlo.{name} {a}, {b} : {t}")) + } + fn widen(&mut self, a: &str) -> String { + let (n, t) = (self.n, self.ty()); + self.op(format!("stablehlo.convert {a} : (tensor<{n}xi32>) -> {t}")) + } + fn splat(&mut self, scalar: &str) -> String { + let t = self.ty(); + self.op(format!("stablehlo.broadcast_in_dim {scalar}, dims = [] : (tensor) -> {t}")) + } + /// `(e·a + round) >> shift`. + fn fx(&mut self, e: &str, a: &str, round: &str, shift: &str) -> String { + let p = self.binary("multiply", e, a); + let r = self.binary("add", &p, round); + self.binary("shift_right_arithmetic", &r, shift) + } +} + +/// The module: `rows` and `probe` over batches of `n` pairs. +pub fn module(n: usize, c: Constants) -> String { + let t = format!("tensor<{n}xi64>"); + let ti = format!("tensor<{n}xi32>"); + let mut out = String::new(); + let _ = writeln!(out, "// The gate's rows and the arithmetic probe, exact in 64-bit integers."); + let _ = writeln!(out, "module @wai_quantum_gate {{"); + // rows + let mut b = Body { text: String::new(), next: 0, n }; + let round = b.constant(c.round); + let shift = b.constant(c.shift); + let e: Vec = (0..8).map(|k| b.splat(&format!("%e{k}"))).collect(); + let a: Vec = ["%a0r", "%a0i", "%a1r", "%a1i"].iter().map(|x| b.widen(x)).collect(); + let last = if c.last_add { "add" } else { "subtract" }; + let terms: [[(usize, usize, &str); 4]; 4] = [ + [(0, 0, "add"), (1, 1, "subtract"), (2, 2, "add"), (3, 3, last)], + [(0, 1, "add"), (1, 0, "add"), (2, 3, "add"), (3, 2, "add")], + [(4, 0, "add"), (5, 1, "subtract"), (6, 2, "add"), (7, 3, "subtract")], + [(4, 1, "add"), (5, 0, "add"), (6, 3, "add"), (7, 2, "add")], + ]; + let mut rows = Vec::new(); + for row in terms { + let mut acc = b.fx(&e[row[0].0], &a[row[0].1], &round, &shift); + for &(k, op, how) in &row[1..] { + let f = b.fx(&e[k], &a[op], &round, &shift); + acc = b.binary(how, &acc, &f); + } + rows.push(acc); + } + let params = (0..8).map(|k| format!("%e{k}: tensor")).chain(["%a0r", "%a0i", "%a1r", "%a1i"].iter().map(|x| format!("{x}: {ti}"))).collect::>().join(", "); + let _ = writeln!(out, " func.func @rows({params}) -> ({t}, {t}, {t}, {t}) {{"); + out.push_str(&b.text); + let _ = writeln!(out, " return {} : {t}, {t}, {t}, {t}", rows.join(", ")); + let _ = writeln!(out, " }}"); + // probe + let mut b = Body { text: String::new(), next: 0, n }; + let round = b.constant(c.round); + let shift = b.constant(c.shift); + let (wa, wb) = (b.widen("%a"), b.widen("%b")); + let prod = b.binary("multiply", &wa, &wb); + let fx = b.fx(&wb, &wa, &round, &shift); + let _ = writeln!(out, " func.func @probe(%a: {ti}, %b: {ti}) -> ({t}, {t}) {{"); + out.push_str(&b.text); + let _ = writeln!(out, " return {prod}, {fx} : {t}, {t}"); + let _ = writeln!(out, " }}"); + let _ = writeln!(out, "}}"); + out +} + +/// The rows the module must give, from the CPU's arithmetic. +pub fn reference_rows(m: &[i64; 8], cols: &[Vec; 4], c: Constants) -> [Vec; 4] { + let fx = |e: i64, a: i32| (e * i64::from(a) + c.round) >> c.shift; + let last = if c.last_add { 1 } else { -1 }; + let mut out: [Vec; 4] = Default::default(); + for (((&a0r, &a0i), &a1r), &a1i) in cols[0].iter().zip(&cols[1]).zip(&cols[2]).zip(&cols[3]) { + out[0].push(fx(m[0], a0r) - fx(m[1], a0i) + fx(m[2], a1r) + last * fx(m[3], a1i)); + out[1].push(fx(m[0], a0i) + fx(m[1], a0r) + fx(m[2], a1i) + fx(m[3], a1r)); + out[2].push(fx(m[4], a0r) - fx(m[5], a0i) + fx(m[6], a1r) - fx(m[7], a1i)); + out[3].push(fx(m[4], a0i) + fx(m[5], a0r) + fx(m[6], a1i) + fx(m[7], a1r)); + } + out +} + +/// A little-endian array in the `.npy` format (version 1.0): `descr` ` Vec { + let dims = match shape { + [] => "()".to_string(), + [n] => format!("({n},)"), + _ => format!("({})", shape.iter().map(|d| d.to_string()).collect::>().join(", ")), + }; + let mut header = format!("{{'descr': '{descr}', 'fortran_order': False, 'shape': {dims}, }}"); + // Magic (6) + version (2) + length (2) + header, padded to 64 bytes, ending in '\n'. + while (10 + header.len() + 1) % 64 != 0 { + header.push(' '); + } + header.push('\n'); + let mut out = b"\x93NUMPY\x01\x00".to_vec(); + out.extend_from_slice(&(header.len() as u16).to_le_bytes()); + out.extend_from_slice(header.as_bytes()); + out.extend_from_slice(data); + out +} + +pub fn npy_i32(v: &[i32]) -> Vec { + npy(">()) +} + +pub fn npy_i64(v: &[i64]) -> Vec { + npy(">()) +} + +pub fn npy_scalar_i64(v: i64) -> Vec { + npy(" Date: Thu, 15 Oct 2026 04:20:41 -0400 Subject: [PATCH 03/19] wai-quantum: the tensor program through a second compiler, held to stated outcomes XLA (through JAX 0.11.2, pinned) now builds and runs the gate's tensor program beside IREE: on the x86-64 CPU (the program and all three mutants), on the discrete GPU through CUDA, and on the macOS runner's arm64 CPU. ci/tensor-program/xla_run.py renames the chosen function to main, runs it and compares every output with the exact one. ci/tensor-program/check.sh runs one build on one device and holds it to the outcome it must have (rows, then probe: exact or differs). A run counts only when it states its outcome: the runtime's SUCCESS or FAILED comparison line, or the driver's exact or differs line. Any other ending (a device that will not open, an exception) fails the check, so it can never read as a mutant caught. The three workflow legs call it instead of repeating the argument lists. Co-Authored-By: Claude Opus 5.5 --- .../workflows/wai-quantum-tensor-program.yml | 131 +++++++++--------- wai/CHANGELOG.md | 18 ++- wai/wai-quantum/README.md | 2 +- wai/wai-quantum/ci/tensor-program/check.sh | 71 ++++++++++ wai/wai-quantum/ci/tensor-program/xla_run.py | 60 ++++++++ 5 files changed, 207 insertions(+), 75 deletions(-) create mode 100755 wai/wai-quantum/ci/tensor-program/check.sh create mode 100644 wai/wai-quantum/ci/tensor-program/xla_run.py diff --git a/.github/workflows/wai-quantum-tensor-program.yml b/.github/workflows/wai-quantum-tensor-program.yml index ce2038b83..aee168283 100644 --- a/.github/workflows/wai-quantum-tensor-program.yml +++ b/.github/workflows/wai-quantum-tensor-program.yml @@ -2,20 +2,24 @@ name: wai-quantum portable tensor program # The gate as a portable tensor program (quantum_fabric::stablehlo, StableHLO # as MLIR text): every value a 64-bit integer that never overflows, so a -# compiler for any device must give the exact rows. An open-source tensor -# compiler and runtime (IREE 3.12.0, Apache-2.0, pinned) builds it for each -# device here, and its runtime compares the outputs with the exact ones the -# crate writes: +# compiler for any device must give the exact rows. Two independent +# open-source tensor compilers, both Apache-2.0 and pinned, build it here: +# IREE (3.12.0), whose runtime compares the outputs with the exact ones the +# crate writes, and XLA (through JAX 0.11.2, the compiler tensor processors +# take this program through), whose outputs ci/tensor-program/xla_run.py +# compares. ci/tensor-program/check.sh holds every run to the outcome it must +# have, and fails a run that ends without stating one. # -# - cpu (the container runner): the x86-64 build; the three mutants, each of -# which must fail where it should; builds for three GPU generations the -# fleet has none of (compiled, not run); +# - cpu (the container runner): both compilers on the x86-64 CPU, the program +# and its three mutants (each must fail where it should); builds for three +# GPU generations the fleet has none of (compiled, not run); # - gpu (the hardware host): every Vulkan device the host has (its discrete # and integrated GPUs and the software one), and the discrete GPU through -# CUDA; -# - metal (the macOS runner): the platform's GPU through Metal. -# Each GPU leg also runs one mutant, so the comparison on that device is -# shown to catch a difference. +# CUDA with each compiler; +# - metal (the macOS runner): the platform's GPU through Metal, and XLA on the +# arm64 CPU. +# Each leg past the first also runs one mutant on each device, so the +# comparison there is shown to catch a difference. concurrency: group: ${{ github.workflow }}-${{ github.ref }} @@ -36,6 +40,7 @@ on: env: IREE_VERSION: 3.12.0 + JAX_VERSION: 0.11.2 jobs: # Does this pull request's own diff touch the paths below? This forge matches a @@ -91,7 +96,7 @@ jobs: [ "$got" = "$want" ] || { echo "::error::the archive is of $got, not $want"; exit 1; } tar -xf src.tar -C src --strip-components=1 rm src.tar.gz src.tar - - name: The program, built and run on the CPU; its mutants; the GPU builds + - name: The program and its mutants on the CPU, both compilers; the GPU builds shell: bash working-directory: src/wai/wai-quantum run: | @@ -100,36 +105,28 @@ jobs: mkdir -p "$cache" exec 9>"$cache/.lock" flock 9 - venv="$cache/iree-$IREE_VERSION" - [ -x "$venv/bin/iree-compile" ] || { rm -rf "$venv"; python3 -m venv "$venv"; "$venv/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } + venv="$cache/compilers-iree-$IREE_VERSION-jax-$JAX_VERSION" + [ -x "$venv/bin/iree-compile" ] || { rm -rf "$venv"; python3 -m venv "$venv"; "$venv/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION" "jax==$JAX_VERSION"; } export RUSTUP_HOME="$cache/rustup" CARGO_HOME="$cache/cargo" CARGO_TARGET_DIR="$cache/target" - export PATH="$venv/bin:$CARGO_HOME/bin:$PATH" + export PATH="$venv/bin:$CARGO_HOME/bin:$PATH" XLA_PYTHON="$venv/bin/python" if ! command -v rustup >/dev/null; then curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --no-modify-path --profile minimal --default-toolchain none fi rustup toolchain install cargo test --locked --features quantum_fabric --lib -- quantum_fabric::stablehlo:: 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; } + check="bash $PWD/ci/tensor-program/check.sh" d="$RUNNER_TEMP/program" cargo run --locked --release -q --example stablehlo_gate --features quantum_fabric -- "$d" cd "$d" - n=4096 - rows="--function=rows --input=@e0.npy --input=@e1.npy --input=@e2.npy --input=@e3.npy --input=@e4.npy --input=@e5.npy --input=@e6.npy --input=@e7.npy --input=@a0r.npy --input=@a0i.npy --input=@a1r.npy --input=@a1i.npy --expected_output=${n}xi64=@r0.bin --expected_output=${n}xi64=@r1.bin --expected_output=${n}xi64=@r2.bin --expected_output=${n}xi64=@r3.bin" - probe="--function=probe --input=@pa.npy --input=@pb.npy --expected_output=${n}xi64=@prod.bin --expected_output=${n}xi64=@fx.bin" - cpu="--iree-hal-target-device=local --iree-hal-local-target-device-backends=llvm-cpu" - # gate: rows and probe exact; m1, m2 (the rounding, the shift): both wrong; - # m3 (row 0's sign): the rows wrong, the probe exact. - for g in gate gate_m1 gate_m2 gate_m3; do - iree-compile $g.mlir $cpu -o $g.vmfb - r=0; iree-run-module --device=local-task --module=$g.vmfb $rows > $g.rows.log 2>&1 || r=1 - p=0; iree-run-module --device=local-task --module=$g.vmfb $probe > $g.probe.log 2>&1 || p=1 - echo "$g on the CPU: rows $([ $r = 0 ] && echo exact || echo differ), probe $([ $p = 0 ] && echo exact || echo differ)" - case "$g/$r$p" in - gate/00|gate_m1/11|gate_m2/11|gate_m3/10) ;; - *) echo "::error::$g gave rows=$r probe=$p"; head -20 $g.rows.log $g.probe.log; exit 1 ;; - esac + # The program exact; m1, m2 (the rounding, the shift) both functions + # wrong; m3 (row 0's last sign) the rows wrong, the probe exact. + for g in 'gate:==' 'gate_m1:!!' 'gate_m2:!!' 'gate_m3:!='; do + m="${g%%:*}" want="${g#*:}" + iree-compile $m.mlir --iree-hal-target-device=local --iree-hal-local-target-device-backends=llvm-cpu --iree-llvmcpu-target-cpu=host -o $m.cpu.vmfb + $check iree local-task $m.cpu.vmfb "$want" + $check xla cpu $m.mlir "$want" done - grep -q "all function outputs matched" gate.rows.log || { echo "::error::the runtime did not compare the rows"; exit 1; } for t in gfx942 gfx1100 gfx1201; do iree-compile gate.mlir --iree-hal-target-backends=rocm --iree-rocm-target=$t -o gate.$t.vmfb echo "built for $t: $(stat -c %s gate.$t.vmfb) bytes" @@ -159,42 +156,40 @@ jobs: git checkout -q FETCH_HEAD got="$(git rev-parse HEAD)" [ "$got" = "$want" ] || { echo "::error::fetched $got, not $want"; exit 1; } - - name: The program on every Vulkan device and through CUDA + - name: The program on every Vulkan device, and through CUDA with each compiler shell: bash working-directory: src/wai/wai-quantum run: | set -euo pipefail export PATH="$HOME/.cargo/bin:$PATH" - venv="$HOME/iree-$IREE_VERSION" - [ -x "$venv/bin/iree-compile" ] || { rm -rf "$venv"; python3 -m venv "$venv"; "$venv/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } - export PATH="$venv/bin:$PATH" + iree="$HOME/iree-$IREE_VERSION" + [ -x "$iree/bin/iree-compile" ] || { rm -rf "$iree"; python3 -m venv "$iree"; "$iree/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } + xla="$HOME/xla-$JAX_VERSION-cuda13" + [ -x "$xla/bin/python" ] && "$xla/bin/python" -I -c 'import jax' 2>/dev/null || { rm -rf "$xla"; python3 -m venv "$xla"; "$xla/bin/pip" install -q "jax[cuda13]==$JAX_VERSION"; } + export PATH="$iree/bin:$PATH" XLA_PYTHON="$xla/bin/python" + check="bash $PWD/ci/tensor-program/check.sh" d="$RUNNER_TEMP/program" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-tensor-program" nice -n 10 cargo run --locked --release -q --example stablehlo_gate --features quantum_fabric -- "$d" cd "$d" - n=4096 - rows="--function=rows --input=@e0.npy --input=@e1.npy --input=@e2.npy --input=@e3.npy --input=@e4.npy --input=@e5.npy --input=@e6.npy --input=@e7.npy --input=@a0r.npy --input=@a0i.npy --input=@a1r.npy --input=@a1i.npy --expected_output=${n}xi64=@r0.bin --expected_output=${n}xi64=@r1.bin --expected_output=${n}xi64=@r2.bin --expected_output=${n}xi64=@r3.bin" - probe="--function=probe --input=@pa.npy --input=@pb.npy --expected_output=${n}xi64=@prod.bin --expected_output=${n}xi64=@fx.bin" - # A capability profile with 64-bit integers; the default has none. - for g in gate gate_m3; do - iree-compile $g.mlir --iree-hal-target-device=vulkan --iree-vulkan-target=ampere -o $g.vulkan.vmfb - iree-compile $g.mlir --iree-hal-target-device=cuda --iree-cuda-target=sm_89 -o $g.cuda.vmfb + # A Vulkan capability profile with 64-bit integers; the default has none. + for m in gate gate_m3; do + iree-compile $m.mlir --iree-hal-target-device=vulkan --iree-vulkan-target=ampere -o $m.vulkan.vmfb + iree-compile $m.mlir --iree-hal-target-device=cuda --iree-cuda-target=sm_89 -o $m.cuda.vmfb done devices="$(iree-run-module --list_devices 2>/dev/null | grep -oE 'vulkan://[^ ]+' | sort -u)" count="$(grep -c . <<<"$devices")" [ "$count" = "$VULKAN_DEVICES" ] || { echo "::error::$count Vulkan devices, not $VULKAN_DEVICES"; iree-run-module --list_devices; exit 1; } - # On each device: the program exact, and a mutant's rows caught (its - # probe is the program's, so exact too). - for dev in $devices cuda; do - api=vulkan; [ "$dev" = cuda ] && api=cuda - iree-run-module --device="$dev" --module=gate.$api.vmfb $rows > rows.log 2>&1 || { echo "::error::$dev: the rows differ"; head -10 rows.log; exit 1; } - iree-run-module --device="$dev" --module=gate.$api.vmfb $probe > probe.log 2>&1 || { echo "::error::$dev: the probe differs"; head -10 probe.log; exit 1; } - if iree-run-module --device="$dev" --module=gate_m3.$api.vmfb $rows > m3.log 2>&1; then echo "::error::$dev: the mutant's rows passed"; exit 1; fi - iree-run-module --device="$dev" --module=gate_m3.$api.vmfb $probe > m3p.log 2>&1 || { echo "::error::$dev: the mutant's probe differs"; head -10 m3p.log; exit 1; } - echo "$dev: rows and probe exact; the mutant's rows caught" + for dev in $devices; do + $check iree "$dev" gate.vulkan.vmfb '==' + $check iree "$dev" gate_m3.vulkan.vmfb '!=' done + $check iree cuda gate.cuda.vmfb '==' + $check iree cuda gate_m3.cuda.vmfb '!=' + $check xla cuda gate.mlir '==' + $check xla cuda gate_m3.mlir '!=' metal: - name: tensor program · Metal (macos) + name: tensor program · Metal and the arm64 CPU (macos) needs: relevant if: needs.relevant.outputs.run == 'true' runs-on: macos @@ -205,29 +200,29 @@ jobs: - name: Put cargo on PATH shell: bash run: echo "$HOME/.cargo/bin" >> "$GITHUB_PATH" - - name: The program on the platform's GPU + - name: The program on the platform's GPU and, through XLA, on its CPU shell: bash working-directory: wai/wai-quantum run: | set -euo pipefail - # The system Python is older than the compiler's packages support. - py="$(command -v /opt/homebrew/bin/python3.12 python3.12 python3 | head -1)" - venv="$HOME/iree-$IREE_VERSION" - [ -x "$venv/bin/iree-compile" ] || { rm -rf "$venv"; "$py" -m venv "$venv"; "$venv/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } - export PATH="$venv/bin:$PATH" + # The system Python is older than the compilers' packages support. + py=/opt/homebrew/bin/python3.12 + [ -x "$py" ] || py="$(command -v python3.12 || command -v python3)" + iree="$HOME/iree-$IREE_VERSION" + [ -x "$iree/bin/iree-compile" ] || { rm -rf "$iree"; "$py" -m venv "$iree"; "$iree/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } + xla="$HOME/xla-$JAX_VERSION" + [ -x "$xla/bin/python" ] && "$xla/bin/python" -I -c 'import jax' 2>/dev/null || { rm -rf "$xla"; "$py" -m venv "$xla"; "$xla/bin/pip" install -q "jax==$JAX_VERSION"; } + export PATH="$iree/bin:$PATH" XLA_PYTHON="$xla/bin/python" + check="bash $PWD/ci/tensor-program/check.sh" d="$RUNNER_TEMP/program" nice -n 10 cargo run --locked --release -q --example stablehlo_gate --features quantum_fabric -- "$d" cd "$d" - n=4096 - rows="--function=rows --input=@e0.npy --input=@e1.npy --input=@e2.npy --input=@e3.npy --input=@e4.npy --input=@e5.npy --input=@e6.npy --input=@e7.npy --input=@a0r.npy --input=@a0i.npy --input=@a1r.npy --input=@a1i.npy --expected_output=${n}xi64=@r0.bin --expected_output=${n}xi64=@r1.bin --expected_output=${n}xi64=@r2.bin --expected_output=${n}xi64=@r3.bin" - probe="--function=probe --input=@pa.npy --input=@pb.npy --expected_output=${n}xi64=@prod.bin --expected_output=${n}xi64=@fx.bin" # Shader source in the module, compiled by the system at load: no # developer toolchain needed. - for g in gate gate_m3; do - iree-compile $g.mlir --iree-hal-target-device=metal --iree-metal-compile-to-metallib=false -o $g.metal.vmfb + for m in gate gate_m3; do + iree-compile $m.mlir --iree-hal-target-device=metal --iree-metal-compile-to-metallib=false -o $m.metal.vmfb done - iree-run-module --device=metal --module=gate.metal.vmfb $rows > rows.log 2>&1 || { echo "::error::the rows differ"; head -10 rows.log; exit 1; } - iree-run-module --device=metal --module=gate.metal.vmfb $probe > probe.log 2>&1 || { echo "::error::the probe differs"; head -10 probe.log; exit 1; } - if iree-run-module --device=metal --module=gate_m3.metal.vmfb $rows > m3.log 2>&1; then echo "::error::the mutant's rows passed"; exit 1; fi - iree-run-module --device=metal --module=gate_m3.metal.vmfb $probe > m3p.log 2>&1 || { echo "::error::the mutant's probe differs"; head -10 m3p.log; exit 1; } - echo "metal: rows and probe exact; the mutant's rows caught" + $check iree metal gate.metal.vmfb '==' + $check iree metal gate_m3.metal.vmfb '!=' + $check xla cpu gate.mlir '==' + $check xla cpu gate_m3.mlir '!=' diff --git a/wai/CHANGELOG.md b/wai/CHANGELOG.md index 93f958108..2958975c7 100644 --- a/wai/CHANGELOG.md +++ b/wai/CHANGELOG.md @@ -32,15 +32,21 @@ Additive. The default features, `full`, and every pinned digest are unchanged. processors, GPUs and CPUs. Every value is a 64-bit integer that never overflows (`|e·a| < 2^62`; a row of four `fx` stays below `2^35`), so a compiler must give the exact rows whatever it does with them. It is not a member: the program leaves the - crate, and the compiler's own runtime compares its outputs with the exact ones. The + crate, a compiler's runtime runs it, and its outputs are compared with the exact ones. The `stablehlo_gate` example writes the module, its three mutants (the same three), a batch of 4096 pairs whose gate and states reach the 32-bit edges, and the outputs a build must give. -- **Measured** with an open-source tensor compiler (IREE 3.12.0): exact on the CPU; on - every Vulkan device of the hardware host (a discrete GPU, an integrated GPU, and a - software one) and through CUDA; through Metal on two generations of the macOS - platform's GPU. Each mutant fails where it should. Builds for three GPU generations - the fleet lacks compile but are not run. CI: `wai-quantum-tensor-program`. +- **Measured** with two independent open-source tensor compilers, IREE 3.12.0 and XLA + (through JAX 0.11.2, the compiler tensor processors take this program through): exact + on the x86-64 CPU with each; on every Vulkan device of the hardware host (a discrete + GPU, an integrated GPU, and a software one); on the discrete GPU through CUDA with + each; through Metal on two generations of the macOS platform's GPU; with XLA on the + arm64 CPU. On the CPU each mutant fails where it should under each compiler; on every + other device, the third mutant's rows do. Builds for three GPU + generations the fleet lacks compile but are not run. `ci/tensor-program/check.sh` + holds each run to its outcome and fails a run that ends without stating one, so a + device that cannot run reads as a failure, never as a mutant caught. CI: + `wai-quantum-tensor-program`. ### wai-quantum 0.5.1: neural-processor, matrix-extension, vector-unit, signal-processor and dataflow members diff --git a/wai/wai-quantum/README.md b/wai/wai-quantum/README.md index 3eddc770d..1931628f8 100644 --- a/wai/wai-quantum/README.md +++ b/wai/wai-quantum/README.md @@ -103,7 +103,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` | **the compute fabric**: members (the CPU and accelerators, quantum processors among them), each kernel implementation certified against the CPU reference in the process that uses it (bit-exact, within a stated tolerance, or statistical by classical fidelity over stated shots), with the certificate saying where it was earned (silicon, an emulator, or a simulator); placement by parity, provenance and metered energy; a signed receipt naming the member and implementation behind a result. Two kernels: apply (a circuit's state) and sample (its measurement counts), the latter with the CPU reference sampler, a corpus shaped for quantum processors, and a virtual quantum processor whose stated noise model lets every step of a processor's path run and be tested without one; its certificates say they were earned in a simulator. `quantum_fabric::stablehlo` writes the gate as a portable tensor program (StableHLO), every value a 64-bit integer that never overflows, for any tensor compiler to build for its own devices; the `stablehlo_gate` example writes it with a batch of inputs and the exact outputs a build must give | +| `quantum_fabric` | **the compute fabric**: members (the CPU and accelerators, quantum processors among them), each kernel implementation certified against the CPU reference in the process that uses it (bit-exact, within a stated tolerance, or statistical by classical fidelity over stated shots), with the certificate saying where it was earned (silicon, an emulator, or a simulator); placement by parity, provenance and metered energy; a signed receipt naming the member and implementation behind a result. Two kernels: apply (a circuit's state) and sample (its measurement counts), the latter with the CPU reference sampler, a corpus shaped for quantum processors, and a virtual quantum processor whose stated noise model lets every step of a processor's path run and be tested without one; its certificates say they were earned in a simulator. `quantum_fabric::stablehlo` writes the gate as a portable tensor program (StableHLO), every value a 64-bit integer that never overflows, for any tensor compiler to build for its own devices; the `stablehlo_gate` example writes it with a batch of inputs and the exact outputs a build must give, and `ci/tensor-program/check.sh` runs a build and holds it to them | | `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 features that reach outside the crate | | `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. With `quantum_sv` too, the **dense f64 state vector** runs on the device bit for bit as on the CPU: the CPU's own launches (diagonal, one-qubit and grouped formulas), each an IEEE binary64 operation with no fused multiply-add, certified against a pinned corpus digest | | `quantum_fabric_opencl` | **the fabric's OpenCL member**, opt-in and not in `full`, native targets only: the same exact 64-bit kernel in OpenCL C, built at run time by each device's own compiler, one launch per gate, on any OpenCL GPU or CPU (an OpenCL accelerator is refused: no fabric class names it). The ICD loader is opened when first used, so a build with the feature runs on a host without OpenCL and finds no member there, saying why. Each device is probed and certified as the CUDA member's are. With `quantum_sv` too, the dense f64 state vector runs bit for bit as on the CPU on every device with binary64 (contraction off); a device without it refuses, saying so | diff --git a/wai/wai-quantum/ci/tensor-program/check.sh b/wai/wai-quantum/ci/tensor-program/check.sh new file mode 100755 index 000000000..e65843ca9 --- /dev/null +++ b/wai/wai-quantum/ci/tensor-program/check.sh @@ -0,0 +1,71 @@ +#!/usr/bin/env bash +# +# check.sh: run one build of the gate's tensor program on one device, both of +# its functions, and hold the outcome to the one it must have. +# +# check.sh iree +# check.sh xla +# +# Run in the directory the stablehlo_gate example wrote. is two +# characters, the rows' outcome then the probe's: `=` exact, `!` differs. The +# conforming program must give `==`; mutants m1 and m2 (the rounding constant, +# the shift) `!!`; m3 (row 0's last sign) `!=`, since its probe is the +# program's. +# +# iree: the module was built by that compiler; its runtime (`iree-run-module`, +# on PATH) runs it on and compares the outputs, and an exact outcome +# must say so in its log. xla: xla_run.py, beside this script, builds and runs +# the module on under ${XLA_PYTHON:-python3}. +# +# Exit status: 0 the outcome is the expected one; 1 it is not; 2 a usage error. + +set -euo pipefail + +[ $# = 4 ] || { sed -n '3,20p' "$0" >&2; exit 2; } +how=$1 dev=$2 module=$3 expect=$4 +case "$expect" in [=!][=!]) ;; *) echo "check.sh: is two of = and !" >&2; exit 2 ;; esac + +n=$(( $(stat -c %s r0.bin 2>/dev/null || stat -f %z r0.bin) / 8 )) +rows_in=(e0.npy e1.npy e2.npy e3.npy e4.npy e5.npy e6.npy e7.npy a0r.npy a0i.npy a1r.npy a1i.npy) +rows_out=(r0.bin r1.bin r2.bin r3.bin) +probe_in=(pa.npy pb.npy) +probe_out=(prod.bin fx.bin) + +# run -- : 0 exact, 1 differs. A +# run that ends any other way (a device that will not open, a module that will +# not load, an exception) is neither, and fails the check: otherwise it would +# read as a mutant caught. +run() { + local fn=$1 log=$2; shift 2 + local status=0 + if [ "$how" = iree ]; then + local args=() seen=0 a + for a in "$@"; do + if [ "$a" = -- ]; then seen=1 + elif [ $seen = 0 ]; then args+=("--input=@$a") + else args+=("--expected_output=${n}xi64=@$a") + fi + done + iree-run-module --device="$dev" --module="$module" --function="$fn" "${args[@]}" > "$log" 2>&1 || status=$? + if [ $status = 0 ] && grep -q '^\[SUCCESS\] all function outputs matched' "$log"; then return 0; fi + if [ $status != 0 ] && grep -q '^\[FAILED\] result\[[0-9]*\]: .* does not match the expected' "$log"; then return 1; fi + else + "${XLA_PYTHON:-python3}" -I "$(dirname "$0")/xla_run.py" "$dev" "$module" "$fn" "$@" > "$log" 2>&1 || status=$? + if [ $status = 0 ] && grep -q "^$fn on .*: exact\$" "$log"; then return 0; fi + if [ $status = 1 ] && grep -q "^$fn on .*: differs\$" "$log"; then return 1; fi + fi + echo "::error::$(basename "$module") $fn on $dev ($how): the run did not complete (exit $status)" + head -30 "$log" + exit 1 +} + +tag="$(basename "$module")" +r='='; run rows "$tag.rows.log" "${rows_in[@]}" -- "${rows_out[@]}" || r='!' +p='='; run probe "$tag.probe.log" "${probe_in[@]}" -- "${probe_out[@]}" || p='!' +word() { [ "$1" = '=' ] && echo exact || echo differ; } +echo "$tag on $dev ($how): rows $(word $r), probe $(word $p)" +if [ "$r$p" != "$expect" ]; then + echo "::error::$tag on $dev ($how): rows $(word $r), probe $(word $p); expected rows $(word "${expect:0:1}"), probe $(word "${expect:1:1}")" + head -20 "$tag.rows.log" "$tag.probe.log" + exit 1 +fi diff --git a/wai/wai-quantum/ci/tensor-program/xla_run.py b/wai/wai-quantum/ci/tensor-program/xla_run.py new file mode 100644 index 000000000..a5d061151 --- /dev/null +++ b/wai/wai-quantum/ci/tensor-program/xla_run.py @@ -0,0 +1,60 @@ +"""Build one function of the gate's tensor program with a second, independent +tensor compiler (XLA, through its Python package), run it on one device, and +compare every output with the exact one. + + python -I xla_run.py ... -- ... + + is the package's name for the device kind (`cpu`, `cuda`). The +module holds two public functions; this compiler builds a module's `main`, so +the chosen function is renamed `main` and the other made private (and unused). +Each expected file is raw little-endian 64-bit integers. + +Exit status: 0 every output exact; 1 an output differs; 2 a usage error. +""" + +import sys + +import numpy as np + + +def main(argv): + if len(argv) < 4 or "--" not in argv or argv[2] not in ("rows", "probe"): + print(__doc__, file=sys.stderr) + return 2 + platform, path, fn = argv[:3] + rest = argv[3:] + k = rest.index("--") + inputs, expected = rest[:k], rest[k + 1:] + + import jax + + jax.config.update("jax_enable_x64", True) + other = "probe" if fn == "rows" else "rows" + text = open(path).read() + for old, new in ((f"func.func @{other}(", f"func.func private @{other}("), (f"func.func @{fn}(", "func.func @main(")): + if text.count(old) != 1: + print(f"{path}: expected one `{old}`", file=sys.stderr) + return 2 + text = text.replace(old, new) + + device = jax.devices(platform)[0] + exe = device.client.compile_and_load(text, [device]) + results = exe.execute([jax.device_put(np.load(p), device) for p in inputs]) + if len(results) != len(expected): + print(f"{fn}: {len(results)} outputs, {len(expected)} expected", file=sys.stderr) + return 1 + bad = 0 + for got, name in zip(results, expected): + got, want = np.asarray(got), np.fromfile(name, dtype=" Date: Thu, 15 Oct 2026 05:01:50 -0400 Subject: [PATCH 04/19] wai-quantum: the HIP member, the CUDA kernel's text on HIP devices quantum_fabric_hip (opt-in, Linux): the gate kernel the CUDA member runs, unchanged, compiled at run time by HIP's runtime compiler for each device's own architecture (--offload-arch from the device), with -ffp-contract=off. Both libraries are loaded when first used, so a build with the feature runs anywhere and finds no member where there is no HIP. The device's properties are read in the runtime's frozen layout (hipDeviceProp_tR0600); properties that do not read as that layout are refused, never guessed at. quantum_fabric::hip holds what names a member (the kernel text, the implementation id hip-lanes64/..., distinct from the CUDA one over the same text, the device identity and the key gpu-hip-); the mutants are the CUDA kernel's, the same edits of the same text. With quantum_sv the dense f64 state vector runs the CUDA kernels' text too (hip-sv64/..., contraction off). Measured on gfx1151: hip_member 9/9 (66020 probe pairs, all four pinned certificates, circuits to 22 qubits, 6 of 6 mutants killed) and hip_sv 3/3 (the state vector's pin b7d129d0..., 16 and 20 qubits bit for bit, 4 of 4 mutants killed). Co-Authored-By: Claude Opus 5.5 --- wai/wai-quantum/Cargo.toml | 10 + wai/wai-quantum/src/lib.rs | 9 +- wai/wai-quantum/src/quantum_fabric.rs | 1 + wai/wai-quantum/src/quantum_fabric/hip.rs | 201 ++++ wai/wai-quantum/src/quantum_fabric/sv64.rs | 31 + wai/wai-quantum/src/quantum_fabric_hip.rs | 1054 ++++++++++++++++++++ wai/wai-quantum/tests/hip_member.rs | 429 ++++++++ wai/wai-quantum/tests/hip_sv.rs | 112 +++ 8 files changed, 1844 insertions(+), 3 deletions(-) create mode 100644 wai/wai-quantum/src/quantum_fabric/hip.rs create mode 100644 wai/wai-quantum/src/quantum_fabric_hip.rs create mode 100644 wai/wai-quantum/tests/hip_member.rs create mode 100644 wai/wai-quantum/tests/hip_sv.rs diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml index 11d067100..e32d114d4 100644 --- a/wai/wai-quantum/Cargo.toml +++ b/wai/wai-quantum/Cargo.toml @@ -162,6 +162,8 @@ quantum_fabric = ["quantum", "quantum_ops", "quantum_toolchain"] 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"] +# The HIP member of the fabric (a GPU the HIP runtime drives, Linux): the CUDA kernel's text, built by HIP's runtime compiler; both libraries loaded at run time. +quantum_fabric_hip = ["quantum_fabric"] quantum_fabric_opencl = ["quantum_fabric", "dep:opencl3"] # The simulated FPGA member: Verilator, invoked at run time, compiles the datapath. quantum_fabric_fpga = ["quantum_fabric"] @@ -205,6 +207,14 @@ required-features = ["quantum_fabric_cuda"] name = "cuda_sv" required-features = ["quantum_fabric_cuda", "quantum_sv"] +[[test]] +name = "hip_member" +required-features = ["quantum_fabric_hip"] + +[[test]] +name = "hip_sv" +required-features = ["quantum_fabric_hip", "quantum_sv"] + [[test]] name = "opencl_member" required-features = ["quantum_fabric_opencl"] diff --git a/wai/wai-quantum/src/lib.rs b/wai/wai-quantum/src/lib.rs index 8330a3059..fbfb541f0 100644 --- a/wai/wai-quantum/src/lib.rs +++ b/wai/wai-quantum/src/lib.rs @@ -10,10 +10,11 @@ //! 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. Nine opt-in features reach outside the crate: +//! member that ran the work. Ten opt-in features reach outside the crate: //! `quantum_fabric_webgpu` (the platform's GPU interface), -//! `quantum_fabric_cuda` (the CUDA driver and runtime compiler) and -//! `quantum_fabric_opencl` (the OpenCL ICD loader), the last two loaded when +//! `quantum_fabric_cuda` (the CUDA driver and runtime compiler), +//! `quantum_fabric_hip` (the HIP runtime and runtime compiler) and +//! `quantum_fabric_opencl` (the OpenCL ICD loader), the last three loaded when //! first used, `quantum_fabric_fpga` (Verilator, run as a program when first //! used), `quantum_fabric_npu` (the platform's model runtime, on macOS), //! `quantum_fabric_sme` (the processor's streaming vector mode), @@ -90,6 +91,8 @@ pub mod quantum_fabric; pub mod quantum_fabric_webgpu; #[cfg(all(feature = "quantum_fabric_cuda", not(target_arch = "wasm32")))] pub mod quantum_fabric_cuda; +#[cfg(all(feature = "quantum_fabric_hip", not(target_arch = "wasm32")))] +pub mod quantum_fabric_hip; #[cfg(all(feature = "quantum_fabric_opencl", not(target_arch = "wasm32")))] pub mod quantum_fabric_opencl; #[cfg(all(feature = "quantum_fabric_fpga", not(target_arch = "wasm32")))] diff --git a/wai/wai-quantum/src/quantum_fabric.rs b/wai/wai-quantum/src/quantum_fabric.rs index 955b716c9..3f15fa170 100644 --- a/wai/wai-quantum/src/quantum_fabric.rs +++ b/wai/wai-quantum/src/quantum_fabric.rs @@ -2048,6 +2048,7 @@ pub mod lanes { } pub mod cuda; +pub mod hip; pub mod fpga; pub mod npu; pub mod opencl; diff --git a/wai/wai-quantum/src/quantum_fabric/hip.rs b/wai/wai-quantum/src/quantum_fabric/hip.rs new file mode 100644 index 000000000..1d21f0022 --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/hip.rs @@ -0,0 +1,201 @@ +//! The gate kernel for HIP devices, and what names a HIP member. The kernel is +//! the CUDA kernel's text, [`cuda::SOURCE`](super::cuda::SOURCE), unchanged: +//! everything it uses (`int2`, `atomicMin`, the launch indices, 64-bit +//! integers) means the same under HIP, and HIP's runtime compiler builds it for +//! each device's own architecture. So the argument that it is bit-exact is the +//! CUDA module's, and the mutants a conforming device must be caught running +//! are the same edits of the same text ([`MUTANTS`]). What differs is the +//! compiler, its options and the device, and the implementation id and the +//! member id bind them. No device is touched here, so all of it is testable on +//! any host; the HIP member itself is `quantum_fabric_hip`. + +use super::FabricClass; +pub use super::cuda::{MUTANTS, mutate}; + +/// The kernel, exactly as it is compiled: the CUDA kernel's text. +pub const SOURCE: &str = super::cuda::SOURCE; + +/// The implementation name the id starts with. +pub const IMPLEMENTATION: &str = "hip-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 contraction of a multiply and +/// an add into one fused operation, though the kernel has no floating point +/// for it to touch. +pub const COMPILE_OPTIONS: [&str; 2] = ["-std=c++17", "-ffp-contract=off"]; + +const DOMAIN_IMPLEMENTATION: &[u8] = b"wai:quantum-fabric-hip\x01"; +const DOMAIN_MEMBER: &[u8] = b"wai:quantum-fabric-hip-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); +} + +/// `"hip-lanes64/<16 hex>"`: the first 8 bytes of +/// BLAKE3(`"wai:quantum-fabric-hip\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, and it differs from the CUDA implementation's id over the same text. +/// 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 HIP member is, as its id hashes it. The member fills it from the +/// runtime; 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 architecture as the runtime names it, with its feature settings + /// (`"gfx1151"`, `"gfx90a:sramecc+:xnack-"`): what the kernel is compiled + /// for. + pub arch: String, + /// `"::"`, in hex. + pub pci: String, + /// The runtime's version, as `hipRuntimeGetVersion` reports it. + pub runtime: u32, + /// The runtime compiler's version, `(major, minor)`. + pub compiler: (u32, u32), + /// Compute units. + pub compute_units: u32, + /// The device's memory, bytes. + pub memory: u64, +} + +impl DeviceIdentity { + /// Always [`FabricClass::Gpu`]. + pub fn class(&self) -> FabricClass { + FabricClass::Gpu + } + + /// The architecture without its feature settings: `"gfx90a"` for + /// `"gfx90a:sramecc+:xnack-"`. + pub fn arch_name(&self) -> &str { + self.arch.split(':').next().unwrap_or("") + } + + /// The option that compiles the kernel for this device: + /// `"--offload-arch="`, the feature settings included. + pub fn arch_option(&self) -> String { + format!("--offload-arch={}", self.arch) + } + + /// `"local/gpu-hip-"`. `d12` is the first 12 hex digits of + /// BLAKE3(`"wai:quantum-fabric-hip-member\x01"` ‖ put(name) ‖ put(uuid) ‖ + /// put(arch) ‖ put(pci) ‖ runtime ‖ compiler major ‖ compiler minor ‖ + /// compute units ‖ 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()); + put(&mut o, self.arch.as_bytes()); + put(&mut o, self.pci.as_bytes()); + o.extend_from_slice(&self.runtime.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.compute_units.to_be_bytes()); + o.extend_from_slice(&self.memory.to_be_bytes()); + format!("local/gpu-hip-{}", hex(&blake3::hash(&o).as_bytes()[..6])) + } + + /// The key a runner declares in CI: `"gpu-hip-"`, for instance + /// `"gpu-hip-gfx1151"`. + pub fn key(&self) -> String { + format!("gpu-hip-{}", self.arch_name()) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + fn identity() -> DeviceIdentity { + DeviceIdentity { + name: "a device".into(), + uuid: "00112233445566778899aabbccddeeff".into(), + arch: "gfx1151".into(), + pci: "0000:c2:00".into(), + runtime: 70_152_801, + compiler: (9, 0), + compute_units: 20, + memory: 64 << 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-hip-") && id.len() == "local/gpu-hip-".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 { arch: "gfx1100".into(), ..identity() }, + DeviceIdentity { arch: "gfx1151:xnack-".into(), ..identity() }, + DeviceIdentity { pci: "0000:c3:00".into(), ..identity() }, + DeviceIdentity { runtime: 60_000_000, ..identity() }, + DeviceIdentity { compiler: (10, 0), ..identity() }, + DeviceIdentity { compiler: (9, 1), ..identity() }, + DeviceIdentity { compute_units: 40, ..identity() }, + DeviceIdentity { memory: 8 << 30, ..identity() }, + ]; + for v in variants { + assert_ne!(v.member_id(), id, "{v:?}"); + } + assert_eq!(base.key(), "gpu-hip-gfx1151"); + assert_eq!(base.arch_option(), "--offload-arch=gfx1151"); + assert_eq!(base.class(), FabricClass::Gpu); + let featured = DeviceIdentity { arch: "gfx90a:sramecc+:xnack-".into(), ..identity() }; + assert_eq!((featured.arch_name(), featured.key().as_str()), ("gfx90a", "gpu-hip-gfx90a")); + assert_eq!(featured.arch_option(), "--offload-arch=gfx90a:sramecc+:xnack-"); + } + + #[test] + fn the_implementation_id_follows_the_kernel_text_and_is_not_the_cuda_one() { + let id = implementation_id(); + assert!(id.starts_with("hip-lanes64/") && id.len() == "hip-lanes64/".len() + 16, "{id}"); + assert_eq!(SOURCE, super::super::cuda::SOURCE); + assert_ne!(id.split('/').nth(1), super::super::cuda::implementation_id().split('/').nth(1)); + for m in &MUTANTS { + assert_ne!(implementation_id_of(&mutate(m)), id, "mutant {}", m.name); + } + } +} diff --git a/wai/wai-quantum/src/quantum_fabric/sv64.rs b/wai/wai-quantum/src/quantum_fabric/sv64.rs index e99fe52d2..fdcc576f1 100644 --- a/wai/wai-quantum/src/quantum_fabric/sv64.rs +++ b/wai/wai-quantum/src/quantum_fabric/sv64.rs @@ -96,6 +96,32 @@ pub fn cl_implementation_id() -> String { cl_implementation_id_of(CL_SOURCE) } +/// The implementation name the HIP id starts with. The HIP kernels are the +/// CUDA kernels' text ([`SOURCE`]), built by HIP's runtime compiler. +pub const HIP_IMPLEMENTATION: &str = "hip-sv64"; + +/// The HIP compile options: no contraction of a multiply and an add into a +/// fused operation, so each operation rounds once. +pub const HIP_COMPILE_OPTIONS: [&str; 2] = ["-std=c++17", "-ffp-contract=off"]; + +const DOMAIN_HIP_IMPLEMENTATION: &[u8] = b"wai:quantum-fabric-sv64-hip\x01"; + +/// `"hip-sv64/<16 hex>"` over the kernel text, the options and the layout. +pub fn hip_implementation_id_of(source: &str) -> String { + let mut o = DOMAIN_HIP_IMPLEMENTATION.to_vec(); + put(&mut o, source.as_bytes()); + for option in HIP_COMPILE_OPTIONS { + put(&mut o, option.as_bytes()); + } + o.extend_from_slice(&HOST_LAYOUT.to_be_bytes()); + let h: String = blake3::hash(&o).as_bytes()[..8].iter().map(|b| format!("{b:02x}")).collect(); + format!("{HIP_IMPLEMENTATION}/{h}") +} + +pub fn hip_implementation_id() -> String { + hip_implementation_id_of(SOURCE) +} + /// One planned gate. #[derive(Clone, Debug, PartialEq)] pub enum Launch { @@ -448,5 +474,10 @@ mod tests { assert_ne!(cl_implementation_id_of(&mutate_cl(m)), cl, "{}", m.name); } assert!(CL_SOURCE.contains("#pragma OPENCL FP_CONTRACT OFF") && CL_SOURCE.contains("#define SV_MAX_D 64")); + let hip = hip_implementation_id(); + assert!(hip.starts_with("hip-sv64/") && hip != id && hip != cl); + for m in &MUTANTS { + assert_ne!(hip_implementation_id_of(&mutate(m)), hip, "{}", m.name); + } } } diff --git a/wai/wai-quantum/src/quantum_fabric_hip.rs b/wai/wai-quantum/src/quantum_fabric_hip.rs new file mode 100644 index 000000000..660514195 --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric_hip.rs @@ -0,0 +1,1054 @@ +//! The fabric's HIP member: the gate kernel of +//! [`quantum_fabric::hip`](crate::quantum_fabric::hip), the CUDA kernel's text, +//! on a GPU the HIP runtime drives, compiled at run time by HIP's runtime +//! compiler for the device's own architecture. See that module, and the CUDA +//! one, for the kernel and the argument that it is bit-exact. +//! +//! - **Discovery** ([`discover`]): the runtime (`libamdhip64`) and the runtime +//! compiler (`libhiprtc`) are loaded when the process first asks for them, +//! so a build with this feature runs on a host with no HIP at all and finds +//! no member there ([`Discovery::unavailable`] says why). Every device the +//! runtime lists is read, 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 [`HipMember`], present as +//! [`Presence::Detected`]. +//! - **Running** ([`HipMember::run_program`]): the state stays on the device; +//! one launch per gate, in op-log order; 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 [`HipError::OutOfRange`], never a state. +//! - **Certification**, **placement**, **receipts** and **metering** are as +//! for the CUDA member. +//! +//! Every call is synchronous; the device is waited on before a call returns. +//! No type of the runtime 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, its architecture, the runtime and the runtime compiler. +//! - **The runtime's device properties are read in its frozen layout** +//! (`hipDeviceProp_tR0600`, the layout every runtime from 6.0 on exports +//! under that name). A device whose properties do not read as that layout +//! (no name, no `gfx` architecture) is refused, never guessed at. +//! - **Plans are per gate.** [`Config::PER_GATE`], as for the CUDA member. +//! - **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::hip::{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 std::ffi::{CStr, CString, c_char, c_int, c_void}; +use std::sync::{Arc, 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 HipError { + /// The runtime or the runtime compiler could not be used, and why. + Unavailable(String), + /// The runtime compiler refused the kernel for this device: its log. + Compile(String), + /// The runtime 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 HipError { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + match self { + HipError::Unavailable(why) => write!(f, "HIP is unavailable: {why}"), + HipError::Compile(e) => write!(f, "the kernel did not compile: {e}"), + HipError::Device(e) => write!(f, "device: {e}"), + HipError::Memory(e) => write!(f, "out of memory: {e}"), + HipError::Circuit(e) => write!(f, "{e}"), + HipError::TooWide { n_qubits, max } => write!(f, "{n_qubits} qubits; this member holds at most {max}"), + HipError::OutOfRange { gate } => write!(f, "the result of op {gate} leaves 32 bits"), + HipError::Internal(why) => write!(f, "host defect: {why}"), + HipError::Probe { a, b } => write!(f, "the device's arithmetic differs from the CPU's at ({a}, {b})"), + } + } +} + +impl std::error::Error for HipError {} + +impl From for MemberError { + fn from(e: HipError) -> MemberError { + match e { + HipError::TooWide { n_qubits, max } => MemberError::TooWide { n_qubits, max }, + HipError::OutOfRange { gate } => MemberError::OutOfRange { gate }, + HipError::Memory(why) => MemberError::Memory(why), + HipError::Device(why) => MemberError::Device(why), + HipError::Circuit(e) => MemberError::Circuit(e), + other => MemberError::Other(other.to_string()), + } + } +} + +/// The runtime's status codes this module tells apart. +const HIP_SUCCESS: c_int = 0; +const HIP_ERROR_OUT_OF_MEMORY: c_int = 2; +const HIP_ERROR_NO_DEVICE: c_int = 100; +const HIP_MEMCPY_HOST_TO_DEVICE: c_int = 1; +const HIP_MEMCPY_DEVICE_TO_HOST: c_int = 2; + +/// `hipDeviceProp_tR0600`: its size, and the offsets of the fields read. +const PROPS_SIZE: usize = 1472; +const PROPS_NAME: usize = 0; +const PROPS_UUID: usize = 256; +const PROPS_TOTAL_MEMORY: usize = 288; +const PROPS_COMPUTE_UNITS: usize = 388; +const PROPS_PCI_BUS: usize = 584; +const PROPS_PCI_DEVICE: usize = 588; +const PROPS_PCI_DOMAIN: usize = 592; +const PROPS_ARCH: usize = 1160; +/// The buffer the properties are written to: more than the layout's size, so +/// a runtime that wrote a larger structure under the same name could not +/// write past it. +const PROPS_BUFFER: usize = 4096; + +#[cfg(target_os = "linux")] +mod dl { + use std::ffi::{c_char, c_int, c_void}; + #[link(name = "dl")] + unsafe extern "C" { + pub fn dlopen(filename: *const c_char, flag: c_int) -> *mut c_void; + pub fn dlsym(handle: *mut c_void, symbol: *const c_char) -> *mut c_void; + pub fn dlerror() -> *const c_char; + } + pub const RTLD_NOW: c_int = 2; +} + +/// The runtime's and the runtime compiler's entry points, each cast to the +/// signature its header declares. +struct Api { + init: unsafe extern "C" fn(u32) -> c_int, + device_count: unsafe extern "C" fn(*mut c_int) -> c_int, + properties: unsafe extern "C" fn(*mut u8, c_int) -> c_int, + set_device: unsafe extern "C" fn(c_int) -> c_int, + mem_info: unsafe extern "C" fn(*mut usize, *mut usize) -> c_int, + malloc: unsafe extern "C" fn(*mut *mut c_void, usize) -> c_int, + free: unsafe extern "C" fn(*mut c_void) -> c_int, + memcpy: unsafe extern "C" fn(*mut c_void, *const c_void, usize, c_int) -> c_int, + module_load_data: unsafe extern "C" fn(*mut *mut c_void, *const c_void) -> c_int, + module_unload: unsafe extern "C" fn(*mut c_void) -> c_int, + module_get_function: unsafe extern "C" fn(*mut *mut c_void, *mut c_void, *const c_char) -> c_int, + #[allow(clippy::type_complexity)] + launch: unsafe extern "C" fn(*mut c_void, u32, u32, u32, u32, u32, u32, u32, *mut c_void, *mut *mut c_void, *mut *mut c_void) -> c_int, + synchronize: unsafe extern "C" fn() -> c_int, + error_string: unsafe extern "C" fn(c_int) -> *const c_char, + runtime_version: unsafe extern "C" fn(*mut c_int) -> c_int, + rtc_create: unsafe extern "C" fn(*mut *mut c_void, *const c_char, *const c_char, c_int, *const *const c_char, *const *const c_char) -> c_int, + rtc_compile: unsafe extern "C" fn(*mut c_void, c_int, *const *const c_char) -> c_int, + rtc_log_size: unsafe extern "C" fn(*mut c_void, *mut usize) -> c_int, + rtc_log: unsafe extern "C" fn(*mut c_void, *mut c_char) -> c_int, + rtc_code_size: unsafe extern "C" fn(*mut c_void, *mut usize) -> c_int, + rtc_code: unsafe extern "C" fn(*mut c_void, *mut c_char) -> c_int, + rtc_destroy: unsafe extern "C" fn(*mut *mut c_void) -> c_int, + rtc_version: unsafe extern "C" fn(*mut c_int, *mut c_int) -> c_int, +} + +static API: OnceLock> = OnceLock::new(); + +/// The first of `names` that loads, or why none did. +#[cfg(target_os = "linux")] +fn load_first(names: &[&str]) -> Result<*mut c_void, String> { + let mut why = Vec::new(); + for name in names { + let c = CString::new(*name).expect("no NUL"); + // SAFETY: dlopen with a valid C string; a failure returns null and + // leaves its reason for dlerror. + let h = unsafe { dl::dlopen(c.as_ptr(), dl::RTLD_NOW) }; + if !h.is_null() { + return Ok(h); + } + // SAFETY: dlerror returns null or a C string valid until the next call. + let e = unsafe { dl::dlerror() }; + why.push(if e.is_null() { format!("{name}: cannot load") } else { unsafe { CStr::from_ptr(e) }.to_string_lossy().into_owned() }); + } + Err(why.join("; ")) +} + +/// `p` as the function pointer type `F`, which the caller names from the +/// symbol's declaration. +/// +/// # Safety +/// +/// `p` must be a symbol of exactly the signature `F` is. +#[cfg(target_os = "linux")] +unsafe fn as_fn(p: *mut c_void) -> F { + assert_eq!(std::mem::size_of::(), std::mem::size_of::<*mut c_void>(), "a function pointer"); + // SAFETY: the sizes agree, and the caller names the symbol's signature. + unsafe { std::mem::transmute_copy(&p) } +} + +#[cfg(target_os = "linux")] +fn open_api() -> Result { + let rt = load_first(&["libamdhip64.so.7", "libamdhip64.so.6", "libamdhip64.so"]).map_err(|e| format!("no HIP runtime library on this host ({e})"))?; + let rtc = load_first(&["libhiprtc.so.7", "libhiprtc.so.6", "libhiprtc.so"]).map_err(|e| format!("no HIP runtime compiler library on this host ({e})"))?; + let sym = |h: *mut c_void, name: &str| -> Result<*mut c_void, String> { + let n = CString::new(name).expect("no NUL"); + // SAFETY: dlsym on a handle dlopen returned, with a valid C string. + let p = unsafe { dl::dlsym(h, n.as_ptr()) }; + if p.is_null() { Err(format!("the HIP libraries have no {name}")) } else { Ok(p) } + }; + // SAFETY: each symbol is cast to the signature its header declares + // (hip_runtime_api.h, hiprtc.h); handles and enums are pointers and ints. + unsafe { + Ok(Api { + init: as_fn(sym(rt, "hipInit")?), + device_count: as_fn(sym(rt, "hipGetDeviceCount")?), + properties: as_fn(sym(rt, "hipGetDevicePropertiesR0600")?), + set_device: as_fn(sym(rt, "hipSetDevice")?), + mem_info: as_fn(sym(rt, "hipMemGetInfo")?), + malloc: as_fn(sym(rt, "hipMalloc")?), + free: as_fn(sym(rt, "hipFree")?), + memcpy: as_fn(sym(rt, "hipMemcpy")?), + module_load_data: as_fn(sym(rt, "hipModuleLoadData")?), + module_unload: as_fn(sym(rt, "hipModuleUnload")?), + module_get_function: as_fn(sym(rt, "hipModuleGetFunction")?), + launch: as_fn(sym(rt, "hipModuleLaunchKernel")?), + synchronize: as_fn(sym(rt, "hipDeviceSynchronize")?), + error_string: as_fn(sym(rt, "hipGetErrorString")?), + runtime_version: as_fn(sym(rt, "hipRuntimeGetVersion")?), + rtc_create: as_fn(sym(rtc, "hiprtcCreateProgram")?), + rtc_compile: as_fn(sym(rtc, "hiprtcCompileProgram")?), + rtc_log_size: as_fn(sym(rtc, "hiprtcGetProgramLogSize")?), + rtc_log: as_fn(sym(rtc, "hiprtcGetProgramLog")?), + rtc_code_size: as_fn(sym(rtc, "hiprtcGetCodeSize")?), + rtc_code: as_fn(sym(rtc, "hiprtcGetCode")?), + rtc_destroy: as_fn(sym(rtc, "hiprtcDestroyProgram")?), + rtc_version: as_fn(sym(rtc, "hiprtcVersion")?), + }) + } +} + +#[cfg(not(target_os = "linux"))] +fn open_api() -> Result { + Err(format!("the HIP runtime is loaded on Linux; this is {}", std::env::consts::OS)) +} + +fn api() -> Result<&'static Api, HipError> { + API.get_or_init(open_api).as_ref().map_err(|e| HipError::Unavailable(e.clone())) +} + +/// Whether the runtime and the runtime compiler can be loaded in this +/// process, or why not. +pub fn available() -> Result<(), String> { + api().map(|_| ()).map_err(|e| match e { + HipError::Unavailable(why) => why, + other => other.to_string(), + }) +} + +/// A runtime status as this module reports it: out of memory as +/// [`HipError::Memory`], anything else but success as [`HipError::Device`]. +fn status(a: &Api, code: c_int, what: &str) -> Result<(), HipError> { + if code == HIP_SUCCESS { + return Ok(()); + } + // SAFETY: the runtime returns a static C string for any code. + let text = unsafe { + let p = (a.error_string)(code); + if p.is_null() { String::new() } else { CStr::from_ptr(p).to_string_lossy().into_owned() } + }; + let why = format!("{what}: {text} ({code})"); + Err(if code == HIP_ERROR_OUT_OF_MEMORY { HipError::Memory(why) } else { HipError::Device(why) }) +} + +/// Device memory, freed when dropped. +struct DeviceBuffer { + ptr: *mut c_void, + bytes: usize, +} + +impl DeviceBuffer { + fn alloc(a: &'static Api, bytes: usize) -> Result { + let mut ptr = std::ptr::null_mut(); + // SAFETY: writes one pointer; the device is current on this thread. + status(a, unsafe { (a.malloc)(&mut ptr, bytes.max(1)) }, "allocating device memory")?; + Ok(DeviceBuffer { ptr, bytes }) + } + + /// A buffer holding `data`. + fn upload(a: &'static Api, data: &[T]) -> Result { + let bytes = std::mem::size_of_val(data); + let b = DeviceBuffer::alloc(a, bytes)?; + // SAFETY: `data` holds `bytes` bytes and the buffer as many. + status(a, unsafe { (a.memcpy)(b.ptr, data.as_ptr().cast(), bytes, HIP_MEMCPY_HOST_TO_DEVICE) }, "uploading")?; + Ok(b) + } + + /// The buffer's contents, as `T`s. + fn download(&self, a: &'static Api) -> Result, HipError> { + let n = self.bytes / std::mem::size_of::(); + let mut out: Vec = Vec::new(); + out.try_reserve_exact(n).map_err(|e| HipError::Memory(e.to_string()))?; + out.resize(n, T::default()); + // SAFETY: `out` holds `n·size_of::()` ≤ `bytes` bytes; the runtime + // waits for the device's work on the buffer first. + status(a, unsafe { (a.memcpy)(out.as_mut_ptr().cast(), self.ptr, n * std::mem::size_of::(), HIP_MEMCPY_DEVICE_TO_HOST) }, "reading back")?; + Ok(out) + } +} + +impl Drop for DeviceBuffer { + fn drop(&mut self) { + if let Ok(a) = api() { + // SAFETY: the pointer came from hipMalloc and is freed once. + unsafe { (a.free)(self.ptr) }; + } + } +} + +/// The kernel compiled for one device and loaded on it: the module and its +/// entry points. +struct Module { + ordinal: c_int, + handle: *mut c_void, + functions: Vec<*mut c_void>, +} + +// SAFETY: the runtime's API is thread-safe, and module and function handles +// are process-wide values any thread may use once its device is current, +// which every use here sets first. +unsafe impl Send for Module {} +// SAFETY: as above; a handle is never mutated through a shared reference. +unsafe impl Sync for Module {} + +impl Drop for Module { + fn drop(&mut self) { + if let Ok(a) = api() { + // SAFETY: the handle came from hipModuleLoadData and is unloaded + // once, on its own device. + unsafe { + (a.set_device)(self.ordinal); + (a.module_unload)(self.handle); + } + } + } +} + +/// `source` compiled with `options` for the device `identity` names, loaded on +/// device `ordinal`, with the entry points `names` in order. +fn compile(a: &'static Api, ordinal: c_int, identity: &DeviceIdentity, source: &str, options: &[&str], names: &[&str]) -> Result { + let src = CString::new(source).map_err(|_| HipError::Internal("a kernel text holding a NUL"))?; + let file = CString::new("wai-quantum.hip").expect("no NUL"); + let mut opts: Vec = vec![CString::new(identity.arch_option()).map_err(|_| HipError::Internal("an architecture holding a NUL"))?]; + for o in options { + opts.push(CString::new(*o).expect("no NUL")); + } + let ptrs: Vec<*const c_char> = opts.iter().map(|o| o.as_ptr()).collect(); + let mut prog = std::ptr::null_mut(); + // SAFETY: valid C strings and an option array of `ptrs.len()` pointers; + // the program is destroyed below on every path. + unsafe { + let r = (a.rtc_create)(&mut prog, src.as_ptr(), file.as_ptr(), 0, std::ptr::null(), std::ptr::null()); + if r != 0 { + return Err(HipError::Compile(format!("the runtime compiler did not create the program ({r})"))); + } + let r = (a.rtc_compile)(prog, ptrs.len() as c_int, ptrs.as_ptr()); + let mut log_size = 0usize; + (a.rtc_log_size)(prog, &mut log_size); + let mut log = vec![0u8; log_size.max(1)]; + (a.rtc_log)(prog, log.as_mut_ptr().cast()); + let log = String::from_utf8_lossy(&log).trim_end_matches('\0').trim().to_string(); + if r != 0 { + (a.rtc_destroy)(&mut prog); + return Err(HipError::Compile(format!("status {r}: {log}"))); + } + let mut size = 0usize; + let r = (a.rtc_code_size)(prog, &mut size); + let mut code = vec![0u8; size.max(1)]; + let r2 = if r == 0 { (a.rtc_code)(prog, code.as_mut_ptr().cast()) } else { r }; + (a.rtc_destroy)(&mut prog); + if r2 != 0 || size == 0 { + return Err(HipError::Compile(format!("no code object ({r2})"))); + } + status(a, (a.set_device)(ordinal), "selecting the device")?; + let mut handle = std::ptr::null_mut(); + status(a, (a.module_load_data)(&mut handle, code.as_ptr().cast()), "loading the code object")?; + let mut module = Module { ordinal, handle, functions: Vec::new() }; + for name in names { + let n = CString::new(*name).expect("no NUL"); + let mut f = std::ptr::null_mut(); + status(a, (a.module_get_function)(&mut f, handle, n.as_ptr()), name)?; + module.functions.push(f); + } + Ok(module) + } +} + +/// One launch argument: a pointer to its value, as the runtime takes it. +fn arg(v: &T) -> *mut c_void { + (v as *const T).cast_mut().cast() +} + +/// Launches `function` over `work` items — [`hip::BLOCK`] threads per block, +/// at most [`hip::GRID_MAX`] blocks, the kernel striding over the rest — and +/// waits for it. +fn launch(a: &'static Api, function: *mut c_void, work: u64, args: &mut [*mut c_void]) -> Result<(), HipError> { + let blocks = blocks_for(work); + // SAFETY: `args` holds a pointer to each of the kernel's parameters, in + // order, each of the type the kernel declares; the caller checks them. + status(a, unsafe { (a.launch)(function, blocks, 1, 1, hip::BLOCK, 1, 1, 0, std::ptr::null_mut(), args.as_mut_ptr(), std::ptr::null_mut()) }, "launching")?; + // SAFETY: waits for the device's work. + status(a, unsafe { (a.synchronize)() }, "running") +} + +/// The blocks a launch over `work` items takes. +fn blocks_for(work: u64) -> u32 { + work.div_ceil(u64::from(hip::BLOCK)).clamp(1, u64::from(hip::GRID_MAX)) as u32 +} + +/// A device that is not a member, and why. +#[derive(Debug)] +pub struct Refusal { + /// The device's name, or its ordinal when the runtime gave none. + pub device: String, + /// Its CI key ([`DeviceIdentity::key`]), when the device could be read. + pub key: String, + pub error: HipError, +} + +/// 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 runtime, no runtime compiler, no + /// device, or the runtime failed to start. `None` when the runtime listed + /// its devices. + pub unavailable: Option, +} + +/// Every device the runtime lists, each read, compiled for and probed in this +/// process before it is listed. On a host without the runtime 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 }; + let a = match api() { + Ok(a) => a, + Err(e) => { + d.unavailable = Some(e.to_string()); + return d; + } + }; + let mut count: c_int = 0; + // SAFETY: plain calls into the loaded runtime; `count` is written once. + let started = unsafe { (a.init)(0) }; + let listed = if started == HIP_SUCCESS { unsafe { (a.device_count)(&mut count) } } else { started }; + if listed == HIP_ERROR_NO_DEVICE { + d.unavailable = Some("the runtime lists no device".into()); + return d; + } + if let Err(e) = status(a, listed, "starting the runtime") { + d.unavailable = Some(format!("the runtime did not start: {e}")); + return d; + } + for ordinal in 0..count { + match open(a, ordinal) { + Ok(m) => d.found.push(m), + Err(r) => d.refused.push(r), + } + } + d +} + +/// 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) } +} + +/// The NUL-terminated text at `offset` in `props`, at most `len` bytes. +fn text_at(props: &[u8], offset: usize, len: usize) -> String { + let field = &props[offset..offset + len]; + let end = field.iter().position(|&b| b == 0).unwrap_or(len); + String::from_utf8_lossy(&field[..end]).into_owned() +} + +fn u32_at(props: &[u8], offset: usize) -> u32 { + u32::from_le_bytes(props[offset..offset + 4].try_into().expect("four bytes")) +} + +/// The device's identity from its properties in the frozen layout, the +/// runtime's version and the compiler's; `None` if the properties do not read +/// as that layout. +fn identity_from(props: &[u8], runtime: u32, compiler: (u32, u32)) -> Option { + let name = text_at(props, PROPS_NAME, 256); + let arch = text_at(props, PROPS_ARCH, 256); + let memory = u64::from_le_bytes(props[PROPS_TOTAL_MEMORY..PROPS_TOTAL_MEMORY + 8].try_into().expect("eight bytes")); + if name.is_empty() || !arch.starts_with("gfx") || !arch.bytes().all(|b| b.is_ascii_alphanumeric() || b":+-".contains(&b)) || memory == 0 { + return None; + } + Some(DeviceIdentity { + name, + uuid: props[PROPS_UUID..PROPS_UUID + 16].iter().map(|b| format!("{b:02x}")).collect(), + arch, + pci: format!("{:04x}:{:02x}:{:02x}", u32_at(props, PROPS_PCI_DOMAIN), u32_at(props, PROPS_PCI_BUS), u32_at(props, PROPS_PCI_DEVICE)), + runtime, + compiler, + compute_units: u32_at(props, PROPS_COMPUTE_UNITS), + memory, + }) +} + +/// Reads the device at `ordinal`, compiles the kernel for it and probes it. +fn open(a: &'static Api, ordinal: c_int) -> Result { + let unread = |error: HipError| Refusal { device: format!("device {ordinal}"), key: String::new(), error }; + let mut props = vec![0u8; PROPS_BUFFER]; + let (mut runtime, mut major, mut minor): (c_int, c_int, c_int) = (0, 0, 0); + // SAFETY: the buffer is larger than the layout the runtime writes; the + // versions are single integers. + unsafe { + status(a, (a.set_device)(ordinal), "selecting the device").map_err(unread)?; + status(a, (a.properties)(props.as_mut_ptr(), ordinal), "reading the device's properties").map_err(unread)?; + (a.runtime_version)(&mut runtime); + (a.rtc_version)(&mut major, &mut minor); + } + let compiler = (u32::try_from(major).unwrap_or(0), u32::try_from(minor).unwrap_or(0)); + let identity = identity_from(&props[..PROPS_SIZE], u32::try_from(runtime).unwrap_or(0), compiler) + .ok_or_else(|| unread(HipError::Device("the device's properties do not read as the layout this member reads (hipDeviceProp_tR0600)".into())))?; + let key = identity.key(); + let refuse = |error: HipError| Refusal { device: identity.name.clone(), key: key.clone(), error }; + let kernel = compile(a, ordinal, &identity, hip::SOURCE, &hip::COMPILE_OPTIONS, &["apply_gate", "probe"]).map_err(refuse)?; + let (mut free, mut total) = (0usize, 0usize); + // SAFETY: writes two sizes; the device is current. + status(a, unsafe { (a.mem_info)(&mut free, &mut total) }, "reading free memory").map_err(refuse)?; + let max_qubits = qubits_for(free as u64); + let detail = format!( + "hip; device {:?}; {}; {} compute units; uuid {}; pci {}; runtime {}; runtime compiler {}.{}; 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.name, identity.arch, identity.compute_units, identity.uuid, identity.pci, identity.runtime, identity.compiler.0, identity.compiler.1, identity.memory, + ); + let member = HipMember { + api: a, + ordinal, + kernel: Arc::new(kernel), + id: identity.member_id(), + key: key.clone(), + implementation: hip::implementation_id(), + detail, + max_qubits, + identity: identity.clone(), + }; + 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(HipError::Probe { a, b })), + }, + Err(e) => Err(refuse(e)), + } +} + +/// One HIP device as a fabric member: the kernel compiled for it and loaded +/// on it, and the arithmetic probe passed in this process. +#[derive(Clone)] +pub struct HipMember { + api: &'static Api, + ordinal: c_int, + kernel: Arc, + identity: DeviceIdentity, + id: String, + key: String, + implementation: String, + detail: String, + max_qubits: u8, +} + +impl std::fmt::Debug for HipMember { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("HipMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).field("max_qubits", &self.max_qubits).finish() + } +} + +impl HipMember { + /// 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, runtime, 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 ([`hip::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 + } + + /// Makes this member's device current on the calling thread. + fn select(&self) -> Result<(), HipError> { + // SAFETY: a plain call into the loaded runtime. + status(self.api, unsafe { (self.api.set_device)(self.ordinal) }, "selecting the device") + } + + /// The device's [`lanes::mul_wide`] and [`lanes::fxmul_wide`] of each + /// pair, in order. + pub fn probe(&self, pairs: &[(i32, i32)]) -> Result, HipError> { + if pairs.is_empty() { + return Ok(Vec::new()); + } + self.select()?; + let flat: Vec = pairs.iter().flat_map(|&(a, b)| [a, b]).collect(); + let cases = DeviceBuffer::upload(self.api, &flat)?; + let out = DeviceBuffer::upload(self.api, &vec![0u32; pairs.len() * 4])?; + let count = pairs.len() as u64; + let mut args = [arg(&cases.ptr), arg(&out.ptr), arg(&count)]; + launch(self.api, self.kernel.functions[1], count, &mut args)?; + let words: Vec = out.download(self.api)?; + 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 [`HipError::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(HipError::TooWide { n_qubits: n, max: self.max_qubits }); + } + if n > MAX_QUBITS || p.start >= 1usize << n { + return Err(HipError::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(HipError::Internal("a tile step; this member runs gate steps only")) }; + let g = p.gates.get(gate as usize).ok_or(HipError::Internal("a step names a gate the program does not hold"))?; + if !g.is_well_formed() { + return Err(HipError::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| HipError::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 [`HipError::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(HipError::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(HipError::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| HipError::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(HipError::Internal("malformed gate words for this state")); + } + } + let pairs = (len / 2) as u64; + self.select()?; + let mut window = Window::new(meter); + window.open(); + let amps = DeviceBuffer::upload(self.api, &init)?; + drop(init); + let report = DeviceBuffer::upload(self.api, &[u32::MAX])?; + for &(w, index) in gates { + let m: [i32; 8] = w.m.map(|x| x as i32); + let mut args = [ + arg(&s.ptr), + arg(&pairs), + arg(&w.tbit), + arg(&w.ctrl), + arg(&m[0]), + arg(&m[1]), + arg(&m[2]), + arg(&m[3]), + arg(&m[4]), + arg(&m[5]), + arg(&m[6]), + arg(&m[7]), + arg(&index), + arg(&report.ptr), + ]; + // The words were checked against the state above, and `amps` + // holds `2·pairs` int2 lanes and `report` one word. + launch(self.api, self.kernel.functions[0], pairs, &mut args)?; + } + let bad: Vec = report.download(self.api)?; + let words: Vec = amps.download(self.api)?; + 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| HipError::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 ([`HipError::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. + pub fn simulate_metered(&self, c: &Circuit, start: usize, meter: &mut dyn RunMeter) -> Result<(StateVector, u64, Option), HipError> { + self.simulate_with(c, start, meter) + } + + fn simulate_with(&self, c: &Circuit, start: usize, meter: &mut M) -> Result<(StateVector, u64, Option), HipError> { + let p = plan::compile_with(c, start, &self.config()).map_err(HipError::Circuit)?; + if p.n_qubits > self.max_qubits { + return Err(HipError::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(HipError::OutOfRange { gate }); + } + let mut amps = Vec::new(); + amps.try_reserve_exact(run.amps.len()).map_err(|e| HipError::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)>), HipError> { + let mut cert = Certification::new(); + let mut calibration = EnergyCalibration::new(); + for c in corpus.circuits(gs) { + let reference = lanes::reference_output(&c).map_err(HipError::Circuit)?; + let (run, (joules_micro, class)) = self.run_program_metered(&plan::compile_with(&c, 0, &self.config()).map_err(HipError::Circuit)?, &mut *meter)?; + if let Some(gate) = run.out_of_range { + return Err(HipError::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) -> HipMember { + HipMember { 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 = compile(self.api, self.ordinal, &self.identity, source, &hip::COMPILE_OPTIONS, &["apply_gate", "probe"])?; + Ok(HipMember { kernel: Arc::new(kernel), implementation: hip::implementation_id_of(source), ..self.clone() }) + } +} + +/// The dense f64 state vector on this device +/// ([`quantum_fabric::sv64`](crate::quantum_fabric::sv64)): the CUDA kernels' +/// text, compiled with contraction off, so each formula is one IEEE binary64 +/// operation at a time and the state comes back with the CPU's bits, or the +/// certificate says it did not. +#[cfg(feature = "quantum_sv")] +#[derive(Clone)] +pub struct HipSv { + member: HipMember, + kernels: Arc, + implementation: String, +} + +#[cfg(feature = "quantum_sv")] +impl std::fmt::Debug for HipSv { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("HipSv").field("member", &self.member.id).field("implementation", &self.implementation).finish() + } +} + +#[cfg(feature = "quantum_sv")] +impl HipMember { + /// The state-vector kernels compiled for this device. + pub fn sv(&self) -> Result { + self.sv_with_source(crate::quantum_fabric::sv64::SOURCE) + } + + /// The same, from another kernel text: for the mutation tests. + #[doc(hidden)] + pub fn sv_with_source(&self, source: &str) -> Result { + use crate::quantum_fabric::sv64; + let kernels = compile(self.api, self.ordinal, &self.identity, source, &sv64::HIP_COMPILE_OPTIONS, &["sv_diag", "sv_one", "sv_many"])?; + Ok(HipSv { member: self.clone(), kernels: Arc::new(kernels), implementation: sv64::hip_implementation_id_of(source) }) + } +} + +#[cfg(feature = "quantum_sv")] +impl HipSv { + /// The implementation id ([`sv64::hip_implementation_id`](crate::quantum_fabric::sv64::hip_implementation_id), or a mutant's). + pub fn implementation(&self) -> &str { + &self.implementation + } + + /// `gates` on `n` qubits from `|0…0⟩`, fused to `width` as the CPU fuses + /// them: the real and imaginary parts. + pub fn run(&self, n: u32, gates: &[crate::gates::Gate], width: usize) -> Result<(Vec, Vec), HipError> { + use crate::quantum_fabric::sv64::{Launch, plan}; + let launches = plan(n, gates, width).map_err(|_| HipError::Internal("a gate the state vector refuses"))?; + // Two f64 arrays of 2^n. + if n == 0 || n > u32::from(self.member.max_qubits.saturating_sub(1)) { + return Err(HipError::TooWide { n_qubits: n.min(255) as u8, max: self.member.max_qubits.saturating_sub(1) }); + } + let a = self.member.api; + self.member.select()?; + let len = 1usize << n; + let mut init = vec![0.0f64; len]; + init[0] = 1.0; + let re = DeviceBuffer::upload(a, &init)?; + init[0] = 0.0; + let im = DeviceBuffer::upload(a, &init)?; + drop(init); + let flat = |m: &[crate::linalg::C]| -> Vec { m.iter().flat_map(|c| [c.re, c.im]).collect() }; + let [diag, one, many] = [0, 1, 2].map(|k| self.kernels.functions[k]); + for l in &launches { + match l { + Launch::Diag { bitpos, diag: d } => { + let bp = DeviceBuffer::upload(a, bitpos)?; + let dg = DeviceBuffer::upload(a, &flat(d))?; + let (count, k) = (len as u64, bitpos.len() as u32); + // `re` and `im` hold `len` doubles, `bp` k positions below + // n and `dg` 2·2^k doubles: the kernel's arguments. + launch(a, diag, count, &mut [arg(&re.ptr), arg(&im.ptr), arg(&count), arg(&bp.ptr), arg(&k), arg(&dg.ptr)])?; + } + Launch::One { s, m } => { + let pairs = (len / 2) as u64; + let v: Vec = flat(m); + let mut args = vec![arg(&re.ptr), arg(&im.ptr), arg(&pairs), arg(s)]; + args.extend(v.iter().map(arg)); + // As above; `s` is a power of two below 2^n. + launch(a, one, pairs, &mut args)?; + } + Launch::Many { sorted, offsets, m } => { + let so = DeviceBuffer::upload(a, sorted)?; + let of = DeviceBuffer::upload(a, offsets)?; + let mm = DeviceBuffer::upload(a, &flat(m))?; + let (groups, k) = ((len >> sorted.len()) as u64, sorted.len() as u32); + // As above; 2^k is at most the kernel's SV_MAX_D, which + // `plan` holds. + launch(a, many, groups, &mut [arg(&re.ptr), arg(&im.ptr), arg(&groups), arg(&so.ptr), arg(&k), arg(&of.ptr), arg(&mm.ptr)])?; + } + } + } + Ok((re.download(a)?, im.download(a)?)) + } + + /// The certificate of [`KERNEL_SV_RUN`](crate::quantum_fabric::sv64::KERNEL_SV_RUN) + /// on this device over `corpus`, against the CPU's state vector. + pub fn certify(&self, corpus: crate::quantum_fabric::sv64::SvCorpus) -> Result { + use crate::quantum_fabric::sv64::{KERNEL_SV_RUN, state_output}; + let mut cert = Certification::new(); + for case in corpus.cases() { + let reference = case.reference().map_err(|_| HipError::Internal("the corpus holds a gate the state vector refuses"))?; + let (re, im) = self.run(case.n, &case.gates, case.width)?; + cert.add(&case.bytes(), &reference, &state_output(&re, &im)); + } + Ok(cert.finish(KERNEL_SV_RUN, &self.implementation, &self.member.id, &corpus.generator())) + } + + /// The member's entry for the state vector, certified in this call. + pub fn entry(&self, corpus: crate::quantum_fabric::sv64::SvCorpus) -> Result { + Ok(KernelEntry { kernel: crate::quantum_fabric::sv64::KERNEL_SV_RUN.into(), implementation: self.implementation.clone(), certificate: self.certify(corpus)?, energy: None }) + } +} + +/// The member on the fabric ([`Fabric`](crate::quantum_fabric::Fabric)): its +/// own implementation, through [`HipMember::simulate_metered`]. Pair it with +/// a [`Member`] whose entries name that implementation, such as +/// [`HipMember::member`]'s; another implementation is +/// [`MemberError::Other`]. +impl Runner for HipMember { + 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 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); + 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!(blocks_for(0), 1); + assert_eq!(blocks_for(1), 1); + assert_eq!(blocks_for(u64::from(hip::BLOCK) + 1), 2); + assert_eq!(blocks_for(1 << 40), hip::GRID_MAX); + } + + /// A properties buffer in the frozen layout. + fn props(name: &str, arch: &str, memory: u64) -> Vec { + let mut p = vec![0u8; PROPS_SIZE]; + p[PROPS_NAME..PROPS_NAME + name.len()].copy_from_slice(name.as_bytes()); + p[PROPS_ARCH..PROPS_ARCH + arch.len()].copy_from_slice(arch.as_bytes()); + p[PROPS_UUID..PROPS_UUID + 16].copy_from_slice(&[0xab; 16]); + p[PROPS_TOTAL_MEMORY..PROPS_TOTAL_MEMORY + 8].copy_from_slice(&memory.to_le_bytes()); + p[PROPS_COMPUTE_UNITS..PROPS_COMPUTE_UNITS + 4].copy_from_slice(&20u32.to_le_bytes()); + p[PROPS_PCI_BUS..PROPS_PCI_BUS + 4].copy_from_slice(&0xc2u32.to_le_bytes()); + p + } + + #[test] + fn properties_read_in_the_frozen_layout_or_not_at_all() { + let id = identity_from(&props("a device", "gfx1151", 64 << 30), 70_152_801, (9, 0)).expect("reads"); + assert_eq!((id.name.as_str(), id.arch.as_str(), id.key().as_str()), ("a device", "gfx1151", "gpu-hip-gfx1151")); + assert_eq!((id.pci.as_str(), id.compute_units, id.memory), ("0000:c2:00", 20, 64 << 30)); + assert_eq!(id.uuid, "ab".repeat(16)); + // Properties that are not that layout are refused, never guessed at. + assert!(identity_from(&props("", "gfx1151", 1), 0, (0, 0)).is_none()); + assert!(identity_from(&props("a device", "", 1), 0, (0, 0)).is_none()); + assert!(identity_from(&props("a device", "sm_89", 1), 0, (0, 0)).is_none()); + assert!(identity_from(&props("a device", "gfx1151", 0), 0, (0, 0)).is_none()); + assert!(identity_from(&props("a device", "gfx\u{7f}", 1), 0, (0, 0)).is_none()); + } + + #[test] + fn the_refusals_that_may_fall_back_do() { + assert!(MemberError::from(HipError::Memory("x".into())).falls_back()); + assert!(MemberError::from(HipError::OutOfRange { gate: 3 }).falls_back()); + assert!(!MemberError::from(HipError::Probe { a: 1, b: 2 }).falls_back()); + assert!(!MemberError::from(HipError::Internal("x")).falls_back()); + } +} diff --git a/wai/wai-quantum/tests/hip_member.rs b/wai/wai-quantum/tests/hip_member.rs new file mode 100644 index 000000000..8a8734928 --- /dev/null +++ b/wai/wai-quantum/tests/hip_member.rs @@ -0,0 +1,429 @@ +//! The fabric's HIP member on the devices of the machine that runs this. +//! +//! The machine declares the devices it must have in `WAI_QUANTUM_HIP_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 runtime 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_HIP_REPORT` names a file, appends the +//! same lines to it, which a CI step can count per device. +//! +//! ```text +//! WAI_QUANTUM_HIP_EXPECT=gpu-hip-gfx1151 \ +//! cargo test --features quantum_fabric_hip --test hip_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::hip::{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_hip::{HipError, HipMember, 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_HIP_REPORT` names. +fn report(line: &str) { + println!("{line}"); + if let Some(path) = std::env::var_os("WAI_QUANTUM_HIP_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_HIP_EXPECT={found_list}"); + let expect = match std::env::var("WAI_QUANTUM_HIP_EXPECT") { + Ok(v) if !v.is_empty() => v, + _ => return Err(format!("WAI_QUANTUM_HIP_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_HIP_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 [HipMember] { + 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!("hip 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-hip-"), "{}", m.id()); + assert_eq!(m.implementation(), hip::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!("hip 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!("hip 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!("hip 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!( + "hip 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!("hip 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(), HipError::Circuit(native), "{c:?}"); + } + assert_eq!(m.simulate(&Circuit::new(MAX_QUBITS + 1)).unwrap_err(), HipError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1))); + let small = m.with_max_qubits(4); + assert_eq!(small.simulate(&Circuit::new(5)).unwrap_err(), HipError::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(), HipError::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(HipError::Internal(_)))); + assert!(matches!(m.run_lanes(&[[0, 0]; 3], &[z_words()]), Err(HipError::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(HipError::Internal(_)))); + // The refusals that may still fall back to another member, do. + assert!(MemberError::from(HipError::TooWide { n_qubits: 5, max: 4 }).falls_back()); + report(&format!("hip refusals: key={} typed", m.key())); + } +} + +/// Whether a member running `mutant`'s text is caught by the mutant's +/// detector. +fn caught(mm: &HipMember, 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!("hip 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!("hip 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 HIP 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(&hip::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!("hip 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!("hip 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!("hip repeat: key={} ok", m.key())); + } +} diff --git a/wai/wai-quantum/tests/hip_sv.rs b/wai/wai-quantum/tests/hip_sv.rs new file mode 100644 index 000000000..9f0d3564c --- /dev/null +++ b/wai/wai-quantum/tests/hip_sv.rs @@ -0,0 +1,112 @@ +//! The dense f64 state vector on the HIP devices of the machine that runs +//! this: bit for bit the CPU's, or the certificate says not. +//! +//! The machine declares its devices in `WAI_QUANTUM_HIP_EXPECT` (comma- +//! separated keys, or `none`), as for `hip_member`; `WAI_QUANTUM_HIP_REPORT` +//! names a file the established lines are appended to. +#![cfg(not(target_arch = "wasm32"))] + +use std::io::Write; +use wai_quantum::gates::Gate; +use wai_quantum::quantum_fabric::sv64::{self, MUTANTS, SV_DEFAULT, SV_DEFAULT_PIN, SvCorpus}; +use wai_quantum::quantum_fabric::{Certification, Parity}; +use wai_quantum::quantum_fabric_hip::{HipMember, discover}; +use wai_quantum::quantum_sv::StateVector; + +fn report(line: &str) { + println!("{line}"); + if let Some(path) = std::env::var_os("WAI_QUANTUM_HIP_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"); + } +} + +fn members() -> Vec { + let d = discover(); + let mut found: Vec = d.found.iter().map(|m| m.key().to_string()).collect(); + found.sort(); + let found_list = if found.is_empty() { "none".to_string() } else { found.join(",") }; + let expect = std::env::var("WAI_QUANTUM_HIP_EXPECT").unwrap_or_else(|_| panic!("WAI_QUANTUM_HIP_EXPECT is not set; found {found_list}")); + assert_eq!(found_list, expect, "the devices found are not the declared ones"); + d.found +} + +fn hex(b: &[u8]) -> String { + b.iter().map(|x| format!("{x:02x}")).collect() +} + +/// The CPU certifying itself on `corpus`: the digest a bit-exact device +/// reproduces. +fn cpu_digest(corpus: SvCorpus) -> String { + let mut cert = Certification::new(); + for case in corpus.cases() { + let r = case.reference().unwrap(); + cert.add(&case.bytes(), &r, &r); + } + hex(&cert.finish(sv64::KERNEL_SV_RUN, "cpu", "cpu", &corpus.generator()).corpus) +} + +#[test] +fn s0_the_state_vector_is_the_cpus_bit_for_bit() { + let cpu = cpu_digest(SV_DEFAULT); + assert_eq!(cpu, SV_DEFAULT_PIN, "this machine's CPU does not reproduce the pin"); + for m in members() { + let sv = m.sv().unwrap_or_else(|e| panic!("{}: {e}", m.key())); + let c = sv.certify(SV_DEFAULT).unwrap_or_else(|e| panic!("{}: {e}", m.key())); + assert_eq!(c.parity, Parity::BitExact, "{}", m.key()); + assert_eq!(c.cases, SV_DEFAULT.cases as u64); + assert_eq!(hex(&c.corpus), cpu, "{}: the digest is not the CPU's", m.key()); + report(&format!("hip sv certified: key={} member={} kernel={} implementation={} parity=bit-exact cases={} digest={}", m.key(), m.id(), c.kernel, c.implementation, c.cases, hex(&c.corpus))); + } +} + +#[test] +fn s1_wide_states_match_the_cpu() { + for m in members() { + let sv = m.sv().unwrap(); + for (n, width) in [(16u32, 0usize), (20, 3)] { + let mut gates = Vec::new(); + let mut a = 0.1f64; + for layer in 0..6u32 { + for q in 0..n { + a += 0.731; + gates.push(match (q + layer) % 5 { + 0 => Gate::rx(q, a), + 1 => Gate::ry(q, a), + 2 => Gate::rz(q, a), + 3 => Gate::sqrt_w(q), + _ => Gate::t(q), + }); + } + let mut q = layer % 2; + while q + 1 < n { + gates.push(if q % 3 == 0 { Gate::fsim(q, q + 1, a, 0.5 * a) } else { Gate::cz(q, q + 1) }); + q += 2; + } + gates.push(Gate::iswap(0, n - 1)); + } + let mut cpu = StateVector::zero(n).unwrap(); + cpu.run(&gates, width).unwrap(); + let (re, im) = sv.run(n, &gates, width).unwrap(); + let same = |a: &[f64], b: &[f64]| a.iter().zip(b).all(|(x, y)| x.to_bits() == y.to_bits()); + assert!(same(&re, &cpu.re) && same(&im, &cpu.im), "{}: n={n} width={width}", m.key()); + } + report(&format!("hip sv wide: key={} 16 and 20 qubits match the CPU bit for bit", m.key())); + } +} + +#[test] +fn s2_mutants_are_caught() { + let small = SvCorpus { cases: 18, ..SV_DEFAULT }; + let cpu = cpu_digest(small); + for m in members() { + let mut killed = 0; + for mutant in &MUTANTS { + let sv = m.sv_with_source(&sv64::mutate(mutant)).unwrap_or_else(|e| panic!("mutant {}: {e}", mutant.name)); + let c = sv.certify(small).unwrap(); + assert!(c.parity != Parity::BitExact || hex(&c.corpus) != cpu, "{}: mutant {} survived", m.key(), mutant.name); + killed += 1; + } + report(&format!("hip sv mutants: key={} killed {killed} of {}", m.key(), MUTANTS.len())); + } +} -- 2.54.0 (Apple Git-157) From 5c8fdaacbb194a0f0aaef260c68f5ca259fb1da9 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 08:52:55 -0400 Subject: [PATCH 05/19] ci: the HIP member's workflow (the AMD host, and a host with no HIP) Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-hip.yml | 153 ++++++++++++++++++++++++++ 1 file changed, 153 insertions(+) create mode 100644 .github/workflows/wai-quantum-hip.yml diff --git a/.github/workflows/wai-quantum-hip.yml b/.github/workflows/wai-quantum-hip.yml new file mode 100644 index 000000000..22ce2f5b6 --- /dev/null +++ b/.github/workflows/wai-quantum-hip.yml @@ -0,0 +1,153 @@ +name: wai-quantum HIP member + +# The fabric's HIP member: the CUDA member's kernel text, unchanged, compiled at +# run time by HIP's runtime compiler for each device's own architecture and +# certified bit-exact against the CPU reference's pinned corpus digests. +# +# - hip-hw (the AMD host, its integrated GPU): 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; and the dense f64 state vector must come back with the CPU's +# bits on its pinned corpus and at 16 and 20 qubits, every state-vector +# mutant caught. +# - hip-absent (a container with no GPU and no HIP 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-hip.yml' + pull_request: + paths: + - 'wai/wai-quantum/**' + - 'rust-toolchain.toml' + - '.github/workflows/wai-quantum-hip.yml' + +jobs: + # Does this pull request's own diff touch the paths below? This forge matches a + # workflow's `paths:` against the push, so a merge of main into a branch fires the + # workflow whatever the pull request changes; every job here waits on this answer. + # A push to main, a schedule, or anything the script cannot determine answers + # `true` (scripts/pr-touches-paths.py). + relevant: + name: does this change touch it + runs-on: linux-arm64 + outputs: + run: ${{ steps.gate.outputs.run }} + steps: + - uses: actions/checkout@v5 + with: + fetch-depth: 0 + - id: gate + run: | + set -euo pipefail + r=$(python3 scripts/pr-touches-paths.py .github/workflows/wai-quantum-hip.yml \ + "${{ github.event.pull_request.base.sha }}" \ + "${{ github.event.pull_request.head.sha }}") + echo "this change touches wai-quantum-hip: $r" + echo "run=$r" >> "$GITHUB_OUTPUT" + + hip-hw: + name: member · HIP (the AMD host) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux-amd + timeout-minutes: 90 + env: + CARGO_PROFILE_DEV_DEBUG: '0' + CARGO_BUILD_JOBS: '12' + 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. The + # host names its own route to the forge (WAI_FORGE_URL). + git fetch -q --depth 1 "${WAI_FORGE_URL:-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-hip" nice -n 10 cargo test --locked --features quantum_fabric_hip,quantum_sv --lib -- quantum_fabric_hip:: quantum_fabric::hip:: quantum_fabric::sv64:: 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" + # The devices this host has, declared here (or by the runner, so a GPU + # and its declaration change in one place). + key="${WAI_QUANTUM_AMD_HIP_EXPECT:-gpu-hip-gfx1151}" + report="$RUNNER_TEMP/hip-report.txt"; : > "$report" + WAI_QUANTUM_HIP_EXPECT="$key" WAI_QUANTUM_HIP_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-hip" nice -n 10 \ + cargo test --locked --release --features quantum_fabric_hip --test hip_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" + grep -qx "hip devices: as declared ($key)" "$report" || { echo "::error::the devices found are not the declared ones"; exit 1; } + n="$(grep -c "^hip certified: key=$key .* parity=bit-exact " "$report")" + [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; } + grep -qx "hip mutants: key=$key killed 6 of 6" "$report" || { echo "::error::a kernel mutant survived"; exit 1; } + - name: The dense f64 state vector, bit for bit + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + export PATH="$HOME/.cargo/bin:$PATH" + key="${WAI_QUANTUM_AMD_HIP_EXPECT:-gpu-hip-gfx1151}" + report="$RUNNER_TEMP/hip-sv-report.txt"; : > "$report" + WAI_QUANTUM_HIP_EXPECT="$key" WAI_QUANTUM_HIP_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-hip" nice -n 10 \ + cargo test --locked --release --features quantum_fabric_hip,quantum_sv --test hip_sv -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/sv.log" + grep -qE '^test result: ok\. 3 passed; 0 failed' "$RUNNER_TEMP/sv.log" || { echo "::error::not every state-vector test ran and passed"; exit 1; } + cat "$report" + grep -q "^hip sv certified: key=$key .* parity=bit-exact cases=48 digest=b7d129d0593adf70dc164477664970c5c0b8c09f79f1a0573036bd60888dbcd1$" "$report" || { echo "::error::the state vector is not the CPU's bit for bit"; exit 1; } + grep -qx "hip sv mutants: key=$key killed 4 of 4" "$report" || { echo "::error::a state-vector mutant survived"; exit 1; } + + hip-absent: + name: member · HIP absent (linux) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux + timeout-minutes: 60 + env: + WAI_QUANTUM_HIP_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 HIP here, and the member says so + shell: bash + working-directory: wai/wai-quantum + run: | + set -euo pipefail + report="$RUNNER_TEMP/hip-report.txt"; : > "$report" + WAI_QUANTUM_HIP_REPORT="$report" cargo test --locked --features quantum_fabric_hip --test hip_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 'hip devices: as declared (none)' "$report" || { echo "::error::the absent case was not established"; exit 1; } + cargo test --locked --features quantum_fabric_hip,quantum_sv --test hip_sv -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/sv.log" + grep -qE '^test result: ok\. 3 passed; 0 failed' "$RUNNER_TEMP/sv.log" || { echo "::error::the state-vector tests did not run over no device"; exit 1; } -- 2.54.0 (Apple Git-157) From d323bbacaebe3aa0811be01744cf47ecc55ef096 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 08:58:32 -0400 Subject: [PATCH 06/19] wai-quantum: OpenCL keys carry the device's vendor id, not a word of its name A key was the vendor's first word, so the GPU vendor whose name begins "Advanced" was keyed opencl-gpu-advanced. A key is now opencl-- from CL_DEVICE_VENDOR_ID, the PCI vendor id for a device on a PCI bus, as the WebGPU member's keys already are; DeviceIdentity gains vendor_id, which the member id binds, and vendor_slug goes. Measured: opencl-cpu-8086, opencl-gpu-10de and opencl-gpu-8086 on the hardware host, opencl-gpu-1027f00 on the macOS runner, opencl-gpu-1002 on the AMD host. The workflow reads the hardware host's declaration from WAI_QUANTUM_HW_OPENCL_KEYS (set beside the old variable, which main's workflow still reads), adds the AMD host's leg (opencl-amd: 9/9 and the state vector, bit-exact there by hand) and names the macOS job by its runner. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-opencl.yml | 91 ++++++++++++++++++-- wai/wai-quantum/src/quantum_fabric/opencl.rs | 34 ++++---- wai/wai-quantum/src/quantum_fabric_opencl.rs | 1 + wai/wai-quantum/tests/opencl_member.rs | 2 +- 4 files changed, 102 insertions(+), 26 deletions(-) diff --git a/.github/workflows/wai-quantum-opencl.yml b/.github/workflows/wai-quantum-opencl.yml index fe3e49fff..cebde7d9b 100644 --- a/.github/workflows/wai-quantum-opencl.yml +++ b/.github/workflows/wai-quantum-opencl.yml @@ -5,13 +5,16 @@ name: wai-quantum OpenCL member # reference's pinned corpus digests. # # - opencl-hw (jetson-hub): the platforms the host declares (by default the -# discrete GPU's and the CPU's). The device +# discrete GPU's and the CPU's), each by a key that carries the device's +# vendor id (opencl--). The device # tests must find exactly the declared devices, agree with the CPU on the # arithmetic probe and on random states (out-of-range runs included), # reproduce all four pinned certificates on each device (two gatesets, # random-v1 and fabric-v1), and catch every kernel mutant on each; and the # dense f64 state vector must come back with the CPU's bits on each. -# - opencl-apple (the macOS runner): the same on the platform's GPU, through its +# - opencl-amd (the AMD host): the same on its integrated GPU, through the +# GPU vendor's OpenCL platform, the state vector included. +# - opencl-macos (the macOS runner): the same on the platform's GPU, through its # OpenCL 1.2; that GPU has no binary64, so the state vector is refused there, # with the reason. # - opencl-absent (a container with no OpenCL): the same build finds no member @@ -67,8 +70,9 @@ jobs: timeout-minutes: 90 env: # The devices are declared by the host, in its runner's environment - # (WAI_QUANTUM_HW_OPENCL_EXPECT); with no declaration, the CPU and the - # discrete GPU. + # (WAI_QUANTUM_HW_OPENCL_KEYS); with no declaration, the CPU and the + # discrete GPU. A key carries the device's vendor id + # (opencl--). CARGO_PROFILE_DEV_DEBUG: '0' CARGO_BUILD_JOBS: '8' steps: @@ -104,7 +108,7 @@ jobs: # skips the platform silently. oneapi=/opt/intel/oneapi export LD_LIBRARY_PATH="$oneapi/compiler/latest/lib:$oneapi/tbb/latest/lib/intel64/gcc4.8:$oneapi/tbb/latest/lib:$oneapi/umf/latest/lib:$oneapi/tcm/latest/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" - export WAI_QUANTUM_OPENCL_EXPECT="${WAI_QUANTUM_HW_OPENCL_EXPECT:-opencl-cpu-intel,opencl-gpu-nvidia}" + export WAI_QUANTUM_OPENCL_EXPECT="${WAI_QUANTUM_HW_OPENCL_KEYS:-opencl-cpu-8086,opencl-gpu-10de}" echo "devices declared: $WAI_QUANTUM_OPENCL_EXPECT" report="$RUNNER_TEMP/opencl-report.txt"; : > "$report" WAI_QUANTUM_OPENCL_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-opencl" nice -n 10 \ @@ -125,7 +129,7 @@ jobs: export PATH="$HOME/.cargo/bin:$PATH" oneapi=/opt/intel/oneapi export LD_LIBRARY_PATH="$oneapi/compiler/latest/lib:$oneapi/tbb/latest/lib/intel64/gcc4.8:$oneapi/tbb/latest/lib:$oneapi/umf/latest/lib:$oneapi/tcm/latest/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" - export WAI_QUANTUM_OPENCL_EXPECT="${WAI_QUANTUM_HW_OPENCL_EXPECT:-opencl-cpu-intel,opencl-gpu-nvidia}" + export WAI_QUANTUM_OPENCL_EXPECT="${WAI_QUANTUM_HW_OPENCL_KEYS:-opencl-cpu-8086,opencl-gpu-10de}" report="$RUNNER_TEMP/opencl-sv-report.txt"; : > "$report" WAI_QUANTUM_OPENCL_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-opencl" nice -n 10 \ cargo test --locked --release --features quantum_fabric_opencl,quantum_sv --test opencl_sv -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/sv.log" @@ -137,14 +141,85 @@ jobs: grep -qx "opencl sv mutants: key=$key killed 4 of 4" "$report" || { echo "::error::$key: a state-vector mutant survived"; exit 1; } done - opencl-apple: + opencl-amd: + name: member · OpenCL GPU (the AMD host) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux-amd + timeout-minutes: 90 + env: + # The devices are declared by the host, in its runner's environment + # (WAI_QUANTUM_AMD_OPENCL_KEYS); with no declaration, its integrated GPU. + 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. The + # host names its own route to the forge (WAI_FORGE_URL). + git fetch -q --depth 1 "${WAI_FORGE_URL:-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-opencl" nice -n 10 cargo test --locked --features quantum_fabric_opencl --lib -- quantum_fabric_opencl:: quantum_fabric::opencl:: 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" + export WAI_QUANTUM_OPENCL_EXPECT="${WAI_QUANTUM_AMD_OPENCL_KEYS:-opencl-gpu-1002}" + echo "devices declared: $WAI_QUANTUM_OPENCL_EXPECT" + report="$RUNNER_TEMP/opencl-report.txt"; : > "$report" + WAI_QUANTUM_OPENCL_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-opencl" nice -n 10 \ + cargo test --locked --release --features quantum_fabric_opencl --test opencl_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" + grep -qx "opencl devices: as declared ($WAI_QUANTUM_OPENCL_EXPECT)" "$report" || { echo "::error::the devices found are not the declared ones"; exit 1; } + for key in ${WAI_QUANTUM_OPENCL_EXPECT//,/ }; do + n="$(grep -c "^opencl certified: key=$key .* parity=bit-exact " "$report")" + [ "$n" = 4 ] || { echo "::error::$key: $n of 4 certificates were earned"; exit 1; } + grep -qx "opencl mutants: key=$key killed 6 of 6" "$report" || { echo "::error::$key: a kernel mutant survived"; exit 1; } + done + - name: The dense f64 state vector, bit for bit + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + export PATH="$HOME/.cargo/bin:$PATH" + export WAI_QUANTUM_OPENCL_EXPECT="${WAI_QUANTUM_AMD_OPENCL_KEYS:-opencl-gpu-1002}" + report="$RUNNER_TEMP/opencl-sv-report.txt"; : > "$report" + WAI_QUANTUM_OPENCL_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-opencl" nice -n 10 \ + cargo test --locked --release --features quantum_fabric_opencl,quantum_sv --test opencl_sv -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/sv.log" + grep -qE '^test result: ok\. 3 passed; 0 failed' "$RUNNER_TEMP/sv.log" || { echo "::error::not every state-vector test ran and passed"; exit 1; } + cat "$report" + for key in ${WAI_QUANTUM_OPENCL_EXPECT//,/ }; do + if grep -qx "opencl sv refused: key=$key the device has no binary64" "$report"; then continue; fi + grep -q "^opencl sv certified: key=$key .* parity=bit-exact cases=48 digest=b7d129d0593adf70dc164477664970c5c0b8c09f79f1a0573036bd60888dbcd1$" "$report" || { echo "::error::$key: the state vector is not the CPU's bit for bit"; exit 1; } + grep -qx "opencl sv mutants: key=$key killed 4 of 4" "$report" || { echo "::error::$key: a state-vector mutant survived"; exit 1; } + done + + opencl-macos: name: member · OpenCL GPU (macos) needs: relevant if: needs.relevant.outputs.run == 'true' runs-on: macos timeout-minutes: 90 env: - WAI_QUANTUM_OPENCL_EXPECT: opencl-gpu-apple + WAI_QUANTUM_OPENCL_EXPECT: opencl-gpu-1027f00 steps: - uses: actions/checkout@v5 # launchd's PATH omits cargo on the macOS host runner. diff --git a/wai/wai-quantum/src/quantum_fabric/opencl.rs b/wai/wai-quantum/src/quantum_fabric/opencl.rs index 2fbc8b1da..28296bbdd 100644 --- a/wai/wai-quantum/src/quantum_fabric/opencl.rs +++ b/wai/wai-quantum/src/quantum_fabric/opencl.rs @@ -113,6 +113,10 @@ pub struct DeviceIdentity { pub version: String, pub driver: String, pub kind: DeviceKind, + /// The device's vendor id (`CL_DEVICE_VENDOR_ID`): for a device on a PCI + /// bus, its PCI vendor id; otherwise whatever unique number the platform + /// reports. + pub vendor_id: u32, /// The device's UUID in hex where the platform reports one /// (`cl_khr_device_uuid`), else empty. pub uuid: String, @@ -133,19 +137,10 @@ impl DeviceIdentity { } } - /// The leading letters and digits of the vendor's first word other than - /// "the", lowercase: `"Maker Corporation"` is `"maker"`, - /// `"Maker(R) Corporation"` is `"maker"`, `"The pocl project"` is `"pocl"`. - pub fn vendor_slug(&self) -> String { - let word = self.vendor.split_whitespace().find(|w| !w.eq_ignore_ascii_case("the")).unwrap_or(""); - let slug: String = word.chars().take_while(|c| c.is_ascii_alphanumeric()).map(|c| c.to_ascii_lowercase()).collect(); - if slug.is_empty() { "unknown".into() } else { slug } - } - /// `"local/opencl--"`. `d12` is the first 12 hex digits of /// BLAKE3(`"wai:quantum-fabric-opencl-member\x01"` ‖ put of the platform, /// its version, the vendor, the name, the version, the driver, the kind and - /// the UUID ‖ compute units ‖ memory), the numbers big-endian. It sorts + /// the UUID ‖ vendor id ‖ compute units ‖ memory), the numbers big-endian. It sorts /// after every `local/cpu-` and `local/gpu-…`, so with no energy /// figures [`place`](super::place) keeps the CPU reference. pub fn member_id(&self) -> String { @@ -155,15 +150,17 @@ impl DeviceIdentity { } put(&mut o, self.kind.name().as_bytes()); put(&mut o, self.uuid.as_bytes()); + o.extend_from_slice(&self.vendor_id.to_be_bytes()); o.extend_from_slice(&self.compute_units.to_be_bytes()); o.extend_from_slice(&self.memory.to_be_bytes()); format!("local/opencl-{}-{}", self.kind.name(), hex(&blake3::hash(&o).as_bytes()[..6])) } - /// The key a runner declares in CI: `"opencl--"`, for - /// instance `"opencl-gpu-maker"`. + /// The key a runner declares in CI: `"opencl--"`, the + /// vendor id in lowercase hex, for instance `"opencl-gpu-10de"`: the same + /// number the WebGPU member's keys carry for a device on a PCI bus. pub fn key(&self) -> String { - format!("opencl-{}-{}", self.kind.name(), self.vendor_slug()) + format!("opencl-{}-{:x}", self.kind.name(), self.vendor_id) } } @@ -206,6 +203,7 @@ mod tests { version: "OpenCL 3.0 CUDA".into(), driver: "595.84".into(), kind: DeviceKind::Gpu, + vendor_id: 0x10de, uuid: "00112233445566778899aabbccddeeff".into(), compute_units: 20, memory: 6 << 30, @@ -232,16 +230,18 @@ mod tests { DeviceIdentity { kind: DeviceKind::Cpu, ..identity() }, DeviceIdentity { kind: DeviceKind::Accelerator, ..identity() }, DeviceIdentity { uuid: String::new(), ..identity() }, + DeviceIdentity { vendor_id: 0x1002, ..identity() }, DeviceIdentity { compute_units: 21, ..identity() }, DeviceIdentity { memory: 8 << 30, ..identity() }, ]; for v in variants { assert_ne!(v.member_id(), id, "{v:?}"); } - assert_eq!(base.key(), "opencl-gpu-graphics"); - assert_eq!(DeviceIdentity { vendor: "Cores(R) Corporation".into(), kind: DeviceKind::Cpu, ..identity() }.key(), "opencl-cpu-cores"); - assert_eq!(DeviceIdentity { vendor: "The pocl project".into(), kind: DeviceKind::Cpu, ..identity() }.key(), "opencl-cpu-pocl"); - assert_eq!(DeviceIdentity { vendor: String::new(), ..identity() }.key(), "opencl-gpu-unknown"); + assert_eq!(base.key(), "opencl-gpu-10de"); + assert_eq!(DeviceIdentity { vendor_id: 0x8086, kind: DeviceKind::Cpu, ..identity() }.key(), "opencl-cpu-8086"); + assert_eq!(DeviceIdentity { vendor_id: 0x0102_7f00, ..identity() }.key(), "opencl-gpu-1027f00"); + // The vendor's name is not in the key: the number is. + assert_eq!(DeviceIdentity { vendor: "another".into(), ..identity() }.key(), base.key()); assert_eq!(base.class(), Some(FabricClass::Gpu)); assert_eq!(DeviceIdentity { kind: DeviceKind::Cpu, ..identity() }.class(), Some(FabricClass::Cpu)); assert_eq!(DeviceIdentity { kind: DeviceKind::Accelerator, ..identity() }.class(), None); diff --git a/wai/wai-quantum/src/quantum_fabric_opencl.rs b/wai/wai-quantum/src/quantum_fabric_opencl.rs index 4b4743509..aae1b7586 100644 --- a/wai/wai-quantum/src/quantum_fabric_opencl.rs +++ b/wai/wai-quantum/src/quantum_fabric_opencl.rs @@ -248,6 +248,7 @@ fn open(platform: &PlatformInfo, position: usize, id: cl_device_id) -> Result,opencl-gpu- \ +//! WAI_QUANTUM_OPENCL_EXPECT=opencl-cpu-,opencl-gpu- \ //! cargo test --features quantum_fabric_opencl --test opencl_member -- --test-threads=1 //! ``` #![cfg(not(target_arch = "wasm32"))] -- 2.54.0 (Apple Git-157) From 99631e2e451511885b547a5aee25a29899753ded Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 08:59:54 -0400 Subject: [PATCH 07/19] ci: legs on framework-128gb for the WebGPU member and the tensor program The WebGPU member on its integrated GPU and the software Vulkan driver (gpu-vulkan-1002, swgpu-vulkan-10005) with the conformance corpus through it; the tensor program through HIP (built for gfx1151) and on both Vulkan devices, the third mutant caught on each. Job names carry the host's name. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-hip.yml | 4 +- .github/workflows/wai-quantum-opencl.yml | 4 +- .../workflows/wai-quantum-tensor-program.yml | 58 ++++++++++++++++++- .github/workflows/wai-quantum-webgpu.yml | 45 +++++++++++++- 4 files changed, 105 insertions(+), 6 deletions(-) diff --git a/.github/workflows/wai-quantum-hip.yml b/.github/workflows/wai-quantum-hip.yml index 22ce2f5b6..accf94484 100644 --- a/.github/workflows/wai-quantum-hip.yml +++ b/.github/workflows/wai-quantum-hip.yml @@ -4,7 +4,7 @@ name: wai-quantum HIP member # run time by HIP's runtime compiler for each device's own architecture and # certified bit-exact against the CPU reference's pinned corpus digests. # -# - hip-hw (the AMD host, its integrated GPU): the device tests must find +# - hip-hw (framework-128gb, its integrated GPU): 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 @@ -57,7 +57,7 @@ jobs: echo "run=$r" >> "$GITHUB_OUTPUT" hip-hw: - name: member · HIP (the AMD host) + name: member · HIP (framework-128gb) needs: relevant if: needs.relevant.outputs.run == 'true' runs-on: linux-amd diff --git a/.github/workflows/wai-quantum-opencl.yml b/.github/workflows/wai-quantum-opencl.yml index cebde7d9b..1e73dd347 100644 --- a/.github/workflows/wai-quantum-opencl.yml +++ b/.github/workflows/wai-quantum-opencl.yml @@ -12,7 +12,7 @@ name: wai-quantum OpenCL member # reproduce all four pinned certificates on each device (two gatesets, # random-v1 and fabric-v1), and catch every kernel mutant on each; and the # dense f64 state vector must come back with the CPU's bits on each. -# - opencl-amd (the AMD host): the same on its integrated GPU, through the +# - opencl-amd (framework-128gb): the same on its integrated GPU, through the # GPU vendor's OpenCL platform, the state vector included. # - opencl-macos (the macOS runner): the same on the platform's GPU, through its # OpenCL 1.2; that GPU has no binary64, so the state vector is refused there, @@ -142,7 +142,7 @@ jobs: done opencl-amd: - name: member · OpenCL GPU (the AMD host) + name: member · OpenCL GPU (framework-128gb) needs: relevant if: needs.relevant.outputs.run == 'true' runs-on: linux-amd diff --git a/.github/workflows/wai-quantum-tensor-program.yml b/.github/workflows/wai-quantum-tensor-program.yml index aee168283..0860ae65d 100644 --- a/.github/workflows/wai-quantum-tensor-program.yml +++ b/.github/workflows/wai-quantum-tensor-program.yml @@ -12,10 +12,12 @@ name: wai-quantum portable tensor program # # - cpu (the container runner): both compilers on the x86-64 CPU, the program # and its three mutants (each must fail where it should); builds for three -# GPU generations the fleet has none of (compiled, not run); +# other GPU architectures (compiled, not run); # - gpu (the hardware host): every Vulkan device the host has (its discrete # and integrated GPUs and the software one), and the discrete GPU through # CUDA with each compiler; +# - amd (framework-128gb): its integrated GPU through HIP, built for the GPU's +# own architecture, and every Vulkan device the host has; # - metal (the macOS runner): the platform's GPU through Metal, and XLA on the # arm64 CPU. # Each leg past the first also runs one mutant on each device, so the @@ -188,6 +190,60 @@ jobs: $check xla cuda gate.mlir '==' $check xla cuda gate_m3.mlir '!=' + amd: + name: tensor program · integrated GPU (framework-128gb) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux-amd + timeout-minutes: 60 + env: + # The Vulkan devices this host has: its integrated GPU and the software + # one. Every one must give the exact rows. + VULKAN_DEVICES: '2' + 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. The host names its own route to + # the forge (WAI_FORGE_URL). + git fetch -q --depth 1 "${WAI_FORGE_URL:-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 program through HIP and on every Vulkan device + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + export PATH="$HOME/.cargo/bin:$PATH" + iree="$HOME/iree-$IREE_VERSION" + [ -x "$iree/bin/iree-compile" ] || { rm -rf "$iree"; python3 -m venv "$iree"; "$iree/bin/pip" install -q "iree-base-compiler==$IREE_VERSION" "iree-base-runtime==$IREE_VERSION"; } + export PATH="$iree/bin:$PATH" + check="bash $PWD/ci/tensor-program/check.sh" + d="$RUNNER_TEMP/program" + CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-tensor-program" nice -n 10 cargo run --locked --release -q --example stablehlo_gate --features quantum_fabric -- "$d" + cd "$d" + # The GPU's own architecture for both: the HIP build's and the Vulkan + # capability profile's (which has 64-bit integers). + arch="${WAI_QUANTUM_AMD_GPU_ARCH:-gfx1151}" + for m in gate gate_m3; do + iree-compile $m.mlir --iree-hal-target-device=hip --iree-hip-target="$arch" -o $m.hip.vmfb + iree-compile $m.mlir --iree-hal-target-device=vulkan --iree-vulkan-target="$arch" -o $m.vulkan.vmfb + done + $check iree hip gate.hip.vmfb '==' + $check iree hip gate_m3.hip.vmfb '!=' + devices="$(iree-run-module --list_devices 2>/dev/null | grep -oE 'vulkan://[^ ]+' | sort -u)" + count="$(grep -c . <<<"$devices")" + [ "$count" = "$VULKAN_DEVICES" ] || { echo "::error::$count Vulkan devices, not $VULKAN_DEVICES"; iree-run-module --list_devices; exit 1; } + for dev in $devices; do + $check iree "$dev" gate.vulkan.vmfb '==' + $check iree "$dev" gate_m3.vulkan.vmfb '!=' + done + metal: name: tensor program · Metal and the arm64 CPU (macos) needs: relevant diff --git a/.github/workflows/wai-quantum-webgpu.yml b/.github/workflows/wai-quantum-webgpu.yml index f78cf37f0..09d881868 100644 --- a/.github/workflows/wai-quantum-webgpu.yml +++ b/.github/workflows/wai-quantum-webgpu.yml @@ -15,7 +15,9 @@ name: wai-quantum WebGPU member # 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. +# libraries that no GPU host has, and the corpus needs none of them. The GPU +# hosts: jetson-hub (a discrete and an integrated GPU) and framework-128gb (an +# integrated GPU). # # The browser job runs the same kernel through a browser's WebGPU, compiled by the # browser's own WGSL compiler: wai-web's `quantum_webgpu` bundle in a headless page @@ -311,6 +313,47 @@ jobs: working-directory: src/wai/wai-rs run: CARGO_TARGET_DIR="$HOME/hw-target/wai-rs" nice -n 10 cargo run --locked --no-default-features --features quantum_fabric_webgpu --bin wai_quantum_vectors -- verify ../quantum-conformance --webgpu + member-amd: + needs: relevant + if: needs.relevant.outputs.run == 'true' + name: member · integrated GPU, Vulkan (framework-128gb) + runs-on: linux-amd + timeout-minutes: 150 + env: + # The adapters are declared by the host, in its runner's environment + # (WAI_QUANTUM_AMD_WEBGPU_EXPECT); with no declaration, the integrated GPU + # and the software Vulkan driver. The wide sizes run on the GPU only. + WAI_QUANTUM_WEBGPU_WIDE: '1' + CARGO_PROFILE_DEV_DEBUG: '0' + CARGO_BUILD_JOBS: '12' + 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. The + # host names its own route to the forge (WAI_FORGE_URL). + git fetch -q --depth 1 "${WAI_FORGE_URL:-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: | + set -euo pipefail + export WAI_QUANTUM_WEBGPU_EXPECT="${WAI_QUANTUM_AMD_WEBGPU_EXPECT:-gpu-vulkan-1002,swgpu-vulkan-10005}" + echo "adapters declared: $WAI_QUANTUM_WEBGPU_EXPECT" + 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 + browser: needs: relevant if: needs.relevant.outputs.run == 'true' -- 2.54.0 (Apple Git-157) From 8432d6115676c518d5c152b3e581c44f7f7cadfa Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 09:03:40 -0400 Subject: [PATCH 08/19] wai-quantum: the avx512 vector-unit member quantum_fabric_vector gains avx512 (x86-64, the foundation instructions): 512-bit vectors, eight 64-bit lanes, each fx an exact 32x32 multiply (vpmuldq), an add and the unit's own 64-bit arithmetic shift (vpsraq), where avx2 needs a logical shift with the sign flipped back. Discovery now opens one member per instruction set the processor reports (native_isas), so a processor with both units has both members. Measured on framework-128gb's processor (silicon): avx2 and avx512 each 9/9 (66020 probe pairs, all four pinned certificates, circuits to 20 qubits, 3 of 3 mutants killed). The workflow's new leg there holds each unit to its vector length, silicon provenance, four certificates and three mutants. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-vector.yml | 49 ++++++++ .../src/quantum_fabric/vector_avx512.rs | 100 ++++++++++++++++ wai/wai-quantum/src/quantum_fabric_vector.rs | 112 +++++++++++++++--- 3 files changed, 247 insertions(+), 14 deletions(-) create mode 100644 wai/wai-quantum/src/quantum_fabric/vector_avx512.rs diff --git a/.github/workflows/wai-quantum-vector.yml b/.github/workflows/wai-quantum-vector.yml index 82e0a01d3..1a0aaee1f 100644 --- a/.github/workflows/wai-quantum-vector.yml +++ b/.github/workflows/wai-quantum-vector.yml @@ -5,6 +5,8 @@ name: wai-quantum CPU vector units # reference's pinned corpus digests. # # - avx2 on silicon: the container runner's x86-64 processor; +# - avx2 and avx512 on silicon together: framework-128gb's processor, which has +# both, one member each; # - sve and rvv under the host's user-mode emulator (its binfmt handler), at # several vector lengths each: the scalable kernels set their length per # pass, and every leg checks the length the member read from the unit is the @@ -171,6 +173,53 @@ jobs: [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; } grep -qx "vector mutants: key=$key killed 3 of 3" "$report" || { echo "::error::a kernel mutant survived"; exit 1; } + vector-amd: + name: member · vector units, avx2 and avx512 (framework-128gb silicon) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux-amd + timeout-minutes: 60 + env: + CARGO_PROFILE_DEV_DEBUG: '0' + CARGO_BUILD_JOBS: '12' + 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. The host names its own route to + # the forge (WAI_FORGE_URL). + git fetch -q --depth 1 "${WAI_FORGE_URL:-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: Both units on silicon + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + export PATH="$HOME/.cargo/bin:$PATH" + export CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-vector" + # The units this host's processor has, declared here (or by the runner). + export WAI_QUANTUM_VECTOR_EXPECT="${WAI_QUANTUM_AMD_VECTOR_EXPECT:-cpu-vector-avx2,cpu-vector-avx512}" + nice -n 10 cargo test --locked --release --features quantum_fabric_vector --lib -- quantum_fabric_vector:: 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; } + report="$RUNNER_TEMP/vector-report.txt"; : > "$report" + WAI_QUANTUM_VECTOR_REPORT="$report" nice -n 10 cargo test --locked --release --features quantum_fabric_vector --test vector_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" + grep -qx "vector devices: as declared ($WAI_QUANTUM_VECTOR_EXPECT)" "$report" || { echo "::error::the units found are not the declared ones"; exit 1; } + for key in ${WAI_QUANTUM_VECTOR_EXPECT//,/ }; do + case "$key" in cpu-vector-avx2) bits=256 ;; cpu-vector-avx512) bits=512 ;; *) echo "::error::no vector length known for $key"; exit 1 ;; esac + grep -q "^vector member: key=$key .* bits=$bits .*; silicon;" "$report" || { echo "::error::$key: not $bits bits on silicon"; exit 1; } + n="$(grep -c "^vector certified: key=$key .* parity=bit-exact " "$report")" + [ "$n" = 4 ] || { echo "::error::$key: $n of 4 certificates were earned"; exit 1; } + grep -qx "vector mutants: key=$key killed 3 of 3" "$report" || { echo "::error::$key: a kernel mutant survived"; exit 1; } + done + vector-absent: name: member · vector unit absent (macos) needs: relevant diff --git a/wai/wai-quantum/src/quantum_fabric/vector_avx512.rs b/wai/wai-quantum/src/quantum_fabric/vector_avx512.rs new file mode 100644 index 000000000..415890772 --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/vector_avx512.rs @@ -0,0 +1,100 @@ +//! The x86-64 vector unit's kernels (AVX-512, its foundation instructions): +//! the gate's fixed-point arithmetic in 512-bit vectors, eight 64-bit lanes, +//! every step exact. +//! +//! This file's text is the implementation id's input. `RND` is the rounding +//! constant, `SH` the shift and `LAST_ADD` whether row 0's last real term is +//! added (a mutant) instead of subtracted. +use std::arch::x86_64::*; + +/// Lanes per vector. +pub const LANES: usize = 8; + +/// Eight operands of `s` from `k`, sign-extended to 64 bits. +#[inline] +#[target_feature(enable = "avx512f")] +unsafe fn load(s: &[i32], k: usize) -> __m512i { + debug_assert!(k + LANES <= s.len()); + // SAFETY: eight values from `k` are in bounds; the load is unaligned. + _mm512_cvtepi32_epi64(unsafe { _mm256_loadu_si256(s.as_ptr().add(k).cast()) }) +} + +/// Writes the eight lanes of `v` to `d` from `k`. +#[inline] +#[target_feature(enable = "avx512f")] +unsafe fn store(d: &mut [i64], k: usize, v: __m512i) { + debug_assert!(k + LANES <= d.len()); + // SAFETY: as for `load`. + unsafe { _mm512_storeu_si512(d.as_mut_ptr().add(k).cast(), v) } +} + +/// `(e·a + RND) >> SH` in each lane. The multiply takes each lane's low 32 +/// bits as signed and forms the exact 64-bit product; the unit's 64-bit +/// arithmetic shift is the shift. +#[inline] +#[target_feature(enable = "avx512f")] +fn fx(e: __m512i, a: __m512i) -> __m512i { + _mm512_srai_epi64::(_mm512_add_epi64(_mm512_mul_epi32(a, e), _mm512_set1_epi64(RND))) +} + +/// A gate's four rows for every pair: `m` the gate's entries, `cols` the pairs' +/// columns (re a0, im a0, re a1, im a1), `out` the rows (0r, 0i, 1r, 1i). +/// +/// # Safety +/// +/// The processor has AVX-512's foundation instructions; every slice has the +/// same length, a multiple of [`LANES`]. +#[target_feature(enable = "avx512f")] +pub unsafe fn rows(m: &[i32; 8], cols: [&[i32]; 4], out: [&mut [i64]; 4]) { + let n = cols[0].len(); + debug_assert!(n.is_multiple_of(LANES) && cols.iter().all(|c| c.len() == n) && out.iter().all(|o| o.len() == n)); + let mut e = [_mm512_setzero_si512(); 8]; + for (v, &x) in e.iter_mut().zip(m) { + *v = _mm512_set1_epi64(i64::from(x)); + } + let [o0r, o0i, o1r, o1i] = out; + macro_rules! f { + ($e:expr, $a:expr) => { + fx::($e, $a) + }; + } + let mut k = 0; + while k < n { + // SAFETY: the caller's slices hold whole vectors. + let (a0r, a0i, a1r, a1i) = unsafe { (load(cols[0], k), load(cols[1], k), load(cols[2], k), load(cols[3], k)) }; + let t = _mm512_add_epi64(_mm512_sub_epi64(f!(e[0], a0r), f!(e[1], a0i)), f!(e[2], a1r)); + let r0r = if LAST_ADD { _mm512_add_epi64(t, f!(e[3], a1i)) } else { _mm512_sub_epi64(t, f!(e[3], a1i)) }; + let r0i = _mm512_add_epi64(_mm512_add_epi64(_mm512_add_epi64(f!(e[0], a0i), f!(e[1], a0r)), f!(e[2], a1i)), f!(e[3], a1r)); + let r1r = _mm512_sub_epi64(_mm512_add_epi64(_mm512_sub_epi64(f!(e[4], a0r), f!(e[5], a0i)), f!(e[6], a1r)), f!(e[7], a1i)); + let r1i = _mm512_add_epi64(_mm512_add_epi64(_mm512_add_epi64(f!(e[4], a0i), f!(e[5], a0r)), f!(e[6], a1i)), f!(e[7], a1r)); + // SAFETY: as for the loads. + unsafe { + store(o0r, k, r0r); + store(o0i, k, r0i); + store(o1r, k, r1r); + store(o1i, k, r1i); + } + k += LANES; + } +} + +/// Each pair's `a·b` and `(a·b + RND) >> SH`. +/// +/// # Safety +/// +/// As for [`rows`]. +#[target_feature(enable = "avx512f")] +pub unsafe fn probe(a: &[i32], b: &[i32], prod: &mut [i64], fixed: &mut [i64]) { + let n = a.len(); + debug_assert!(n.is_multiple_of(LANES) && b.len() == n && prod.len() == n && fixed.len() == n); + let mut k = 0; + while k < n { + // SAFETY: the caller's slices hold whole vectors. + unsafe { + let (va, vb) = (load(a, k), load(b, k)); + store(prod, k, _mm512_mul_epi32(va, vb)); + store(fixed, k, fx::(vb, va)); + } + k += LANES; + } +} diff --git a/wai/wai-quantum/src/quantum_fabric_vector.rs b/wai/wai-quantum/src/quantum_fabric_vector.rs index ca255b6dc..883efffb5 100644 --- a/wai/wai-quantum/src/quantum_fabric_vector.rs +++ b/wai/wai-quantum/src/quantum_fabric_vector.rs @@ -2,6 +2,9 @@ //! the processor's own vector extension, one kernel per instruction set: //! //! - `avx2` (x86-64): 256-bit vectors, four 64-bit lanes ([`AVX2_TEXT`]); +//! - `avx512` (x86-64, the foundation instructions): 512-bit vectors, eight +//! 64-bit lanes, with the unit's own 64-bit arithmetic shift +//! ([`AVX512_TEXT`]); //! - `sve` (aarch64, the scalable vector extension): 128 to 2048-bit vectors, //! the length the processor's ([`SVE_TEXT`]); //! - `rvv` (riscv64, the vector extension 1.0): any vector length @@ -30,6 +33,9 @@ use crate::quantum_fabric::{Certificate, Certification, Corpus, FabricClass, Ker #[cfg(target_arch = "x86_64")] #[path = "quantum_fabric/vector_avx2.rs"] mod avx2; +#[cfg(target_arch = "x86_64")] +#[path = "quantum_fabric/vector_avx512.rs"] +mod avx512; /// One probe case: [`lanes::mul_wide`] and [`lanes::fxmul_wide`] as /// `(high, low)` words. @@ -77,17 +83,19 @@ impl From for MemberError { #[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] pub enum Isa { Avx2, + Avx512, Sve, Rvv, } impl Isa { - pub const ALL: [Isa; 3] = [Isa::Avx2, Isa::Sve, Isa::Rvv]; + pub const ALL: [Isa; 4] = [Isa::Avx2, Isa::Avx512, Isa::Sve, Isa::Rvv]; /// The toolchain's name for the extension. pub fn label(self) -> &'static str { match self { Isa::Avx2 => "avx2", + Isa::Avx512 => "avx512", Isa::Sve => "sve", Isa::Rvv => "rvv", } @@ -97,6 +105,7 @@ impl Isa { pub fn key(self) -> &'static str { match self { Isa::Avx2 => "cpu-vector-avx2", + Isa::Avx512 => "cpu-vector-avx512", Isa::Sve => "cpu-vector-sve", Isa::Rvv => "cpu-vector-rvv", } @@ -105,7 +114,7 @@ impl Isa { /// The processor architecture the extension belongs to. pub fn arch(self) -> &'static str { match self { - Isa::Avx2 => "x86_64", + Isa::Avx2 | Isa::Avx512 => "x86_64", Isa::Sve => "aarch64", Isa::Rvv => "riscv64", } @@ -477,6 +486,9 @@ rvv_probe_fn!(rvv_probe_m2, "0x20000000", "29"); /// The text the conforming x86-64 kernels are compiled from. pub const AVX2_TEXT: &str = include_str!("quantum_fabric/vector_avx2.rs"); +/// The text the conforming 512-bit x86-64 kernels are compiled from. +pub const AVX512_TEXT: &str = include_str!("quantum_fabric/vector_avx512.rs"); + /// The text the conforming scalable kernels are assembled from. pub const SVE_TEXT: &str = concat!(sve_rows_text!("0x20000000", "30", "sub"), sve_probe_text!("0x20000000", "30")); @@ -487,6 +499,7 @@ pub const RVV_TEXT: &str = concat!(rvv_rows_text!("0x20000000", "30", "vsub.vv") fn text_of(isa: Isa, v: Variant) -> &'static str { match (isa, v) { (Isa::Avx2, _) => AVX2_TEXT, + (Isa::Avx512, _) => AVX512_TEXT, (Isa::Sve, Variant::Conforming) => SVE_TEXT, (Isa::Sve, Variant::M1) => concat!(sve_rows_text!("0x10000000", "30", "sub"), sve_probe_text!("0x10000000", "30")), (Isa::Sve, Variant::M2) => concat!(sve_rows_text!("0x20000000", "29", "sub"), sve_probe_text!("0x20000000", "29")), @@ -522,9 +535,15 @@ pub fn implementation_id_of(isa: Isa, v: Variant) -> String { format!("{}-lanes64/{h}", isa.label()) } -/// The instruction set this processor's architecture would use. +/// The first instruction set written for this processor's architecture. pub fn native_isa() -> Option { - Isa::ALL.into_iter().find(|i| i.arch() == std::env::consts::ARCH) + native_isas().into_iter().next() +} + +/// Every instruction set written for this processor's architecture, in +/// [`Isa::ALL`]'s order; which of them the processor has is [`has`]'s answer. +pub fn native_isas() -> Vec { + Isa::ALL.into_iter().filter(|i| i.arch() == std::env::consts::ARCH).collect() } /// The kernel's `AT_HWCAP` on Linux: the word whose bits name the base @@ -559,6 +578,16 @@ pub fn has(isa: Isa) -> Result<(), String> { unreachable!("the architecture was checked") } } + Isa::Avx512 => { + #[cfg(target_arch = "x86_64")] + { + if std::arch::is_x86_feature_detected!("avx512f") { Ok(()) } else { Err("the processor reports no avx512f".into()) } + } + #[cfg(not(target_arch = "x86_64"))] + { + unreachable!("the architecture was checked") + } + } Isa::Sve => { #[cfg(target_arch = "aarch64")] { @@ -592,6 +621,7 @@ pub fn has(isa: Isa) -> Result<(), String> { fn vector_bits(isa: Isa) -> u32 { match isa { Isa::Avx2 => 256, + Isa::Avx512 => 512, Isa::Sve => { #[cfg(target_arch = "aarch64")] { @@ -631,20 +661,28 @@ pub struct Discovery { pub unavailable: Option, } -/// This processor's vector unit, probed in this process, or why there is none. +/// This processor's vector units, one member per instruction set it reports, +/// each probed in this process; or why there is none. pub fn discover() -> Discovery { let mut d = Discovery { found: Vec::new(), refused: Vec::new(), unavailable: None }; - let Some(isa) = native_isa() else { + let isas = native_isas(); + if isas.is_empty() { d.unavailable = Some(format!("no vector kernel is written for {}", std::env::consts::ARCH)); return d; - }; - if let Err(why) = has(isa) { - d.unavailable = Some(why); - return d; } - match VectorMember::open(isa, Variant::Conforming) { - Ok(m) => d.found.push(m), - Err(e) => d.refused.push((isa.key().into(), e)), + let mut missing = Vec::new(); + for isa in isas { + if let Err(why) = has(isa) { + missing.push(why); + continue; + } + match VectorMember::open(isa, Variant::Conforming) { + Ok(m) => d.found.push(m), + Err(e) => d.refused.push((isa.key().into(), e)), + } + } + if d.found.is_empty() && d.refused.is_empty() { + d.unavailable = Some(missing.join("; ")); } d } @@ -796,6 +834,25 @@ impl VectorMember { return Ok((prod, fx)); } } + Isa::Avx512 => { + #[cfg(target_arch = "x86_64")] + { + let (a, b) = (padded(a, avx512::LANES), padded(b, avx512::LANES)); + let (mut prod, mut fx) = (vec![0i64; a.len()], vec![0i64; a.len()]); + // SAFETY: the member exists only where the processor + // reports avx512f; every slice holds whole vectors. + unsafe { + match self.variant { + Variant::M1 => avx512::probe::<{ 1 << 28 }, 30>(&a, &b, &mut prod, &mut fx), + Variant::M2 => avx512::probe::<{ 1 << 29 }, 29>(&a, &b, &mut prod, &mut fx), + Variant::Conforming | Variant::M3 => avx512::probe::<{ 1 << 29 }, 30>(&a, &b, &mut prod, &mut fx), + } + } + prod.truncate(n); + fx.truncate(n); + return Ok((prod, fx)); + } + } Isa::Sve => { #[cfg(target_arch = "aarch64")] { @@ -862,6 +919,32 @@ impl VectorMember { return Ok(o); } } + Isa::Avx512 => { + #[cfg(target_arch = "x86_64")] + { + let c = [padded(&cols[0], avx512::LANES), padded(&cols[1], avx512::LANES), padded(&cols[2], avx512::LANES), padded(&cols[3], avx512::LANES)]; + let len = c[0].len(); + let mut o = [vec![0i64; len], vec![0i64; len], vec![0i64; len], vec![0i64; len]]; + { + let [o0, o1, o2, o3] = &mut o; + let (c, out) = ([&c[0][..], &c[1][..], &c[2][..], &c[3][..]], [&mut o0[..], &mut o1[..], &mut o2[..], &mut o3[..]]); + // SAFETY: the member exists only where the processor + // reports avx512f; every slice holds whole vectors. + unsafe { + match self.variant { + Variant::Conforming => avx512::rows::<{ 1 << 29 }, 30, false>(m, c, out), + Variant::M1 => avx512::rows::<{ 1 << 28 }, 30, false>(m, c, out), + Variant::M2 => avx512::rows::<{ 1 << 29 }, 29, false>(m, c, out), + Variant::M3 => avx512::rows::<{ 1 << 29 }, 30, true>(m, c, out), + } + } + } + for v in &mut o { + v.truncate(n); + } + return Ok(o); + } + } Isa::Sve => { #[cfg(target_arch = "aarch64")] { @@ -1050,11 +1133,12 @@ mod tests { .into_iter() .flat_map(|isa| [Variant::Conforming, Variant::M1, Variant::M2, Variant::M3].into_iter().map(move |v| implementation_id_of(isa, v))) .collect(); - assert_eq!(ids.len(), 12); + assert_eq!(ids.len(), 16); for isa in Isa::ALL { assert!(implementation_id_of(isa, Variant::Conforming).starts_with(&format!("{}-lanes64/", isa.label()))); } assert!(AVX2_TEXT.contains("_mm256_mul_epi32")); + assert!(AVX512_TEXT.contains("_mm512_mul_epi32") && AVX512_TEXT.contains("_mm512_srai_epi64")); assert!(SVE_TEXT.contains("whilelt") && !SVE_TEXT.contains("smstart")); assert!(RVV_TEXT.contains("vsetvli") && RVV_TEXT.contains("vsext.vf2")); } -- 2.54.0 (Apple Git-157) From 54fc4ed52903ab97fe415805260a8b8c130a4555 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 09:17:34 -0400 Subject: [PATCH 09/19] wai-quantum: the npu-cores member, the gate on a neural processor's own core quantum_fabric_npu_cores (opt-in, Linux): the gate's fixed-point arithmetic on one compute core of a neural processor's tile array, programmed directly rather than through a model runtime. The kernel (npu_cores.cc) is the same 64-bit transcription of lanes::exact_row as the other members'; the core forms each gathered pair's four rows, the host checks their range and keeps the low 32 bits. A line carries a 16-word header (the mode and the gate's words), so one program serves every gate and the arithmetic probe; each kernel variant (the conforming one and the three mutants) is its own program. At run time the member finds the device in the accelerator class by its driver, builds the program with the tile array's open toolchain (npu_cores_design.py, run by the Python WAI_QUANTUM_NPU_CORES_PYTHON names; cached per variant, generation and toolchain), compiles the host side (npu_cores_host.cpp) with the system C++ compiler against the device runtime and loads it. The key is npu-cores--; the member id binds the PCI ids, revision, runtime name, firmware and toolchain versions. Measured on the XDNA 2 processor (silicon), npu_cores_member 9/9 in 100 s: 66020 probe pairs, all four pinned certificates, circuits to 20 qubits, 3 of 3 mutants killed. Also: the tensor program's HIP build is --iree-hal-target-backends=rocm --iree-rocm-target (the compiler has no --iree-hip-target), found by the leg's dry run on the CI user. Co-Authored-By: Claude Opus 5.5 --- .../workflows/wai-quantum-tensor-program.yml | 2 +- wai/wai-quantum/Cargo.toml | 6 + wai/wai-quantum/src/lib.rs | 7 +- wai/wai-quantum/src/quantum_fabric.rs | 1 + .../src/quantum_fabric/npu_cores.cc | 65 ++ .../src/quantum_fabric/npu_cores.rs | 273 +++++++ .../src/quantum_fabric/npu_cores_design.py | 98 +++ .../src/quantum_fabric/npu_cores_host.cpp | 125 +++ .../src/quantum_fabric_npu_cores.rs | 742 ++++++++++++++++++ wai/wai-quantum/tests/npu_cores_member.rs | 414 ++++++++++ 10 files changed, 1731 insertions(+), 2 deletions(-) create mode 100644 wai/wai-quantum/src/quantum_fabric/npu_cores.cc create mode 100644 wai/wai-quantum/src/quantum_fabric/npu_cores.rs create mode 100644 wai/wai-quantum/src/quantum_fabric/npu_cores_design.py create mode 100644 wai/wai-quantum/src/quantum_fabric/npu_cores_host.cpp create mode 100644 wai/wai-quantum/src/quantum_fabric_npu_cores.rs create mode 100644 wai/wai-quantum/tests/npu_cores_member.rs diff --git a/.github/workflows/wai-quantum-tensor-program.yml b/.github/workflows/wai-quantum-tensor-program.yml index 0860ae65d..8d1e7cc1b 100644 --- a/.github/workflows/wai-quantum-tensor-program.yml +++ b/.github/workflows/wai-quantum-tensor-program.yml @@ -231,7 +231,7 @@ jobs: # capability profile's (which has 64-bit integers). arch="${WAI_QUANTUM_AMD_GPU_ARCH:-gfx1151}" for m in gate gate_m3; do - iree-compile $m.mlir --iree-hal-target-device=hip --iree-hip-target="$arch" -o $m.hip.vmfb + iree-compile $m.mlir --iree-hal-target-backends=rocm --iree-rocm-target="$arch" -o $m.hip.vmfb iree-compile $m.mlir --iree-hal-target-device=vulkan --iree-vulkan-target="$arch" -o $m.vulkan.vmfb done $check iree hip gate.hip.vmfb '==' diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml index e32d114d4..489e4be88 100644 --- a/wai/wai-quantum/Cargo.toml +++ b/wai/wai-quantum/Cargo.toml @@ -169,6 +169,8 @@ quantum_fabric_opencl = ["quantum_fabric", "dep:opencl3"] quantum_fabric_fpga = ["quantum_fabric"] # The neural-processor member: exact fp16 limb products on the platform's neural processor. quantum_fabric_npu = ["quantum_fabric", "dep:objc2", "dep:block2", "dep:objc2-foundation", "dep:objc2-core-ml"] +# The npu-cores member: the gate's arithmetic on a neural processor's compute core, programmed directly (Linux): the program built by the tile array's open toolchain, the host side compiled against the device runtime, both at run time. +quantum_fabric_npu_cores = ["quantum_fabric"] # The matrix-extension member: the gate's arithmetic in the processor's streaming vector mode (SME). quantum_fabric_sme = ["quantum_fabric"] # The DSP vector-unit member: the gate's arithmetic in the signal processor's vector unit. @@ -215,6 +217,10 @@ required-features = ["quantum_fabric_hip"] name = "hip_sv" required-features = ["quantum_fabric_hip", "quantum_sv"] +[[test]] +name = "npu_cores_member" +required-features = ["quantum_fabric_npu_cores"] + [[test]] name = "opencl_member" required-features = ["quantum_fabric_opencl"] diff --git a/wai/wai-quantum/src/lib.rs b/wai/wai-quantum/src/lib.rs index fbfb541f0..41f23b688 100644 --- a/wai/wai-quantum/src/lib.rs +++ b/wai/wai-quantum/src/lib.rs @@ -10,13 +10,16 @@ //! 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. Ten opt-in features reach outside the crate: +//! member that ran the work. Eleven opt-in features reach outside the crate: //! `quantum_fabric_webgpu` (the platform's GPU interface), //! `quantum_fabric_cuda` (the CUDA driver and runtime compiler), //! `quantum_fabric_hip` (the HIP runtime and runtime compiler) and //! `quantum_fabric_opencl` (the OpenCL ICD loader), the last three loaded when //! first used, `quantum_fabric_fpga` (Verilator, run as a program when first //! used), `quantum_fabric_npu` (the platform's model runtime, on macOS), +//! `quantum_fabric_npu_cores` (a neural processor's compute cores, through +//! the device runtime and the tile array's open toolchain, both used when +//! first needed), //! `quantum_fabric_sme` (the processor's streaming vector mode), //! `quantum_fabric_vector` (the processor's vector extension), //! `quantum_fabric_dsp` (the signal processor's vector unit, in a build for @@ -99,6 +102,8 @@ pub mod quantum_fabric_opencl; pub mod quantum_fabric_fpga; #[cfg(all(feature = "quantum_fabric_npu", not(target_arch = "wasm32")))] pub mod quantum_fabric_npu; +#[cfg(all(feature = "quantum_fabric_npu_cores", not(target_arch = "wasm32")))] +pub mod quantum_fabric_npu_cores; #[cfg(all(feature = "quantum_fabric_sme", not(target_arch = "wasm32")))] pub mod quantum_fabric_sme; #[cfg(all(feature = "quantum_fabric_dsp", not(target_arch = "wasm32")))] diff --git a/wai/wai-quantum/src/quantum_fabric.rs b/wai/wai-quantum/src/quantum_fabric.rs index 3f15fa170..55d2e130a 100644 --- a/wai/wai-quantum/src/quantum_fabric.rs +++ b/wai/wai-quantum/src/quantum_fabric.rs @@ -2051,6 +2051,7 @@ pub mod cuda; pub mod hip; pub mod fpga; pub mod npu; +pub mod npu_cores; pub mod opencl; pub mod plan; #[cfg(feature = "quantum_fabric_dataflow")] diff --git a/wai/wai-quantum/src/quantum_fabric/npu_cores.cc b/wai/wai-quantum/src/quantum_fabric/npu_cores.cc new file mode 100644 index 000000000..51050fadf --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/npu_cores.cc @@ -0,0 +1,65 @@ +// wai-quantum fabric: the gate's rows on one compute core of a neural +// processor's tile array, compiled for the core by the array's open compiler. +// +// A transcription of `lanes::exact_row` in 64-bit integers: each +// fx(e, a) = (e·a + RND) >> SH is an exact 64-bit product of two 32-bit +// values (at most 2^62 in size), an add and an arithmetic shift, and each row +// is the sum of four such terms (at most 2^34 in size). The core writes the +// four rows whole; the host checks each row's range and keeps its low 32 +// bits, as `lanes::apply_words` does. Nothing here is floating point or +// depends on the order anything runs in. +// +// RND, SH and LAST_ADD (whether row 0's last real term is added instead of +// subtracted, a mutant) are given by the build, so each variant of the kernel +// is its own compiled program. +// +// One line: a 16-word header (word 0 the mode, words 8..15 the gate's +// entries m0..m7), then `n` cases of four words. Mode 0: each case is a pair's +// columns (re a0, im a0, re a1, im a1), and its output the four rows +// (0r, 0i, 1r, 1i). Mode 1, the arithmetic probe: each case is (a, b, 0, 0), +// and its output a·b and fx(b, a), then two zeros. Every output value is 64 +// bits, written as two words, the low one first. +#include + +#if !defined(RND) || !defined(SH) || !defined(LAST_ADD) +#error "RND, SH and LAST_ADD are given by the build" +#endif + +static inline int64_t fx(int32_t e, int32_t a) { + return ((int64_t)e * (int64_t)a + (int64_t)RND) >> SH; +} + +static inline void put(int32_t *o, int64_t v) { + uint64_t u = (uint64_t)v; + o[0] = (int32_t)(uint32_t)u; + o[1] = (int32_t)(uint32_t)(u >> 32); +} + +extern "C" void wq_line(const int32_t *in, int32_t *out, int32_t n) { + const int32_t mode = in[0]; + const int32_t *m = in + 8; + const int32_t *p = in + 16; + if (mode == 1) { + for (int32_t k = 0; k < n; k++) { + const int32_t a = p[4 * k], b = p[4 * k + 1]; + put(out + 8 * k + 0, (int64_t)a * (int64_t)b); + put(out + 8 * k + 2, fx(b, a)); + put(out + 8 * k + 4, 0); + put(out + 8 * k + 6, 0); + } + return; + } + const int32_t m0 = m[0], m1 = m[1], m2 = m[2], m3 = m[3], m4 = m[4], m5 = m[5], m6 = m[6], m7 = m[7]; + for (int32_t k = 0; k < n; k++) { + const int32_t a0r = p[4 * k], a0i = p[4 * k + 1], a1r = p[4 * k + 2], a1i = p[4 * k + 3]; + int64_t r0r = fx(m0, a0r) - fx(m1, a0i) + fx(m2, a1r); + r0r = LAST_ADD ? r0r + fx(m3, a1i) : r0r - fx(m3, a1i); + const int64_t r0i = fx(m0, a0i) + fx(m1, a0r) + fx(m2, a1i) + fx(m3, a1r); + const int64_t r1r = fx(m4, a0r) - fx(m5, a0i) + fx(m6, a1r) - fx(m7, a1i); + const int64_t r1i = fx(m4, a0i) + fx(m5, a0r) + fx(m6, a1i) + fx(m7, a1r); + put(out + 8 * k + 0, r0r); + put(out + 8 * k + 2, r0i); + put(out + 8 * k + 4, r1r); + put(out + 8 * k + 6, r1i); + } +} diff --git a/wai/wai-quantum/src/quantum_fabric/npu_cores.rs b/wai/wai-quantum/src/quantum_fabric/npu_cores.rs new file mode 100644 index 000000000..fc85bf262 --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/npu_cores.rs @@ -0,0 +1,273 @@ +//! What names the npu-cores member, and the line layout it sends: the +//! kernel one compute core of a neural processor's tile array runs +//! ([`KERNEL`], `npu_cores.cc`), the program around it ([`DESIGN`], built by +//! the array's open toolchain), the host side ([`HOST`]), the 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 member itself is `quantum_fabric_npu_cores`. +//! +//! **Why it is bit-exact.** The kernel is the same transcription of +//! [`lanes::exact_row`](super::lanes) as the other members': each operand a +//! 32-bit value, each `fx(a, b) = (a·b + 2^29) >> 30` an exact 64-bit +//! product, an add and an arithmetic shift, each row the sum of four. The core +//! writes the rows whole; the host checks each row's range and keeps its low +//! 32 bits, as [`lanes::apply_words`](super::lanes::apply_words) does. The +//! member's probe checks the core's products and shifts against +//! [`lanes::mul_wide`](super::lanes::mul_wide) and +//! [`lanes::fxmul_wide`](super::lanes::fxmul_wide) before the device is a +//! member, and its certificates check whole runs against the CPU's pinned +//! corpus digests. + +/// The kernel, exactly as it is compiled for the core. +pub const KERNEL: &str = include_str!("npu_cores.cc"); + +/// The program around the kernel, run by the toolchain's Python to build it. +pub const DESIGN: &str = include_str!("npu_cores_design.py"); + +/// The host side, compiled against the device runtime and loaded. +pub const HOST: &str = include_str!("npu_cores_host.cpp"); + +/// The implementation name the id starts with. +pub const IMPLEMENTATION: &str = "npu-cores-lanes64"; + +/// Cases per line. +pub const LINE: usize = 256; + +/// Lines per dispatch. +pub const LINES: usize = 4; + +/// Cases per dispatch. +pub const CASES: usize = LINE * LINES; + +/// Words before a line's cases: word 0 the mode, words 8..15 the gate's +/// entries. +pub const HEADER: usize = 16; + +/// Words per case in, and out (four 64-bit values). +pub const CASE_IN: usize = 4; +pub const CASE_OUT: usize = 8; + +/// The modes a line's header names. +pub const MODE_ROWS: i32 = 0; +pub const MODE_PROBE: i32 = 1; + +/// Version of the host layout: the line format above, the gather of a gate's +/// pairs, the 64-bit rows and the host's range check. Part of the +/// implementation id. +pub const HOST_LAYOUT: u32 = 1; + +/// The kernels a member runs: the conforming one, or a mutant. Each is its +/// own compiled program. +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub enum Variant { + Conforming, + /// The rounding constant is half what it should be. + M1, + /// The product is shifted by one bit too few. + M2, + /// Row 0's real part adds its last term instead of subtracting it. + M3, +} + +impl Variant { + pub const MUTANTS: [Variant; 3] = [Variant::M1, Variant::M2, Variant::M3]; + + /// The kernel's constants: the rounding constant, the shift, and whether + /// row 0's last real term is added. + pub fn constants(self) -> (i64, u32, bool) { + match self { + Variant::Conforming => (1 << 29, 30, false), + Variant::M1 => (1 << 28, 30, false), + Variant::M2 => (1 << 29, 29, false), + Variant::M3 => (1 << 29, 30, true), + } + } + + pub fn label(self) -> &'static str { + match self { + Variant::Conforming => "conforming", + Variant::M1 => "m1", + Variant::M2 => "m2", + Variant::M3 => "m3", + } + } +} + +fn put(o: &mut Vec, bytes: &[u8]) { + o.extend_from_slice(&(bytes.len() as u64).to_be_bytes()); + o.extend_from_slice(bytes); +} + +fn hex(bytes: &[u8]) -> String { + bytes.iter().map(|b| format!("{b:02x}")).collect() +} + +/// `"npu-cores-lanes64/<16 hex>"`: the first 8 bytes of +/// BLAKE3(`"wai:quantum-fabric-npu-cores\x01"` ‖ put([`KERNEL`]) ‖ +/// put([`DESIGN`]) ‖ put([`HOST`]) ‖ the variant's constants ‖ [`LINE`] ‖ +/// [`LINES`] ‖ [`HOST_LAYOUT`]), the numbers big-endian. The toolchain that +/// builds the program is the device's and is in the member id instead. +pub fn implementation_id_of(v: Variant) -> String { + let (rnd, sh, last_add) = v.constants(); + let mut o = b"wai:quantum-fabric-npu-cores\x01".to_vec(); + put(&mut o, KERNEL.as_bytes()); + put(&mut o, DESIGN.as_bytes()); + put(&mut o, HOST.as_bytes()); + o.extend_from_slice(&rnd.to_be_bytes()); + o.extend_from_slice(&sh.to_be_bytes()); + o.push(u8::from(last_add)); + o.extend_from_slice(&(LINE as u64).to_be_bytes()); + o.extend_from_slice(&(LINES as u64).to_be_bytes()); + o.extend_from_slice(&HOST_LAYOUT.to_be_bytes()); + format!("{IMPLEMENTATION}/{}", hex(&blake3::hash(&o).as_bytes()[..8])) +} + +pub fn implementation_id() -> String { + implementation_id_of(Variant::Conforming) +} + +/// The toolchain's name for a device's tile-array generation, by its PCI +/// device id, for the ones this member's program is written for. +pub fn generation(device: u16) -> Option<&'static str> { + match device { + // The first generation, and the second (512-bit cores, eight columns). + 0x1502 => Some("npu1"), + 0x17f0 => Some("npu2"), + _ => None, + } +} + +/// What an npu-cores member is, as its id hashes it. The member fills it from +/// the device and the toolchain; the fields are plain values here so that the +/// id is testable without a device. +#[derive(Clone, Debug, PartialEq, Eq)] +pub struct DeviceIdentity { + /// PCI vendor and device ids, and the revision. + pub vendor: u16, + pub device: u16, + pub revision: u8, + /// The device's name as its runtime gives it. + pub name: String, + /// The device's firmware version, as its driver reports it. + pub firmware: String, + /// The toolchain that built the program, as it reports itself (one line + /// per package). + pub toolchain: String, +} + +impl DeviceIdentity { + /// `"local/npu-cores-"`: the first 12 hex digits of + /// BLAKE3(`"wai:quantum-fabric-npu-cores-member\x01"` ‖ vendor ‖ device ‖ + /// revision ‖ put(name) ‖ put(firmware) ‖ put(toolchain)), the numbers + /// big-endian. + pub fn member_id(&self) -> String { + let mut o = b"wai:quantum-fabric-npu-cores-member\x01".to_vec(); + o.extend_from_slice(&self.vendor.to_be_bytes()); + o.extend_from_slice(&self.device.to_be_bytes()); + o.push(self.revision); + put(&mut o, self.name.as_bytes()); + put(&mut o, self.firmware.as_bytes()); + put(&mut o, self.toolchain.as_bytes()); + format!("local/npu-cores-{}", hex(&blake3::hash(&o).as_bytes()[..6])) + } + + /// The key a runner declares in CI: `"npu-cores--"`, the + /// PCI ids in hex. + pub fn key(&self) -> String { + format!("npu-cores-{:04x}-{:04x}", self.vendor, self.device) + } +} + +/// One dispatch's input: [`LINES`] lines, each a header naming `mode` and +/// `m`, then [`LINE`] cases from `cases` (zeros past its end). At most +/// [`CASES`] cases. +pub fn pack(mode: i32, m: &[i32; 8], cases: &[[i32; 4]]) -> Vec { + assert!(cases.len() <= CASES, "{} cases; a dispatch holds {CASES}", cases.len()); + let mut v = Vec::with_capacity(LINES * (HEADER + CASE_IN * LINE)); + for l in 0..LINES { + let mut head = [0i32; HEADER]; + head[0] = mode; + head[8..16].copy_from_slice(m); + v.extend_from_slice(&head); + for k in 0..LINE { + v.extend_from_slice(&cases.get(l * LINE + k).copied().unwrap_or([0; 4])); + } + } + v +} + +/// The words one dispatch's output holds. +pub const OUT_WORDS: usize = CASES * CASE_OUT; + +/// A dispatch's output as each case's four 64-bit values. +pub fn unpack(out: &[i32]) -> Vec<[i64; 4]> { + assert_eq!(out.len(), OUT_WORDS); + out.as_chunks::<8>() + .0 + .iter() + .map(|w| { + let v = |k: usize| (u64::from(w[2 * k] as u32) | (u64::from(w[2 * k + 1] as u32) << 32)) as i64; + [v(0), v(1), v(2), v(3)] + }) + .collect() +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn every_variant_has_its_own_implementation_id() { + let ids: std::collections::BTreeSet = [Variant::Conforming, Variant::M1, Variant::M2, Variant::M3].into_iter().map(implementation_id_of).collect(); + assert_eq!(ids.len(), 4); + assert!(implementation_id().starts_with("npu-cores-lanes64/") && implementation_id().len() == "npu-cores-lanes64/".len() + 16); + for v in Variant::MUTANTS { + assert_ne!(v.constants(), Variant::Conforming.constants()); + } + assert!(KERNEL.contains("extern \"C\" void wq_line(") && DESIGN.contains("\"wq_line\"") && HOST.contains("wq_npu_run")); + } + + #[test] + fn every_field_of_the_identity_is_in_the_member_id() { + let base = DeviceIdentity { vendor: 0x1022, device: 0x17f0, revision: 0x11, name: "a device".into(), firmware: "1.1.2.65".into(), toolchain: "a 1\nb 2".into() }; + let id = base.member_id(); + assert!(id.starts_with("local/npu-cores-") && id.len() == "local/npu-cores-".len() + 12, "{id}"); + assert!(id.as_str() > "local/cpu-x86_64", "sorts after the CPU reference"); + for v in [ + DeviceIdentity { vendor: 0x1023, ..base.clone() }, + DeviceIdentity { device: 0x1502, ..base.clone() }, + DeviceIdentity { revision: 0x10, ..base.clone() }, + DeviceIdentity { name: "another".into(), ..base.clone() }, + DeviceIdentity { firmware: "1.1.2.66".into(), ..base.clone() }, + DeviceIdentity { toolchain: "a 1\nb 3".into(), ..base.clone() }, + ] { + assert_ne!(v.member_id(), id, "{v:?}"); + } + assert_eq!(base.key(), "npu-cores-1022-17f0"); + assert_eq!((generation(0x17f0), generation(0x1502), generation(0x1234)), (Some("npu2"), Some("npu1"), None)); + } + + #[test] + fn a_dispatch_is_lines_of_a_header_and_cases() { + let m = [1, 2, 3, 4, 5, 6, 7, 8]; + let cases: Vec<[i32; 4]> = (0..LINE as i32 + 3).map(|k| [k, -k, 2 * k, -2 * k]).collect(); + let v = pack(MODE_ROWS, &m, &cases); + assert_eq!(v.len(), LINES * (HEADER + CASE_IN * LINE)); + let stride = HEADER + CASE_IN * LINE; + for l in 0..LINES { + assert_eq!(v[l * stride], MODE_ROWS); + assert_eq!(&v[l * stride + 8..l * stride + 16], &m); + } + // The second line starts with case LINE; past the cases, zeros. + assert_eq!(&v[stride + HEADER..stride + HEADER + 4], &[LINE as i32, -(LINE as i32), 2 * LINE as i32, -2 * LINE as i32]); + assert_eq!(&v[stride + HEADER + 4 * 3..stride + HEADER + 4 * 4], &[0, 0, 0, 0]); + assert_eq!(pack(MODE_PROBE, &m, &[])[0], MODE_PROBE); + let mut out = vec![0i32; OUT_WORDS]; + out[0] = -1; + out[1] = 0x7fff_ffff; + out[2] = 5; + let u = unpack(&out); + assert_eq!(u.len(), CASES); + assert_eq!(u[0], [i64::MAX, 5, 0, 0]); + } +} diff --git a/wai/wai-quantum/src/quantum_fabric/npu_cores_design.py b/wai/wai-quantum/src/quantum_fabric/npu_cores_design.py new file mode 100644 index 000000000..ae6ad3e8b --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/npu_cores_design.py @@ -0,0 +1,98 @@ +"""The tile array's program for wai-quantum's npu-cores member. + +One compute core runs the kernel (npu_cores.cc) over a dispatch of LINES +lines of LINE cases each: the array's data movers bring each line into the +core's memory and take its outputs back to the host's buffer. The member runs +this script with the open toolchain's Python to build the program, then loads +it itself; nothing here runs on the device. + + python npu_cores_design.py --kernel npu_cores.cc --rnd 536870912 --sh 30 --last-add 0 \\ + --line 256 --lines 4 --device npu2 --xclbin final.xclbin --insts insts.bin + +With --versions it prints the toolchain's versions, one per line, and exits. +""" + +import argparse +import importlib.metadata as md +import os +import site +import sys + +HEADER = 16 # words before a line's cases +CASE_IN = 4 # words per case in +CASE_OUT = 8 # words per case out: four 64-bit values + + +def versions(): + for name in ("mlir_aie", "llvm-aie"): + try: + print(f"{name} {md.version(name)}") + except md.PackageNotFoundError: + print(f"{name} absent") + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("--versions", action="store_true") + ap.add_argument("--kernel") + ap.add_argument("--rnd", type=int) + ap.add_argument("--sh", type=int) + ap.add_argument("--last-add", type=int, choices=[0, 1]) + ap.add_argument("--line", type=int) + ap.add_argument("--lines", type=int) + ap.add_argument("--device", default="npu2") + ap.add_argument("--xclbin") + ap.add_argument("--insts") + a = ap.parse_args() + if a.versions: + versions() + return + # The core compiler ships in its own package beside this interpreter's. + if "PEANO_INSTALL_DIR" not in os.environ: + for d in site.getsitepackages(): + if os.path.isdir(os.path.join(d, "llvm-aie", "bin")): + os.environ["PEANO_INSTALL_DIR"] = os.path.join(d, "llvm-aie") + break + + import numpy as np + import aie.iron as iron + from aie.iron import CompileTime, In, ObjectFifo, Out, Program, Runtime, Worker + from aie.iron.device import from_name + from aie.iron.kernel import ExternalFunction + + source = open(a.kernel).read() + line, lines = a.line, a.lines + flags = [f"-DRND={a.rnd}", f"-DSH={a.sh}", f"-DLAST_ADD={a.last_add}"] + + @iron.jit + def design(inp: In, out: Out, *, n: CompileTime[int]): + lin = np.ndarray[(HEADER + CASE_IN * line,), np.dtype[np.int32]] + lout = np.ndarray[(CASE_OUT * line,), np.dtype[np.int32]] + vin = np.ndarray[(n * (HEADER + CASE_IN * line),), np.dtype[np.int32]] + vout = np.ndarray[(n * CASE_OUT * line,), np.dtype[np.int32]] + of_in = ObjectFifo(lin, name="in") + of_out = ObjectFifo(lout, name="out") + kernel = ExternalFunction("wq_line", source_string=source, arg_types=[lin, lout, np.int32], compile_flags=flags) + + def core(of_in, of_out, kernel): + o = of_out.acquire(1) + i = of_in.acquire(1) + kernel(i, o, line) + of_in.release(1) + of_out.release(1) + + worker = Worker(core, [of_in.cons(), of_out.prod(), kernel]) + + def sequence(x, y, in_h, out_h): + in_h.fill(x) + out_h.drain(y, wait=True) + + rt = Runtime(sequence, [vin, vout, of_in.prod(), of_out.cons()]) + return Program(iron.get_current_device(), rt, workers=[worker]).resolve_program() + + iron.set_current_device(from_name(a.device, n_cols=1)) + design.specialize(n=lines).compile(xclbin_path=a.xclbin, inst_path=a.insts) + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/wai/wai-quantum/src/quantum_fabric/npu_cores_host.cpp b/wai/wai-quantum/src/quantum_fabric/npu_cores_host.cpp new file mode 100644 index 000000000..135f58c6f --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/npu_cores_host.cpp @@ -0,0 +1,125 @@ +// wai-quantum fabric: the host side of the npu-cores member. The device +// runtime's interface for a neural processor (a hardware context over a +// registered program) is C++, so the member compiles this file with the +// system's C++ compiler against that runtime and loads it; this is the whole +// of what it calls. Every function returns 0 on success, or 1 with the +// runtime's message in `err`. +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace { + +struct Npu { + xrt::device device; + xrt::hw_context context; + xrt::kernel kernel; + xrt::bo instructions; + uint32_t count = 0; +}; + +void say(char *err, size_t len, const std::string &what) { + if (len == 0) return; + size_t n = what.size() < len - 1 ? what.size() : len - 1; + std::memcpy(err, what.data(), n); + err[n] = 0; +} + +std::vector words_of(const char *path) { + std::ifstream f(path, std::ios::binary); + if (!f) throw std::runtime_error(std::string("cannot read ") + path); + std::vector bytes((std::istreambuf_iterator(f)), std::istreambuf_iterator()); + if (bytes.empty() || bytes.size() % 4 != 0) throw std::runtime_error(std::string(path) + " is not a whole number of words"); + std::vector w(bytes.size() / 4); + std::memcpy(w.data(), bytes.data(), bytes.size()); + return w; +} + +} // namespace + +// The device's name as the runtime gives it. +extern "C" int wq_npu_name(char *name, size_t len, char *err, size_t errlen) { + try { + xrt::device d(0); + say(name, len, d.get_info()); + return 0; + } catch (const std::exception &e) { + say(err, errlen, e.what()); + return 1; + } +} + +// Opens device 0 with the program `xclbin` and its instruction stream +// `insts`, kernel `kernel` (the first whose name starts with it). +extern "C" int wq_npu_open(const char *xclbin, const char *insts, const char *kernel, void **out, char *err, size_t errlen) { + try { + auto *n = new Npu(); + n->device = xrt::device(0); + xrt::xclbin x{std::string(xclbin)}; + std::string name; + for (auto &k : x.get_kernels()) { + if (k.get_name().rfind(kernel, 0) == 0) { + name = k.get_name(); + break; + } + } + if (name.empty()) { + delete n; + say(err, errlen, std::string("the program has no kernel named ") + kernel); + return 1; + } + n->device.register_xclbin(x); + n->context = xrt::hw_context(n->device, x.get_uuid()); + n->kernel = xrt::kernel(n->context, name); + auto w = words_of(insts); + n->count = static_cast(w.size()); + n->instructions = xrt::bo(n->device, w.size() * sizeof(uint32_t), XCL_BO_FLAGS_CACHEABLE, n->kernel.group_id(1)); + std::memcpy(n->instructions.map(), w.data(), w.size() * sizeof(uint32_t)); + n->instructions.sync(XCL_BO_SYNC_BO_TO_DEVICE); + *out = n; + return 0; + } catch (const std::exception &e) { + say(err, errlen, e.what()); + return 1; + } +} + +// One dispatch: `in` (in_words) to the program's first buffer, its second +// buffer (out_words) back into `out`. +extern "C" int wq_npu_run(void *h, const int32_t *in, size_t in_words, int32_t *out, size_t out_words, char *err, size_t errlen) { + try { + auto *n = static_cast(h); + xrt::bo a(n->device, in_words * 4, XRT_BO_FLAGS_HOST_ONLY, n->kernel.group_id(3)); + xrt::bo b(n->device, out_words * 4, XRT_BO_FLAGS_HOST_ONLY, n->kernel.group_id(4)); + std::memcpy(a.map(), in, in_words * 4); + std::memset(b.map(), 0, out_words * 4); + a.sync(XCL_BO_SYNC_BO_TO_DEVICE); + b.sync(XCL_BO_SYNC_BO_TO_DEVICE); + unsigned int opcode = 3; + auto run = n->kernel(opcode, n->instructions, n->count, a, b); + auto state = run.wait(); + if (state != ERT_CMD_STATE_COMPLETED) { + say(err, errlen, "the dispatch ended in state " + std::to_string(static_cast(state))); + return 1; + } + b.sync(XCL_BO_SYNC_BO_FROM_DEVICE); + std::memcpy(out, b.map(), out_words * 4); + return 0; + } catch (const std::exception &e) { + say(err, errlen, e.what()); + return 1; + } +} + +extern "C" void wq_npu_close(void *h) { + delete static_cast(h); +} diff --git a/wai/wai-quantum/src/quantum_fabric_npu_cores.rs b/wai/wai-quantum/src/quantum_fabric_npu_cores.rs new file mode 100644 index 000000000..55aaa0b47 --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric_npu_cores.rs @@ -0,0 +1,742 @@ +//! The fabric's npu-cores member: the gate's fixed-point arithmetic on one +//! compute core of a neural processor's tile array, programmed directly +//! rather than through a model runtime. See +//! [`quantum_fabric::npu_cores`](crate::quantum_fabric::npu_cores) for the +//! kernel, the line layout and the argument that it is bit-exact. +//! +//! - **Discovery** ([`discover`]): the device is found in the kernel's +//! accelerator class (`/sys/class/accel`), by the driver this member's +//! runtime drives. Its program is built with the tile array's open +//! toolchain, whose Python the host names (`WAI_QUANTUM_NPU_CORES_PYTHON`), +//! into a cache directory keyed by the implementation id and the toolchain +//! (`WAI_QUANTUM_NPU_CORES_CACHE`, else the system's temporary directory). +//! The host side ([`npu_cores::HOST`]) is compiled with the system's C++ +//! compiler (`CXX`, else `c++`) against the device runtime and loaded. The +//! device's arithmetic is then probed against [`lanes`] in this process. +//! A host with no such device finds no member and says why +//! ([`Discovery::unavailable`]); a device that cannot be built for or +//! disagrees is refused with the reason. +//! - **Running**: the host gathers the pairs a gate selects, the core forms +//! their four rows a dispatch of [`npu_cores::CASES`] pairs at a time, and +//! the host checks each row's range and keeps its low 32 bits. +//! - **Certification, placement and receipts** are as for the other members; +//! every certificate's provenance is silicon. +//! +//! # Honest boundaries +//! +//! - **One core.** The program uses one compute core of one column; the array +//! has many. Placement over the columns is not here yet. +//! - **Unmetered**, as every member is by default. + +use crate::quantum::{Amp, Circuit, Gateset, MAX_QUBITS, ONE, QuantumError, StateVector}; +use crate::quantum_fabric::lanes::{self, GateWords, Lane}; +use crate::quantum_fabric::npu_cores::{self, DeviceIdentity, Variant}; +use crate::quantum_fabric::plan::{self, Config, LaneRun, Program, Step}; +use crate::quantum_fabric::{Certificate, Certification, Corpus, FabricClass, KernelEntry, Member, MemberError, Presence, Provenance, RunFuture, RunMeter, Runner, apply_kernel}; +use std::ffi::{CStr, CString, c_char, c_int, c_void}; +use std::path::{Path, PathBuf}; +use std::process::Command; +use std::sync::{Arc, Mutex, OnceLock}; + +/// One probe case: [`lanes::mul_wide`] and [`lanes::fxmul_wide`] as +/// `(high, low)` words. +pub type ProbeOutput = ((u32, u32), (u32, u32)); + +/// Why a device is not a member, or a run has no state. +#[derive(Debug, PartialEq, Eq)] +#[non_exhaustive] +pub enum NpuCoresError { + /// No device, no toolchain, or no runtime, and why. + Unavailable(String), + /// The program or the host side did not build: the tool's output. + Build(String), + /// The runtime reported an error. + Device(String), + /// Exactly what the CPU reference refuses. + Circuit(QuantumError), + /// More qubits than this member 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: 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 NpuCoresError { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + match self { + NpuCoresError::Unavailable(why) => write!(f, "no tile array to drive: {why}"), + NpuCoresError::Build(why) => write!(f, "the build failed: {why}"), + NpuCoresError::Device(why) => write!(f, "device: {why}"), + NpuCoresError::Circuit(e) => write!(f, "{e}"), + NpuCoresError::TooWide { n_qubits, max } => write!(f, "{n_qubits} qubits; this member holds at most {max}"), + NpuCoresError::OutOfRange { gate } => write!(f, "the result of op {gate} leaves 32 bits"), + NpuCoresError::Internal(why) => write!(f, "host defect: {why}"), + NpuCoresError::Probe { a, b } => write!(f, "the device's arithmetic differs from the CPU's at ({a}, {b})"), + } + } +} + +impl std::error::Error for NpuCoresError {} + +impl From for MemberError { + fn from(e: NpuCoresError) -> MemberError { + match e { + NpuCoresError::TooWide { n_qubits, max } => MemberError::TooWide { n_qubits, max }, + NpuCoresError::OutOfRange { gate } => MemberError::OutOfRange { gate }, + NpuCoresError::Device(why) => MemberError::Device(why), + NpuCoresError::Circuit(e) => MemberError::Circuit(e), + other => MemberError::Other(other.to_string()), + } + } +} + +/// The driver whose devices this member's runtime drives. +const DRIVER: &str = "amdxdna"; // allow-name: the kernel driver's name, as sysfs gives it + +/// A device in the accelerator class, as sysfs describes it. +#[derive(Clone, Debug, PartialEq, Eq)] +pub struct Found { + pub path: PathBuf, + pub vendor: u16, + pub device: u16, + pub revision: u8, + pub firmware: String, +} + +/// The accelerator-class devices under `root` (normally `/sys/class/accel`) +/// whose driver is the one this member's runtime drives. +pub fn scan(root: &Path) -> Vec { + let read = |p: PathBuf| std::fs::read_to_string(p).map(|s| s.trim().to_string()).unwrap_or_default(); + let num = |s: String| u32::from_str_radix(s.trim_start_matches("0x"), 16).ok(); + let mut out = Vec::new(); + let Ok(entries) = std::fs::read_dir(root) else { return out }; + let mut names: Vec = entries.filter_map(|e| e.ok().map(|e| e.path())).collect(); + names.sort(); + for p in names { + let dev = p.join("device"); + let driver = std::fs::read_link(dev.join("driver")).ok().and_then(|l| l.file_name().map(|n| n.to_string_lossy().into_owned())).unwrap_or_default(); + if driver != DRIVER { + continue; + } + let (Some(vendor), Some(device)) = (num(read(dev.join("vendor"))), num(read(dev.join("device")))) else { continue }; + out.push(Found { + path: p.clone(), + vendor: vendor as u16, + device: device as u16, + revision: num(read(dev.join("revision"))).unwrap_or(0) as u8, + firmware: read(dev.join("fw_version")), + }); + } + out +} + +/// The host side's entry points. +struct Host { + name: unsafe extern "C" fn(*mut c_char, usize, *mut c_char, usize) -> c_int, + open: unsafe extern "C" fn(*const c_char, *const c_char, *const c_char, *mut *mut c_void, *mut c_char, usize) -> c_int, + run: unsafe extern "C" fn(*mut c_void, *const i32, usize, *mut i32, usize, *mut c_char, usize) -> c_int, + close: unsafe extern "C" fn(*mut c_void), +} + +#[cfg(target_os = "linux")] +mod dl { + use std::ffi::{c_char, c_int, c_void}; + #[link(name = "dl")] + unsafe extern "C" { + pub fn dlopen(filename: *const c_char, flag: c_int) -> *mut c_void; + pub fn dlsym(handle: *mut c_void, symbol: *const c_char) -> *mut c_void; + pub fn dlerror() -> *const c_char; + } + pub const RTLD_NOW: c_int = 2; +} + +/// `p` as the function pointer type `F`, which the caller names from the +/// symbol's declaration. +/// +/// # Safety +/// +/// `p` must be a symbol of exactly the signature `F` is. +#[cfg(target_os = "linux")] +unsafe fn as_fn(p: *mut c_void) -> F { + assert_eq!(std::mem::size_of::(), std::mem::size_of::<*mut c_void>(), "a function pointer"); + // SAFETY: the sizes agree, and the caller names the symbol's signature. + unsafe { std::mem::transmute_copy(&p) } +} + +#[cfg(target_os = "linux")] +fn load_host(lib: &Path) -> Result { + let c = CString::new(lib.as_os_str().as_encoded_bytes()).map_err(|_| "a path holding a NUL".to_string())?; + // SAFETY: dlopen and dlsym with valid C strings; each symbol is cast to the + // signature the host side declares. + unsafe { + let h = dl::dlopen(c.as_ptr(), dl::RTLD_NOW); + if h.is_null() { + let e = dl::dlerror(); + return Err(if e.is_null() { format!("{}: cannot load", lib.display()) } else { CStr::from_ptr(e).to_string_lossy().into_owned() }); + } + let sym = |name: &str| -> Result<*mut c_void, String> { + let n = CString::new(name).expect("no NUL"); + let p = dl::dlsym(h, n.as_ptr()); + if p.is_null() { Err(format!("the host side has no {name}")) } else { Ok(p) } + }; + Ok(Host { name: as_fn(sym("wq_npu_name")?), open: as_fn(sym("wq_npu_open")?), run: as_fn(sym("wq_npu_run")?), close: as_fn(sym("wq_npu_close")?) }) + } +} + +#[cfg(not(target_os = "linux"))] +fn load_host(_lib: &Path) -> Result { + Err(format!("the device runtime is loaded on Linux; this is {}", std::env::consts::OS)) +} + +/// The system's C++ compiler and its version line. +fn cxx() -> Result<(String, String), String> { + let cxx = std::env::var("CXX").unwrap_or_else(|_| "c++".into()); + let out = Command::new(&cxx).arg("--version").output().map_err(|e| format!("no C++ compiler {cxx}: {e}"))?; + Ok((cxx, String::from_utf8_lossy(&out.stdout).lines().next().unwrap_or("").to_string())) +} + +/// The toolchain's Python, and its versions as it reports them. +fn toolchain() -> Result<(PathBuf, String), String> { + let py = std::env::var_os("WAI_QUANTUM_NPU_CORES_PYTHON").filter(|s| !s.is_empty()).map(PathBuf::from).ok_or("no toolchain named (WAI_QUANTUM_NPU_CORES_PYTHON, the Python of an environment with the tile array's toolchain)")?; + let dir = cache_base().join("design"); + std::fs::create_dir_all(&dir).map_err(|e| format!("{}: {e}", dir.display()))?; + let script = dir.join(format!("design-{}.py", hex(&blake3::hash(npu_cores::DESIGN.as_bytes()).as_bytes()[..6]))); + if !script.exists() { + std::fs::write(&script, npu_cores::DESIGN).map_err(|e| format!("{}: {e}", script.display()))?; + } + let out = Command::new(&py).arg(&script).arg("--versions").output().map_err(|e| format!("{}: {e}", py.display()))?; + let versions = String::from_utf8_lossy(&out.stdout).trim().to_string(); + if !out.status.success() || versions.is_empty() || versions.contains("absent") { + return Err(format!("{} does not have the toolchain: {versions} {}", py.display(), String::from_utf8_lossy(&out.stderr).trim())); + } + Ok((py, versions)) +} + +fn hex(bytes: &[u8]) -> String { + bytes.iter().map(|b| format!("{b:02x}")).collect() +} + +fn cache_base() -> PathBuf { + std::env::var_os("WAI_QUANTUM_NPU_CORES_CACHE").map(PathBuf::from).unwrap_or_else(|| std::env::temp_dir().join("wai-quantum-npu-cores")) +} + +/// The cache directory for a variant's program built by a toolchain for a +/// generation. +fn program_dir(v: Variant, generation: &str, toolchain: &str) -> PathBuf { + let mut o = npu_cores::implementation_id_of(v).into_bytes(); + o.push(0); + o.extend_from_slice(generation.as_bytes()); + o.push(0); + o.extend_from_slice(toolchain.as_bytes()); + cache_base().join(format!("program-{}-{}", v.label(), hex(&blake3::hash(&o).as_bytes()[..8]))) +} + +/// Builds into a staging directory beside `dir` with `build`, then moves it +/// into place, so that a concurrent build never sees half a directory. +fn build_into(dir: &Path, build: impl FnOnce(&Path) -> Result<(), String>) -> Result<(), String> { + if dir.exists() { + return Ok(()); + } + let staging = dir.with_extension(format!("building-{}", std::process::id())); + let _ = std::fs::remove_dir_all(&staging); + std::fs::create_dir_all(&staging).map_err(|e| format!("{}: {e}", staging.display()))?; + if let Err(e) = build(&staging) { + let _ = std::fs::remove_dir_all(&staging); + return Err(e); + } + match std::fs::rename(&staging, dir) { + Ok(()) => Ok(()), + // Another process put it in place first. + Err(_) if dir.exists() => { + let _ = std::fs::remove_dir_all(&staging); + Ok(()) + } + Err(e) => Err(format!("{} into {}: {e}", staging.display(), dir.display())), + } +} + +/// The variant's program (`final.xclbin`, `insts.bin`), from the cache or +/// built into it. +fn program(py: &Path, toolchain: &str, generation: &str, v: Variant) -> Result { + let dir = program_dir(v, generation, toolchain); + let (rnd, sh, last_add) = v.constants(); + build_into(&dir, |s| { + std::fs::write(s.join("npu_cores.cc"), npu_cores::KERNEL).map_err(|e| e.to_string())?; + std::fs::write(s.join("design.py"), npu_cores::DESIGN).map_err(|e| e.to_string())?; + let out = Command::new(py) + .current_dir(s) + .args(["design.py", "--kernel", "npu_cores.cc", "--device", generation, "--xclbin", "final.xclbin", "--insts", "insts.bin"]) + .args(["--rnd", &rnd.to_string(), "--sh", &sh.to_string(), "--last-add", &u8::from(last_add).to_string()]) + .args(["--line", &npu_cores::LINE.to_string(), "--lines", &npu_cores::LINES.to_string()]) + .output() + .map_err(|e| format!("{}: {e}", py.display()))?; + if !out.status.success() || !s.join("final.xclbin").exists() || !s.join("insts.bin").exists() { + let mut log = String::from_utf8_lossy(&out.stderr).into_owned(); + log.push_str(&String::from_utf8_lossy(&out.stdout)); + return Err(log.lines().rev().take(20).collect::>().into_iter().rev().collect::>().join("\n")); + } + Ok(()) + }) + .map_err(NpuCoresError::Build)?; + Ok(dir) +} + +/// The host side, compiled once per process into the cache and loaded. +fn host() -> Result<&'static Host, NpuCoresError> { + static HOST: OnceLock> = OnceLock::new(); + HOST.get_or_init(|| { + let (cxx, version) = cxx()?; + let mut o = npu_cores::HOST.as_bytes().to_vec(); + o.extend_from_slice(version.as_bytes()); + let dir = cache_base().join(format!("host-{}", hex(&blake3::hash(&o).as_bytes()[..8]))); + build_into(&dir, |s| { + std::fs::write(s.join("host.cpp"), npu_cores::HOST).map_err(|e| e.to_string())?; + let out = Command::new(&cxx) + .current_dir(s) + .args(["-O2", "-shared", "-fPIC", "-std=c++17", "host.cpp", "-o", "libwqnpu.so", "-lxrt_coreutil"]) + .output() + .map_err(|e| format!("{cxx}: {e}"))?; + if !out.status.success() { + return Err(format!("the host side did not compile against the device runtime: {}", String::from_utf8_lossy(&out.stderr).lines().take(12).collect::>().join("\n"))); + } + Ok(()) + })?; + load_host(&dir.join("libwqnpu.so")) + }) + .as_ref() + .map_err(|e| NpuCoresError::Unavailable(e.clone())) +} + +fn message(buf: &[u8]) -> String { + let end = buf.iter().position(|&b| b == 0).unwrap_or(buf.len()); + String::from_utf8_lossy(&buf[..end]).into_owned() +} + +/// A program open on the device: a hardware context over it, which the +/// runtime holds until it is dropped. +struct Context { + host: &'static Host, + handle: Mutex<*mut c_void>, +} + +// SAFETY: the handle is only used under the mutex, and the runtime's +// objects behind it may be used from any thread one at a time. +unsafe impl Send for Context {} +// SAFETY: as above. +unsafe impl Sync for Context {} + +impl Context { + fn open(host: &'static Host, dir: &Path) -> Result { + let c = |p: PathBuf| CString::new(p.as_os_str().as_encoded_bytes()).map_err(|_| NpuCoresError::Internal("a path holding a NUL")); + let (xclbin, insts) = (c(dir.join("final.xclbin"))?, c(dir.join("insts.bin"))?); + let kernel = CString::new("MLIR_AIE").expect("no NUL"); + let mut err = [0u8; 1024]; + let mut h = std::ptr::null_mut(); + // SAFETY: valid C strings and an error buffer of its stated length. + let r = unsafe { (host.open)(xclbin.as_ptr(), insts.as_ptr(), kernel.as_ptr(), &mut h, err.as_mut_ptr().cast(), err.len()) }; + if r != 0 { + return Err(NpuCoresError::Device(message(&err))); + } + Ok(Context { host, handle: Mutex::new(h) }) + } + + /// One dispatch: `input` in, [`npu_cores::OUT_WORDS`] words out. + fn run(&self, input: &[i32]) -> Result, NpuCoresError> { + let mut out = vec![0i32; npu_cores::OUT_WORDS]; + let mut err = [0u8; 1024]; + let h = self.handle.lock().unwrap_or_else(|p| p.into_inner()); + // SAFETY: the handle is open; the buffers hold the stated words. + let r = unsafe { (self.host.run)(*h, input.as_ptr(), input.len(), out.as_mut_ptr(), out.len(), err.as_mut_ptr().cast(), err.len()) }; + if r != 0 { + return Err(NpuCoresError::Device(message(&err))); + } + Ok(out) + } +} + +impl Drop for Context { + fn drop(&mut self) { + let h = *self.handle.get_mut().unwrap_or_else(|p| p.into_inner()); + // SAFETY: the handle came from open and is closed once. + unsafe { (self.host.close)(h) }; + } +} + +/// A device that is not a member, and why. +#[derive(Debug)] +pub struct Refusal { + pub device: String, + pub key: String, + pub error: NpuCoresError, +} + +/// What [`discover`] found. +#[derive(Debug)] +pub struct Discovery { + pub found: Vec, + pub refused: Vec, + pub unavailable: Option, +} + +/// The tile arrays on this host, each built for, opened and probed in this +/// process; or why there is none. +pub fn discover() -> Discovery { + let mut d = Discovery { found: Vec::new(), refused: Vec::new(), unavailable: None }; + let devices = scan(Path::new("/sys/class/accel")); + if devices.is_empty() { + d.unavailable = Some(format!("no accelerator-class device whose driver is {DRIVER}")); + return d; + } + // The runtime opens one device; a second is left alone and said so. + for (k, f) in devices.into_iter().enumerate() { + let key = format!("npu-cores-{:04x}-{:04x}", f.vendor, f.device); + if k > 0 { + d.refused.push(Refusal { device: f.path.display().to_string(), key, error: NpuCoresError::Unavailable("the runtime opens the first device only".into()) }); + continue; + } + match NpuCoresMember::open(&f, Variant::Conforming) { + Ok(m) => d.found.push(m), + Err(error) => d.refused.push(Refusal { device: f.path.display().to_string(), key, error }), + } + } + d +} + +/// The most qubits a run holds. +pub const MAX_NPU_CORES_QUBITS: u8 = if 22 < MAX_QUBITS { 22 } else { MAX_QUBITS }; + +/// A neural processor's compute core as a fabric member. +#[derive(Clone)] +pub struct NpuCoresMember { + generation: &'static str, + python: PathBuf, + context: Arc, + variant: Variant, + identity: DeviceIdentity, + id: String, + key: String, + implementation: String, + detail: String, + max_qubits: u8, +} + +impl std::fmt::Debug for NpuCoresMember { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("NpuCoresMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).finish() + } +} + +impl NpuCoresMember { + fn open(found: &Found, variant: Variant) -> Result { + let generation = npu_cores::generation(found.device).ok_or_else(|| NpuCoresError::Unavailable(format!("no program is written for device {:04x}:{:04x}", found.vendor, found.device)))?; + let (python, versions) = toolchain().map_err(NpuCoresError::Unavailable)?; + let host = host()?; + let mut name = [0u8; 256]; + let mut err = [0u8; 1024]; + // SAFETY: buffers of their stated lengths. + if unsafe { (host.name)(name.as_mut_ptr().cast(), name.len(), err.as_mut_ptr().cast(), err.len()) } != 0 { + return Err(NpuCoresError::Device(message(&err))); + } + let identity = DeviceIdentity { vendor: found.vendor, device: found.device, revision: found.revision, name: message(&name), firmware: found.firmware.clone(), toolchain: versions.clone() }; + let dir = program(&python, &versions, generation, variant)?; + let context = Arc::new(Context::open(host, &dir)?); + let member = NpuCoresMember { + generation, + python, + context, + variant, + id: identity.member_id(), + key: identity.key(), + implementation: npu_cores::implementation_id_of(variant), + detail: format!( + "npu; device {:?} ({:04x}:{:04x} rev {:02x}), the tile array's {generation} generation, firmware {}; one compute core; toolchain {}; up to {MAX_NPU_CORES_QUBITS} qubits; silicon; unmetered", + identity.name, + found.vendor, + found.device, + found.revision, + found.firmware, + versions.replace('\n', ", ") + ), + identity, + max_qubits: MAX_NPU_CORES_QUBITS, + }; + if variant == Variant::Conforming { + let pairs = lanes::probe_pairs(4096, 0x5eed); + let out = member.probe(&pairs)?; + if let Some((a, b)) = lanes::first_mismatch(&pairs, &out) { + return Err(NpuCoresError::Probe { a, b }); + } + } + Ok(member) + } + + /// The same device running a mutant's program: for the mutation tests. + #[doc(hidden)] + pub fn with_variant(&self, v: Variant) -> Result { + let dir = program(&self.python, &self.identity.toolchain, self.generation, v)?; + let context = Arc::new(Context::open(host()?, &dir)?); + Ok(NpuCoresMember { context, variant: v, implementation: npu_cores::implementation_id_of(v), ..self.clone() }) + } + + #[doc(hidden)] + pub fn with_max_qubits(&self, max_qubits: u8) -> NpuCoresMember { + NpuCoresMember { max_qubits: max_qubits.min(self.max_qubits), ..self.clone() } + } + + pub fn id(&self) -> &str { + &self.id + } + + pub fn key(&self) -> &str { + &self.key + } + + pub fn class(&self) -> FabricClass { + FabricClass::Npu + } + + pub fn detail(&self) -> &str { + &self.detail + } + + pub fn implementation(&self) -> &str { + &self.implementation + } + + pub fn variant(&self) -> Variant { + self.variant + } + + pub fn identity(&self) -> &DeviceIdentity { + &self.identity + } + + pub fn provenance(&self) -> Provenance { + Provenance::Silicon + } + + pub fn max_qubits(&self) -> u8 { + self.max_qubits + } + + pub fn config(&self) -> Config { + Config::PER_GATE + } + + /// The device's [`lanes::mul_wide`] and [`lanes::fxmul_wide`] of each pair. + pub fn probe(&self, pairs: &[(i32, i32)]) -> Result, NpuCoresError> { + let mut out = Vec::with_capacity(pairs.len()); + let split = |v: i64| ((v as u64 >> 32) as u32, v as u64 as u32); + for chunk in pairs.chunks(npu_cores::CASES) { + let cases: Vec<[i32; 4]> = chunk.iter().map(|&(a, b)| [a, b, 0, 0]).collect(); + let got = npu_cores::unpack(&self.context.run(&npu_cores::pack(npu_cores::MODE_PROBE, &[0; 8], &cases))?); + out.extend(got[..chunk.len()].iter().map(|v| (split(v[0]), split(v[1])))); + } + Ok(out) + } + + /// The four rows of each pair, from the device. + fn rows(&self, m: &[i32; 8], cases: &[[i32; 4]]) -> Result, NpuCoresError> { + let mut out = Vec::with_capacity(cases.len()); + for chunk in cases.chunks(npu_cores::CASES) { + let got = npu_cores::unpack(&self.context.run(&npu_cores::pack(npu_cores::MODE_ROWS, m, chunk))?); + out.extend_from_slice(&got[..chunk.len()]); + } + Ok(out) + } + + /// Gates in order from `state`, as [`lanes::run_words`] runs them. + pub fn run_lanes(&self, state: &[Lane], gates: &[GateWords]) -> Result { + let len = state.len(); + if len < 2 || !len.is_power_of_two() { + return Err(NpuCoresError::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(NpuCoresError::TooWide { n_qubits: n, max: self.max_qubits }); + } + let planned: Vec<(&GateWords, u32)> = gates.iter().zip(0u32..).collect(); + self.execute(state.to_vec(), &planned) + } + + fn execute(&self, mut amps: Vec, gates: &[(&GateWords, u32)]) -> Result { + let len = amps.len(); + 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(NpuCoresError::Internal("malformed gate words for this state")); + } + } + let mut report: Option = None; + for &(w, index) in gates { + let m: [i32; 8] = w.m.map(|x| x as i32); + let (tbit, ctrl) = (w.tbit as usize, w.ctrl as usize); + let mut at = Vec::new(); + let mut cases = Vec::new(); + for p in 0..len / 2 { + let low = p & (tbit - 1); + let i = ((p - low) << 1) | low; + if i & ctrl != ctrl { + continue; + } + at.push(i); + let (a0, a1) = (amps[i], amps[i | tbit]); + cases.push([a0[0], a0[1], a1[0], a1[1]]); + } + let rows = self.rows(&m, &cases)?; + for (&i, r) in at.iter().zip(&rows) { + if r.iter().any(|&v| i32::try_from(v).is_err()) { + report = Some(report.map_or(index, |x| x.min(index))); + } + let low = |v: i64| v as u64 as u32 as i32; + amps[i] = [low(r[0]), low(r[1])]; + amps[i | tbit] = [low(r[2]), low(r[3])]; + } + } + Ok(LaneRun { amps, out_of_range: report }) + } + + pub fn run_program(&self, p: &Program) -> Result { + let n = p.n_qubits; + if n == 0 { + return Ok(LaneRun { amps: vec![[ONE as i32, 0]], out_of_range: None }); + } + if n > self.max_qubits { + return Err(NpuCoresError::TooWide { n_qubits: n, max: self.max_qubits }); + } + if p.start >= 1usize << n { + return Err(NpuCoresError::Internal("the start is outside the state")); + } + let mut gates = Vec::with_capacity(p.steps.len()); + for step in &p.steps { + let Step::Gate { gate } = *step else { return Err(NpuCoresError::Internal("a tile step; this member runs gate steps only")) }; + let g = p.gates.get(gate as usize).ok_or(NpuCoresError::Internal("a step names a gate the program does not hold"))?; + if !g.is_well_formed() { + return Err(NpuCoresError::Internal("malformed gate words for this state")); + } + gates.push((&g.words, g.index)); + } + let mut amps = vec![[0i32, 0]; 1usize << n]; + amps[p.start] = [ONE as i32, 0]; + self.execute(amps, &gates) + } + + pub fn simulate_from(&self, c: &Circuit, start: usize) -> Result { + let p = plan::compile_with(c, start, &self.config()).map_err(NpuCoresError::Circuit)?; + if p.n_qubits > self.max_qubits { + return Err(NpuCoresError::TooWide { n_qubits: p.n_qubits, max: self.max_qubits }); + } + let run = self.run_program(&p)?; + if let Some(gate) = run.out_of_range { + return Err(NpuCoresError::OutOfRange { gate }); + } + Ok(StateVector { n_qubits: p.n_qubits, amps: run.amps.iter().map(|l| Amp { re: i64::from(l[0]), im: i64::from(l[1]) }).collect() }) + } + + pub fn simulate(&self, c: &Circuit) -> Result { + self.simulate_from(c, 0) + } + + /// The certificate of [`apply_kernel`]`(gs)` on this member, with its + /// provenance. + pub fn certify(&self, gs: Gateset, corpus: Corpus) -> Result { + let mut cert = Certification::new(); + for c in corpus.circuits(gs) { + let reference = lanes::reference_output(&c).map_err(NpuCoresError::Circuit)?; + let run = self.run_program(&plan::compile_with(&c, 0, &self.config()).map_err(NpuCoresError::Circuit)?)?; + if let Some(gate) = run.out_of_range { + return Err(NpuCoresError::OutOfRange { gate }); + } + cert.add(&lanes::circuit_bytes(&c), &reference, &lanes::output(&run.amps)); + } + Ok(cert.finish(apply_kernel(gs), &self.implementation, &self.id, &corpus.generator(gs)).with_provenance(Provenance::Silicon)) + } + + pub fn member(&self, corpus: Corpus) -> Result { + let mut kernels = Vec::new(); + for gs in Gateset::ALL { + kernels.push(KernelEntry { kernel: apply_kernel(gs).into(), implementation: self.implementation.clone(), certificate: self.certify(gs, corpus)?, energy: None }); + } + Ok(Member { id: self.id.clone(), class: FabricClass::Npu, presence: Presence::Detected, kernels, detail: self.detail.clone() }) + } +} + +impl Runner for NpuCoresMember { + 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_from(c, start)?, 0, None)) + }) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + /// A sysfs tree under `root`: `accel` entries, each a `device` + /// directory with the files given and a `driver` link to a directory of + /// that name. + /// An entry's name, its driver's name and its files. + type Entry<'a> = (&'a str, &'a str, &'a [(&'a str, &'a str)]); + + fn tree(root: &Path, entries: &[Entry]) { + for (name, driver, files) in entries { + let dev = root.join("class").join(name).join("device"); + std::fs::create_dir_all(&dev).unwrap(); + let target = root.join("drivers").join(driver); + std::fs::create_dir_all(&target).unwrap(); + #[cfg(unix)] + std::os::unix::fs::symlink(&target, dev.join("driver")).unwrap(); + for (f, v) in *files { + std::fs::write(dev.join(f), format!("{v}\n")).unwrap(); + } + } + } + + #[test] + #[cfg(unix)] + fn the_scan_finds_the_runtimes_devices_and_nothing_else() { + let root = std::env::temp_dir().join(format!("wq-npu-cores-scan-{}", std::process::id())); + let _ = std::fs::remove_dir_all(&root); + tree( + &root, + &[ + ("accel0", DRIVER, &[("vendor", "0x1022"), ("device", "0x17f0"), ("revision", "0x11"), ("fw_version", "1.1.2.65")]), + ("accel1", "another", &[("vendor", "0x1234"), ("device", "0x0001")]), + ("accel2", DRIVER, &[("vendor", "not hex")]), + ], + ); + let found = scan(&root.join("class")); + assert_eq!(found.len(), 1, "{found:?}"); + assert_eq!((found[0].vendor, found[0].device, found[0].revision, found[0].firmware.as_str()), (0x1022, 0x17f0, 0x11, "1.1.2.65")); + assert!(scan(&root.join("absent")).is_empty()); + let _ = std::fs::remove_dir_all(&root); + } + + #[test] + fn a_program_is_cached_per_variant_generation_and_toolchain() { + let a = program_dir(Variant::Conforming, "npu2", "t 1"); + assert_ne!(a, program_dir(Variant::M1, "npu2", "t 1")); + assert_ne!(a, program_dir(Variant::Conforming, "npu1", "t 1")); + assert_ne!(a, program_dir(Variant::Conforming, "npu2", "t 2")); + assert_eq!(a, program_dir(Variant::Conforming, "npu2", "t 1")); + } + + #[test] + fn the_refusals_that_may_fall_back_do() { + assert!(MemberError::from(NpuCoresError::OutOfRange { gate: 3 }).falls_back()); + assert!(!MemberError::from(NpuCoresError::Probe { a: 1, b: 2 }).falls_back()); + assert!(!MemberError::from(NpuCoresError::Internal("x")).falls_back()); + } +} diff --git a/wai/wai-quantum/tests/npu_cores_member.rs b/wai/wai-quantum/tests/npu_cores_member.rs new file mode 100644 index 000000000..032e27329 --- /dev/null +++ b/wai/wai-quantum/tests/npu_cores_member.rs @@ -0,0 +1,414 @@ +//! The fabric's npu-cores member (a neural processor's compute core) on the +//! machine that runs this. +//! +//! The machine declares the devices it must have in `WAI_QUANTUM_NPU_CORES_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_NPU_CORES_REPORT` names a file, appends the +//! same lines to it, which a CI step can count per device. +//! +//! ```text +//! WAI_QUANTUM_NPU_CORES_EXPECT=npu-cores-1022-17f0 \ +//! WAI_QUANTUM_NPU_CORES_PYTHON= \ +//! cargo test --release --features quantum_fabric_npu_cores --test npu_cores_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::lanes::{self, GateWords, Lane}; +use wai_quantum::quantum_fabric::plan::{Program, Step}; +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::npu_cores::{Variant, implementation_id_of}; +use wai_quantum::quantum_fabric_npu_cores::{Discovery, NpuCoresError, NpuCoresMember, 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_NPU_CORES_REPORT` names. +fn report(line: &str) { + println!("{line}"); + if let Some(path) = std::env::var_os("WAI_QUANTUM_NPU_CORES_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_NPU_CORES_EXPECT={found_list}"); + let expect = match std::env::var("WAI_QUANTUM_NPU_CORES_EXPECT") { + Ok(v) if !v.is_empty() => v, + _ => return Err(format!("WAI_QUANTUM_NPU_CORES_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_NPU_CORES_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 [NpuCoresMember] { + 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!("npu-cores devices: as declared ({expect})")), + Err(why) => panic!("{why}"), + } + for m in &d.found { + assert_eq!(m.class(), FabricClass::Npu); + assert!(m.id().starts_with("local/npu-cores-"), "{}", m.id()); + assert_eq!(m.implementation(), implementation_id_of(Variant::Conforming)); + assert_eq!(m.identity().member_id(), m.id()); + assert!(m.max_qubits() >= 20, "{}: holds only {} qubits", m.key(), m.max_qubits()); + report(&format!("npu-cores 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!("npu-cores 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!("npu-cores 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!( + "npu-cores 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(20); + let mut s = 0x0077_6964_6500_u64; + for (k, n) in [12u8, 16, 18, 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!("npu-cores 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(), NpuCoresError::Circuit(native), "{c:?}"); + } + assert_eq!(m.simulate(&Circuit::new(MAX_QUBITS + 1)).unwrap_err(), NpuCoresError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1))); + let small = m.with_max_qubits(4); + assert_eq!(small.simulate(&Circuit::new(5)).unwrap_err(), NpuCoresError::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(), NpuCoresError::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(NpuCoresError::Internal(_)))); + assert!(matches!(m.run_lanes(&[[0, 0]; 3], &[z_words()]), Err(NpuCoresError::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(NpuCoresError::Internal(_)))); + // The refusals that may still fall back to another member, do. + assert!(MemberError::from(NpuCoresError::TooWide { n_qubits: 5, max: 4 }).falls_back()); + report(&format!("npu-cores refusals: key={} typed", m.key())); + } +} + +/// The kernel mutants, compiled in: each must be caught by the probe or by a +/// certificate. +#[test] +fn c6_mutants_are_killed() { + let edges = lanes::probe_pairs(0, 0); + for m in members() { + let mut killed = 0; + for v in Variant::MUTANTS { + let mm = m.with_variant(v).unwrap_or_else(|e| panic!("{}: mutant {v:?} did not build: {e}", m.key())); + assert_ne!(mm.implementation(), m.implementation()); + 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(); + let caught = differing > 0 + || Gateset::ALL.into_iter().all(|gs| match mm.certify(gs, RANDOM_V1) { + Err(_) => true, + Ok(c) => c.parity != Parity::BitExact || c.corpus != lanes::certify_corpus(gs, "local/test", RANDOM_V1).corpus, + }); + println!("npu-cores mutant {v:?} on {}: {differing} of {} edge pairs differ", m.key(), edges.len()); + assert!(caught, "{}: mutant {v:?} survived", m.key()); + killed += 1; + } + report(&format!("npu-cores mutants: key={} killed {killed} of 3", m.key())); + } +} + +#[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::Npu, 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!("npu-cores 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!("npu-cores repeat: key={} ok", m.key())); + } +} -- 2.54.0 (Apple Git-157) From 4ea41b67d9923d9f2bc6916382400d9af75b6dc8 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 09:21:04 -0400 Subject: [PATCH 10/19] ci: the npu-cores member's workflow (framework-128gb's neural processor, and a host with none) Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-npu-cores.yml | 160 ++++++++++++++++++++ 1 file changed, 160 insertions(+) create mode 100644 .github/workflows/wai-quantum-npu-cores.yml diff --git a/.github/workflows/wai-quantum-npu-cores.yml b/.github/workflows/wai-quantum-npu-cores.yml new file mode 100644 index 000000000..9478b11ea --- /dev/null +++ b/.github/workflows/wai-quantum-npu-cores.yml @@ -0,0 +1,160 @@ +name: wai-quantum npu-cores member + +# The fabric's npu-cores member: the gate's fixed-point arithmetic on one +# compute core of a neural processor's tile array, programmed directly. The +# member builds the core's program with the tile array's open toolchain +# (Apache-2.0 with the LLVM exception, pinned below) and its host side with +# the system's C++ compiler against the device runtime, both at run time. +# +# - npu-cores-hw (framework-128gb, its neural processor): 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) on +# silicon, and catch every kernel mutant, each its own program. +# - npu-cores-absent (a container with no such device): 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-npu-cores.yml' + pull_request: + paths: + - 'wai/wai-quantum/**' + - 'rust-toolchain.toml' + - '.github/workflows/wai-quantum-npu-cores.yml' + +env: + # The tile array's toolchain: the compiler and its core back end, the + # build the compiler's release pins (its utils/peano-requirements.txt). + TOOLCHAIN_VERSION: 1.4.3 + CORE_COMPILER: llvm-aie==22.0.0.2026090701+3e93bf7b + +jobs: + # Does this pull request's own diff touch the paths below? This forge matches a + # workflow's `paths:` against the push, so a merge of main into a branch fires the + # workflow whatever the pull request changes; every job here waits on this answer. + # A push to main, a schedule, or anything the script cannot determine answers + # `true` (scripts/pr-touches-paths.py). + relevant: + name: does this change touch it + runs-on: linux-arm64 + outputs: + run: ${{ steps.gate.outputs.run }} + steps: + - uses: actions/checkout@v5 + with: + fetch-depth: 0 + - id: gate + run: | + set -euo pipefail + r=$(python3 scripts/pr-touches-paths.py .github/workflows/wai-quantum-npu-cores.yml \ + "${{ github.event.pull_request.base.sha }}" \ + "${{ github.event.pull_request.head.sha }}") + echo "this change touches wai-quantum-npu-cores: $r" + echo "run=$r" >> "$GITHUB_OUTPUT" + + npu-cores-hw: + name: member · npu-cores, neural processor (framework-128gb silicon) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux-amd + timeout-minutes: 90 + env: + CARGO_PROFILE_DEV_DEBUG: '0' + CARGO_BUILD_JOBS: '12' + 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. The host names its own route to + # the forge (WAI_FORGE_URL). + git fetch -q --depth 1 "${WAI_FORGE_URL:-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 toolchain (cached) + shell: bash + run: | + set -euo pipefail + venv="$HOME/aie-$TOOLCHAIN_VERSION" + if ! "$venv/bin/python" -I -c 'import aie.iron' 2>/dev/null; then + rm -rf "$venv" + python3 -m venv "$venv" + tag="$(python3 -c 'import sys; print(f"cp{sys.version_info.major}{sys.version_info.minor}")')" + "$venv/bin/pip" install -q "https://github.com/Xilinx/mlir-aie/releases/download/v$TOOLCHAIN_VERSION/mlir_aie-$TOOLCHAIN_VERSION-$tag-$tag-manylinux_2_35_x86_64.whl" # allow-name: the toolchain's release URL + "$venv/bin/pip" install -q -f https://github.com/Xilinx/llvm-aie/releases/expanded_assets/nightly "$CORE_COMPILER" # allow-name: the core compiler's release index + "$venv/bin/pip" install -q "numpy>=2.5.2,<3.0" rich ml_dtypes cloudpickle aiofiles + "$venv/bin/pip" install -q -f https://llvm.github.io/eudsl "eudsl-python-extras==0.1.0.20260801.905+68a0d7a" --config-settings=EUDSL_PYTHON_EXTRAS_HOST_PACKAGE_PREFIX=aie + fi + echo "WAI_QUANTUM_NPU_CORES_PYTHON=$venv/bin/python" >> "$GITHUB_ENV" + - 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-npu-cores" nice -n 10 cargo test --locked --features quantum_fabric_npu_cores --lib -- quantum_fabric_npu_cores:: quantum_fabric::npu_cores:: 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" + export WAI_QUANTUM_NPU_CORES_CACHE="$HOME/npu-cores-cache" + # The device this host has, declared here (or by the runner). + key="${WAI_QUANTUM_AMD_NPU_CORES_EXPECT:-npu-cores-1022-17f0}" + report="$RUNNER_TEMP/npu-cores-report.txt"; : > "$report" + WAI_QUANTUM_NPU_CORES_EXPECT="$key" WAI_QUANTUM_NPU_CORES_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-npu-cores" nice -n 10 \ + cargo test --locked --release --features quantum_fabric_npu_cores --test npu_cores_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" + grep -qx "npu-cores devices: as declared ($key)" "$report" || { echo "::error::the device found is not the declared one"; exit 1; } + grep -q "^npu-cores member: key=$key .*; silicon;" "$report" || { echo "::error::the member does not say it ran on silicon"; exit 1; } + n="$(grep -c "^npu-cores certified: key=$key .* parity=bit-exact " "$report")" + [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; } + grep -qx "npu-cores mutants: key=$key killed 3 of 3" "$report" || { echo "::error::a kernel mutant survived"; exit 1; } + + npu-cores-absent: + name: member · npu-cores absent (linux) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux + timeout-minutes: 60 + env: + WAI_QUANTUM_NPU_CORES_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 neural processor here, and the member says so + shell: bash + working-directory: wai/wai-quantum + run: | + set -euo pipefail + report="$RUNNER_TEMP/npu-cores-report.txt"; : > "$report" + cargo test --locked --features quantum_fabric_npu_cores --lib -- quantum_fabric_npu_cores:: quantum_fabric::npu_cores:: 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; } + WAI_QUANTUM_NPU_CORES_REPORT="$report" cargo test --locked --features quantum_fabric_npu_cores --test npu_cores_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 'npu-cores devices: as declared (none)' "$report" || { echo "::error::the absent case was not established"; exit 1; } -- 2.54.0 (Apple Git-157) From 0176aab894d4cab48524c2c95af95d2592be6f10 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 09:22:27 -0400 Subject: [PATCH 11/19] wai-quantum: npu-cores imports CStr only where it is used (Linux) Co-Authored-By: Claude Opus 5.5 --- wai/wai-quantum/src/quantum_fabric_npu_cores.rs | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/wai/wai-quantum/src/quantum_fabric_npu_cores.rs b/wai/wai-quantum/src/quantum_fabric_npu_cores.rs index 55aaa0b47..a1c4483d3 100644 --- a/wai/wai-quantum/src/quantum_fabric_npu_cores.rs +++ b/wai/wai-quantum/src/quantum_fabric_npu_cores.rs @@ -33,7 +33,7 @@ use crate::quantum_fabric::lanes::{self, GateWords, Lane}; use crate::quantum_fabric::npu_cores::{self, DeviceIdentity, Variant}; use crate::quantum_fabric::plan::{self, Config, LaneRun, Program, Step}; use crate::quantum_fabric::{Certificate, Certification, Corpus, FabricClass, KernelEntry, Member, MemberError, Presence, Provenance, RunFuture, RunMeter, Runner, apply_kernel}; -use std::ffi::{CStr, CString, c_char, c_int, c_void}; +use std::ffi::{CString, c_char, c_int, c_void}; use std::path::{Path, PathBuf}; use std::process::Command; use std::sync::{Arc, Mutex, OnceLock}; @@ -176,7 +176,7 @@ fn load_host(lib: &Path) -> Result { let h = dl::dlopen(c.as_ptr(), dl::RTLD_NOW); if h.is_null() { let e = dl::dlerror(); - return Err(if e.is_null() { format!("{}: cannot load", lib.display()) } else { CStr::from_ptr(e).to_string_lossy().into_owned() }); + return Err(if e.is_null() { format!("{}: cannot load", lib.display()) } else { std::ffi::CStr::from_ptr(e).to_string_lossy().into_owned() }); } let sym = |name: &str| -> Result<*mut c_void, String> { let n = CString::new(name).expect("no NUL"); -- 2.54.0 (Apple Git-157) From d73a7990d93d802e08e9f0b2da9996ccdfdef4c1 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 09:39:45 -0400 Subject: [PATCH 12/19] wai-quantum 0.6.0: the release that breaks, so the next field does not The branch carries three API breaks since 0.5.1 (OpenCL keys and DeviceIdentity::vendor_id, Isa::Avx512, Chip::vector), so it releases as 0.6.0. The device identities (cuda, hip, opencl, npu_cores), the dataflow Chip and the vector Isa are #[non_exhaustive] now, so adding a field or an instruction set is not a break again. wai-quantum-qpu becomes 0.2.0 against wai-quantum 0.6.0, and wai-rs requires 0.6.0. README rows for the HIP, npu-cores and avx512 members and the OpenCL key; the CHANGELOG entry. Verified on the build host: wai-quantum --features full --lib 597 passed; every member feature checks with its tests and examples; wai-quantum-qpu's tests pass; wai-rs builds with its quantum features. Co-Authored-By: Claude Opus 5.5 --- wai/CHANGELOG.md | 49 +++++++++++++++++-- wai/tools/quantum-meter/Cargo.lock | 2 +- wai/tools/quantum-wgsl-check/Cargo.lock | 2 +- wai/wai-pack/Cargo.lock | 2 +- wai/wai-quantum-qpu/Cargo.lock | 4 +- wai/wai-quantum-qpu/Cargo.toml | 4 +- wai/wai-quantum-serve/Cargo.lock | 2 +- wai/wai-quantum/Cargo.lock | 2 +- wai/wai-quantum/Cargo.toml | 2 +- wai/wai-quantum/README.md | 6 ++- wai/wai-quantum/src/quantum_fabric/cuda.rs | 1 + wai/wai-quantum/src/quantum_fabric/hip.rs | 1 + .../src/quantum_fabric/npu_cores.rs | 1 + wai/wai-quantum/src/quantum_fabric/opencl.rs | 1 + .../src/quantum_fabric_dataflow.rs | 1 + wai/wai-quantum/src/quantum_fabric_vector.rs | 1 + wai/wai-rs/Cargo.lock | 2 +- wai/wai-rs/Cargo.toml | 2 +- wai/wai-web/Cargo.lock | 2 +- 19 files changed, 69 insertions(+), 18 deletions(-) diff --git a/wai/CHANGELOG.md b/wai/CHANGELOG.md index 2958975c7..0617eee72 100644 --- a/wai/CHANGELOG.md +++ b/wai/CHANGELOG.md @@ -6,9 +6,39 @@ wire format may change and such changes are recorded here. ## v0.4.0 — current -### wai-quantum 0.5.2: the dataflow device's vector engine, and the gate as a portable tensor program +### wai-quantum 0.6.0: the HIP, npu-cores and avx512 members, the dataflow device's vector engine, and the gate as a portable tensor program -Additive. The default features, `full`, and every pinned digest are unchanged. +Breaking in three places (the last bullet). The default features, `full`, and every pinned +digest are unchanged. + +- **The HIP member** (`quantum_fabric_hip`, opt-in, Linux): the CUDA member's kernel text, + unchanged, compiled at run time by HIP's runtime compiler for each device's own + architecture, with contraction off. The runtime and its compiler are loaded when first + used; the device's properties are read in the runtime's frozen layout, or the device is + refused. Its own implementation id (`hip-lanes64/…`), its own key (`gpu-hip-`); the + mutants are the CUDA kernel's. With `quantum_sv`, the dense f64 state vector runs the CUDA + kernels' text too (`hip-sv64/…`). +- **The npu-cores member** (`quantum_fabric_npu_cores`, opt-in, Linux): the gate's + arithmetic on one compute core of a neural processor's tile array, programmed directly + rather than through a model runtime. The core forms each gathered pair's four rows in + exact 64-bit integers; the host checks their range and keeps the low 32 bits. The program + (`npu_cores.cc` in `npu_cores_design.py`) is built at run time with the tile array's open + toolchain and cached per variant, generation and toolchain; the host side + (`npu_cores_host.cpp`) is compiled against the device runtime and loaded. One program + serves every gate and the arithmetic probe; each mutant is its own program. Key + `npu-cores--`. +- **The avx512 vector unit** (`quantum_fabric_vector`): eight 64-bit lanes, each `fx` an + exact 32×32 multiply, an add and the unit's own 64-bit arithmetic shift. Discovery opens + one member per instruction set the processor reports, so a processor with both units has + both members. +- **Measured** on an integrated GPU (gfx1151), a neural processor's tile array and the + processor beside them, all silicon: the HIP member 9/9 (66020 probe pairs, all four pinned + certificates, circuits to 22 qubits, 6 of 6 mutants) and its state vector 3/3 (the pin + `b7d129d0…`, 4 of 4 mutants); the npu-cores member 9/9 (66020 probe pairs, all four + certificates, circuits to 20 qubits, 3 of 3 mutants); avx2 and avx512 9/9 each; and, + unchanged, the WebGPU member (`gpu-vulkan-1002`) and the OpenCL member + (`opencl-gpu-1002`, the state vector's pin included) bit-exact there. CI runs each of them + on that host. - **The vector engine as a member** (`quantum_fabric_dataflow`, key `dataflow-1e52-b140-vector`): on the device generation whose vector engine multiplies 23-bit numbers, the gate's @@ -46,7 +76,20 @@ Additive. The default features, `full`, and every pinned digest are unchanged. generations the fleet lacks compile but are not run. `ci/tensor-program/check.sh` holds each run to its outcome and fails a run that ends without stating one, so a device that cannot run reads as a failure, never as a mutant caught. CI: - `wai-quantum-tensor-program`. + `wai-quantum-tensor-program`. That host's GPU runs it through HIP, built for the GPU's own + architecture, and on both of its Vulkan devices. +- **Breaking:** + - an OpenCL key is `opencl--`, the device's `CL_DEVICE_VENDOR_ID` in hex + (`opencl-cpu-8086`, `opencl-gpu-10de`, `opencl-gpu-1027f00`, `opencl-gpu-1002`), not a word + of the vendor's name; `opencl::DeviceIdentity` gains `vendor_id`, which the member id + binds, and `vendor_slug` is gone; + - `quantum_fabric_vector::Isa` gains `Avx512` (`Isa::ALL` has four), and `native_isas` lists + every set written for the architecture; + - `quantum_fabric_dataflow::Chip` gains `vector`. + +### wai-quantum-qpu 0.2.0 + +Builds against wai-quantum 0.6.0; nothing of its own changes. ### wai-quantum 0.5.1: neural-processor, matrix-extension, vector-unit, signal-processor and dataflow members diff --git a/wai/tools/quantum-meter/Cargo.lock b/wai/tools/quantum-meter/Cargo.lock index cd24b9ed9..1f4435f1b 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.5.2" +version = "0.6.0" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/tools/quantum-wgsl-check/Cargo.lock b/wai/tools/quantum-wgsl-check/Cargo.lock index 60e84a45d..ac48ce6b7 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.5.2" +version = "0.6.0" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-pack/Cargo.lock b/wai/wai-pack/Cargo.lock index 55ab937cf..ff12a4851 100644 --- a/wai/wai-pack/Cargo.lock +++ b/wai/wai-pack/Cargo.lock @@ -703,7 +703,7 @@ dependencies = [ [[package]] name = "wai-quantum" -version = "0.5.2" +version = "0.6.0" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-quantum-qpu/Cargo.lock b/wai/wai-quantum-qpu/Cargo.lock index 1b3261977..8864c9a27 100644 --- a/wai/wai-quantum-qpu/Cargo.lock +++ b/wai/wai-quantum-qpu/Cargo.lock @@ -422,7 +422,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" [[package]] name = "wai-quantum" -version = "0.5.2" +version = "0.6.0" dependencies = [ "blake3", "ed25519-dalek", @@ -431,7 +431,7 @@ dependencies = [ [[package]] name = "wai-quantum-qpu" -version = "0.1.0" +version = "0.2.0" dependencies = [ "blake3", "serde_json", diff --git a/wai/wai-quantum-qpu/Cargo.toml b/wai/wai-quantum-qpu/Cargo.toml index eeba1828a..d068933f9 100644 --- a/wai/wai-quantum-qpu/Cargo.toml +++ b/wai/wai-quantum-qpu/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "wai-quantum-qpu" -version = "0.1.0" +version = "0.2.0" edition = "2024" rust-version = "1.97" license = "Apache-2.0" @@ -10,6 +10,6 @@ keywords = ["quantum", "qpu", "openqasm", "verification", "receipts"] categories = ["science", "api-bindings"] [dependencies] -wai-quantum = { version = "0.5.0", path = "../wai-quantum", default-features = false, features = ["quantum_fabric"] } +wai-quantum = { version = "0.6.0", path = "../wai-quantum", default-features = false, features = ["quantum_fabric"] } serde_json = "1" blake3 = "1" diff --git a/wai/wai-quantum-serve/Cargo.lock b/wai/wai-quantum-serve/Cargo.lock index 15d5d9cbe..1dc465b84 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.5.2" +version = "0.6.0" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-quantum/Cargo.lock b/wai/wai-quantum/Cargo.lock index c4a208df5..ad3d6b295 100644 --- a/wai/wai-quantum/Cargo.lock +++ b/wai/wai-quantum/Cargo.lock @@ -1079,7 +1079,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a" [[package]] name = "wai-quantum" -version = "0.5.2" +version = "0.6.0" dependencies = [ "blake3", "block2", diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml index 489e4be88..c184371c1 100644 --- a/wai/wai-quantum/Cargo.toml +++ b/wai/wai-quantum/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "wai-quantum" -version = "0.5.2" +version = "0.6.0" edition = "2024" rust-version = "1.97" license = "Apache-2.0" diff --git a/wai/wai-quantum/README.md b/wai/wai-quantum/README.md index 1931628f8..d9233e101 100644 --- a/wai/wai-quantum/README.md +++ b/wai/wai-quantum/README.md @@ -106,12 +106,14 @@ never creates information. | `quantum_fabric` | **the compute fabric**: members (the CPU and accelerators, quantum processors among them), each kernel implementation certified against the CPU reference in the process that uses it (bit-exact, within a stated tolerance, or statistical by classical fidelity over stated shots), with the certificate saying where it was earned (silicon, an emulator, or a simulator); placement by parity, provenance and metered energy; a signed receipt naming the member and implementation behind a result. Two kernels: apply (a circuit's state) and sample (its measurement counts), the latter with the CPU reference sampler, a corpus shaped for quantum processors, and a virtual quantum processor whose stated noise model lets every step of a processor's path run and be tested without one; its certificates say they were earned in a simulator. `quantum_fabric::stablehlo` writes the gate as a portable tensor program (StableHLO), every value a 64-bit integer that never overflows, for any tensor compiler to build for its own devices; the `stablehlo_gate` example writes it with a batch of inputs and the exact outputs a build must give, and `ci/tensor-program/check.sh` runs a build and holds it to them | | `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 features that reach outside the crate | | `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. With `quantum_sv` too, the **dense f64 state vector** runs on the device bit for bit as on the CPU: the CPU's own launches (diagonal, one-qubit and grouped formulas), each an IEEE binary64 operation with no fused multiply-add, certified against a pinned corpus digest | -| `quantum_fabric_opencl` | **the fabric's OpenCL member**, opt-in and not in `full`, native targets only: the same exact 64-bit kernel in OpenCL C, built at run time by each device's own compiler, one launch per gate, on any OpenCL GPU or CPU (an OpenCL accelerator is refused: no fabric class names it). The ICD loader is opened when first used, so a build with the feature runs on a host without OpenCL and finds no member there, saying why. Each device is probed and certified as the CUDA member's are. With `quantum_sv` too, the dense f64 state vector runs bit for bit as on the CPU on every device with binary64 (contraction off); a device without it refuses, saying so | +| `quantum_fabric_hip` | **the fabric's HIP member**, opt-in and not in `full`, Linux: the CUDA member's kernel text, unchanged, compiled at run time by HIP's runtime compiler for each device's own architecture, one launch per gate. The runtime and its compiler are loaded when first used, so a build with the feature runs on a host without them and finds no member there, saying why; the device's properties are read in the runtime's frozen layout or the device is refused. Probed and certified as the CUDA member is, under its own implementation id, and with `quantum_sv` the dense f64 state vector runs bit for bit as on the CPU | +| `quantum_fabric_opencl` | **the fabric's OpenCL member**, opt-in and not in `full`, native targets only: the same exact 64-bit kernel in OpenCL C, built at run time by each device's own compiler, one launch per gate, on any OpenCL GPU or CPU (an OpenCL accelerator is refused: no fabric class names it). The ICD loader is opened when first used, so a build with the feature runs on a host without OpenCL and finds no member there, saying why. Each device is probed and certified as the CUDA member's are, and keyed by its vendor id (`opencl--`). With `quantum_sv` too, the dense f64 state vector runs bit for bit as on the CPU on every device with binary64 (contraction off); a device without it refuses, saying so | | `quantum_fabric_fpga` | **the fabric's FPGA member, simulated**, opt-in and not in `full`, native targets only: the gate datapath in Verilog (each 32×32 product from four 16-bit partial products, the size of an FPGA's multiplier blocks; a state RAM; one pair per two cycles), compiled at run time by Verilator into a cycle-accurate simulation driven over a pipe. Probed and certified as the other members are, and every certificate says it was earned in a simulator, so placement ranks it after silicon. A host without Verilator finds no member and says why | | `quantum_fabric_npu` | **the fabric's neural-processor member**, opt-in and not in `full`, macOS only: each 32-bit operand split into eight 4-bit limbs, every limb product and limb sum formed in fp16 on the platform's neural processor (each value an integer fp16 holds exactly, whatever order it is added in), the host assembling each product and applying the fixed-point arithmetic. The member exists only where the platform model runtime's compute plan places every operation of the model on the neural processor; a model it would run elsewhere is refused with the placement it reported. Probed and certified as every member is | +| `quantum_fabric_npu_cores` | **the fabric's npu-cores member**, opt-in and not in `full`, Linux: the gate's fixed-point arithmetic on one compute core of a neural processor's tile array, programmed directly rather than through a model runtime: exact 64-bit products on the core, the host gathering each gate's pairs and checking each row's range. At run time the member finds the device by its driver, builds the core's program with the tile array's open toolchain (`WAI_QUANTUM_NPU_CORES_PYTHON` names its Python; the program is cached per variant, generation and toolchain) and compiles its host side against the device runtime. Each of the three kernel mutants is its own program and must be caught. Probed and certified as the other members are; its certificates' provenance is silicon | | `quantum_fabric_sme` | **the fabric's matrix-extension member**, opt-in and not in `full`, aarch64 only: the gate's fixed-point arithmetic in the processor's streaming vector mode (SME), 64-bit lanes, the host gathering each gate's pairs and checking each row's range. Every instruction it uses exists only in that mode, so a kernel that ran, ran there. Detected by the processor's feature flags (on Linux, the kernel's `AT_HWCAP2`, which a user-mode emulator supplies); under an emulator the runner says so (`WAI_QUANTUM_SME_EMULATOR`) and every certificate carries it. Three kernel mutants are compiled in and must be caught | | `quantum_fabric_dsp` | **the fabric's DSP vector-unit member**, opt-in and not in `full`, in a build for the signal processor with its vector unit in 128-byte mode: the gate's fixed-point arithmetic in 32-bit lanes, each 64-bit value a pair of words, every product the vector unit's exact 32×32 multiply and every carry and borrow kept; the host gathers each gate's pairs and checks each row's range. Its instructions exist only on that unit, so a kernel that ran, ran there. Under an emulator the runner says so (`WAI_QUANTUM_DSP_EMULATOR`) and every certificate carries it. Three kernel mutants are compiled in and must be caught. Its class is `dsp` | -| `quantum_fabric_vector` | **the fabric's CPU vector-unit members**, opt-in and not in `full`: the gate's fixed-point arithmetic in the processor's own vector extension, one kernel per instruction set — `avx2` (x86-64, 256-bit), `sve` (aarch64, the scalable vector extension, any length) and `rvv` (riscv64, the vector extension 1.0, any length) — each operand sign-extended to 64-bit lanes, the host gathering each gate's pairs and checking each row's range. The scalable kernels set their length per pass; the member reads the length from the unit and names it in its id and detail. Detected by the processor's feature flags (on Linux, `AT_HWCAP` for the vector extension, which a user-mode emulator supplies); under an emulator the runner says so (`WAI_QUANTUM_VECTOR_EMULATOR`). Its id extends the CPU reference's, so placement with no energy figures keeps the reference. Three kernel mutants per instruction set are compiled in and must be caught | +| `quantum_fabric_vector` | **the fabric's CPU vector-unit members**, opt-in and not in `full`: the gate's fixed-point arithmetic in the processor's own vector extension, one kernel per instruction set, and one member per set the processor has — `avx2` (x86-64, 256-bit), `avx512` (x86-64, 512-bit, with the unit's own 64-bit arithmetic shift), `sve` (aarch64, the scalable vector extension, any length) and `rvv` (riscv64, the vector extension 1.0, any length) — each operand sign-extended to 64-bit lanes, the host gathering each gate's pairs and checking each row's range. The scalable kernels set their length per pass; the member reads the length from the unit and names it in its id and detail. Detected by the processor's feature flags (on Linux, `AT_HWCAP` for the vector extension, which a user-mode emulator supplies); under an emulator the runner says so (`WAI_QUANTUM_VECTOR_EMULATOR`). Its id extends the CPU reference's, so placement with no energy figures keeps the reference. Three kernel mutants per instruction set are compiled in and must be caught | | `quantum_fabric_dataflow` | **the fabric's dataflow member**, opt-in and not in `full`, Linux: the gate's fixed-point arithmetic on the first RISC-V core of each of 32 compute tiles of a dataflow accelerator, each over a contiguous slice of a gate's pairs. The program is built in this crate (`quantum_fabric::rv32`, an RV32IM encoder whose listings an independent assembler reproduces word for word): every 64-bit value a pair of words, `mul`/`mulh` for products, `sltu` for carries and borrows. The device's library (`WAI_QUANTUM_DATAFLOW_LIB`, loaded at run time) presents the device as a PCI endpoint; the host opens two memory windows on each tile, holds its cores in reset, writes the program at the reset address and that tile's parameters and slice beside it, releases the first cores and clocks the device until every tile reports done. The device is named by its PCI ids (`dataflow-1e52-`); the library is a simulator, and every certificate says so (`WAI_QUANTUM_DATAFLOW_SIMULATOR`). Three kernel mutants are compiled in. On the device generation whose vector engine multiplies 23-bit numbers, a second member (`-vector`) runs the gate on that engine, 32 lanes at a time: the first core configures the tile's unpacker and packer and hands the tile its instructions inline (`quantum_fabric::tensix`, transcribed from the simulator's own table and checked against it; `quantum_fabric::vector_engine`), each 32×32 product from four 23-bit partial products | | `full` | everything | diff --git a/wai/wai-quantum/src/quantum_fabric/cuda.rs b/wai/wai-quantum/src/quantum_fabric/cuda.rs index ffb0dd41f..67eaec1e0 100644 --- a/wai/wai-quantum/src/quantum_fabric/cuda.rs +++ b/wai/wai-quantum/src/quantum_fabric/cuda.rs @@ -98,6 +98,7 @@ pub fn implementation_id_of(source: &str) -> String { /// driver; the fields are plain values here so that the id is testable /// without a device. #[derive(Clone, Debug, PartialEq, Eq)] +#[non_exhaustive] pub struct DeviceIdentity { pub name: String, /// The device's UUID, in hex. diff --git a/wai/wai-quantum/src/quantum_fabric/hip.rs b/wai/wai-quantum/src/quantum_fabric/hip.rs index 1d21f0022..d3281a3da 100644 --- a/wai/wai-quantum/src/quantum_fabric/hip.rs +++ b/wai/wai-quantum/src/quantum_fabric/hip.rs @@ -77,6 +77,7 @@ pub fn implementation_id_of(source: &str) -> String { /// runtime; the fields are plain values here so that the id is testable /// without a device. #[derive(Clone, Debug, PartialEq, Eq)] +#[non_exhaustive] pub struct DeviceIdentity { pub name: String, /// The device's UUID, in hex. diff --git a/wai/wai-quantum/src/quantum_fabric/npu_cores.rs b/wai/wai-quantum/src/quantum_fabric/npu_cores.rs index fc85bf262..dcd924b6a 100644 --- a/wai/wai-quantum/src/quantum_fabric/npu_cores.rs +++ b/wai/wai-quantum/src/quantum_fabric/npu_cores.rs @@ -141,6 +141,7 @@ pub fn generation(device: u16) -> Option<&'static str> { /// the device and the toolchain; the fields are plain values here so that the /// id is testable without a device. #[derive(Clone, Debug, PartialEq, Eq)] +#[non_exhaustive] pub struct DeviceIdentity { /// PCI vendor and device ids, and the revision. pub vendor: u16, diff --git a/wai/wai-quantum/src/quantum_fabric/opencl.rs b/wai/wai-quantum/src/quantum_fabric/opencl.rs index 28296bbdd..1898eaf01 100644 --- a/wai/wai-quantum/src/quantum_fabric/opencl.rs +++ b/wai/wai-quantum/src/quantum_fabric/opencl.rs @@ -103,6 +103,7 @@ impl DeviceKind { /// runtime; the fields are plain values here so that the id is testable /// without a device. #[derive(Clone, Debug, PartialEq, Eq)] +#[non_exhaustive] pub struct DeviceIdentity { /// The platform's name, version and vendor, as it reports them. pub platform: String, diff --git a/wai/wai-quantum/src/quantum_fabric_dataflow.rs b/wai/wai-quantum/src/quantum_fabric_dataflow.rs index 1f0fd5ebc..b29290466 100644 --- a/wai/wai-quantum/src/quantum_fabric_dataflow.rs +++ b/wai/wai-quantum/src/quantum_fabric_dataflow.rs @@ -192,6 +192,7 @@ pub const TILES: usize = 32; /// The device generations this member knows: the PCI device id, the compute /// tiles it uses, and how a memory window is set up. #[derive(Clone, Copy, Debug, PartialEq, Eq)] +#[non_exhaustive] pub struct Chip { pub device: u16, /// The compute tiles' (x, y) on the device's network, as the host diff --git a/wai/wai-quantum/src/quantum_fabric_vector.rs b/wai/wai-quantum/src/quantum_fabric_vector.rs index 883efffb5..5666c6298 100644 --- a/wai/wai-quantum/src/quantum_fabric_vector.rs +++ b/wai/wai-quantum/src/quantum_fabric_vector.rs @@ -81,6 +81,7 @@ impl From for MemberError { /// The instruction sets a member runs. #[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +#[non_exhaustive] pub enum Isa { Avx2, Avx512, diff --git a/wai/wai-rs/Cargo.lock b/wai/wai-rs/Cargo.lock index 5206f40a0..c68a2c75f 100644 --- a/wai/wai-rs/Cargo.lock +++ b/wai/wai-rs/Cargo.lock @@ -4300,7 +4300,7 @@ version = "0.1.0" [[package]] name = "wai-quantum" -version = "0.5.2" +version = "0.6.0" dependencies = [ "blake3", "ed25519-dalek", diff --git a/wai/wai-rs/Cargo.toml b/wai/wai-rs/Cargo.toml index 1e71b366c..84ebc361e 100644 --- a/wai/wai-rs/Cargo.toml +++ b/wai/wai-rs/Cargo.toml @@ -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.5.0", path = "../wai-quantum", optional = true, default-features = false } +wai-quantum = { version = "0.6.0", path = "../wai-quantum", optional = true, default-features = false } # ---- the integer operator set (behind `neural_int`) ---- # `wai-ops/1` (extensions/integer-operators.md): an integer decoder that decodes diff --git a/wai/wai-web/Cargo.lock b/wai/wai-web/Cargo.lock index 0c3b53fde..fd194df41 100644 --- a/wai/wai-web/Cargo.lock +++ b/wai/wai-web/Cargo.lock @@ -1231,7 +1231,7 @@ version = "0.1.0" [[package]] name = "wai-quantum" -version = "0.5.2" +version = "0.6.0" dependencies = [ "blake3", "ed25519-dalek", -- 2.54.0 (Apple Git-157) From 06263e507ec027118d1b5fee60c047e91695015a Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 09:59:37 -0400 Subject: [PATCH 13/19] wai-quantum: the tensor-engine member, the gate's products on a matrix engine quantum_fabric_tensor_engine (opt-in): the gate's sixteen products per pair as exact limb convolutions in one bf16 matrix multiply per tile of 512 pairs, accumulated in fp32, on a tensor processor's matrix engine. Each 32-bit value is four 8-bit limbs (bf16 holds each exactly), every limb product is at most 2^16 and every coefficient a sum of at most four, so fp32 holds every partial sum exactly in any order: the argument holds for any engine that multiplies bf16 and accumulates in fp32. The host checks every coefficient is an integer below 2^24, assembles the products in 64 bits and applies the fixed-point arithmetic, as the neural-processor member's host does. The probe uses its own weight layout, sixteen independent pairs a multiply. The kernel (tensor_engine_helper.py) runs as a long-lived helper under the Python WAI_QUANTUM_TENSOR_ENGINE_PYTHON names; today its target is the kernel toolchain's simulator (Apache-2.0), so every certificate's provenance is simulator:. Four mutants: the three host constants, and Radix9, nine-bit limbs the engine's bf16 inputs cannot hold, which the engine gets wrong. Measured on the simulator: tensor_engine_member 9/9 in 340 s (66020 probe pairs, all four pinned certificates, circuits to 20 qubits, 4 of 4 mutants killed). Workflow: the simulator on the container runner, and the absent case on linux-arm64, where the HIP and npu-cores absent legs move too. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-hip.yml | 4 +- .github/workflows/wai-quantum-npu-cores.yml | 4 +- .../workflows/wai-quantum-tensor-engine.yml | 156 ++++++ wai/wai-quantum/Cargo.toml | 6 + wai/wai-quantum/src/lib.rs | 8 +- wai/wai-quantum/src/quantum_fabric.rs | 1 + .../src/quantum_fabric/tensor_engine.rs | 380 +++++++++++++ .../quantum_fabric/tensor_engine_helper.py | 82 +++ .../src/quantum_fabric_tensor_engine.rs | 505 ++++++++++++++++++ wai/wai-quantum/tests/tensor_engine_member.rs | 414 ++++++++++++++ 10 files changed, 1554 insertions(+), 6 deletions(-) create mode 100644 .github/workflows/wai-quantum-tensor-engine.yml create mode 100644 wai/wai-quantum/src/quantum_fabric/tensor_engine.rs create mode 100644 wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py create mode 100644 wai/wai-quantum/src/quantum_fabric_tensor_engine.rs create mode 100644 wai/wai-quantum/tests/tensor_engine_member.rs diff --git a/.github/workflows/wai-quantum-hip.yml b/.github/workflows/wai-quantum-hip.yml index accf94484..760a6ba86 100644 --- a/.github/workflows/wai-quantum-hip.yml +++ b/.github/workflows/wai-quantum-hip.yml @@ -122,10 +122,10 @@ jobs: grep -qx "hip sv mutants: key=$key killed 4 of 4" "$report" || { echo "::error::a state-vector mutant survived"; exit 1; } hip-absent: - name: member · HIP absent (linux) + name: member · HIP absent (linux-arm64) needs: relevant if: needs.relevant.outputs.run == 'true' - runs-on: linux + runs-on: linux-arm64 timeout-minutes: 60 env: WAI_QUANTUM_HIP_EXPECT: none diff --git a/.github/workflows/wai-quantum-npu-cores.yml b/.github/workflows/wai-quantum-npu-cores.yml index 9478b11ea..405140f3e 100644 --- a/.github/workflows/wai-quantum-npu-cores.yml +++ b/.github/workflows/wai-quantum-npu-cores.yml @@ -129,10 +129,10 @@ jobs: grep -qx "npu-cores mutants: key=$key killed 3 of 3" "$report" || { echo "::error::a kernel mutant survived"; exit 1; } npu-cores-absent: - name: member · npu-cores absent (linux) + name: member · npu-cores absent (linux-arm64) needs: relevant if: needs.relevant.outputs.run == 'true' - runs-on: linux + runs-on: linux-arm64 timeout-minutes: 60 env: WAI_QUANTUM_NPU_CORES_EXPECT: none diff --git a/.github/workflows/wai-quantum-tensor-engine.yml b/.github/workflows/wai-quantum-tensor-engine.yml new file mode 100644 index 000000000..2937b6da7 --- /dev/null +++ b/.github/workflows/wai-quantum-tensor-engine.yml @@ -0,0 +1,156 @@ +name: wai-quantum tensor-engine member + +# The fabric's tensor-engine member: the gate's products as exact bf16 limb +# convolutions on a tensor processor's matrix engine, accumulated in fp32. The +# kernel runs under the engine's kernel toolchain (Apache-2.0, pinned below), +# whose simulator is the target here: every certificate says it was earned in +# a simulator. +# +# - tensor-engine-sim (the container runner): the device tests must find the +# declared target, agree with the CPU on the arithmetic probe and on random +# states (out-of-range runs included), reproduce all four pinned +# certificates, and catch every mutant, the limb size the engine's inputs +# cannot hold among them. +# - tensor-engine-absent: with no kernel toolchain named, the same build finds +# no member and says why; 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-tensor-engine.yml' + pull_request: + paths: + - 'wai/wai-quantum/**' + - 'rust-toolchain.toml' + - '.github/workflows/wai-quantum-tensor-engine.yml' + +env: + KERNEL_TOOLCHAIN: nki==0.6.0+31049202112.g85070674 + +jobs: + # Does this pull request's own diff touch the paths below? This forge matches a + # workflow's `paths:` against the push, so a merge of main into a branch fires the + # workflow whatever the pull request changes; every job here waits on this answer. + # A push to main, a schedule, or anything the script cannot determine answers + # `true` (scripts/pr-touches-paths.py). + relevant: + name: does this change touch it + runs-on: linux-arm64 + outputs: + run: ${{ steps.gate.outputs.run }} + steps: + - uses: actions/checkout@v5 + with: + fetch-depth: 0 + - id: gate + run: | + set -euo pipefail + r=$(python3 scripts/pr-touches-paths.py .github/workflows/wai-quantum-tensor-engine.yml \ + "${{ github.event.pull_request.base.sha }}" \ + "${{ github.event.pull_request.head.sha }}") + echo "this change touches wai-quantum-tensor-engine: $r" + echo "run=$r" >> "$GITHUB_OUTPUT" + + tensor-engine-sim: + name: member · tensor engine, the kernel toolchain's simulator (linux) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux + timeout-minutes: 60 + container: + image: debian@sha256:913f6706df59a68922d1dd08f78c2476560a8d367897200a6005b00e5f67c2d5 + options: --memory=6g --memory-swap=6g + env: + WAI_QUANTUM_TENSOR_ENGINE_EXPECT: tensor-engine-simulator + steps: + - name: Tools + shell: bash + run: | + set -euo pipefail + apt-get update -qq + apt-get install -y -qq --no-install-recommends ca-certificates curl git build-essential python3 python3-venv >/dev/null + # No Node for a checkout action: the forge's archive of the exact commit, + # its embedded commit id checked. + - name: Fetch the source + shell: bash + run: | + set -euo pipefail + want="${{ github.event.pull_request.head.sha || github.sha }}" + rm -rf src src.tar.gz src.tar && mkdir src + curl -sSf -o src.tar.gz "http://100.64.224.0:3400/Transaction-Science/open-standards/archive/$want.tar.gz" + gzip -dc src.tar.gz > src.tar + got="$(git get-tar-commit-id < src.tar)" + [ "$got" = "$want" ] || { echo "::error::the archive is of $got, not $want"; exit 1; } + tar -xf src.tar -C src --strip-components=1 + rm src.tar.gz src.tar + - name: The member's tests on the simulator + shell: bash + working-directory: src/wai/wai-quantum + run: | + set -euo pipefail + cache=/ci-cache/wai-quantum-tensor-engine + mkdir -p "$cache" + exec 9>"$cache/.lock" + flock 9 + venv="$cache/kernel-toolchain" + if ! "$venv/bin/python" -I -c 'import nki, ml_dtypes' 2>/dev/null || [ "$(cat "$venv/pinned" 2>/dev/null)" != "$KERNEL_TOOLCHAIN" ]; then + rm -rf "$venv" + python3 -m venv "$venv" + "$venv/bin/pip" install -q --extra-index-url https://pip.repos.neuron.amazonaws.com "$KERNEL_TOOLCHAIN" numpy ml_dtypes # allow-name: the kernel toolchain's package index + echo "$KERNEL_TOOLCHAIN" > "$venv/pinned" + fi + export RUSTUP_HOME="$cache/rustup" CARGO_HOME="$cache/cargo" CARGO_TARGET_DIR="$cache/target" + export PATH="$CARGO_HOME/bin:$PATH" WAI_QUANTUM_TENSOR_ENGINE_PYTHON="$venv/bin/python" WAI_QUANTUM_TENSOR_ENGINE_CACHE="$RUNNER_TEMP/tensor-engine" + if ! command -v rustup >/dev/null; then + curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y --no-modify-path --profile minimal --default-toolchain none + fi + rustup toolchain install + cargo test --locked --features quantum_fabric_tensor_engine --lib -- quantum_fabric_tensor_engine:: quantum_fabric::tensor_engine:: 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; } + key="$WAI_QUANTUM_TENSOR_ENGINE_EXPECT" + report="$RUNNER_TEMP/tensor-engine-report.txt"; : > "$report" + WAI_QUANTUM_TENSOR_ENGINE_REPORT="$report" cargo test --locked --release --features quantum_fabric_tensor_engine --test tensor_engine_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" + grep -qx "tensor-engine devices: as declared ($key)" "$report" || { echo "::error::the target found is not the declared one"; exit 1; } + grep -q "^tensor-engine member: key=$key .*; simulator:" "$report" || { echo "::error::the member does not say it ran in a simulator"; exit 1; } + n="$(grep -c "^tensor-engine certified: key=$key .* parity=bit-exact " "$report")" + [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; } + grep -qx "tensor-engine mutants: key=$key killed 4 of 4" "$report" || { echo "::error::a mutant survived"; exit 1; } + + tensor-engine-absent: + name: member · tensor engine absent (linux-arm64) + needs: relevant + if: needs.relevant.outputs.run == 'true' + runs-on: linux-arm64 + timeout-minutes: 60 + env: + WAI_QUANTUM_TENSOR_ENGINE_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 kernel toolchain here, and the member says so + shell: bash + working-directory: wai/wai-quantum + run: | + set -euo pipefail + report="$RUNNER_TEMP/tensor-engine-report.txt"; : > "$report" + WAI_QUANTUM_TENSOR_ENGINE_REPORT="$report" cargo test --locked --features quantum_fabric_tensor_engine --test tensor_engine_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 'tensor-engine devices: as declared (none)' "$report" || { echo "::error::the absent case was not established"; exit 1; } diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml index c184371c1..26a026636 100644 --- a/wai/wai-quantum/Cargo.toml +++ b/wai/wai-quantum/Cargo.toml @@ -171,6 +171,8 @@ quantum_fabric_fpga = ["quantum_fabric"] quantum_fabric_npu = ["quantum_fabric", "dep:objc2", "dep:block2", "dep:objc2-foundation", "dep:objc2-core-ml"] # The npu-cores member: the gate's arithmetic on a neural processor's compute core, programmed directly (Linux): the program built by the tile array's open toolchain, the host side compiled against the device runtime, both at run time. quantum_fabric_npu_cores = ["quantum_fabric"] +# The tensor-engine member: the gate's products as exact bf16 limb convolutions on a tensor processor's matrix engine, through its kernel toolchain's Python run as a helper. +quantum_fabric_tensor_engine = ["quantum_fabric"] # The matrix-extension member: the gate's arithmetic in the processor's streaming vector mode (SME). quantum_fabric_sme = ["quantum_fabric"] # The DSP vector-unit member: the gate's arithmetic in the signal processor's vector unit. @@ -221,6 +223,10 @@ required-features = ["quantum_fabric_hip", "quantum_sv"] name = "npu_cores_member" required-features = ["quantum_fabric_npu_cores"] +[[test]] +name = "tensor_engine_member" +required-features = ["quantum_fabric_tensor_engine"] + [[test]] name = "opencl_member" required-features = ["quantum_fabric_opencl"] diff --git a/wai/wai-quantum/src/lib.rs b/wai/wai-quantum/src/lib.rs index 41f23b688..7dd1c366c 100644 --- a/wai/wai-quantum/src/lib.rs +++ b/wai/wai-quantum/src/lib.rs @@ -10,7 +10,7 @@ //! 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. Eleven opt-in features reach outside the crate: +//! member that ran the work. Twelve opt-in features reach outside the crate: //! `quantum_fabric_webgpu` (the platform's GPU interface), //! `quantum_fabric_cuda` (the CUDA driver and runtime compiler), //! `quantum_fabric_hip` (the HIP runtime and runtime compiler) and @@ -19,7 +19,9 @@ //! used), `quantum_fabric_npu` (the platform's model runtime, on macOS), //! `quantum_fabric_npu_cores` (a neural processor's compute cores, through //! the device runtime and the tile array's open toolchain, both used when -//! first needed), +//! first needed), `quantum_fabric_tensor_engine` (a tensor processor's +//! matrix engine, through its kernel toolchain's Python, run as a helper when +//! first used), //! `quantum_fabric_sme` (the processor's streaming vector mode), //! `quantum_fabric_vector` (the processor's vector extension), //! `quantum_fabric_dsp` (the signal processor's vector unit, in a build for @@ -104,6 +106,8 @@ pub mod quantum_fabric_fpga; pub mod quantum_fabric_npu; #[cfg(all(feature = "quantum_fabric_npu_cores", not(target_arch = "wasm32")))] pub mod quantum_fabric_npu_cores; +#[cfg(all(feature = "quantum_fabric_tensor_engine", not(target_arch = "wasm32")))] +pub mod quantum_fabric_tensor_engine; #[cfg(all(feature = "quantum_fabric_sme", not(target_arch = "wasm32")))] pub mod quantum_fabric_sme; #[cfg(all(feature = "quantum_fabric_dsp", not(target_arch = "wasm32")))] diff --git a/wai/wai-quantum/src/quantum_fabric.rs b/wai/wai-quantum/src/quantum_fabric.rs index 55d2e130a..eb9022544 100644 --- a/wai/wai-quantum/src/quantum_fabric.rs +++ b/wai/wai-quantum/src/quantum_fabric.rs @@ -2062,6 +2062,7 @@ pub mod tensix; pub mod vector_engine; pub mod sampling; pub mod stablehlo; +pub mod tensor_engine; #[cfg(feature = "quantum_sv")] pub mod sv64; pub mod vqpu; diff --git a/wai/wai-quantum/src/quantum_fabric/tensor_engine.rs b/wai/wai-quantum/src/quantum_fabric/tensor_engine.rs new file mode 100644 index 000000000..53379bc4c --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/tensor_engine.rs @@ -0,0 +1,380 @@ +//! What names the tensor-engine member, and the arithmetic it relies on: the +//! gate's products as exact limb convolutions on a tensor processor's matrix +//! engine, which multiplies bf16 values and accumulates in fp32. No device is +//! touched here, so all of it is testable on any host; the member itself is +//! `quantum_fabric_tensor_engine`, and the kernel is [`HELPER`]. +//! +//! **Why it is bit-exact on any such engine.** Each 32-bit value is split into +//! four limbs of [`Variant::bits`] (8) bits, the lower three in `[0, 255]` and +//! the top one signed in `[-128, 127]`; bf16 holds every one of them exactly. +//! A product `e·a` is `Σ c_s·2^(8s)` with `c_s = Σ_{i+j=s} e_i·a_j`, seven +//! coefficients of at most four limb products each, so every product the +//! engine forms is an integer at most `2^16` in size and every coefficient +//! one below `2^18`: fp32 holds each partial sum exactly, in whatever order the +//! engine adds them. The weights put each of a gate's sixteen products per pair +//! (four terms in each of four rows) in its own seven columns, so one matrix +//! multiply forms them all. The host checks every coefficient is an integer +//! below `2^24`, assembles each product exactly in 64 bits, and applies the +//! fixed-point arithmetic of [`lanes::apply_words`](super::lanes::apply_words): +//! `fx = (product + 2^29) >> 30`, each row the sum of four, the range checked +//! and the low 32 bits kept. +//! +//! So the engine does every multiplication and the host none; the host does +//! the carries, the rounding and the sums, as the neural-processor member's +//! host does. + +/// The kernel and its driver, run by the member with the matrix engine's +/// toolchain's Python. +pub const HELPER: &str = include_str!("tensor_engine_helper.py"); + +/// The implementation name the id starts with. +pub const IMPLEMENTATION: &str = "tensor-engine-limbs8"; + +/// Pairs per matrix multiply: the moving matrix's columns. +pub const TILE: usize = 512; + +/// The contraction: four components of a pair, four limbs each. +pub const K: usize = 16; + +/// The outputs per pair: sixteen products, seven coefficients each. +pub const PRODUCTS: usize = 16; +pub const COEFFICIENTS: usize = 7; +pub const M: usize = PRODUCTS * COEFFICIENTS; + +/// The largest coefficient magnitude fp32 holds exactly. +pub const EXACT: f32 = 16_777_216.0; + +/// The sixteen products, as `(entry, component)`: rows 0r, 0i, 1r and 1i, +/// four terms each, in the order [`lanes::exact_row`](super::lanes) adds them. +pub const TERMS: [(usize, usize); PRODUCTS] = + [(0, 0), (1, 1), (2, 2), (3, 3), (0, 1), (1, 0), (2, 3), (3, 2), (4, 0), (5, 1), (6, 2), (7, 3), (4, 1), (5, 0), (6, 3), (7, 2)]; + +/// Version of the host layout: the weights' and the limbs' layout, the +/// assembly and the host's range check. Part of the implementation id. +pub const HOST_LAYOUT: u32 = 1; + +/// The kernels a member runs: the conforming one, or a mutant. +#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] +pub enum Variant { + Conforming, + /// The rounding constant is half what it should be. + M1, + /// The product is shifted by one bit too few. + M2, + /// Row 0's real part adds its last term instead of subtracting it. + M3, + /// Limbs of nine bits, which bf16 cannot hold: the engine rounds them, + /// and the products are wrong. The exactness argument rests on the limb + /// size, and this mutant shows the engine is where it would fail. + Radix9, +} + +impl Variant { + pub const MUTANTS: [Variant; 4] = [Variant::M1, Variant::M2, Variant::M3, Variant::Radix9]; + + /// The rounding constant, the shift, and whether row 0's last real term + /// is added. + pub fn constants(self) -> (i64, u32, bool) { + match self { + Variant::Conforming | Variant::Radix9 => (1 << 29, 30, false), + Variant::M1 => (1 << 28, 30, false), + Variant::M2 => (1 << 29, 29, false), + Variant::M3 => (1 << 29, 30, true), + } + } + + /// Bits per limb. + pub fn bits(self) -> u32 { + if self == Variant::Radix9 { 9 } else { 8 } + } + + pub fn label(self) -> &'static str { + match self { + Variant::Conforming => "conforming", + Variant::M1 => "m1", + Variant::M2 => "m2", + Variant::M3 => "m3", + Variant::Radix9 => "radix9", + } + } +} + +fn hex(bytes: &[u8]) -> String { + bytes.iter().map(|b| format!("{b:02x}")).collect() +} + +/// `"tensor-engine-limbs8/<16 hex>"` over the kernel's text, the variant's +/// constants and limb size, the tile and the host layout. +pub fn implementation_id_of(v: Variant) -> String { + let (rnd, sh, last_add) = v.constants(); + let mut o = b"wai:quantum-fabric-tensor-engine\x01".to_vec(); + o.extend_from_slice(&(HELPER.len() as u64).to_be_bytes()); + o.extend_from_slice(HELPER.as_bytes()); + o.extend_from_slice(&rnd.to_be_bytes()); + o.extend_from_slice(&sh.to_be_bytes()); + o.push(u8::from(last_add)); + o.extend_from_slice(&v.bits().to_be_bytes()); + o.extend_from_slice(&(TILE as u64).to_be_bytes()); + o.extend_from_slice(&HOST_LAYOUT.to_be_bytes()); + format!("{IMPLEMENTATION}/{}", hex(&blake3::hash(&o).as_bytes()[..8])) +} + +pub fn implementation_id() -> String { + implementation_id_of(Variant::Conforming) +} + +/// `x`'s four limbs of `bits` bits: the lower three unsigned, the top one the +/// signed rest, so that `x = Σ limb_i·2^(bits·i)`. +pub fn limbs(x: i32, bits: u32) -> [i32; 4] { + let mask = (1i64 << bits) - 1; + let v = i64::from(x); + [(v & mask) as i32, ((v >> bits) & mask) as i32, ((v >> (2 * bits)) & mask) as i32, (v >> (3 * bits)) as i32] +} + +/// `v` as bf16 bits, rounded to nearest with ties to even, as a conversion to +/// the engine's input type rounds. Exact for every integer of at most 256 in +/// size. +pub fn bf16(v: i32) -> u16 { + let bits = (v as f32).to_bits(); + let round = 0x7fff + ((bits >> 16) & 1); + ((bits + round) >> 16) as u16 +} + +/// The value bf16 bits stand for. +pub fn from_bf16(b: u16) -> f32 { + f32::from_bits(u32::from(b) << 16) +} + +/// The stationary matrix for gate entries `m`: `K × M` bf16 bits, row-major. +/// Row `c·4 + j` is limb `j` of component `c`; column `p·7 + s` coefficient +/// `s` of product `p`, which takes limb `s − j` of its entry where it exists. +pub fn weights(m: &[i32; 8], bits: u32) -> Vec { + let mut w = vec![0u16; K * M]; + let el: Vec<[i32; 4]> = m.iter().map(|&e| limbs(e, bits)).collect(); + for (p, &(k, c)) in TERMS.iter().enumerate() { + for j in 0..4 { + for (i, &l) in el[k].iter().enumerate() { + w[(c * 4 + j) * M + p * COEFFICIENTS + i + j] = bf16(l); + } + } + } + w +} + +/// Pairs per probe request: one per product slot. +pub const PROBE_PAIRS: usize = PRODUCTS; + +/// The stationary matrix for the arithmetic probe: product `p`'s weights are +/// the limbs of `bs[p]`, taken against component `p % 4`. With the probe's +/// moving matrix ([`probe_moving`]), product `q` of row `q / 4` is +/// `a_q · bs[q]`: sixteen independent pairs in one multiply. +pub fn probe_weights(bs: &[i32; PROBE_PAIRS], bits: u32) -> Vec { + let mut w = vec![0u16; K * M]; + for (p, &b) in bs.iter().enumerate() { + let c = p % 4; + for j in 0..4 { + for (i, &l) in limbs(b, bits).iter().enumerate() { + w[(c * 4 + j) * M + p * COEFFICIENTS + i + j] = bf16(l); + } + } + } + w +} + +/// The probe's cases: pair `q`'s `a` at row `q / 4`, component `q % 4`. +pub fn probe_moving(a: &[i32; PROBE_PAIRS]) -> Vec<[i32; 4]> { + (0..PROBE_PAIRS / 4).map(|r| [a[4 * r], a[4 * r + 1], a[4 * r + 2], a[4 * r + 3]]).collect() +} + +/// The moving matrix for `cases` (each a pair's four components): `K × n` +/// bf16 bits, row-major, `n` the cases rounded up to a whole number of +/// [`TILE`]s, the rest zeros. +pub fn moving(cases: &[[i32; 4]], bits: u32) -> (Vec, usize) { + let n = cases.len().div_ceil(TILE).max(1) * TILE; + let mut a = vec![0u16; K * n]; + for (q, case) in cases.iter().enumerate() { + for (c, &x) in case.iter().enumerate() { + for (j, &l) in limbs(x, bits).iter().enumerate() { + a[(c * 4 + j) * n + q] = bf16(l); + } + } + } + (a, n) +} + +/// Why the engine's coefficients are not a product's. +#[derive(Debug, PartialEq, Eq)] +pub struct Inexact { + pub product: usize, + pub case: usize, +} + +/// The sixteen products of each of the first `count` cases, from the engine's +/// coefficients (`M × n` f32, row-major): each coefficient must be an integer +/// below [`EXACT`] in size, and each product is `Σ c_s·2^(bits·s)`, exact in +/// 64 bits. +pub fn assemble(coefficients: &[f32], n: usize, count: usize, bits: u32) -> Result, Inexact> { + assert_eq!(coefficients.len(), M * n); + let mut out = Vec::with_capacity(count); + for q in 0..count { + let mut ps = [0i64; PRODUCTS]; + for (p, slot) in ps.iter_mut().enumerate() { + let mut v = 0i64; + for s in 0..COEFFICIENTS { + let c = coefficients[(p * COEFFICIENTS + s) * n + q]; + if c.is_nan() || c.abs() >= EXACT || c.fract() != 0.0 { + return Err(Inexact { product: p, case: q }); + } + v += (c as i64) << (bits * s as u32); + } + *slot = v; + } + out.push(ps); + } + Ok(out) +} + +/// A pair's four rows from its sixteen products, with the variant's +/// constants. +pub fn rows(p: &[i64; PRODUCTS], rnd: i64, sh: u32, last_add: bool) -> [i64; 4] { + let fx = |v: i64| (v + rnd) >> sh; + let last = if last_add { fx(p[3]) } else { -fx(p[3]) }; + [ + fx(p[0]) - fx(p[1]) + fx(p[2]) + last, + fx(p[4]) + fx(p[5]) + fx(p[6]) + fx(p[7]), + fx(p[8]) - fx(p[9]) + fx(p[10]) - fx(p[11]), + fx(p[12]) + fx(p[13]) + fx(p[14]) + fx(p[15]), + ] +} + +/// What a tensor-engine member is, as its id hashes it. +#[derive(Clone, Debug, PartialEq, Eq)] +#[non_exhaustive] +pub struct DeviceIdentity { + /// Where the kernel runs, as the driver reports it: `"simulator"` or a + /// device's name. + pub target: String, + /// The driver's report of its packages and versions. + pub versions: String, +} + +impl DeviceIdentity { + pub fn new(target: String, versions: String) -> DeviceIdentity { + DeviceIdentity { target, versions } + } + + /// `"local/tensor-engine-"` over the target and the versions. + pub fn member_id(&self) -> String { + let mut o = b"wai:quantum-fabric-tensor-engine-member\x01".to_vec(); + o.extend_from_slice(&(self.target.len() as u64).to_be_bytes()); + o.extend_from_slice(self.target.as_bytes()); + o.extend_from_slice(&(self.versions.len() as u64).to_be_bytes()); + o.extend_from_slice(self.versions.as_bytes()); + format!("local/tensor-engine-{}", hex(&blake3::hash(&o).as_bytes()[..6])) + } + + /// The key a runner declares: `"tensor-engine-"`. + pub fn key(&self) -> String { + format!("tensor-engine-{}", self.target) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + /// The engine's arithmetic, as it is: bf16 inputs, products and sums in + /// f32, in the given order. + fn engine(w: &[u16], a: &[u16], n: usize) -> Vec { + let mut c = vec![0f32; M * n]; + for col in 0..M { + for q in 0..n { + let mut acc = 0f32; + for k in 0..K { + acc += from_bf16(w[k * M + col]) * from_bf16(a[k * n + q]); + } + c[col * n + q] = acc; + } + } + c + } + + #[test] + fn limbs_rebuild_every_value_and_bf16_holds_them() { + for x in [0, 1, -1, i32::MAX, i32::MIN, 0x1234_5678, -0x0765_4321, 1 << 29, -(1 << 29) - 1] { + for bits in [8, 9] { + let l = limbs(x, bits); + let back: i64 = l.iter().enumerate().map(|(i, &v)| i64::from(v) << (bits * i as u32)).sum(); + assert_eq!(back, i64::from(x), "{x} in {bits}-bit limbs"); + } + for v in limbs(x, 8) { + assert_eq!(from_bf16(bf16(v)), v as f32, "{v} in bf16"); + } + } + // Nine-bit limbs are past what bf16 holds. + assert_ne!(from_bf16(bf16(257)), 257.0); + assert_eq!(from_bf16(bf16(256)), 256.0); + } + + #[test] + fn the_engines_arithmetic_gives_every_product_exactly() { + let mut s = 0x7e57u64; + 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)) as u32 as i32 + }; + let m: [i32; 8] = [i32::MAX, i32::MIN, 1 << 30, -(1 << 30), 0x1234_5678, next(), next(), -1]; + let mut cases: Vec<[i32; 4]> = (0..300).map(|_| [next(), next(), next(), next()]).collect(); + cases[0] = [i32::MIN, i32::MAX, 0, -1]; + let (a, n) = moving(&cases, 8); + assert_eq!(n, TILE); + let c = engine(&weights(&m, 8), &a, n); + let ps = assemble(&c, n, cases.len(), 8).expect("exact"); + for (q, case) in cases.iter().enumerate() { + for (p, &(k, comp)) in TERMS.iter().enumerate() { + assert_eq!(ps[q][p], i64::from(m[k]) * i64::from(case[comp]), "case {q} product {p}"); + } + } + // The probe's layout: sixteen independent pairs in one multiply. + let bs: [i32; PROBE_PAIRS] = std::array::from_fn(|k| if k == 0 { i32::MIN } else { next() }); + let az: [i32; PROBE_PAIRS] = std::array::from_fn(|k| if k == 0 { i32::MIN } else { next() }); + let (a, n) = moving(&probe_moving(&az), 8); + let ps = assemble(&engine(&probe_weights(&bs, 8), &a, n), n, PROBE_PAIRS / 4, 8).unwrap(); + for q in 0..PROBE_PAIRS { + assert_eq!(ps[q / 4][q], i64::from(az[q]) * i64::from(bs[q]), "probe pair {q}"); + } + // The rows are the gate's. + let one = 1i32 << 30; + let id = [one, 0, 0, 0, 0, 0, one, 0]; + let (a, n) = moving(&[[5, -6, 7, -8]], 8); + let ps = assemble(&engine(&weights(&id, 8), &a, n), n, 1, 8).unwrap(); + assert_eq!(rows(&ps[0], 1 << 29, 30, false), [5, -6, 7, -8]); + } + + #[test] + fn nine_bit_limbs_come_back_wrong_from_the_engine() { + let m = [0x1234_5ff7, 0, 0, 0, 0, 0, 0, 0]; + let cases = [[0x7fff_ffff, 0, 0, 0]]; + let (a, n) = moving(&cases, 9); + let got = match assemble(&engine(&weights(&m, 9), &a, n), n, 1, 9) { + Ok(ps) => ps[0][0], + Err(_) => return, + }; + assert_ne!(got, i64::from(m[0]) * i64::from(cases[0][0])); + } + + #[test] + fn ids_follow_the_variant_and_the_identity() { + let ids: std::collections::BTreeSet = [Variant::Conforming, Variant::M1, Variant::M2, Variant::M3, Variant::Radix9].into_iter().map(implementation_id_of).collect(); + assert_eq!(ids.len(), 5); + let a = DeviceIdentity::new("simulator".into(), "x 1".into()); + assert_eq!(a.key(), "tensor-engine-simulator"); + assert!(a.member_id().starts_with("local/tensor-engine-") && a.member_id().as_str() > "local/cpu-x86_64"); + assert_ne!(a.member_id(), DeviceIdentity::new("simulator".into(), "x 2".into()).member_id()); + assert_ne!(a.member_id(), DeviceIdentity::new("device".into(), "x 1".into()).member_id()); + assert!(HELPER.contains("nc_matmul") && HELPER.contains("wq-tensor-engine 1")); + } +} diff --git a/wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py b/wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py new file mode 100644 index 000000000..da16ec1d1 --- /dev/null +++ b/wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py @@ -0,0 +1,82 @@ +"""The device side of wai-quantum's tensor-engine member: one kernel on a +tensor processor's matrix engine, driven over a pipe. + +The kernel is one bf16 matrix multiply per tile of TILE pairs, accumulated in +fp32: the gate's limb weights (16 x 112, the stationary matrix) times the +pairs' limbs (16 x TILE, the moving matrix). Each output is one limb +coefficient of one of the gate's sixteen products per pair; the member +assembles the products and does everything else. Every input is an integer +bf16 holds exactly and every output an integer below 2^24, so fp32 holds each +partial sum exactly in any order (see quantum_fabric::tensor_engine). + + python -I tensor_engine_helper.py + +Protocol (little-endian). On start, one line: "wq-tensor-engine 1 + ...". Then requests until end of input: u32 0x45545157, +u32 n (a multiple of TILE), the weights as 16*112 u16 (bf16 bits, row-major), +the pairs' limbs as 16*n u16 (row-major); each answered by u32 0 and the +coefficients as 112*n f32 (row-major), or by u32 1, u32 length and a message. +""" + +import struct +import sys + +TILE = 512 +K = 16 +M = 112 +MAGIC = 0x45545157 + + +def main(): + import numpy as np + import ml_dtypes + import nki + import nki.isa as nisa + import nki.language as nl + + @nki.jit + def coefficients(w, a): + tiles = a.shape[1] // TILE + out = nl.ndarray((M, a.shape[1]), dtype=nl.float32, buffer=nl.shared_hbm) + ws = nl.load(w) + for t in nl.affine_range(tiles): + xs = nl.load(a[:, t * TILE:(t + 1) * TILE]) + acc = nl.ndarray((M, TILE), dtype=nl.float32, buffer=nl.psum) + nisa.nc_matmul(acc, ws, xs) + nl.store(out[:, t * TILE:(t + 1) * TILE], nl.copy(acc)) + return out + + inp, out = sys.stdin.buffer, sys.stdout.buffer + version = getattr(nki, "__version__", "unknown") + out.write(f"wq-tensor-engine 1 simulator nki {version} numpy {np.__version__} ml_dtypes {ml_dtypes.__version__}\n".encode()) + out.flush() + + def read(n): + b = inp.read(n) + if len(b) != n: + raise EOFError + return b + + while True: + try: + magic, n = struct.unpack("`), and +//! placement ranks it after silicon. The same kernel and the same argument +//! are what a device run would use. +//! - **The engine multiplies; the host assembles**, as the neural-processor +//! member's host does. +//! - **Unmetered**, as every member is by default. + +use crate::quantum::{Amp, Circuit, Gateset, MAX_QUBITS, ONE, QuantumError, StateVector}; +use crate::quantum_fabric::lanes::{self, GateWords, Lane}; +use crate::quantum_fabric::plan::{self, Config, LaneRun, Program, Step}; +use crate::quantum_fabric::tensor_engine::{self, DeviceIdentity, Variant}; +use crate::quantum_fabric::{Certificate, Certification, Corpus, FabricClass, KernelEntry, Member, MemberError, Presence, Provenance, RunFuture, RunMeter, Runner, apply_kernel}; +use std::io::{BufRead, BufReader, Read, Write}; +use std::path::PathBuf; +use std::process::{Child, ChildStdin, ChildStdout, Command, Stdio}; +use std::sync::{Arc, Mutex}; + +/// One probe case: [`lanes::mul_wide`] and [`lanes::fxmul_wide`] as +/// `(high, low)` words. +pub type ProbeOutput = ((u32, u32), (u32, u32)); + +/// Why the engine is not a member, or a run has no state. +#[derive(Debug, PartialEq, Eq)] +#[non_exhaustive] +pub enum TensorEngineError { + /// No Python named, or it has no kernel toolchain, and why. + Unavailable(String), + /// The helper failed or said why it could not run a request. + Device(String), + /// A coefficient came back that is not an exact integer below 2^24. + Inexact { product: usize, case: usize }, + /// Exactly what the CPU reference refuses. + Circuit(QuantumError), + /// More qubits than this member 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: a defect of the host, never a result. + Internal(&'static str), + /// The engine's products disagreed with the CPU's on this pair. + Probe { a: i32, b: i32 }, +} + +impl std::fmt::Display for TensorEngineError { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + match self { + TensorEngineError::Unavailable(why) => write!(f, "no matrix engine to drive: {why}"), + TensorEngineError::Device(why) => write!(f, "the engine's driver: {why}"), + TensorEngineError::Inexact { product, case } => write!(f, "coefficient of product {product} for case {case} is not an exact integer below 2^24"), + TensorEngineError::Circuit(e) => write!(f, "{e}"), + TensorEngineError::TooWide { n_qubits, max } => write!(f, "{n_qubits} qubits; this member holds at most {max}"), + TensorEngineError::OutOfRange { gate } => write!(f, "the result of op {gate} leaves 32 bits"), + TensorEngineError::Internal(why) => write!(f, "host defect: {why}"), + TensorEngineError::Probe { a, b } => write!(f, "the engine's products differ from the CPU's at ({a}, {b})"), + } + } +} + +impl std::error::Error for TensorEngineError {} + +impl From for MemberError { + fn from(e: TensorEngineError) -> MemberError { + match e { + TensorEngineError::TooWide { n_qubits, max } => MemberError::TooWide { n_qubits, max }, + TensorEngineError::OutOfRange { gate } => MemberError::OutOfRange { gate }, + TensorEngineError::Device(why) => MemberError::Device(why), + TensorEngineError::Circuit(e) => MemberError::Circuit(e), + other => MemberError::Other(other.to_string()), + } + } +} + +/// The request tag the helper checks. +const MAGIC: u32 = 0x4554_5157; + +/// The helper process: its pipes, used one request at a time. +struct Helper { + child: Child, + stdin: ChildStdin, + stdout: BufReader, +} + +impl Drop for Helper { + fn drop(&mut self) { + let _ = self.child.kill(); + let _ = self.child.wait(); + } +} + +/// The helper, started with `python`, and its first line: the target and the +/// versions it reports. +fn start(python: &PathBuf) -> Result<(Helper, String, String), String> { + let dir = std::env::var_os("WAI_QUANTUM_TENSOR_ENGINE_CACHE").map(PathBuf::from).unwrap_or_else(|| std::env::temp_dir().join("wai-quantum-tensor-engine")); + std::fs::create_dir_all(&dir).map_err(|e| format!("{}: {e}", dir.display()))?; + let hash: String = blake3::hash(tensor_engine::HELPER.as_bytes()).as_bytes()[..6].iter().map(|b| format!("{b:02x}")).collect(); + let script = dir.join(format!("helper-{hash}.py")); + if !script.exists() { + let staging = dir.join(format!("helper-{hash}.py.{}", std::process::id())); + std::fs::write(&staging, tensor_engine::HELPER).map_err(|e| format!("{}: {e}", staging.display()))?; + std::fs::rename(&staging, &script).map_err(|e| format!("{}: {e}", script.display()))?; + } + let mut child = Command::new(python) + .arg("-I") + .arg(&script) + .stdin(Stdio::piped()) + .stdout(Stdio::piped()) + .stderr(Stdio::null()) + .spawn() + .map_err(|e| format!("{}: {e}", python.display()))?; + let stdin = child.stdin.take().ok_or("no pipe to the helper")?; + let mut stdout = BufReader::new(child.stdout.take().ok_or("no pipe from the helper")?); + let mut line = String::new(); + stdout.read_line(&mut line).map_err(|e| format!("the helper's first line: {e}"))?; + let mut words = line.split_whitespace(); + if (words.next(), words.next()) != (Some("wq-tensor-engine"), Some("1")) { + let _ = child.kill(); + return Err(format!("{} did not start the kernel's driver (it said {:?}); does it have the kernel toolchain?", python.display(), line.trim())); + } + let target = words.next().unwrap_or("").to_string(); + let versions = words.collect::>().join(" "); + if target.is_empty() || !target.bytes().all(|b| b.is_ascii_alphanumeric() || b == b'-') { + let _ = child.kill(); + return Err(format!("the helper named no target ({:?})", line.trim())); + } + Ok((Helper { child, stdin, stdout }, target, versions)) +} + +impl Helper { + /// One request: the coefficients of `M × n` outputs. + fn coefficients(&mut self, weights: &[u16], moving: &[u16], n: usize) -> Result, TensorEngineError> { + let io = |e: std::io::Error| TensorEngineError::Device(format!("the pipe to the helper: {e}")); + let mut req = Vec::with_capacity(8 + 2 * (weights.len() + moving.len())); + req.extend_from_slice(&MAGIC.to_le_bytes()); + req.extend_from_slice(&(n as u32).to_le_bytes()); + for w in weights.iter().chain(moving) { + req.extend_from_slice(&w.to_le_bytes()); + } + self.stdin.write_all(&req).map_err(io)?; + self.stdin.flush().map_err(io)?; + let mut status = [0u8; 4]; + self.stdout.read_exact(&mut status).map_err(io)?; + if u32::from_le_bytes(status) != 0 { + let mut len = [0u8; 4]; + self.stdout.read_exact(&mut len).map_err(io)?; + let mut msg = vec![0u8; u32::from_le_bytes(len).min(1 << 16) as usize]; + self.stdout.read_exact(&mut msg).map_err(io)?; + return Err(TensorEngineError::Device(String::from_utf8_lossy(&msg).into_owned())); + } + let mut bytes = vec![0u8; tensor_engine::M * n * 4]; + self.stdout.read_exact(&mut bytes).map_err(io)?; + Ok(bytes.as_chunks::<4>().0.iter().map(|b| f32::from_le_bytes(*b)).collect()) + } +} + +/// What [`discover`] found. +#[derive(Debug)] +pub struct Discovery { + pub found: Vec, + pub refused: Vec<(String, TensorEngineError)>, + pub unavailable: Option, +} + +/// The matrix engine the named Python drives, probed in this process; or why +/// there is none. +pub fn discover() -> Discovery { + let mut d = Discovery { found: Vec::new(), refused: Vec::new(), unavailable: None }; + let Some(python) = std::env::var_os("WAI_QUANTUM_TENSOR_ENGINE_PYTHON").filter(|s| !s.is_empty()).map(PathBuf::from) else { + d.unavailable = Some("no kernel toolchain named (WAI_QUANTUM_TENSOR_ENGINE_PYTHON, the Python of an environment with the matrix engine's kernel toolchain)".into()); + return d; + }; + match start(&python) { + Err(why) => d.unavailable = Some(why), + Ok((helper, target, versions)) => { + let identity = DeviceIdentity::new(target, versions); + let key = identity.key(); + match TensorEngineMember::open(helper, identity) { + Ok(m) => d.found.push(m), + Err(e) => d.refused.push((key, e)), + } + } + } + d +} + +/// The most qubits a run holds. +pub const MAX_TENSOR_ENGINE_QUBITS: u8 = if 22 < MAX_QUBITS { 22 } else { MAX_QUBITS }; + +/// A tensor processor's matrix engine as a fabric member. +#[derive(Clone)] +pub struct TensorEngineMember { + helper: Arc>, + variant: Variant, + identity: DeviceIdentity, + provenance: Provenance, + id: String, + key: String, + implementation: String, + detail: String, + max_qubits: u8, +} + +impl std::fmt::Debug for TensorEngineMember { + fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result { + f.debug_struct("TensorEngineMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).finish() + } +} + +impl TensorEngineMember { + fn open(helper: Helper, identity: DeviceIdentity) -> Result { + // Anything but a device the helper names is a simulation of the engine. + let provenance = if identity.target == "simulator" { Provenance::Simulator(identity.versions.clone()) } else { Provenance::Silicon }; + let member = TensorEngineMember { + helper: Arc::new(Mutex::new(helper)), + variant: Variant::Conforming, + provenance: provenance.clone(), + id: identity.member_id(), + key: identity.key(), + implementation: tensor_engine::implementation_id(), + detail: format!( + "tensor; the matrix engine, bf16 limb products accumulated in fp32, {} pairs a multiply; target {}; {}; up to {MAX_TENSOR_ENGINE_QUBITS} qubits; {}; unmetered", + tensor_engine::TILE, + identity.target, + identity.versions, + provenance.label() + ), + identity, + max_qubits: MAX_TENSOR_ENGINE_QUBITS, + }; + let pairs = lanes::probe_pairs(4096, 0x5eed); + let out = member.probe(&pairs)?; + if let Some((a, b)) = lanes::first_mismatch(&pairs, &out) { + return Err(TensorEngineError::Probe { a, b }); + } + Ok(member) + } + + /// The same engine running a mutant: for the mutation tests. + #[doc(hidden)] + pub fn with_variant(&self, v: Variant) -> TensorEngineMember { + TensorEngineMember { variant: v, implementation: tensor_engine::implementation_id_of(v), ..self.clone() } + } + + #[doc(hidden)] + pub fn with_max_qubits(&self, max_qubits: u8) -> TensorEngineMember { + TensorEngineMember { max_qubits: max_qubits.min(self.max_qubits), ..self.clone() } + } + + pub fn id(&self) -> &str { + &self.id + } + + pub fn key(&self) -> &str { + &self.key + } + + pub fn class(&self) -> FabricClass { + FabricClass::Tensor + } + + pub fn detail(&self) -> &str { + &self.detail + } + + pub fn implementation(&self) -> &str { + &self.implementation + } + + pub fn variant(&self) -> Variant { + self.variant + } + + pub fn identity(&self) -> &DeviceIdentity { + &self.identity + } + + pub fn provenance(&self) -> &Provenance { + &self.provenance + } + + pub fn max_qubits(&self) -> u8 { + self.max_qubits + } + + pub fn config(&self) -> Config { + Config::PER_GATE + } + + /// The engine's sixteen products of each case, for stationary weights `w`. + fn run(&self, w: &[u16], cases: &[[i32; 4]]) -> Result, TensorEngineError> { + let bits = self.variant.bits(); + let (a, n) = tensor_engine::moving(cases, bits); + let c = self.helper.lock().unwrap_or_else(|p| p.into_inner()).coefficients(w, &a, n)?; + tensor_engine::assemble(&c, n, cases.len(), bits).map_err(|e| TensorEngineError::Inexact { product: e.product, case: e.case }) + } + + /// Each case's sixteen products under gate entries `m`, from the engine. + fn products(&self, m: &[i32; 8], cases: &[[i32; 4]]) -> Result, TensorEngineError> { + self.run(&tensor_engine::weights(m, self.variant.bits()), cases) + } + + /// The engine's [`lanes::mul_wide`] of each pair, and the host's + /// [`lanes::fxmul_wide`] of that product: sixteen pairs a multiply + /// ([`tensor_engine::probe_weights`]). + pub fn probe(&self, pairs: &[(i32, i32)]) -> Result, TensorEngineError> { + let (rnd, sh, _) = self.variant.constants(); + let split = |v: i64| ((v as u64 >> 32) as u32, v as u64 as u32); + let mut out = Vec::with_capacity(pairs.len()); + for chunk in pairs.chunks(tensor_engine::PROBE_PAIRS) { + let a: [i32; tensor_engine::PROBE_PAIRS] = std::array::from_fn(|q| chunk.get(q).map_or(0, |p| p.0)); + let b: [i32; tensor_engine::PROBE_PAIRS] = std::array::from_fn(|q| chunk.get(q).map_or(0, |p| p.1)); + let ps = self.run(&tensor_engine::probe_weights(&b, self.variant.bits()), &tensor_engine::probe_moving(&a))?; + for q in 0..chunk.len() { + let v = ps[q / 4][q]; + out.push((split(v), split((v + rnd) >> sh))); + } + } + Ok(out) + } + + /// Gates in order from `state`, as [`lanes::run_words`] runs them. + pub fn run_lanes(&self, state: &[Lane], gates: &[GateWords]) -> Result { + let len = state.len(); + if len < 2 || !len.is_power_of_two() { + return Err(TensorEngineError::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(TensorEngineError::TooWide { n_qubits: n, max: self.max_qubits }); + } + let planned: Vec<(&GateWords, u32)> = gates.iter().zip(0u32..).collect(); + self.execute(state.to_vec(), &planned) + } + + fn execute(&self, mut amps: Vec, gates: &[(&GateWords, u32)]) -> Result { + let len = amps.len(); + 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(TensorEngineError::Internal("malformed gate words for this state")); + } + } + let (rnd, sh, last_add) = self.variant.constants(); + let mut report: Option = None; + for &(w, index) in gates { + let m: [i32; 8] = w.m.map(|x| x as i32); + let (tbit, ctrl) = (w.tbit as usize, w.ctrl as usize); + let mut at = Vec::new(); + let mut cases = Vec::new(); + for p in 0..len / 2 { + let low = p & (tbit - 1); + let i = ((p - low) << 1) | low; + if i & ctrl != ctrl { + continue; + } + at.push(i); + let (a0, a1) = (amps[i], amps[i | tbit]); + cases.push([a0[0], a0[1], a1[0], a1[1]]); + } + let products = self.products(&m, &cases)?; + for (&i, p) in at.iter().zip(&products) { + let r = tensor_engine::rows(p, rnd, sh, last_add); + if r.iter().any(|&v| i32::try_from(v).is_err()) { + report = Some(report.map_or(index, |x| x.min(index))); + } + let low = |v: i64| v as u64 as u32 as i32; + amps[i] = [low(r[0]), low(r[1])]; + amps[i | tbit] = [low(r[2]), low(r[3])]; + } + } + Ok(LaneRun { amps, out_of_range: report }) + } + + pub fn run_program(&self, p: &Program) -> Result { + let n = p.n_qubits; + if n == 0 { + return Ok(LaneRun { amps: vec![[ONE as i32, 0]], out_of_range: None }); + } + if n > self.max_qubits { + return Err(TensorEngineError::TooWide { n_qubits: n, max: self.max_qubits }); + } + if p.start >= 1usize << n { + return Err(TensorEngineError::Internal("the start is outside the state")); + } + let mut gates = Vec::with_capacity(p.steps.len()); + for step in &p.steps { + let Step::Gate { gate } = *step else { return Err(TensorEngineError::Internal("a tile step; this member runs gate steps only")) }; + let g = p.gates.get(gate as usize).ok_or(TensorEngineError::Internal("a step names a gate the program does not hold"))?; + if !g.is_well_formed() { + return Err(TensorEngineError::Internal("malformed gate words for this state")); + } + gates.push((&g.words, g.index)); + } + let mut amps = vec![[0i32, 0]; 1usize << n]; + amps[p.start] = [ONE as i32, 0]; + self.execute(amps, &gates) + } + + pub fn simulate_from(&self, c: &Circuit, start: usize) -> Result { + let p = plan::compile_with(c, start, &self.config()).map_err(TensorEngineError::Circuit)?; + if p.n_qubits > self.max_qubits { + return Err(TensorEngineError::TooWide { n_qubits: p.n_qubits, max: self.max_qubits }); + } + let run = self.run_program(&p)?; + if let Some(gate) = run.out_of_range { + return Err(TensorEngineError::OutOfRange { gate }); + } + Ok(StateVector { n_qubits: p.n_qubits, amps: run.amps.iter().map(|l| Amp { re: i64::from(l[0]), im: i64::from(l[1]) }).collect() }) + } + + pub fn simulate(&self, c: &Circuit) -> Result { + self.simulate_from(c, 0) + } + + /// The certificate of [`apply_kernel`]`(gs)` on this member, with its + /// provenance. + pub fn certify(&self, gs: Gateset, corpus: Corpus) -> Result { + let mut cert = Certification::new(); + for c in corpus.circuits(gs) { + let reference = lanes::reference_output(&c).map_err(TensorEngineError::Circuit)?; + let run = self.run_program(&plan::compile_with(&c, 0, &self.config()).map_err(TensorEngineError::Circuit)?)?; + if let Some(gate) = run.out_of_range { + return Err(TensorEngineError::OutOfRange { gate }); + } + cert.add(&lanes::circuit_bytes(&c), &reference, &lanes::output(&run.amps)); + } + Ok(cert.finish(apply_kernel(gs), &self.implementation, &self.id, &corpus.generator(gs)).with_provenance(self.provenance.clone())) + } + + pub fn member(&self, corpus: Corpus) -> Result { + let mut kernels = Vec::new(); + for gs in Gateset::ALL { + kernels.push(KernelEntry { kernel: apply_kernel(gs).into(), implementation: self.implementation.clone(), certificate: self.certify(gs, corpus)?, energy: None }); + } + Ok(Member { id: self.id.clone(), class: FabricClass::Tensor, presence: Presence::Detected, kernels, detail: self.detail.clone() }) + } +} + +impl Runner for TensorEngineMember { + 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_from(c, start)?, 0, None)) + }) + } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn the_refusals_that_may_fall_back_do() { + assert!(MemberError::from(TensorEngineError::OutOfRange { gate: 3 }).falls_back()); + assert!(!MemberError::from(TensorEngineError::Probe { a: 1, b: 2 }).falls_back()); + assert!(!MemberError::from(TensorEngineError::Inexact { product: 0, case: 0 }).falls_back()); + } + + #[test] + fn a_python_that_is_not_the_driver_is_refused_with_its_words() { + // `true` reads its arguments, prints nothing and exits: no driver line. + let Ok(t) = which("true") else { return }; + let err = start(&t).err().expect("refused"); + assert!(err.contains("did not start the kernel's driver"), "{err}"); + } + + fn which(name: &str) -> Result { + for d in std::env::var_os("PATH").map(|p| std::env::split_paths(&p).collect::>()).unwrap_or_default() { + let c = d.join(name); + if c.is_file() { + return Ok(c); + } + } + Err(()) + } +} diff --git a/wai/wai-quantum/tests/tensor_engine_member.rs b/wai/wai-quantum/tests/tensor_engine_member.rs new file mode 100644 index 000000000..d5c27c7d4 --- /dev/null +++ b/wai/wai-quantum/tests/tensor_engine_member.rs @@ -0,0 +1,414 @@ +//! The fabric's tensor-engine member (a tensor processor's matrix engine) on +//! the machine that runs this. +//! +//! The machine declares the devices it must have in `WAI_QUANTUM_TENSOR_ENGINE_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_TENSOR_ENGINE_REPORT` names a file, appends the +//! same lines to it, which a CI step can count per device. +//! +//! ```text +//! WAI_QUANTUM_TENSOR_ENGINE_EXPECT=tensor-engine-simulator \ +//! WAI_QUANTUM_TENSOR_ENGINE_PYTHON= \ +//! cargo test --release --features quantum_fabric_tensor_engine --test tensor_engine_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::lanes::{self, GateWords, Lane}; +use wai_quantum::quantum_fabric::plan::{Program, Step}; +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::tensor_engine::{Variant, implementation_id_of}; +use wai_quantum::quantum_fabric_tensor_engine::{Discovery, TensorEngineError, TensorEngineMember, 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_TENSOR_ENGINE_REPORT` names. +fn report(line: &str) { + println!("{line}"); + if let Some(path) = std::env::var_os("WAI_QUANTUM_TENSOR_ENGINE_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(|(k, e)| format!("{k}: {e}")).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_TENSOR_ENGINE_EXPECT={found_list}"); + let expect = match std::env::var("WAI_QUANTUM_TENSOR_ENGINE_EXPECT") { + Ok(v) if !v.is_empty() => v, + _ => return Err(format!("WAI_QUANTUM_TENSOR_ENGINE_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_TENSOR_ENGINE_EXPECT={expect:?} is not a sorted, comma-separated list of keys, nor none. {how}")); + } + for k in &keys { + if !found.contains(k) + && let Some((_, e)) = d.refused.iter().find(|(key, _)| key == k) + { + return Err(format!("the declared device {k} was refused: {e}. {how}")); + } + } + 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 [TensorEngineMember] { + 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!("tensor-engine devices: as declared ({expect})")), + Err(why) => panic!("{why}"), + } + for m in &d.found { + assert_eq!(m.class(), FabricClass::Tensor); + assert!(m.id().starts_with("local/tensor-engine-"), "{}", m.id()); + assert_eq!(m.implementation(), implementation_id_of(Variant::Conforming)); + assert_eq!(m.identity().member_id(), m.id()); + assert!(m.max_qubits() >= 20, "{}: holds only {} qubits", m.key(), m.max_qubits()); + report(&format!("tensor-engine 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!("tensor-engine 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!("tensor-engine 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!( + "tensor-engine 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(20); + let mut s = 0x0077_6964_6500_u64; + for (k, n) in [12u8, 16, 18, 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!("tensor-engine 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(), TensorEngineError::Circuit(native), "{c:?}"); + } + assert_eq!(m.simulate(&Circuit::new(MAX_QUBITS + 1)).unwrap_err(), TensorEngineError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1))); + let small = m.with_max_qubits(4); + assert_eq!(small.simulate(&Circuit::new(5)).unwrap_err(), TensorEngineError::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(), TensorEngineError::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(TensorEngineError::Internal(_)))); + assert!(matches!(m.run_lanes(&[[0, 0]; 3], &[z_words()]), Err(TensorEngineError::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(TensorEngineError::Internal(_)))); + // The refusals that may still fall back to another member, do. + assert!(MemberError::from(TensorEngineError::TooWide { n_qubits: 5, max: 4 }).falls_back()); + report(&format!("tensor-engine refusals: key={} typed", m.key())); + } +} + +/// The kernel mutants, compiled in: each must be caught by the probe or by a +/// certificate. +#[test] +fn c6_mutants_are_killed() { + let edges = lanes::probe_pairs(0, 0); + for m in members() { + let mut killed = 0; + for v in Variant::MUTANTS { + let mm = m.with_variant(v); + assert_ne!(mm.implementation(), m.implementation()); + 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(); + let caught = differing > 0 + || Gateset::ALL.into_iter().all(|gs| match mm.certify(gs, RANDOM_V1) { + Err(_) => true, + Ok(c) => c.parity != Parity::BitExact || c.corpus != lanes::certify_corpus(gs, "local/test", RANDOM_V1).corpus, + }); + println!("tensor-engine mutant {v:?} on {}: {differing} of {} edge pairs differ", m.key(), edges.len()); + assert!(caught, "{}: mutant {v:?} survived", m.key()); + killed += 1; + } + report(&format!("tensor-engine mutants: key={} killed {killed} of {}", m.key(), Variant::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::Tensor, 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!("tensor-engine 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!("tensor-engine repeat: key={} ok", m.key())); + } +} -- 2.54.0 (Apple Git-157) From 3346bb2cd696703cba5a5d2d611a10bb74f1c5d9 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 10:00:30 -0400 Subject: [PATCH 14/19] wai-quantum 0.6.0: the tensor-engine member in the CHANGELOG and README Co-Authored-By: Claude Opus 5.5 --- wai/CHANGELOG.md | 15 +++++++++++++-- wai/wai-quantum/README.md | 1 + 2 files changed, 14 insertions(+), 2 deletions(-) diff --git a/wai/CHANGELOG.md b/wai/CHANGELOG.md index 0e33b6276..b4cc07d74 100644 --- a/wai/CHANGELOG.md +++ b/wai/CHANGELOG.md @@ -6,7 +6,7 @@ wire format may change and such changes are recorded here. ## v0.4.0 — current -### wai-quantum 0.6.0: the HIP, npu-cores and avx512 members, the dataflow device's vector engine, and the gate as a portable tensor program +### wai-quantum 0.6.0: the HIP, npu-cores, tensor-engine and avx512 members, the dataflow device's vector engine, and the gate as a portable tensor program Breaking in three places (the last bullet). The default features, `full`, and every pinned digest are unchanged. @@ -27,6 +27,15 @@ digest are unchanged. (`npu_cores_host.cpp`) is compiled against the device runtime and loaded. One program serves every gate and the arithmetic probe; each mutant is its own program. Key `npu-cores--`. +- **The tensor-engine member** (`quantum_fabric_tensor_engine`, opt-in): the gate's sixteen + products per pair as exact limb convolutions in one bf16 matrix multiply per tile of pairs, + accumulated in fp32, on a tensor processor's matrix engine. Each value is four 8-bit limbs, + which bf16 holds exactly; every limb product is at most `2^16` and every coefficient a sum of + at most four, so fp32 holds every partial sum exactly in any order, on any engine that + multiplies bf16 and accumulates in fp32. The host checks each coefficient is an integer + below `2^24`, assembles the products and does the rest. The kernel runs as a helper under + the engine's kernel toolchain; its target here is that toolchain's simulator, and every + certificate says so. One mutant is nine-bit limbs, which the engine's inputs cannot hold. - **The avx512 vector unit** (`quantum_fabric_vector`): eight 64-bit lanes, each `fx` an exact 32×32 multiply, an add and the unit's own 64-bit arithmetic shift. Discovery opens one member per instruction set the processor reports, so a processor with both units has @@ -35,7 +44,9 @@ digest are unchanged. processor beside them, all silicon: the HIP member 9/9 (66020 probe pairs, all four pinned certificates, circuits to 22 qubits, 6 of 6 mutants) and its state vector 3/3 (the pin `b7d129d0…`, 4 of 4 mutants); the npu-cores member 9/9 (66020 probe pairs, all four - certificates, circuits to 20 qubits, 3 of 3 mutants); avx2 and avx512 9/9 each; and, + certificates, circuits to 20 qubits, 3 of 3 mutants); avx2 and avx512 9/9 each; the + tensor-engine member on its simulator 9/9 (66020 probe pairs, all four certificates, 4 of 4 + mutants); and, unchanged, the WebGPU member (`gpu-vulkan-1002`) and the OpenCL member (`opencl-gpu-1002`, the state vector's pin included) bit-exact there. CI runs each of them on that host. diff --git a/wai/wai-quantum/README.md b/wai/wai-quantum/README.md index d9233e101..b5df4ad0e 100644 --- a/wai/wai-quantum/README.md +++ b/wai/wai-quantum/README.md @@ -111,6 +111,7 @@ never creates information. | `quantum_fabric_fpga` | **the fabric's FPGA member, simulated**, opt-in and not in `full`, native targets only: the gate datapath in Verilog (each 32×32 product from four 16-bit partial products, the size of an FPGA's multiplier blocks; a state RAM; one pair per two cycles), compiled at run time by Verilator into a cycle-accurate simulation driven over a pipe. Probed and certified as the other members are, and every certificate says it was earned in a simulator, so placement ranks it after silicon. A host without Verilator finds no member and says why | | `quantum_fabric_npu` | **the fabric's neural-processor member**, opt-in and not in `full`, macOS only: each 32-bit operand split into eight 4-bit limbs, every limb product and limb sum formed in fp16 on the platform's neural processor (each value an integer fp16 holds exactly, whatever order it is added in), the host assembling each product and applying the fixed-point arithmetic. The member exists only where the platform model runtime's compute plan places every operation of the model on the neural processor; a model it would run elsewhere is refused with the placement it reported. Probed and certified as every member is | | `quantum_fabric_npu_cores` | **the fabric's npu-cores member**, opt-in and not in `full`, Linux: the gate's fixed-point arithmetic on one compute core of a neural processor's tile array, programmed directly rather than through a model runtime: exact 64-bit products on the core, the host gathering each gate's pairs and checking each row's range. At run time the member finds the device by its driver, builds the core's program with the tile array's open toolchain (`WAI_QUANTUM_NPU_CORES_PYTHON` names its Python; the program is cached per variant, generation and toolchain) and compiles its host side against the device runtime. Each of the three kernel mutants is its own program and must be caught. Probed and certified as the other members are; its certificates' provenance is silicon | +| `quantum_fabric_tensor_engine` | **the fabric's tensor-engine member**, opt-in and not in `full`: the gate's products as exact limb convolutions on a tensor processor's matrix engine, one bf16 matrix multiply per tile of pairs accumulated in fp32 (four 8-bit limbs per value, which bf16 holds exactly; every partial sum an integer below `2^24`, which fp32 holds exactly in any order), the host assembling the products and applying the fixed-point arithmetic. The kernel runs as a helper under the engine's kernel toolchain (`WAI_QUANTUM_TENSOR_ENGINE_PYTHON`); where its target is that toolchain's simulator, every certificate says so. Four mutants must be caught, nine-bit limbs the engine cannot hold among them | | `quantum_fabric_sme` | **the fabric's matrix-extension member**, opt-in and not in `full`, aarch64 only: the gate's fixed-point arithmetic in the processor's streaming vector mode (SME), 64-bit lanes, the host gathering each gate's pairs and checking each row's range. Every instruction it uses exists only in that mode, so a kernel that ran, ran there. Detected by the processor's feature flags (on Linux, the kernel's `AT_HWCAP2`, which a user-mode emulator supplies); under an emulator the runner says so (`WAI_QUANTUM_SME_EMULATOR`) and every certificate carries it. Three kernel mutants are compiled in and must be caught | | `quantum_fabric_dsp` | **the fabric's DSP vector-unit member**, opt-in and not in `full`, in a build for the signal processor with its vector unit in 128-byte mode: the gate's fixed-point arithmetic in 32-bit lanes, each 64-bit value a pair of words, every product the vector unit's exact 32×32 multiply and every carry and borrow kept; the host gathers each gate's pairs and checks each row's range. Its instructions exist only on that unit, so a kernel that ran, ran there. Under an emulator the runner says so (`WAI_QUANTUM_DSP_EMULATOR`) and every certificate carries it. Three kernel mutants are compiled in and must be caught. Its class is `dsp` | | `quantum_fabric_vector` | **the fabric's CPU vector-unit members**, opt-in and not in `full`: the gate's fixed-point arithmetic in the processor's own vector extension, one kernel per instruction set, and one member per set the processor has — `avx2` (x86-64, 256-bit), `avx512` (x86-64, 512-bit, with the unit's own 64-bit arithmetic shift), `sve` (aarch64, the scalable vector extension, any length) and `rvv` (riscv64, the vector extension 1.0, any length) — each operand sign-extended to 64-bit lanes, the host gathering each gate's pairs and checking each row's range. The scalable kernels set their length per pass; the member reads the length from the unit and names it in its id and detail. Detected by the processor's feature flags (on Linux, `AT_HWCAP` for the vector extension, which a user-mode emulator supplies); under an emulator the runner says so (`WAI_QUANTUM_VECTOR_EMULATOR`). Its id extends the CPU reference's, so placement with no energy figures keeps the reference. Three kernel mutants per instruction set are compiled in and must be caught | -- 2.54.0 (Apple Git-157) From 2a30dcd360e14a321c328d877858f80e56e2911e Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 10:06:50 -0400 Subject: [PATCH 15/19] wai-quantum: the gate's products as a portable matrix-unit program stablehlo::limb_module: the tensor-engine member's limb layout as one StableHLO dot_general of bf16 operands, accumulated and returned in f32, the program a compiler maps to a matrix unit. stablehlo_gate writes it with the batch's limb weights and limbs and the exact coefficients, and a nine-bit-limb data set whose exact coefficients bf16 cannot reproduce (a run that matches them is not comparing). check.sh holds the limb module to "=!" (exact on eight-bit limbs, different on nine); both drivers read typed raw files, and floats are compared exactly (IREE's --expected_f32_threshold=0). Measured: exact under both compilers on the x86-64 and arm64 CPUs and on the discrete GPU through CUDA; the nine-bit limbs differ everywhere. Not exact: the IREE 3.12 HIP build for gfx1151, whose bf16 multiply returns non-integers (-14.0002 for -14), so the workflow shows that step and does not require it. Not built: the IREE Vulkan and Metal builds do not lower bf16 in this version. Co-Authored-By: Claude Opus 5.5 --- .../workflows/wai-quantum-tensor-program.yml | 25 +++++++++ wai/wai-quantum/ci/tensor-program/check.sh | 39 ++++++++----- wai/wai-quantum/ci/tensor-program/xla_run.py | 55 +++++++++++++------ wai/wai-quantum/examples/stablehlo_gate.rs | 35 +++++++++++- .../src/quantum_fabric/stablehlo.rs | 38 +++++++++++++ 5 files changed, 160 insertions(+), 32 deletions(-) diff --git a/.github/workflows/wai-quantum-tensor-program.yml b/.github/workflows/wai-quantum-tensor-program.yml index 8d1e7cc1b..43b544238 100644 --- a/.github/workflows/wai-quantum-tensor-program.yml +++ b/.github/workflows/wai-quantum-tensor-program.yml @@ -22,6 +22,12 @@ name: wai-quantum portable tensor program # arm64 CPU. # Each leg past the first also runs one mutant on each device, so the # comparison there is shown to catch a difference. +# +# A second program, the limb module, is the gate's products as one bf16 +# matrix multiply accumulated in f32 (quantum_fabric::stablehlo::limb_module): +# what a compiler maps to a matrix unit. It must be exact on eight-bit limbs +# and differ on nine-bit ones, wherever the compiler builds it exactly: the +# CPUs under both compilers, and the discrete GPU through CUDA under both. concurrency: group: ${{ github.workflow }}-${{ github.ref }} @@ -129,6 +135,12 @@ jobs: $check iree local-task $m.cpu.vmfb "$want" $check xla cpu $m.mlir "$want" done + # The products on a matrix unit: one bf16 multiply accumulated in f32, + # exact on eight-bit limbs; nine-bit limbs, which bf16 cannot hold, + # must differ (quantum_fabric::stablehlo::limb_module). + iree-compile limbs.mlir --iree-hal-target-device=local --iree-hal-local-target-device-backends=llvm-cpu --iree-llvmcpu-target-cpu=host -o limbs.cpu.vmfb + $check iree local-task limbs.cpu.vmfb '=!' + $check xla cpu limbs.mlir '=!' for t in gfx942 gfx1100 gfx1201; do iree-compile gate.mlir --iree-hal-target-backends=rocm --iree-rocm-target=$t -o gate.$t.vmfb echo "built for $t: $(stat -c %s gate.$t.vmfb) bytes" @@ -189,6 +201,10 @@ jobs: $check iree cuda gate_m3.cuda.vmfb '!=' $check xla cuda gate.mlir '==' $check xla cuda gate_m3.mlir '!=' + # The products on the GPU's matrix path, through both compilers. + iree-compile limbs.mlir --iree-hal-target-device=cuda --iree-cuda-target=sm_89 -o limbs.cuda.vmfb + $check iree cuda limbs.cuda.vmfb '=!' + $check xla cuda limbs.mlir '=!' amd: name: tensor program · integrated GPU (framework-128gb) @@ -243,6 +259,12 @@ jobs: $check iree "$dev" gate.vulkan.vmfb '==' $check iree "$dev" gate_m3.vulkan.vmfb '!=' done + # The products through this compiler's HIP build: shown, not + # required. Its bf16 multiply for this GPU returns values that are + # not integers (measured 2026-10-15), so the coefficients differ; the + # step says what it found, and a build that becomes exact shows here. + iree-compile limbs.mlir --iree-hal-target-backends=rocm --iree-rocm-target="$arch" -o limbs.hip.vmfb + $check iree hip limbs.hip.vmfb '=!' || echo "(the HIP build's limb products are not exact; not required)" metal: name: tensor program · Metal and the arm64 CPU (macos) @@ -282,3 +304,6 @@ jobs: $check iree metal gate_m3.metal.vmfb '!=' $check xla cpu gate.mlir '==' $check xla cpu gate_m3.mlir '!=' + # The products on the arm64 CPU through XLA. (The Metal build of + # this compiler version does not lower bf16.) + $check xla cpu limbs.mlir '=!' diff --git a/wai/wai-quantum/ci/tensor-program/check.sh b/wai/wai-quantum/ci/tensor-program/check.sh index e65843ca9..1af3d7330 100755 --- a/wai/wai-quantum/ci/tensor-program/check.sh +++ b/wai/wai-quantum/ci/tensor-program/check.sh @@ -7,10 +7,12 @@ # check.sh xla # # Run in the directory the stablehlo_gate example wrote. is two -# characters, the rows' outcome then the probe's: `=` exact, `!` differs. The -# conforming program must give `==`; mutants m1 and m2 (the rounding constant, -# the shift) `!!`; m3 (row 0's last sign) `!=`, since its probe is the -# program's. +# characters, `=` exact or `!` differs. For a gate module, the rows' outcome +# then the probe's: the conforming program must give `==`; mutants m1 and m2 +# (the rounding constant, the shift) `!!`; m3 (row 0's last sign) `!=`, since +# its probe is the program's. For the limb module (a name starting `limbs`), +# the coefficients of eight-bit limbs, then of nine-bit limbs, which bf16 +# cannot hold: a conforming build gives `=!`. # # iree: the module was built by that compiler; its runtime (`iree-run-module`, # on PATH) runs it on and compares the outputs, and an exact outcome @@ -25,7 +27,9 @@ set -euo pipefail how=$1 dev=$2 module=$3 expect=$4 case "$expect" in [=!][=!]) ;; *) echo "check.sh: is two of = and !" >&2; exit 2 ;; esac -n=$(( $(stat -c %s r0.bin 2>/dev/null || stat -f %z r0.bin) / 8 )) +size() { stat -c %s "$1" 2>/dev/null || stat -f %z "$1"; } +n=$(( $(size r0.bin) / 8 )) +case "$(basename "$module")" in limbs*) kind=limbs ;; *) kind=gate ;; esac rows_in=(e0.npy e1.npy e2.npy e3.npy e4.npy e5.npy e6.npy e7.npy a0r.npy a0i.npy a1r.npy a1i.npy) rows_out=(r0.bin r1.bin r2.bin r3.bin) probe_in=(pa.npy pb.npy) @@ -42,11 +46,12 @@ run() { local args=() seen=0 a for a in "$@"; do if [ "$a" = -- ]; then seen=1 - elif [ $seen = 0 ]; then args+=("--input=@$a") - else args+=("--expected_output=${n}xi64=@$a") + elif [ $seen = 0 ]; then case "$a" in *=@*) args+=("--input=$a") ;; *) args+=("--input=@$a") ;; esac + else case "$a" in *=@*) args+=("--expected_output=$a") ;; *) args+=("--expected_output=${n}xi64=@$a") ;; esac fi done - iree-run-module --device="$dev" --module="$module" --function="$fn" "${args[@]}" > "$log" 2>&1 || status=$? + # Floats are compared exactly too. + iree-run-module --device="$dev" --module="$module" --function="$fn" --expected_f32_threshold=0 "${args[@]}" > "$log" 2>&1 || status=$? if [ $status = 0 ] && grep -q '^\[SUCCESS\] all function outputs matched' "$log"; then return 0; fi if [ $status != 0 ] && grep -q '^\[FAILED\] result\[[0-9]*\]: .* does not match the expected' "$log"; then return 1; fi else @@ -60,12 +65,20 @@ run() { } tag="$(basename "$module")" -r='='; run rows "$tag.rows.log" "${rows_in[@]}" -- "${rows_out[@]}" || r='!' -p='='; run probe "$tag.probe.log" "${probe_in[@]}" -- "${probe_out[@]}" || p='!' word() { [ "$1" = '=' ] && echo exact || echo differ; } -echo "$tag on $dev ($how): rows $(word $r), probe $(word $p)" +if [ "$kind" = limbs ]; then + ln=$(( $(size c.bin) / 4 / 112 )) + r='='; run coefficients "$tag.limbs8.log" "16x112xbf16=@w.bin" "16x${ln}xbf16=@a.bin" -- "112x${ln}xf32=@c.bin" || r='!' + p='='; run coefficients "$tag.limbs9.log" "16x112xbf16=@w9.bin" "16x${ln}xbf16=@a9.bin" -- "112x${ln}xf32=@c9.bin" || p='!' + first="eight-bit limbs" second="nine-bit limbs" logs=("$tag.limbs8.log" "$tag.limbs9.log") +else + r='='; run rows "$tag.rows.log" "${rows_in[@]}" -- "${rows_out[@]}" || r='!' + p='='; run probe "$tag.probe.log" "${probe_in[@]}" -- "${probe_out[@]}" || p='!' + first=rows second=probe logs=("$tag.rows.log" "$tag.probe.log") +fi +echo "$tag on $dev ($how): $first $(word $r), $second $(word $p)" if [ "$r$p" != "$expect" ]; then - echo "::error::$tag on $dev ($how): rows $(word $r), probe $(word $p); expected rows $(word "${expect:0:1}"), probe $(word "${expect:1:1}")" - head -20 "$tag.rows.log" "$tag.probe.log" + echo "::error::$tag on $dev ($how): $first $(word $r), $second $(word $p); expected $first $(word "${expect:0:1}"), $second $(word "${expect:1:1}")" + head -20 "${logs[@]}" exit 1 fi diff --git a/wai/wai-quantum/ci/tensor-program/xla_run.py b/wai/wai-quantum/ci/tensor-program/xla_run.py index a5d061151..b68027ebc 100644 --- a/wai/wai-quantum/ci/tensor-program/xla_run.py +++ b/wai/wai-quantum/ci/tensor-program/xla_run.py @@ -2,12 +2,15 @@ tensor compiler (XLA, through its Python package), run it on one device, and compare every output with the exact one. - python -I xla_run.py ... -- ... + python -I xla_run.py ... -- ... - is the package's name for the device kind (`cpu`, `cuda`). The -module holds two public functions; this compiler builds a module's `main`, so -the chosen function is renamed `main` and the other made private (and unused). -Each expected file is raw little-endian 64-bit integers. + is the package's name for the device kind (`cpu`, `cuda`, `tpu`). +This compiler builds a module's `main`, so the chosen function is renamed +`main` and every other public one made private (and unused). An input is a +`.npy` file, or `x=@` (`16x112xbf16=@w.bin`); an +expected output is a raw file of little-endian 64-bit integers, or +`x=@`. Every output must equal its expected values +exactly, floats included. Exit status: 0 every output exact; 1 an output differs; 2 a usage error. """ @@ -18,7 +21,7 @@ import numpy as np def main(argv): - if len(argv) < 4 or "--" not in argv or argv[2] not in ("rows", "probe"): + if len(argv) < 4 or "--" not in argv: print(__doc__, file=sys.stderr) return 2 platform, path, fn = argv[:3] @@ -29,27 +32,45 @@ def main(argv): import jax jax.config.update("jax_enable_x64", True) - other = "probe" if fn == "rows" else "rows" + import re + text = open(path).read() - for old, new in ((f"func.func @{other}(", f"func.func private @{other}("), (f"func.func @{fn}(", "func.func @main(")): - if text.count(old) != 1: - print(f"{path}: expected one `{old}`", file=sys.stderr) - return 2 - text = text.replace(old, new) + if text.count(f"func.func @{fn}(") != 1: + print(f"{path}: expected one public function @{fn}", file=sys.stderr) + return 2 + text = re.sub(r"func\.func @(\w+)\(", lambda g: "func.func @main(" if g.group(1) == fn else f"func.func private @{g.group(1)}(", text) + + import ml_dtypes + + types = {"bf16": ml_dtypes.bfloat16, "f32": "x=@` as an array, or None for a plain path.""" + if "=@" not in spec: + return None + head, file = spec.split("=@", 1) + *dims, ty = head.split("x") + return np.fromfile(file, dtype=types[ty]).reshape([int(d) for d in dims]) device = jax.devices(platform)[0] exe = device.client.compile_and_load(text, [device]) - results = exe.execute([jax.device_put(np.load(p), device) for p in inputs]) + args = [typed(p) if typed(p) is not None else np.load(p) for p in inputs] + results = exe.execute([jax.device_put(x, device) for x in args]) if len(results) != len(expected): print(f"{fn}: {len(results)} outputs, {len(expected)} expected", file=sys.stderr) return 1 bad = 0 for got, name in zip(results, expected): - got, want = np.asarray(got), np.fromfile(name, dtype=">())); write(&dir, "fx.bin", &raw(&pairs.iter().map(|&(a, b)| (i64::from(a) * i64::from(b) + (1 << 29)) >> 30).collect::>())); - println!("wrote the module, its mutants, and a batch of {N} pairs to {}", dir.display()); + // The products on a matrix unit: the gate's limb weights against the + // batch's limbs, and the exact coefficients. Then nine-bit limbs, which + // bf16 cannot hold: their exact coefficients are not what any bf16 unit + // gives, so a run that matches them is not comparing. + write(&dir, "limbs.mlir", limb_module(N).as_bytes()); + let cases: Vec<[i32; 4]> = (0..N).map(|q| [cols[0][q], cols[1][q], cols[2][q], cols[3][q]]).collect(); + let m32: [i32; 8] = m.map(|e| e as i32); + for (bits, suffix) in [(8, ""), (9, "9")] { + let int_limbs = |x: i32| tensor_engine::limbs(x, bits); + let mut w = vec![0i32; tensor_engine::K * tensor_engine::M]; + for (p, &(k, c)) in tensor_engine::TERMS.iter().enumerate() { + for j in 0..4 { + for (i, &l) in int_limbs(m32[k]).iter().enumerate() { + w[(c * 4 + j) * tensor_engine::M + p * tensor_engine::COEFFICIENTS + i + j] = l; + } + } + } + let mut a = vec![0i32; tensor_engine::K * N]; + for (q, case) in cases.iter().enumerate() { + for (c, &x) in case.iter().enumerate() { + for (j, &l) in int_limbs(x).iter().enumerate() { + a[(c * 4 + j) * N + q] = l; + } + } + } + let bf = |v: &[i32]| v.iter().flat_map(|&x| tensor_engine::bf16(x).to_le_bytes()).collect::>(); + write(&dir, &format!("w{suffix}.bin"), &bf(&w)); + write(&dir, &format!("a{suffix}.bin"), &bf(&a)); + write(&dir, &format!("c{suffix}.bin"), &reference_coefficients(&w, &a, N).iter().flat_map(|x| x.to_le_bytes()).collect::>()); + } + println!("wrote the module, its mutants, the limb module, and a batch of {N} pairs to {}", dir.display()); } diff --git a/wai/wai-quantum/src/quantum_fabric/stablehlo.rs b/wai/wai-quantum/src/quantum_fabric/stablehlo.rs index 2d679efcb..d71337f48 100644 --- a/wai/wai-quantum/src/quantum_fabric/stablehlo.rs +++ b/wai/wai-quantum/src/quantum_fabric/stablehlo.rs @@ -113,6 +113,41 @@ pub fn module(n: usize, c: Constants) -> String { out } +/// The gate's products on a matrix unit, as a second module: one +/// `dot_general` of the gate's limb weights (`16 × 112`, bf16) and a batch's +/// limbs (`16 × n`, bf16), accumulated and returned in f32. The layout and the +/// argument that every coefficient is exact on any unit that multiplies bf16 +/// and accumulates in f32 are [`tensor_engine`](super::tensor_engine)'s; a +/// compiler that maps this to a matrix unit must give the exact coefficients. +pub fn limb_module(n: usize) -> String { + use super::tensor_engine::{K, M}; + let (w, a, c) = (format!("tensor<{K}x{M}xbf16>"), format!("tensor<{K}x{n}xbf16>"), format!("tensor<{M}x{n}xf32>")); + format!( + "// The gate's products as limb coefficients: one bf16 matrix multiply, f32 results.\n\ + module @wai_quantum_limbs {{\n\ + \x20 func.func @coefficients(%w: {w}, %a: {a}) -> {c} {{\n\ + \x20 %c = stablehlo.dot_general %w, %a, contracting_dims = [0] x [0] : ({w}, {a}) -> {c}\n\ + \x20 return %c : {c}\n\ + \x20 }}\n\ + }}\n" + ) +} + +/// The exact coefficients of [`limb_module`] for limbs `w` (`16 × 112`) and +/// `a` (`16 × n`), given as the integers the limbs are (not as bf16 holds +/// them): `112 × n`, row-major. +pub fn reference_coefficients(w: &[i32], a: &[i32], n: usize) -> Vec { + use super::tensor_engine::{K, M}; + let mut c = vec![0f32; M * n]; + for col in 0..M { + for q in 0..n { + let v: i64 = (0..K).map(|k| i64::from(w[k * M + col]) * i64::from(a[k * n + q])).sum(); + c[col * n + q] = v as f32; + } + } + c +} + /// The rows the module must give, from the CPU's arithmetic. pub fn reference_rows(m: &[i64; 8], cols: &[Vec; 4], c: Constants) -> [Vec; 4] { let fx = |e: i64, a: i32| (e * i64::from(a) + c.round) >> c.shift; @@ -169,6 +204,9 @@ mod tests { let m = module(64, CONFORMING); assert!(m.contains("func.func @rows(") && m.contains("func.func @probe(")); assert_eq!(m.matches("stablehlo.shift_right_arithmetic").count(), 16 + 1); + let l = limb_module(512); + assert!(l.contains("func.func @coefficients(%w: tensor<16x112xbf16>, %a: tensor<16x512xbf16>) -> tensor<112x512xf32>"), "{l}"); + assert!(l.contains("stablehlo.dot_general %w, %a, contracting_dims = [0] x [0]")); for c in [Constants { round: 1 << 28, ..CONFORMING }, Constants { shift: 29, ..CONFORMING }, Constants { last_add: true, ..CONFORMING }] { assert_ne!(module(64, c), m); } -- 2.54.0 (Apple Git-157) From 5f6a55ec8d5622f4f0583b481ef28b09a024d567 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 10:07:13 -0400 Subject: [PATCH 16/19] wai-quantum 0.6.0: the limb program in the CHANGELOG Co-Authored-By: Claude Opus 5.5 --- wai/CHANGELOG.md | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/wai/CHANGELOG.md b/wai/CHANGELOG.md index b4cc07d74..34ff18026 100644 --- a/wai/CHANGELOG.md +++ b/wai/CHANGELOG.md @@ -89,6 +89,14 @@ digest are unchanged. device that cannot run reads as a failure, never as a mutant caught. CI: `wai-quantum-tensor-program`. That host's GPU runs it through HIP, built for the GPU's own architecture, and on both of its Vulkan devices. +- **The products on a matrix unit, as a portable program** (`stablehlo::limb_module`): the + tensor-engine member's limb layout as one StableHLO `dot_general` of bf16 operands + accumulated in f32, the program a compiler maps to a matrix unit. With it, the example writes + a nine-bit-limb data set whose exact coefficients bf16 cannot reproduce, so a run that + matches them is not comparing. Exact under both compilers on the x86-64 and arm64 CPUs and + on the discrete GPU through CUDA. Not exact through one compiler's HIP build for the + integrated GPU, whose bf16 multiply returns values that are not integers; CI shows that + step and does not require it. - **Breaking:** - an OpenCL key is `opencl--`, the device's `CL_DEVICE_VENDOR_ID` in hex (`opencl-cpu-8086`, `opencl-gpu-10de`, `opencl-gpu-1027f00`, `opencl-gpu-1002`), not a word -- 2.54.0 (Apple Git-157) From cec78c1f374c9fde9ac49599b8dd67de0816c4f3 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 10:12:55 -0400 Subject: [PATCH 17/19] wai-quantum: the tensor-engine helper runs on a named device platform too WAI_QUANTUM_TENSOR_ENGINE_TARGET names a device platform (trn1, trn2, ...): the helper then runs the same kernel compiled for that device (nki.jit on NumPy inputs), reports that target, and the member's key is tensor-engine- with silicon provenance. Unset, the target is the toolchain's simulator, as before. Co-Authored-By: Claude Opus 5.5 --- .../quantum_fabric/tensor_engine_helper.py | 25 ++++++++++++++++--- .../src/quantum_fabric_tensor_engine.rs | 12 ++++++++- 2 files changed, 33 insertions(+), 4 deletions(-) diff --git a/wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py b/wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py index da16ec1d1..6a3da3308 100644 --- a/wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py +++ b/wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py @@ -9,7 +9,11 @@ assembles the products and does everything else. Every input is an integer bf16 holds exactly and every output an integer below 2^24, so fp32 holds each partial sum exactly in any order (see quantum_fabric::tensor_engine). - python -I tensor_engine_helper.py + python -I tensor_engine_helper.py [--target simulator|] + +The target is the kernel toolchain's simulator (the default), or a device +platform the toolchain names (trn1, trn2, ...): then the kernel is compiled +for that device and runs on it. Protocol (little-endian). On start, one line: "wq-tensor-engine 1 ...". Then requests until end of input: u32 0x45545157, @@ -28,6 +32,20 @@ MAGIC = 0x45545157 def main(): + import os + + target = "simulator" + if len(sys.argv) == 3 and sys.argv[1] == "--target": + target = sys.argv[2] + elif len(sys.argv) != 1: + print(__doc__, file=sys.stderr) + return 2 + if not target.replace("-", "").isalnum(): + print(f"a target is letters, digits and hyphens: {target!r}", file=sys.stderr) + return 2 + if target != "simulator": + os.environ.setdefault("NEURON_PLATFORM_TARGET_OVERRIDE", target) + import numpy as np import ml_dtypes import nki @@ -48,7 +66,8 @@ def main(): inp, out = sys.stdin.buffer, sys.stdout.buffer version = getattr(nki, "__version__", "unknown") - out.write(f"wq-tensor-engine 1 simulator nki {version} numpy {np.__version__} ml_dtypes {ml_dtypes.__version__}\n".encode()) + run = nki.simulate(coefficients) if target == "simulator" else coefficients + out.write(f"wq-tensor-engine 1 {target} nki {version} numpy {np.__version__} ml_dtypes {ml_dtypes.__version__}\n".encode()) out.flush() def read(n): @@ -67,7 +86,7 @@ def main(): raise ValueError(f"a request of {n} pairs (magic {magic:#x})") w = np.frombuffer(read(K * M * 2), " Result<(Helper, String, String), String> { std::fs::write(&staging, tensor_engine::HELPER).map_err(|e| format!("{}: {e}", staging.display()))?; std::fs::rename(&staging, &script).map_err(|e| format!("{}: {e}", script.display()))?; } + // The simulator unless the host names a device platform. + let target = std::env::var("WAI_QUANTUM_TENSOR_ENGINE_TARGET").ok().filter(|t| !t.is_empty()).unwrap_or_else(|| "simulator".into()); + if !target.bytes().all(|b| b.is_ascii_alphanumeric() || b == b'-') { + return Err(format!("WAI_QUANTUM_TENSOR_ENGINE_TARGET={target:?} is not a platform name")); + } let mut child = Command::new(python) .arg("-I") .arg(&script) + .args(["--target", &target]) .stdin(Stdio::piped()) .stdout(Stdio::piped()) .stderr(Stdio::null()) -- 2.54.0 (Apple Git-157) From 05b8cca5753020bdc086c20bd78e7de3e686cc19 Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 20:48:13 -0400 Subject: [PATCH 18/19] ci: the dataflow workflow checks the instruction table on the vector generation only, and keeps it in the cache The table test failed on both legs of #266: the cached simulator had been built before the table was copied beside it, so the table was missing. A cache entry without its table is rebuilt now. And the encoder is a transcription of the b140 generation's table, which the 401e table does not match (its PACR has other fields; measured), so only the b140 leg names the table and requires its check. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-dataflow.yml | 17 ++++++++++++++--- wai/wai-quantum/src/quantum_fabric/tensix.rs | 6 ++++-- 2 files changed, 18 insertions(+), 5 deletions(-) diff --git a/.github/workflows/wai-quantum-dataflow.yml b/.github/workflows/wai-quantum-dataflow.yml index 0873f83c1..5e5082e5c 100644 --- a/.github/workflows/wai-quantum-dataflow.yml +++ b/.github/workflows/wai-quantum-dataflow.yml @@ -63,12 +63,17 @@ jobs: max-parallel: 1 matrix: include: + # `vector`: the generation whose vector engine the vector programs + # drive, and whose instruction table their encoder is checked against. - pci: '1e52:b140' build: bh key: dataflow-1e52-b140,dataflow-1e52-b140-vector + vector: 'true' - pci: '1e52:401e' build: wh key: dataflow-1e52-401e + vector: '' + container: image: debian@sha256:913f6706df59a68922d1dd08f78c2476560a8d367897200a6005b00e5f67c2d5 options: --memory=6g --memory-swap=6g @@ -108,7 +113,9 @@ jobs: exec 9>"$cache/.lock" flock 9 lib="$cache/sim-$SIM_COMMIT/${{ matrix.build }}/libttsim.so" - if [ ! -f "$lib" ]; then + # The table is part of the cache entry: an entry built before it was + # kept is rebuilt, not used without it. + if [ ! -f "$lib" ] || [ ! -f "$(dirname "$lib")/tensix_isa.json" ]; then b="$cache/sim-build" rm -rf "$b" && git init -q "$b" git -C "$b" fetch -q --depth 1 "$SIM_REPO" "$SIM_COMMIT" @@ -161,12 +168,16 @@ jobs: export RUSTUP_HOME="$cache/rustup" CARGO_HOME="$cache/cargo" CARGO_TARGET_DIR="$cache/target" export PATH="$CARGO_HOME/bin:$PATH" export WAI_QUANTUM_DATAFLOW_LIB="$cache/sim-$SIM_COMMIT/${{ matrix.build }}/libttsim.so" - export WAI_QUANTUM_TENSIX_ISA="$cache/sim-$SIM_COMMIT/${{ matrix.build }}/tensix_isa.json" + # The vector programs' encoder is a transcription of this generation's + # table; another generation lays some instructions out differently. + if [ -n "${{ matrix.vector }}" ]; then export WAI_QUANTUM_TENSIX_ISA="$cache/sim-$SIM_COMMIT/${{ matrix.build }}/tensix_isa.json"; fi cargo test --locked --release -j 6 --features quantum_fabric_dataflow --lib -- quantum_fabric_dataflow:: quantum_fabric::rv32:: quantum_fabric::tensix:: quantum_fabric::vector_engine:: --nocapture 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; } # The vector engine's encoder against the simulator's own table (the # test returns early without the table, so require its line). - grep -qE '^tensix layouts: [0-9]+ instructions match the table' "$RUNNER_TEMP/unit.log" || { echo "::error::the instruction table was not checked"; exit 1; } + if [ -n "${{ matrix.vector }}" ]; then + grep -qE '^tensix layouts: [0-9]+ instructions match the table' "$RUNNER_TEMP/unit.log" || { echo "::error::the instruction table was not checked"; exit 1; } + fi report="$RUNNER_TEMP/dataflow-report.txt"; : > "$report" WAI_QUANTUM_DATAFLOW_REPORT="$report" cargo test --locked --release -j 6 --features quantum_fabric_dataflow --test dataflow_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; } diff --git a/wai/wai-quantum/src/quantum_fabric/tensix.rs b/wai/wai-quantum/src/quantum_fabric/tensix.rs index f7be541d0..54c7c63cd 100644 --- a/wai/wai-quantum/src/quantum_fabric/tensix.rs +++ b/wai/wai-quantum/src/quantum_fabric/tensix.rs @@ -127,8 +127,10 @@ mod tests { assert_eq!(inline(0xA280_0002).rotate_right(2), 0xA280_0002); } - /// With `WAI_QUANTUM_TENSIX_ISA` naming the simulator's table - /// (`data/bh/tensix_isa.json`), every transcribed layout must be the table's. + /// With `WAI_QUANTUM_TENSIX_ISA` naming the simulator's table for the + /// generation the vector programs drive (`data/bh/tensix_isa.json`), every + /// transcribed layout must be the table's. (Another generation's table lays + /// some instructions out differently: `PACR` has other fields.) #[test] fn layouts_are_the_simulators_table() { let Some(path) = std::env::var_os("WAI_QUANTUM_TENSIX_ISA") else { return }; -- 2.54.0 (Apple Git-157) From 9225edce712b5b12e1ad82018a6937c19d72fdad Mon Sep 17 00:00:00 2001 From: Transaction Science Date: Thu, 15 Oct 2026 20:55:59 -0400 Subject: [PATCH 19/19] ci: the dataflow workflow checks each declared member, not the list as one key The b140 leg declares two members (the cores and the vector engine), and the checks used the comma-joined list as one key, so no report line matched. They loop over the declared keys now. Found by running the leg by hand in its CI image; with this change all three members pass there (the simulator named, four certificates, three of three mutants each). Co-Authored-By: Claude Opus 5.5 --- .github/workflows/wai-quantum-dataflow.yml | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/.github/workflows/wai-quantum-dataflow.yml b/.github/workflows/wai-quantum-dataflow.yml index 5e5082e5c..4aaeceb62 100644 --- a/.github/workflows/wai-quantum-dataflow.yml +++ b/.github/workflows/wai-quantum-dataflow.yml @@ -182,12 +182,14 @@ jobs: WAI_QUANTUM_DATAFLOW_REPORT="$report" cargo test --locked --release -j 6 --features quantum_fabric_dataflow --test dataflow_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_DATAFLOW_EXPECT" - grep -qx "dataflow devices: as declared ($key)" "$report" || { echo "::error::the device found is not the declared one"; exit 1; } - grep -q "^dataflow member: key=$key .*simulator:device simulator 283f621e" "$report" || { echo "::error::the member does not say it ran on the pinned simulator"; exit 1; } - n="$(grep -c "^dataflow certified: key=$key .* parity=bit-exact " "$report")" - [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; } - grep -qx "dataflow mutants: key=$key killed 3 of 3" "$report" || { echo "::error::a kernel mutant survived"; exit 1; } + grep -qx "dataflow devices: as declared ($WAI_QUANTUM_DATAFLOW_EXPECT)" "$report" || { echo "::error::the devices found are not the declared ones"; exit 1; } + # Each declared member (a generation with a vector engine has two). + for key in ${WAI_QUANTUM_DATAFLOW_EXPECT//,/ }; do + grep -q "^dataflow member: key=$key .*simulator:device simulator 283f621e" "$report" || { echo "::error::$key: the member does not say it ran on the pinned simulator"; exit 1; } + n="$(grep -c "^dataflow certified: key=$key .* parity=bit-exact " "$report")" + [ "$n" = 4 ] || { echo "::error::$key: $n of 4 certificates were earned"; exit 1; } + grep -qx "dataflow mutants: key=$key killed 3 of 3" "$report" || { echo "::error::$key: a kernel mutant survived"; exit 1; } + done dataflow-absent: name: member · dataflow device absent (macos) -- 2.54.0 (Apple Git-157)