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

wai-quantum 0.6.0: HIP, npu-cores, tensor-engine and avx512 members; the dataflow vector engine; the gate as portable tensor programs #266

Merged
dcharlot merged 24 commits from wai/quantum-dataflow-vector into main 2026-10-15 20:58:49 -04:00
Conversation 0 Commits 24 Files changed 54 +7516 -108
54 changed files with 7516 additions and 108 deletions
Download patch file Download diff file Expand all files Collapse all files

42
.github/workflows/wai-quantum-dataflow.yml vendored
View file

@ -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
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,14 +113,19 @@ 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"
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 +155,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,18 +168,28 @@ 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"
# 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).
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; }
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)

153
.github/workflows/wai-quantum-hip.yml vendored Normal file
View file

@ -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 (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
# 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 (framework-128gb)
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-arm64)
needs: relevant
if: needs.relevant.outputs.run == 'true'
runs-on: linux-arm64
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; }

160
.github/workflows/wai-quantum-npu-cores.yml vendored Normal file
View file

@ -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-arm64)
needs: relevant
if: needs.relevant.outputs.run == 'true'
runs-on: linux-arm64
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; }

91
.github/workflows/wai-quantum-opencl.yml vendored
View file

@ -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-<kind>-<vendor id, hex>). 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 (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,
# 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-<kind>-<vendor id, hex>).
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 (framework-128gb)
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.

156
.github/workflows/wai-quantum-tensor-engine.yml vendored Normal file
View file

@ -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; }

309
.github/workflows/wai-quantum-tensor-program.yml vendored Normal file
View file

@ -0,0 +1,309 @@
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. 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): both compilers on the x86-64 CPU, the program
# and its three mutants (each must fail where it should); builds for three
# 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
# 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 }}
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
JAX_VERSION: 0.11.2
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 and its mutants on the CPU, both compilers; 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/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" 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"
# 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
# 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"
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 with each compiler
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"; }
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"
# 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; }
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 '!='
# 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)
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-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 '=='
$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
# 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)
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 and, through XLA, on its CPU
shell: bash
working-directory: wai/wai-quantum
run: |
set -euo pipefail
# 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"
# Shader source in the module, compiled by the system at load: no
# developer toolchain needed.
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
$check iree metal gate.metal.vmfb '=='
$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 '=!'

49
.github/workflows/wai-quantum-vector.yml vendored
View file

@ -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

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

@ -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'

104
wai/CHANGELOG.md
View file

@ -6,6 +6,110 @@ wire format may change and such changes are recorded here.
## v0.4.0 — current
### 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.
- **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-<arch>`); 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-<pci vendor>-<device>`.
- **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
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; 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.
- **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.
- **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, 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 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`. 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-<kind>-<vendor id>`, 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.
### Parameter sets kept loaded between decodes
- **The store.** A decode through the integer operator set lowers its parameter set and

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

@ -1017,7 +1017,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
[[package]]
name = "wai-quantum"
version = "0.5.1"
version = "0.6.0"
dependencies = [
"blake3",
"ed25519-dalek",

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

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

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

@ -703,7 +703,7 @@ dependencies = [
[[package]]
name = "wai-quantum"
version = "0.5.1"
version = "0.6.0"
dependencies = [
"blake3",
"ed25519-dalek",

4
wai/wai-quantum-qpu/Cargo.lock generated
View file

@ -422,7 +422,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
[[package]]
name = "wai-quantum"
version = "0.5.1"
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",

4
wai/wai-quantum-qpu/Cargo.toml
View file

@ -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"

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

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

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

@ -1079,7 +1079,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
[[package]]
name = "wai-quantum"
version = "0.5.1"
version = "0.6.0"
dependencies = [
"blake3",
"block2",

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

@ -1,6 +1,6 @@
[package]
name = "wai-quantum"
version = "0.5.1"
version = "0.6.0"
edition = "2024"
rust-version = "1.97"
license = "Apache-2.0"
@ -162,11 +162,17 @@ 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"]
# 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 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.
@ -205,6 +211,22 @@ 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 = "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"]
@ -241,6 +263,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"]

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

@ -103,16 +103,19 @@ 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, 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-<kind>-<vendor id>`). 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_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 — `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-<device>`); the library is a simulator, and every certificate says so (`WAI_QUANTUM_DATAFLOW_SIMULATOR`). Three kernel mutants are compiled in |
| `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-<device>`); 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`.

84
wai/wai-quantum/ci/tensor-program/check.sh Executable file
View file

@ -0,0 +1,84 @@
#!/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 <device> <module.vmfb> <expect>
# check.sh xla <platform> <module.mlir> <expect>
#
# Run in the directory the stablehlo_gate example wrote. <expect> is two
# 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 <device> 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 <platform> 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: <expect> is two of = and !" >&2; exit 2 ;; esac
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)
probe_out=(prod.bin fx.bin)
# run <function> <log> <inputs...> -- <expected...>: 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 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
# 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
"${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")"
word() { [ "$1" = '=' ] && echo exact || echo differ; }
if [ "$kind" = limbs ]; then
ln=$(( $(size c.bin) / 4 / 112 ))
r='='; run coefficients "$tag.limbs8.log" "[email protected]" "16x${ln}[email protected]" -- "112x${ln}[email protected]" || r='!'
p='='; run coefficients "$tag.limbs9.log" "[email protected]" "16x${ln}[email protected]" -- "112x${ln}[email protected]" || 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): $first $(word $r), $second $(word $p); expected $first $(word "${expect:0:1}"), $second $(word "${expect:1:1}")"
head -20 "${logs[@]}"
exit 1
fi

81
wai/wai-quantum/ci/tensor-program/xla_run.py Normal file
View file

@ -0,0 +1,81 @@
"""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 <platform> <module.mlir> <function> <input>... -- <expected>...
<platform> 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 `<shape>x<type>=@<raw file>` (`16x112xbf16=@w.bin`); an
expected output is a raw file of little-endian 64-bit integers, or
`<shape>x<type>=@<raw file>`. Every output must equal its expected values
exactly, floats included.
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:
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)
import re
text = open(path).read()
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": "<f4", "i32": "<i4", "i64": "<i8"}
def typed(spec):
"""`<shape>x<type>=@<file>` 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])
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):
want = typed(name)
if want is None:
want = np.fromfile(name, dtype="<i8")
got = np.asarray(got)
if got.dtype != want.dtype or got.size != want.size:
print(f"{name}: {got.dtype}{list(got.shape)}, not {want.dtype}{list(want.shape)}")
bad += 1
continue
got = got.reshape(want.shape)
if not np.array_equal(got, want):
print(f"{name}: {int(np.sum(got != want))} of {want.size} values differ")
bad += 1
print(f"{fn} on {device.platform} ({device.device_kind}): {'exact' if bad == 0 else 'differs'}")
return 1 if bad else 0
if __name__ == "__main__":
sys.exit(main(sys.argv[1:]))

11
wai/wai-quantum/examples/rv32_listing.rs
View file

@ -9,6 +9,7 @@
//!
//! In `<dir>`, for each program `<name>`: `<name>.s` (the listing, with a
//! `.text` header) and `<name>.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::<Vec<u8>>()).expect("the words");

100
wai/wai-quantum/examples/stablehlo_gate.rs Normal file
View file

@ -0,0 +1,100 @@
//! 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 -- <dir>
//! ```
//!
//! In `<dir>`: `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, limb_module, module, npy_i32, npy_scalar_i64, reference_coefficients, reference_rows};
use wai_quantum::quantum_fabric::tensor_engine;
/// Pairs per batch.
const N: usize = 4096;
fn raw(v: &[i64]) -> Vec<u8> {
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<i32>; 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<i32>, Vec<i32>) = 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::<Vec<i64>>()));
write(&dir, "fx.bin", &raw(&pairs.iter().map(|&(a, b)| (i64::from(a) * i64::from(b) + (1 << 29)) >> 30).collect::<Vec<i64>>()));
// 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::<Vec<u8>>();
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::<Vec<u8>>());
}
println!("wrote the module, its mutants, the limb module, and a batch of {N} pairs to {}", dir.display());
}

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

@ -10,12 +10,18 @@
//! 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. 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) 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_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_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
@ -90,12 +96,18 @@ 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")))]
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_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")))]

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

@ -2048,13 +2048,21 @@ pub mod lanes {
}
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")]
pub mod rv32;
#[cfg(feature = "quantum_fabric_dataflow")]
pub mod tensix;
#[cfg(feature = "quantum_fabric_dataflow")]
pub mod vector_engine;
pub mod sampling;
pub mod stablehlo;
pub mod tensor_engine;
#[cfg(feature = "quantum_sv")]
pub mod sv64;
pub mod vqpu;

1
wai/wai-quantum/src/quantum_fabric/cuda.rs
View file

@ -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.

202
wai/wai-quantum/src/quantum_fabric/hip.rs Normal file
View file

@ -0,0 +1,202 @@
//! 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<u8>, 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)]
#[non_exhaustive]
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,
/// `"<domain>:<bus>:<device>"`, 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=<arch>"`, the feature settings included.
pub fn arch_option(&self) -> String {
format!("--offload-arch={}", self.arch)
}
/// `"local/gpu-hip-<d12>"`. `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-<arch>`, 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-<arch name>"`, 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);
}
}
}

65
wai/wai-quantum/src/quantum_fabric/npu_cores.cc Normal file
View file

@ -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 <stdint.h>
#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);
}
}

274
wai/wai-quantum/src/quantum_fabric/npu_cores.rs Normal file
View file

@ -0,0 +1,274 @@
//! 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<u8>, 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)]
#[non_exhaustive]
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-<d12>"`: 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-<vendor>-<device>"`, 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<i32> {
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<String> = [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]);
}
}

98
wai/wai-quantum/src/quantum_fabric/npu_cores_design.py Normal file
View file

@ -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())

125
wai/wai-quantum/src/quantum_fabric/npu_cores_host.cpp Normal file
View file

@ -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 <xrt/xrt_bo.h>
#include <xrt/xrt_device.h>
#include <xrt/xrt_hw_context.h>
#include <xrt/xrt_kernel.h>
#include <xrt/experimental/xrt_xclbin.h>
#include <cstdint>
#include <cstring>
#include <exception>
#include <fstream>
#include <string>
#include <vector>
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<uint32_t> 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<char> bytes((std::istreambuf_iterator<char>(f)), std::istreambuf_iterator<char>());
if (bytes.empty() || bytes.size() % 4 != 0) throw std::runtime_error(std::string(path) + " is not a whole number of words");
std::vector<uint32_t> 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<xrt::info::device::name>());
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<uint32_t>(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<void *>(), 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<Npu *>(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<void *>(), in, in_words * 4);
std::memset(b.map<void *>(), 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<int>(state)));
return 1;
}
b.sync(XCL_BO_SYNC_BO_FROM_DEVICE);
std::memcpy(out, b.map<void *>(), 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<Npu *>(h);
}

35
wai/wai-quantum/src/quantum_fabric/opencl.rs
View file

@ -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,
@ -113,6 +114,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 +138,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-<kind>-<d12>"`. `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-<arch>` and `local/gpu-…`, so with no energy
/// figures [`place`](super::place) keeps the CPU reference.
pub fn member_id(&self) -> String {
@ -155,15 +151,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-<kind>-<vendor>"`, for
/// instance `"opencl-gpu-maker"`.
/// The key a runner declares in CI: `"opencl-<kind>-<vendor id>"`, 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 +204,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 +231,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);

15
wai/wai-quantum/src/quantum_fabric/rv32.rs
View file

@ -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 });

228
wai/wai-quantum/src/quantum_fabric/stablehlo.rs Normal file
View file

@ -0,0 +1,228 @@
//! 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<i64>) -> {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<String> = (0..8).map(|k| b.splat(&format!("%e{k}"))).collect();
let a: Vec<String> = ["%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<i64>")).chain(["%a0r", "%a0i", "%a1r", "%a1i"].iter().map(|x| format!("{x}: {ti}"))).collect::<Vec<_>>().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 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<f32> {
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<i32>; 4], c: Constants) -> [Vec<i64>; 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<i64>; 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` `<i4`
/// or `<i8`, `shape` empty for a scalar.
pub fn npy(descr: &str, shape: &[usize], data: &[u8]) -> Vec<u8> {
let dims = match shape {
[] => "()".to_string(),
[n] => format!("({n},)"),
_ => format!("({})", shape.iter().map(|d| d.to_string()).collect::<Vec<_>>().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<u8> {
npy("<i4", &[v.len()], &v.iter().flat_map(|x| x.to_le_bytes()).collect::<Vec<u8>>())
}
pub fn npy_i64(v: &[i64]) -> Vec<u8> {
npy("<i8", &[v.len()], &v.iter().flat_map(|x| x.to_le_bytes()).collect::<Vec<u8>>())
}
pub fn npy_scalar_i64(v: i64) -> Vec<u8> {
npy("<i8", &[], &v.to_le_bytes())
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn the_module_has_both_functions_and_its_mutants_differ() {
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);
}
}
#[test]
fn npy_headers_are_aligned_and_reference_rows_are_the_gate() {
for v in [npy_i32(&[1, 2, 3]), npy_i64(&[7; 512]), npy_scalar_i64(-5)] {
let hl = u16::from_le_bytes([v[8], v[9]]) as usize;
assert_eq!((10 + hl) % 64, 0);
assert_eq!(v[10 + hl - 1], b'\n');
}
// One pair, the identity gate (1.0 = 2^30): the rows are the inputs.
let one = 1i64 << 30;
let m = [one, 0, 0, 0, 0, 0, one, 0];
let r = reference_rows(&m, &[vec![5], vec![-6], vec![7], vec![-8]], CONFORMING);
assert_eq!(r.map(|x| x[0]), [5, -6, 7, -8]);
}
}

31
wai/wai-quantum/src/quantum_fabric/sv64.rs
View file

@ -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);
}
}
}

149
wai/wai-quantum/src/quantum_fabric/tensix.rs Normal file
View file

@ -0,0 +1,149 @@
//! 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 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 };
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());
}
}

380
wai/wai-quantum/src/quantum_fabric/tensor_engine.rs Normal file
View file

@ -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<u16> {
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<u16> {
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<u16>, 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<Vec<[i64; PRODUCTS]>, 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-<d12>"` 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-<target>"`.
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<f32> {
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<String> = [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"));
}
}

101
wai/wai-quantum/src/quantum_fabric/tensor_engine_helper.py Normal file
View file

@ -0,0 +1,101 @@
"""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 [--target simulator|<device platform>]
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 <target>
<package> <version> ...". 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 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
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")
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):
b = inp.read(n)
if len(b) != n:
raise EOFError
return b
while True:
try:
magic, n = struct.unpack("<II", read(8))
except EOFError:
return
try:
if magic != MAGIC or n == 0 or n % TILE:
raise ValueError(f"a request of {n} pairs (magic {magic:#x})")
w = np.frombuffer(read(K * M * 2), "<u2").view(ml_dtypes.bfloat16).reshape(K, M)
a = np.frombuffer(read(K * n * 2), "<u2").view(ml_dtypes.bfloat16).reshape(K, n)
c = np.asarray(run(w, a), dtype=np.float32)
out.write(struct.pack("<I", 0))
out.write(np.ascontiguousarray(c, dtype="<f4").tobytes())
except EOFError:
return
except Exception as e: # report, keep serving
msg = f"{type(e).__name__}: {e}".encode()
out.write(struct.pack("<II", 1, len(msg)) + msg)
out.flush()
if __name__ == "__main__":
sys.exit(main())

100
wai/wai-quantum/src/quantum_fabric/vector_avx512.rs Normal file
View file

@ -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<const RND: i64, const SH: u32>(e: __m512i, a: __m512i) -> __m512i {
_mm512_srai_epi64::<SH>(_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<const RND: i64, const SH: u32, const LAST_ADD: bool>(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::<RND, SH>($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<const RND: i64, const SH: u32>(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::<RND, SH>(vb, va));
}
k += LANES;
}
}

315
wai/wai-quantum/src/quantum_fabric/vector_engine.rs Normal file
View file

@ -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<u32>, 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<u32>, 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<u32>, 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);
}
}

229
wai/wai-quantum/src/quantum_fabric_dataflow.rs
View file

@ -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<u32>, String), (Vec<u32>, 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.
@ -144,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
@ -152,6 +201,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 +235,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 +385,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 +398,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<usize> = (0..counts.len()).collect();
let mut step = expected.clamp(1, budget.min(u64::from(u32::MAX)));
let mut waiting: Vec<usize> = (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 +424,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<Device>, String> {
@ -403,21 +462,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-<vendor>-<device>`.
/// The key a runner declares for the cores: `dataflow-<vendor>-<device>`.
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 +508,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 +519,10 @@ pub struct DataflowMember {
}
impl DataflowMember {
fn open(chip: Chip, variant: Variant) -> Result<DataflowMember, DataflowError> {
fn open(chip: Chip, engine: Engine, variant: Variant) -> Result<DataflowMember, DataflowError> {
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 +531,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::<String>());
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 +563,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 +628,9 @@ impl DataflowMember {
/// The kernel's [`lanes::mul_wide`] and [`lanes::fxmul_wide`] of each pair.
pub fn probe(&self, pairs: &[(i32, i32)]) -> Result<Vec<ProbeOutput>, 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 +641,7 @@ impl DataflowMember {
dev.l1_write_words(j, LAYOUT.column(1), &slice.iter().map(|p| p.1 as u32).collect::<Vec<u32>>());
}
let counts: Vec<u32> = 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 +652,9 @@ impl DataflowMember {
/// The four rows of each pair, from the kernel.
fn rows(&self, m: &[i32; 8], cols: &[Vec<i32>; 4]) -> Result<[Vec<i64>; 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 +671,7 @@ impl DataflowMember {
}
}
let counts: Vec<u32> = 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 +684,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<Vec<ProbeOutput>, 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::<Vec<u32>>());
dev.l1_write_words(j, vector_engine::IN + 4 * n as u32, &slice.iter().map(|p| p.1 as u32).collect::<Vec<u32>>());
}
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<i32>; 4]) -> Result<[Vec<i64>; 4], DataflowError> {
let n = vector_engine::PAIRS;
let mut dev = self.device()?;
let len = cols[0].len();
let mut o: [Vec<i64>; 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::<Vec<u32>>());
}
}
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<LaneRun, DataflowError> {
let len = state.len();
@ -759,6 +910,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");

1054
wai/wai-quantum/src/quantum_fabric_hip.rs Normal file
View file

File diff suppressed because it is too large Load diff

742
wai/wai-quantum/src/quantum_fabric_npu_cores.rs Normal file
View file

@ -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::{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<NpuCoresError> 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<Found> {
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<PathBuf> = 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<F: Copy>(p: *mut c_void) -> F {
assert_eq!(std::mem::size_of::<F>(), 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<Host, String> {
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 { 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");
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<Host, String> {
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<PathBuf, NpuCoresError> {
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::<Vec<_>>().into_iter().rev().collect::<Vec<_>>().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<Result<Host, String>> = 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::<Vec<_>>().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<Context, NpuCoresError> {
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<Vec<i32>, 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<NpuCoresMember>,
pub refused: Vec<Refusal>,
pub unavailable: Option<String>,
}
/// 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<Context>,
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<NpuCoresMember, NpuCoresError> {
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<NpuCoresMember, NpuCoresError> {
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<Vec<ProbeOutput>, 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<Vec<[i64; 4]>, 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<LaneRun, NpuCoresError> {
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<Lane>, gates: &[(&GateWords, u32)]) -> Result<LaneRun, NpuCoresError> {
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<u32> = 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<LaneRun, NpuCoresError> {
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<StateVector, NpuCoresError> {
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<StateVector, NpuCoresError> {
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<Certificate, NpuCoresError> {
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<Member, NpuCoresError> {
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<k>` 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());
}
}

1
wai/wai-quantum/src/quantum_fabric_opencl.rs
View file

@ -248,6 +248,7 @@ fn open(platform: &PlatformInfo, position: usize, id: cl_device_id) -> Result<Op
version: device.version().unwrap_or_default(),
driver: device.driver_version().unwrap_or_default(),
kind,
vendor_id: device.vendor_id().unwrap_or(0),
uuid,
compute_units: device.max_compute_units().unwrap_or(0),
memory,

515
wai/wai-quantum/src/quantum_fabric_tensor_engine.rs Normal file
View file

@ -0,0 +1,515 @@
//! The fabric's tensor-engine member: the gate's products on a tensor
//! processor's matrix engine, as exact limb convolutions in one bf16 matrix
//! multiply per tile of pairs, accumulated in fp32. See
//! [`quantum_fabric::tensor_engine`](crate::quantum_fabric::tensor_engine)
//! for the layout and the argument that it is bit-exact on any engine that
//! multiplies bf16 and accumulates in fp32.
//!
//! - **Discovery** ([`discover`]): the kernel ([`tensor_engine::HELPER`]) is
//! run as a long-lived helper by the Python the host names
//! (`WAI_QUANTUM_TENSOR_ENGINE_PYTHON`, an environment with the engine's
//! kernel toolchain). The helper says where the kernel runs (its target)
//! and what it runs with; its arithmetic is then probed against [`lanes`]
//! in this process. The target is the toolchain's simulator unless the
//! host names a device platform (`WAI_QUANTUM_TENSOR_ENGINE_TARGET`, for
//! instance `trn1`), where the kernel is compiled for the device and runs
//! on it, and the provenance is silicon. A host with no such Python finds no
//! member and says why.
//! - **Running**: the host gathers the pairs a gate selects and sends their
//! limbs and the gate's limb weights; the engine returns every product's
//! coefficients; the host checks each is an exact integer, assembles the
//! products and applies the fixed-point arithmetic, checking each row's
//! range and keeping its low 32 bits.
//!
//! # Honest boundaries
//!
//! - **The target is the helper's.** Today it is the kernel toolchain's
//! simulator, whose arithmetic is a model of the engine's, not the engine:
//! every certificate's provenance says so (`simulator:<versions>`), 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<TensorEngineError> 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<ChildStdout>,
}
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()))?;
}
// 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())
.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::<Vec<_>>().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<Vec<f32>, 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<TensorEngineMember>,
pub refused: Vec<(String, TensorEngineError)>,
pub unavailable: Option<String>,
}
/// 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<Mutex<Helper>>,
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<TensorEngineMember, TensorEngineError> {
// 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<Vec<[i64; tensor_engine::PRODUCTS]>, 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<Vec<[i64; tensor_engine::PRODUCTS]>, 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<Vec<ProbeOutput>, 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<LaneRun, TensorEngineError> {
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<Lane>, gates: &[(&GateWords, u32)]) -> Result<LaneRun, TensorEngineError> {
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<u32> = 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<LaneRun, TensorEngineError> {
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<StateVector, TensorEngineError> {
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<StateVector, TensorEngineError> {
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<Certificate, TensorEngineError> {
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<Member, TensorEngineError> {
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<PathBuf, ()> {
for d in std::env::var_os("PATH").map(|p| std::env::split_paths(&p).collect::<Vec<_>>()).unwrap_or_default() {
let c = d.join(name);
if c.is_file() {
return Ok(c);
}
}
Err(())
}
}

113
wai/wai-quantum/src/quantum_fabric_vector.rs
View file

@ -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.
@ -75,19 +81,22 @@ impl From<VectorError> for MemberError {
/// The instruction sets a member runs.
#[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)]
#[non_exhaustive]
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 +106,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 +115,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 +487,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 +500,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 +536,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> {
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> {
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 +579,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 +622,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 +662,28 @@ pub struct Discovery {
pub unavailable: Option<String>,
}
/// 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 +835,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 +920,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 +1134,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"));
}

6
wai/wai-quantum/tests/dataflow_member.rs
View file

@ -9,7 +9,7 @@
//! same lines to it, which a CI step can count per device.
//!
//! ```text
//! WAI_QUANTUM_DATAFLOW_LIB=<the device library> WAI_QUANTUM_DATAFLOW_EXPECT=dataflow-1e52-b140 \
//! WAI_QUANTUM_DATAFLOW_LIB=<the device library> 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()));
}
}

429
wai/wai-quantum/tests/hip_member.rs Normal file
View file

@ -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<Discovery> = 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<String, String> {
let mut found: Vec<&str> = d.found.iter().map(|m| m.key()).collect();
found.sort_unstable();
let found_list = if found.is_empty() { "none".to_string() } else { found.join(",") };
let refused: Vec<String> = d.refused.iter().map(|r| format!("{} ({:?}): {}", r.key, r.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<Circuit> {
let gate = |base, controls: Vec<u8>, target, param| Gate { base, controls, target, param };
let bad = [
(1, vec![gate(BaseGate::H, vec![], 0, 0), gate(BaseGate::P, vec![], 0, 0)]),
(1, vec![gate(BaseGate::P, vec![], 0, 33)]),
(1, vec![gate(BaseGate::X, vec![], 3, 0)]),
(2, vec![gate(BaseGate::X, vec![0], 0, 0)]),
(2, vec![gate(BaseGate::X, vec![5], 1, 0)]),
];
let mut out = Vec::new();
for gs in Gateset::ALL {
for (n, ops) in bad.clone() {
let mut c = Circuit::with_gateset(n, gs);
c.ops = ops;
out.push(c);
}
}
out
}
/// `Z` on qubit 0: on `[(0, 0), (i32::MIN, 0)]` its exact result is `2^31`,
/// which leaves 32 bits.
fn z_words() -> GateWords {
lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 0, param: 0 }, Gateset::V2)
}
const Z_STATE: [Lane; 2] = [[0, 0], [i32::MIN, 0]];
#[test]
fn 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<Lane> = (0..1usize << n).map(|_| [value(&mut s), value(&mut s)]).collect();
let gates: Vec<GateWords> = (0..1 + next(&mut s) % 8)
.map(|_| {
let t = (next(&mut s) % u64::from(n)) as u32;
let ctrl = (0..n).filter(|&b| b != t && next(&mut s).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()));
}
}

112
wai/wai-quantum/tests/hip_sv.rs Normal file
View file

@ -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<HipMember> {
let d = discover();
let mut found: Vec<String> = 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()));
}
}

414
wai/wai-quantum/tests/npu_cores_member.rs Normal file
View file

@ -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=<the toolchain's 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<Discovery> = 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<String, String> {
let mut found: Vec<&str> = d.found.iter().map(|m| m.key()).collect();
found.sort_unstable();
let found_list = if found.is_empty() { "none".to_string() } else { found.join(",") };
let refused: Vec<String> = d.refused.iter().map(|r| format!("{} ({}): {}", r.key, r.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<Circuit> {
let gate = |base, controls: Vec<u8>, target, param| Gate { base, controls, target, param };
let bad = [
(1, vec![gate(BaseGate::H, vec![], 0, 0), gate(BaseGate::P, vec![], 0, 0)]),
(1, vec![gate(BaseGate::P, vec![], 0, 33)]),
(1, vec![gate(BaseGate::X, vec![], 3, 0)]),
(2, vec![gate(BaseGate::X, vec![0], 0, 0)]),
(2, vec![gate(BaseGate::X, vec![5], 1, 0)]),
];
let mut out = Vec::new();
for gs in Gateset::ALL {
for (n, ops) in bad.clone() {
let mut c = Circuit::with_gateset(n, gs);
c.ops = ops;
out.push(c);
}
}
out
}
/// `Z` on qubit 0: on `[(0, 0), (i32::MIN, 0)]` its exact result is `2^31`,
/// which leaves 32 bits.
fn z_words() -> GateWords {
lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 0, param: 0 }, Gateset::V2)
}
const Z_STATE: [Lane; 2] = [[0, 0], [i32::MIN, 0]];
#[test]
fn 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<Lane> = (0..1usize << n).map(|_| [value(&mut s), value(&mut s)]).collect();
let gates: Vec<GateWords> = (0..1 + next(&mut s) % 8)
.map(|_| {
let t = (next(&mut s) % u64::from(n)) as u32;
let ctrl = (0..n).filter(|&b| b != t && next(&mut s).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
wai/wai-quantum/tests/opencl_member.rs
View file

@ -9,7 +9,7 @@
//! same lines to it, which a CI step can count per device.
//!
//! ```text
//! WAI_QUANTUM_OPENCL_EXPECT=opencl-cpu-<vendor>,opencl-gpu-<vendor> \
//! WAI_QUANTUM_OPENCL_EXPECT=opencl-cpu-<vendor id>,opencl-gpu-<vendor id> \
//! cargo test --features quantum_fabric_opencl --test opencl_member -- --test-threads=1
//! ```
#![cfg(not(target_arch = "wasm32"))]

414
wai/wai-quantum/tests/tensor_engine_member.rs Normal file
View file

@ -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=<the kernel toolchain's 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<Discovery> = 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<String, String> {
let mut found: Vec<&str> = d.found.iter().map(|m| m.key()).collect();
found.sort_unstable();
let found_list = if found.is_empty() { "none".to_string() } else { found.join(",") };
let refused: Vec<String> = d.refused.iter().map(|(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<Circuit> {
let gate = |base, controls: Vec<u8>, target, param| Gate { base, controls, target, param };
let bad = [
(1, vec![gate(BaseGate::H, vec![], 0, 0), gate(BaseGate::P, vec![], 0, 0)]),
(1, vec![gate(BaseGate::P, vec![], 0, 33)]),
(1, vec![gate(BaseGate::X, vec![], 3, 0)]),
(2, vec![gate(BaseGate::X, vec![0], 0, 0)]),
(2, vec![gate(BaseGate::X, vec![5], 1, 0)]),
];
let mut out = Vec::new();
for gs in Gateset::ALL {
for (n, ops) in bad.clone() {
let mut c = Circuit::with_gateset(n, gs);
c.ops = ops;
out.push(c);
}
}
out
}
/// `Z` on qubit 0: on `[(0, 0), (i32::MIN, 0)]` its exact result is `2^31`,
/// which leaves 32 bits.
fn z_words() -> GateWords {
lanes::gate_words(&Gate { base: BaseGate::Z, controls: vec![], target: 0, param: 0 }, Gateset::V2)
}
const Z_STATE: [Lane; 2] = [[0, 0], [i32::MIN, 0]];
#[test]
fn 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<Lane> = (0..1usize << n).map(|_| [value(&mut s), value(&mut s)]).collect();
let gates: Vec<GateWords> = (0..1 + next(&mut s) % 8)
.map(|_| {
let t = (next(&mut s) % u64::from(n)) as u32;
let ctrl = (0..n).filter(|&b| b != t && next(&mut s).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
wai/wai-rs/Cargo.lock generated
View file

@ -4300,7 +4300,7 @@ version = "0.1.0"
[[package]]
name = "wai-quantum"
version = "0.5.1"
version = "0.6.0"
dependencies = [
"blake3",
"ed25519-dalek",

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

@ -100,7 +100,7 @@ curve25519-dalek = { version = "4", optional = true }
# ---- quantum (optional, behind the `quantum*` features) ----
# The quantum stack now lives in its own dependency-light crate so it can be used
# without the media/codec chain. `wai` re-exports it unchanged.
wai-quantum = { version = "0.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

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

@ -1231,7 +1231,7 @@ version = "0.1.0"
[[package]]
name = "wai-quantum"
version = "0.5.1"
version = "0.6.0"
dependencies = [
"blake3",
"ed25519-dalek",