diff --git a/.github/workflows/wai-quantum-cuda.yml b/.github/workflows/wai-quantum-cuda.yml
index c7b1718ab..df041e5a1 100644
--- a/.github/workflows/wai-quantum-cuda.yml
+++ b/.github/workflows/wai-quantum-cuda.yml
@@ -8,8 +8,11 @@ name: wai-quantum CUDA member
 #   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.
+#   the tensor-core kernel, every product formed on the tensor cores from 8-bit
+#   limbs, must do all of the same under its own implementation id, with its
+#   own mutants caught; 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.
 # - cuda-absent (a container with no GPU and no CUDA libraries): the same build
 #   finds no member and says why, rather than failing; the device tests run over
 #   no member with the declaration `none`.
@@ -103,6 +106,23 @@ jobs:
           n="$(grep -c "^cuda certified: key=$key .* parity=bit-exact " "$report")"
           [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; }
           grep -qx "cuda mutants: key=$key killed 6 of 6" "$report" || { echo "::error::a kernel mutant survived"; exit 1; }
+      - name: The tensor-core kernel, every product on the tensor cores
+        shell: bash
+        working-directory: src/wai/wai-quantum
+        run: |
+          set -euo pipefail
+          export PATH="$HOME/.cargo/bin:$PATH"
+          report="$RUNNER_TEMP/cuda-tensor-report.txt"; : > "$report"
+          WAI_QUANTUM_CUDA_REPORT="$report" CARGO_TARGET_DIR="$HOME/hw-target/wai-quantum-cuda" nice -n 10 \
+            cargo test --locked --release --features quantum_fabric_cuda --test cuda_tensor -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/tensor.log"
+          grep -qE '^test result: ok\. 9 passed; 0 failed' "$RUNNER_TEMP/tensor.log" || { echo "::error::not every tensor-core test ran and passed"; exit 1; }
+          cat "$report"
+          key="$WAI_QUANTUM_CUDA_EXPECT"
+          grep -qx "cuda tensor devices: as declared ($key)" "$report" || { echo "::error::the devices found are not the declared ones"; exit 1; }
+          grep -q "^cuda tensor member: key=$key .* implementation=cuda-tensor-limbs8/[0-9a-f]\{16\} " "$report" || { echo "::error::the device ran no tensor-core member"; exit 1; }
+          n="$(grep -c "^cuda tensor certified: key=$key .* implementation=cuda-tensor-limbs8/[0-9a-f]\{16\} parity=bit-exact " "$report")"
+          [ "$n" = 4 ] || { echo "::error::$n of 4 tensor-core certificates were earned"; exit 1; }
+          grep -qx "cuda tensor mutants: key=$key killed 7 of 7" "$report" || { echo "::error::a tensor-core mutant survived"; exit 1; }
       - name: The dense f64 state vector, bit for bit
         shell: bash
         working-directory: src/wai/wai-quantum
@@ -146,5 +166,8 @@ jobs:
           WAI_QUANTUM_CUDA_REPORT="$report" cargo test --locked --features quantum_fabric_cuda --test cuda_member -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/device.log"
           grep -qE '^test result: ok\. 9 passed; 0 failed' "$RUNNER_TEMP/device.log" || { echo "::error::not every device test ran and passed"; exit 1; }
           grep -qx 'cuda devices: as declared (none)' "$report" || { echo "::error::the absent case was not established"; exit 1; }
+          WAI_QUANTUM_CUDA_REPORT="$report" cargo test --locked --features quantum_fabric_cuda --test cuda_tensor -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/tensor.log"
+          grep -qE '^test result: ok\. 9 passed; 0 failed' "$RUNNER_TEMP/tensor.log" || { echo "::error::the tensor-core tests did not run over no device"; exit 1; }
+          grep -qx 'cuda tensor devices: as declared (none)' "$report" || { echo "::error::the tensor-core absent case was not established"; exit 1; }
           cargo test --locked --features quantum_fabric_cuda,quantum_sv --test cuda_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; }
diff --git a/.github/workflows/wai-quantum-metal.yml b/.github/workflows/wai-quantum-metal.yml
new file mode 100644
index 000000000..559064178
--- /dev/null
+++ b/.github/workflows/wai-quantum-metal.yml
@@ -0,0 +1,139 @@
+name: wai-quantum Metal member
+
+# The fabric's Metal member: the gate kernels in the Metal shading language,
+# compiled at run time by the operating system's own compiler for the device,
+# and certified bit-exact against the CPU reference's pinned corpus digests.
+#
+# - metal-hw (the macOS runner's 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 matrix kernel, every product formed by the GPU's SIMD-group matrix
+#   multiply from 8-bit limbs, must do all of the same under its own
+#   implementation id, with its own mutants caught.
+# - metal-absent (a container with no Metal): 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-metal.yml'
+  pull_request:
+    paths:
+      - 'wai/wai-quantum/**'
+      - 'rust-toolchain.toml'
+      - '.github/workflows/wai-quantum-metal.yml'
+
+jobs:
+  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-metal.yml \
+                "${{ github.event.pull_request.base.sha }}" \
+                "${{ github.event.pull_request.head.sha }}")
+          echo "this change touches wai-quantum-metal: $r"
+          echo "run=$r" >> "$GITHUB_OUTPUT"
+
+  metal-hw:
+    name: member · Metal (macos)
+    needs: relevant
+    if: needs.relevant.outputs.run == 'true'
+    runs-on: macos
+    timeout-minutes: 90
+    env:
+      WAI_QUANTUM_METAL_EXPECT: gpu-msl-1009
+    steps:
+      - uses: actions/checkout@v5
+      # launchd's PATH omits cargo on the macOS host runner.
+      - name: Put cargo on PATH
+        shell: bash
+        run: echo "$HOME/.cargo/bin" >> "$GITHUB_PATH"
+      - name: The member's unit tests
+        shell: bash
+        working-directory: wai/wai-quantum
+        run: |
+          set -euo pipefail
+          nice -n 10 cargo test --locked --features quantum_fabric_metal --lib -- quantum_fabric_metal:: quantum_fabric::metal:: 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: wai/wai-quantum
+        run: |
+          set -euo pipefail
+          report="$RUNNER_TEMP/metal-report.txt"; : > "$report"
+          WAI_QUANTUM_METAL_REPORT="$report" nice -n 10 \
+            cargo test --locked --release --features quantum_fabric_metal --test metal_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_METAL_EXPECT"
+          grep -qx "metal devices: as declared ($key)" "$report" || { echo "::error::the devices found are not the declared ones"; exit 1; }
+          n="$(grep -c "^metal certified: key=$key .* implementation=metal-lanes64/[0-9a-f]\{16\} parity=bit-exact " "$report")"
+          [ "$n" = 4 ] || { echo "::error::$n of 4 certificates were earned"; exit 1; }
+          grep -qx "metal mutants: key=$key killed 6 of 6" "$report" || { echo "::error::a kernel mutant survived"; exit 1; }
+      - name: The matrix kernel, every product by the SIMD-group matrix multiply
+        shell: bash
+        working-directory: wai/wai-quantum
+        run: |
+          set -euo pipefail
+          report="$RUNNER_TEMP/metal-matrix-report.txt"; : > "$report"
+          WAI_QUANTUM_METAL_REPORT="$report" nice -n 10 \
+            cargo test --locked --release --features quantum_fabric_metal --test metal_matrix -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/matrix.log"
+          grep -qE '^test result: ok\. 9 passed; 0 failed' "$RUNNER_TEMP/matrix.log" || { echo "::error::not every matrix test ran and passed"; exit 1; }
+          cat "$report"
+          key="$WAI_QUANTUM_METAL_EXPECT"
+          grep -qx "metal matrix devices: as declared ($key)" "$report" || { echo "::error::the devices found are not the declared ones"; exit 1; }
+          grep -q "^metal matrix member: key=$key .* implementation=metal-matrix-limbs8/[0-9a-f]\{16\} " "$report" || { echo "::error::the device ran no matrix member"; exit 1; }
+          n="$(grep -c "^metal matrix certified: key=$key .* implementation=metal-matrix-limbs8/[0-9a-f]\{16\} parity=bit-exact " "$report")"
+          [ "$n" = 4 ] || { echo "::error::$n of 4 matrix certificates were earned"; exit 1; }
+          grep -qx "metal matrix mutants: key=$key killed 7 of 7" "$report" || { echo "::error::a matrix mutant survived"; exit 1; }
+
+  metal-absent:
+    name: member · Metal absent (linux)
+    needs: relevant
+    if: needs.relevant.outputs.run == 'true'
+    runs-on: linux
+    timeout-minutes: 60
+    env:
+      WAI_QUANTUM_METAL_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 Metal here, and the member says so
+        shell: bash
+        working-directory: wai/wai-quantum
+        run: |
+          set -euo pipefail
+          cargo test --locked --features quantum_fabric_metal --lib -- quantum_fabric_metal:: quantum_fabric::metal:: 2>&1 | tee "$RUNNER_TEMP/unit.log"
+          grep -qE '^test quantum_fabric_metal::tests::off_macos_nothing_is_found_and_discovery_says_why \.\.\. ok$' "$RUNNER_TEMP/unit.log" || { echo "::error::the off-platform discovery test did not run"; exit 1; }
+          report="$RUNNER_TEMP/metal-report.txt"; : > "$report"
+          for t in metal_member metal_matrix; do
+            WAI_QUANTUM_METAL_REPORT="$report" cargo test --locked --features quantum_fabric_metal --test "$t" -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/$t.log"
+            grep -qE '^test result: ok\. 9 passed; 0 failed' "$RUNNER_TEMP/$t.log" || { echo "::error::$t did not run over no device"; exit 1; }
+          done
+          grep -qx 'metal devices: as declared (none)' "$report" || { echo "::error::the absent case was not established"; exit 1; }
+          grep -qx 'metal matrix devices: as declared (none)' "$report" || { echo "::error::the matrix absent case was not established"; exit 1; }
diff --git a/wai/CHANGELOG.md b/wai/CHANGELOG.md
index 4b8cbc504..831ec7541 100644
--- a/wai/CHANGELOG.md
+++ b/wai/CHANGELOG.md
@@ -6,6 +6,60 @@ wire format may change and such changes are recorded here.
 
 ## v0.4.0 — current
 
+### wai-quantum 0.6.1: the CUDA member's tensor-core kernel, and the Metal member
+
+Additive. Every pinned digest is unchanged.
+
+- **Every product on the tensor cores** (`quantum_fabric::cuda::TENSOR_SOURCE`,
+  `tensor_limbs.cu`; `CudaMember::tensor`): a second CUDA kernel for the same gate, for
+  devices of compute capability 8.0 or later (`DeviceIdentity::has_tensor_cores`). Each
+  32-bit value is four 8-bit limbs in bf16, the lower three unsigned and the top one
+  signed. One warp multiplies sixteen pairs' limbs by the gate's limb weights (sixteen
+  products per pair, seven coefficients each) in fourteen `m16n8k16` steps, accumulated in
+  f32. Every limb product is at most `2^16` and every coefficient sums at most four, so f32
+  holds every partial sum exactly in any order. The kernel adds the coefficients, shifted
+  into place, in 64-bit integers; the rest is the other kernel's arithmetic on exact
+  products. The member id is the device's; the implementation id is the kernel's own
+  (`cuda-tensor-limbs8/…`), and the member is probed before it is returned. A device below
+  8.0 is `CudaError::Unsupported` (new). Seven mutants (`TENSOR_MUTANTS`), one of which is
+  nine-bit limbs, which bf16 does not hold.
+- **Measured** on the discrete GPU (compute capability 8.9, silicon): 66020 probe pairs
+  agree, all four pinned certificates, circuits to 22 qubits match native, 7 of 7 mutants
+  killed. The other kernel and the state vector are unchanged there (all four pins, 6 of 6
+  and 4 of 4 mutants). CI: a step of `wai-quantum-cuda`; its absent leg runs the new tests
+  over no device.
+- **The Metal member** (`quantum_fabric_metal`, opt-in, macOS): the gate on the platform's
+  GPU through its own interface, compiled at run time by the operating system's compiler.
+  Two kernels, each with its own implementation id and mutants:
+  - `lanes64.metal` (`metal-lanes64/…`): the CUDA kernel's exact 64-bit integer arithmetic
+    in the Metal shading language, one serial dispatch per gate in one command buffer, the
+    state in memory the GPU and the host share;
+  - `matrix_limbs.metal` (`metal-matrix-limbs8/…`, `MetalMember::matrix`): every product
+    formed by the GPU's SIMD-group matrix multiply, from 8-bit limbs in bf16 accumulated in
+    f32, by the tensor-core kernel's argument. One SIMD group multiplies eight pairs' limbs
+    by the gate's limb weights in fourteen 8 × 8 column blocks, and each thread joins its
+    four products and forms one row. A device below GPU family 1007 is
+    `MetalError::Unsupported`.
+
+  A device is keyed by its GPU family number (`gpu-msl-<family>`), and its member id binds
+  the device, its registry id, family and memory, and the operating system's version. Off
+  macOS the member builds, finds nothing and says why.
+- **Measured** on the macOS runner's GPU (family 1009, silicon): both kernels 9/9, each
+  with 66020 probe pairs, all four pinned certificates and circuits to 22 qubits; 6 of 6
+  and 7 of 7 mutants killed. Nine-bit limbs break 297 of the 484 edge pairs, so the
+  multiply takes bf16 inputs. CI: `wai-quantum-metal`, on that runner, and its absent leg
+  on Linux.
+- **Measured by hand, outside CI:**
+  - the gate as a portable tensor program, through XLA on a cloud tensor processor: the
+    gate's module exact and its three mutants failing where they should, and the limb
+    program (`stablehlo::limb_module`) exact on the processor's matrix unit, with the
+    nine-bit set differing as it must;
+  - `wai-quantum-qpu`'s `native_turns` against the service's simulator, no hardware used:
+    with no noise, statistical parity at fidelity 1.0 on both gatesets' corpora (16
+    circuits of 5 qubits, 1000 shots each), which checks the compilation to the service's
+    native gates and the wire contract; with one of the service's device noise models,
+    fidelity 0.909 (V1) and 0.927 (V2).
+
 ### `wai.world.model`'s integer renderer through the integer operator set
 
 - **The lowering, from the models the renderer holds.** `wai-ops` gains
diff --git a/wai/tools/quantum-meter/Cargo.lock b/wai/tools/quantum-meter/Cargo.lock
index 1f4435f1b..2f4f18e60 100644
--- a/wai/tools/quantum-meter/Cargo.lock
+++ b/wai/tools/quantum-meter/Cargo.lock
@@ -1017,7 +1017,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "ed25519-dalek",
diff --git a/wai/tools/quantum-wgsl-check/Cargo.lock b/wai/tools/quantum-wgsl-check/Cargo.lock
index ac48ce6b7..3ca21b7b3 100644
--- a/wai/tools/quantum-wgsl-check/Cargo.lock
+++ b/wai/tools/quantum-wgsl-check/Cargo.lock
@@ -599,7 +599,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "ed25519-dalek",
diff --git a/wai/wai-pack/Cargo.lock b/wai/wai-pack/Cargo.lock
index d8c710032..69503345d 100644
--- a/wai/wai-pack/Cargo.lock
+++ b/wai/wai-pack/Cargo.lock
@@ -713,7 +713,7 @@ dependencies = [
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "ed25519-dalek",
diff --git a/wai/wai-quantum-qpu/Cargo.lock b/wai/wai-quantum-qpu/Cargo.lock
index 8864c9a27..f9ac5d194 100644
--- a/wai/wai-quantum-qpu/Cargo.lock
+++ b/wai/wai-quantum-qpu/Cargo.lock
@@ -422,7 +422,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "ed25519-dalek",
diff --git a/wai/wai-quantum-serve/Cargo.lock b/wai/wai-quantum-serve/Cargo.lock
index ee89e0f5b..817232a82 100644
--- a/wai/wai-quantum-serve/Cargo.lock
+++ b/wai/wai-quantum-serve/Cargo.lock
@@ -599,7 +599,7 @@ dependencies = [
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "ed25519-dalek",
diff --git a/wai/wai-quantum/Cargo.lock b/wai/wai-quantum/Cargo.lock
index ad3d6b295..e5fb779b2 100644
--- a/wai/wai-quantum/Cargo.lock
+++ b/wai/wai-quantum/Cargo.lock
@@ -1079,7 +1079,7 @@ checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "block2",
@@ -1088,6 +1088,7 @@ dependencies = [
  "objc2",
  "objc2-core-ml",
  "objc2-foundation",
+ "objc2-metal",
  "opencl3",
  "serde_json",
  "wgpu",
diff --git a/wai/wai-quantum/Cargo.toml b/wai/wai-quantum/Cargo.toml
index 26a026636..b60b10869 100644
--- a/wai/wai-quantum/Cargo.toml
+++ b/wai/wai-quantum/Cargo.toml
@@ -1,6 +1,6 @@
 [package]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 edition = "2024"
 rust-version = "1.97"
 license = "Apache-2.0"
@@ -47,13 +47,19 @@ opencl3 = { version = "0.12.3", optional = true, default-features = false, featu
 [target.'cfg(target_os = "macos")'.dependencies]
 objc2 = { version = "0.6.5", optional = true }
 block2 = { version = "0.6.2", optional = true }
-objc2-foundation = { version = "0.3.2", optional = true, default-features = false, features = ["std", "NSArray", "NSDictionary", "NSError", "NSString", "NSURL", "NSValue"] }
+objc2-foundation = { version = "0.3.2", optional = true, default-features = false, features = ["std", "NSArray", "NSDictionary", "NSError", "NSProcessInfo", "NSString", "NSURL", "NSValue"] }
 objc2-core-ml = { version = "0.3.2", optional = true, default-features = false, features = [
     "std", "block2", "MLModel", "MLModel_MLModelCompilation", "MLModelConfiguration", "MLMultiArray", "MLDictionaryFeatureProvider",
     "MLFeatureProvider", "MLFeatureValue", "MLComputePlan", "MLComputePlanDeviceUsage", "MLComputeDeviceProtocol",
     "MLNeuralEngineComputeDevice", "MLModelStructure", "MLModelStructureProgram", "MLModelStructureProgramFunction",
     "MLModelStructureProgramBlock", "MLModelStructureProgramOperation",
 ] }
+# The Metal member: the platform's GPU interface, its kernels compiled at run
+# time by the system's own compiler.
+objc2-metal = { version = "0.3.2", optional = true, default-features = false, features = [
+    "std", "MTLAllocation", "MTLBuffer", "MTLCommandBuffer", "MTLCommandEncoder", "MTLCommandQueue", "MTLComputeCommandEncoder",
+    "MTLComputePipeline", "MTLDevice", "MTLLibrary", "MTLPipeline", "MTLResource", "MTLTypes",
+] }
 
 [target.'cfg(all(target_arch = "wasm32", target_os = "unknown"))'.dependencies]
 wgpu = { version = "30.0.1", optional = true, default-features = false, features = ["std", "wgsl", "webgpu"] }
@@ -169,6 +175,7 @@ quantum_fabric_opencl = ["quantum_fabric", "dep:opencl3"]
 quantum_fabric_fpga = ["quantum_fabric"]
 # The neural-processor member: exact fp16 limb products on the platform's neural processor.
 quantum_fabric_npu = ["quantum_fabric", "dep:objc2", "dep:block2", "dep:objc2-foundation", "dep:objc2-core-ml"]
+quantum_fabric_metal = ["quantum_fabric", "dep:objc2", "dep:objc2-foundation", "dep:objc2-metal"]
 # 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.
@@ -207,6 +214,18 @@ required-features = ["quantum_fabric_webgpu"]
 name = "cuda_member"
 required-features = ["quantum_fabric_cuda"]
 
+[[test]]
+name = "cuda_tensor"
+required-features = ["quantum_fabric_cuda"]
+
+[[test]]
+name = "metal_member"
+required-features = ["quantum_fabric_metal"]
+
+[[test]]
+name = "metal_matrix"
+required-features = ["quantum_fabric_metal"]
+
 [[test]]
 name = "cuda_sv"
 required-features = ["quantum_fabric_cuda", "quantum_sv"]
diff --git a/wai/wai-quantum/README.md b/wai/wai-quantum/README.md
index b5df4ad0e..a0ade27d6 100644
--- a/wai/wai-quantum/README.md
+++ b/wai/wai-quantum/README.md
@@ -105,7 +105,8 @@ never creates information.
 | `quantum_vml`, `quantum_phasor`, `quantum_qfhrr`, `quantum_kernel`, `quantum_qdata`, `quantum_phasor_meter` | the interference / phasor ML layer |
 | `quantum_fabric` | **the compute fabric**: members (the CPU and accelerators, quantum processors among them), each kernel implementation certified against the CPU reference in the process that uses it (bit-exact, within a stated tolerance, or statistical by classical fidelity over stated shots), with the certificate saying where it was earned (silicon, an emulator, or a simulator); placement by parity, provenance and metered energy; a signed receipt naming the member and implementation behind a result. Two kernels: apply (a circuit's state) and sample (its measurement counts), the latter with the CPU reference sampler, a corpus shaped for quantum processors, and a virtual quantum processor whose stated noise model lets every step of a processor's path run and be tested without one; its certificates say they were earned in a simulator. `quantum_fabric::stablehlo` writes the gate as a portable tensor program (StableHLO), every value a 64-bit integer that never overflows, for any tensor compiler to build for its own devices; the `stablehlo_gate` example writes it with a batch of inputs and the exact outputs a build must give, 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_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. On a device of compute capability 8.0 or later, a **second kernel runs every product on the tensor cores** (`CudaMember::tensor`): each value as four 8-bit limbs in bf16, accumulated in f32, which holds every partial sum exactly; probed and certified under its own implementation id, with its own mutants. 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_metal` | **the fabric's Metal member**, opt-in and not in `full`, macOS: the gate kernel in the Metal shading language with exact 64-bit integers, compiled at run time by the operating system's own compiler, one serial dispatch per gate in one command buffer, the state in memory the GPU and the host share. Elsewhere the build finds no member and says 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, keyed by its GPU family number (`gpu-msl-<family>`). On a GPU with SIMD-group matrices, a **second kernel forms every product with the matrix multiply** (`MetalMember::matrix`): each value as four 8-bit limbs in bf16, accumulated in f32, which holds every partial sum exactly; probed and certified under its own implementation id, with its own mutants |
 | `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 |
diff --git a/wai/wai-quantum/src/lib.rs b/wai/wai-quantum/src/lib.rs
index 7dd1c366c..1d56e98c3 100644
--- a/wai/wai-quantum/src/lib.rs
+++ b/wai/wai-quantum/src/lib.rs
@@ -10,13 +10,15 @@
 //! 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. Twelve opt-in features reach outside the crate:
+//! member that ran the work. These opt-in features reach outside the crate:
 //! `quantum_fabric_webgpu` (the platform's GPU interface),
 //! `quantum_fabric_cuda` (the CUDA driver and runtime compiler),
 //! `quantum_fabric_hip` (the HIP runtime and runtime compiler) and
 //! `quantum_fabric_opencl` (the OpenCL ICD loader), the last three loaded when
 //! first used, `quantum_fabric_fpga` (Verilator, run as a program when first
 //! used), `quantum_fabric_npu` (the platform's model runtime, on macOS),
+//! `quantum_fabric_metal` (the platform's GPU interface and its runtime
+//! compiler, 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
@@ -104,6 +106,8 @@ pub mod quantum_fabric_opencl;
 pub mod quantum_fabric_fpga;
 #[cfg(all(feature = "quantum_fabric_npu", not(target_arch = "wasm32")))]
 pub mod quantum_fabric_npu;
+#[cfg(all(feature = "quantum_fabric_metal", not(target_arch = "wasm32")))]
+pub mod quantum_fabric_metal;
 #[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")))]
diff --git a/wai/wai-quantum/src/quantum_fabric.rs b/wai/wai-quantum/src/quantum_fabric.rs
index eb9022544..eea7dc12f 100644
--- a/wai/wai-quantum/src/quantum_fabric.rs
+++ b/wai/wai-quantum/src/quantum_fabric.rs
@@ -2048,6 +2048,7 @@ pub mod lanes {
 }
 
 pub mod cuda;
+pub mod metal;
 pub mod hip;
 pub mod fpga;
 pub mod npu;
diff --git a/wai/wai-quantum/src/quantum_fabric/cuda.rs b/wai/wai-quantum/src/quantum_fabric/cuda.rs
index 67eaec1e0..ccdd88761 100644
--- a/wai/wai-quantum/src/quantum_fabric/cuda.rs
+++ b/wai/wai-quantum/src/quantum_fabric/cuda.rs
@@ -24,6 +24,29 @@
 //! device is a member, and its certificates check whole runs against the CPU's
 //! pinned corpus digests.
 //!
+//! **The tensor-core kernel** ([`TENSOR_SOURCE`], `tensor_limbs.cu`) is a
+//! second implementation of the same gate, for devices of compute capability
+//! [`TENSOR_MIN_COMPUTE`] or later, whose every product is formed on the
+//! tensor cores, bf16 inputs accumulated in f32. Each 32-bit value is split
+//! into four 8-bit limbs, the lower three unsigned and the top one signed, so:
+//!
+//! - every limb is an integer of at most `2^8` in size, which bf16 holds
+//!   exactly (8 significant bits);
+//! - every limb product is at most `2^16` in size, and every coefficient of a
+//!   product, `c_s = Σ_{i+j=s} e_i·a_j`, sums at most four of them, so it is
+//!   an integer below `2^18`: f32 holds every partial sum exactly, in whatever
+//!   order the tensor core adds them;
+//! - the host side of the kernel adds each coefficient, shifted by `8s` bits,
+//!   in 64-bit two's complement, which gives the exact product `e·a`; the
+//!   rest is [`SOURCE`]'s arithmetic on exact products.
+//!
+//! One warp multiplies sixteen pairs' limbs (16 × 16) by the gate's limb
+//! weights (16 × 112: sixteen products, seven coefficients each) in fourteen
+//! `m16n8k16` steps; the layout is the tensor engine's
+//! ([`tensor_engine`](super::tensor_engine)). It has its own implementation id
+//! ([`tensor_implementation_id`]) and its own mutants ([`TENSOR_MUTANTS`]),
+//! one of which, nine-bit limbs, bf16 cannot hold.
+//!
 //! [`lanes::exact_row`]: super::lanes
 //! [`lanes::apply_words`]: super::lanes::apply_words
 //! [`lanes::run_words`]: super::lanes::run_words
@@ -183,6 +206,87 @@ pub fn mutate(m: &Mutant) -> String {
     SOURCE.replacen(m.find, m.replace, 1)
 }
 
+/// The tensor-core kernel, `tensor_limbs.cu`, exactly as it is compiled.
+pub const TENSOR_SOURCE: &str = include_str!("tensor_limbs.cu");
+
+/// The tensor-core kernel's implementation name.
+pub const TENSOR_IMPLEMENTATION: &str = "cuda-tensor-limbs8";
+
+/// Version of the tensor-core kernel's host-side layout: [`HOST_LAYOUT`]'s,
+/// except that a warp takes sixteen pairs (or probe cases), so a launch has
+/// [`TENSOR_THREADS_PER_ITEM`] threads per item. Part of its implementation
+/// id.
+pub const TENSOR_HOST_LAYOUT: u32 = 1;
+
+/// Threads per pair (or probe case) of a tensor-core launch: 32 lanes take 16.
+pub const TENSOR_THREADS_PER_ITEM: u64 = 2;
+
+/// The compute capability the tensor-core kernel needs at least: bf16
+/// `mma.sync` at `m16n8k16`.
+pub const TENSOR_MIN_COMPUTE: (u32, u32) = (8, 0);
+
+const DOMAIN_TENSOR_IMPLEMENTATION: &[u8] = b"wai:quantum-fabric-cuda-tensor\x01";
+
+/// `"cuda-tensor-limbs8/<16 hex>"`: as [`implementation_id`], with the domain
+/// `"wai:quantum-fabric-cuda-tensor\x01"`, [`TENSOR_SOURCE`] and
+/// [`TENSOR_HOST_LAYOUT`].
+pub fn tensor_implementation_id() -> String {
+    tensor_implementation_id_of(TENSOR_SOURCE)
+}
+
+/// [`tensor_implementation_id`] for another kernel text, such as a mutant's.
+pub fn tensor_implementation_id_of(source: &str) -> String {
+    let mut o = DOMAIN_TENSOR_IMPLEMENTATION.to_vec();
+    put(&mut o, source.as_bytes());
+    for option in COMPILE_OPTIONS {
+        put(&mut o, option.as_bytes());
+    }
+    o.extend_from_slice(&TENSOR_HOST_LAYOUT.to_be_bytes());
+    format!("{TENSOR_IMPLEMENTATION}/{}", hex(&blake3::hash(&o).as_bytes()[..8]))
+}
+
+impl DeviceIdentity {
+    /// Whether the device can run the tensor-core kernel: compute capability
+    /// [`TENSOR_MIN_COMPUTE`] or later.
+    pub fn has_tensor_cores(&self) -> bool {
+        self.compute >= TENSOR_MIN_COMPUTE
+    }
+}
+
+/// The tensor-core kernel's mutants, as [`MUTANTS`] are the other's. Each
+/// `find` occurs exactly once in [`TENSOR_SOURCE`].
+pub const TENSOR_MUTANTS: [Mutant; 7] = [
+    // The rounding constant is half what it should be.
+    Mutant { name: "T1", find: "return (v + ((i64)1 << 29)) >> 30;", replace: "return (v + ((i64)1 << 28)) >> 30;", detector: Detector::Probe },
+    // The product is shifted by one bit too few.
+    Mutant { name: "T2", find: "return (v + ((i64)1 << 29)) >> 30;", replace: "return (v + ((i64)1 << 29)) >> 29;", detector: Detector::Probe },
+    // Nine-bit limbs: a limb above 2^8 has nine significant bits, which bf16
+    // does not hold, so the products are not exact.
+    Mutant { name: "T3", find: "#define LIMB 8", replace: "#define LIMB 9", detector: Detector::Probe },
+    // The second half of the B fragment is loaded with its two rows swapped.
+    Mutant {
+        name: "T4",
+        find: "u32 b1 = pack(W[2 * t + 8][col], W[2 * t + 9][col]);",
+        replace: "u32 b1 = pack(W[2 * t + 9][col], W[2 * t + 8][col]);",
+        detector: Detector::Probe,
+    },
+    // The control test is inverted.
+    Mutant { name: "T5", find: "if ((i & (u64)ctrl) == (u64)ctrl) {", replace: "if ((i & (u64)ctrl) != (u64)ctrl) {", detector: Detector::Certificate },
+    // A sign of row 0's real part is wrong, on the imaginary part of the
+    // top-right entry (as C5).
+    Mutant { name: "T6", find: "fx(pr[2]) - fx(pr[3]);", replace: "fx(pr[2]) + fx(pr[3]);", detector: Detector::Certificate },
+    // The range guard never fires.
+    Mutant { name: "T7", find: "if (!(fits32(o0r) && fits32(o0i) && fits32(o1r) && fits32(o1i))) {", replace: "if (false) {", detector: Detector::Range },
+];
+
+/// [`TENSOR_SOURCE`] with `m` applied. Panics unless `m.find` occurs exactly
+/// once.
+pub fn tensor_mutate(m: &Mutant) -> String {
+    let n = TENSOR_SOURCE.matches(m.find).count();
+    assert_eq!(n, 1, "mutant {}: `{}` occurs {n} times in the tensor-core kernel", m.name, m.find);
+    TENSOR_SOURCE.replacen(m.find, m.replace, 1)
+}
+
 #[cfg(test)]
 mod tests {
     use super::*;
@@ -245,4 +349,32 @@ mod tests {
             assert!(MUTANTS.iter().any(|m| m.detector == d), "no mutant for {d:?}");
         }
     }
+
+    #[test]
+    fn the_tensor_kernel_has_its_own_id_and_mutants() {
+        let id = tensor_implementation_id();
+        assert!(id.starts_with("cuda-tensor-limbs8/") && id.len() == "cuda-tensor-limbs8/".len() + 16, "{id}");
+        assert_ne!(tensor_implementation_id_of(SOURCE), implementation_id_of(SOURCE), "the domains differ");
+        let mut seen = std::collections::BTreeSet::new();
+        for m in &TENSOR_MUTANTS {
+            let edited = tensor_mutate(m);
+            assert_ne!(edited, TENSOR_SOURCE, "mutant {} changes nothing", m.name);
+            assert_ne!(tensor_implementation_id_of(&edited), id, "mutant {}", m.name);
+            assert!(seen.insert(edited), "mutant {} duplicates another", m.name);
+        }
+        for d in [Detector::Probe, Detector::Certificate, Detector::Range] {
+            assert!(TENSOR_MUTANTS.iter().any(|m| m.detector == d), "no tensor mutant for {d:?}");
+        }
+        // The kernel's entry points and parameters are the other kernel's.
+        for entry in ["extern \"C\" __global__ void apply_gate(", "extern \"C\" __global__ void probe("] {
+            assert_eq!((SOURCE.matches(entry).count(), TENSOR_SOURCE.matches(entry).count()), (1, 1), "{entry}");
+        }
+    }
+
+    #[test]
+    fn tensor_cores_start_at_compute_8_0() {
+        for (compute, has) in [((7, 5), false), ((7, 0), false), ((8, 0), true), ((8, 6), true), ((8, 9), true), ((9, 0), true), ((12, 0), true)] {
+            assert_eq!(DeviceIdentity { compute, ..identity() }.has_tensor_cores(), has, "{compute:?}");
+        }
+    }
 }
diff --git a/wai/wai-quantum/src/quantum_fabric/lanes64.metal b/wai/wai-quantum/src/quantum_fabric/lanes64.metal
new file mode 100644
index 000000000..80032f848
--- /dev/null
+++ b/wai/wai-quantum/src/quantum_fabric/lanes64.metal
@@ -0,0 +1,104 @@
+// wai-quantum fabric: the gate kernel for the platform's GPU, in the Metal
+// shading language.
+//
+// A transcription of `lanes::exact_row` and `lanes::apply_words` in 64-bit
+// integers, as lanes64.cu is. Amplitudes are pairs of 32-bit fixed-point
+// values (30 fractional bits). Each product `fx(a, b) = (a·b + 2^29) >> 30`
+// is exact in 64 bits (a 32×32 product is at most 2^62 in size), and so is
+// each row's sum of four products. The lane stores the low 32 bits of each
+// sum; a sum outside 32 bits records the gate's op-log index in `report`, the
+// lowest such index winning, so the host refuses the run with the gate the
+// CPU reference names.
+//
+// Nothing here is floating point, depends on the order threads run in, or
+// reads an amplitude another thread of the same dispatch writes: each thread
+// owns one pair (i, j), and a dispatch is one gate. The host encodes a run's
+// gates as serial dispatches of one encoder, so each sees the last one's
+// writes.
+//
+// `>>` on a negative signed value is arithmetic here, as Rust's is on i64;
+// the host's probe checks this, among the rest of the arithmetic, before a
+// device is a member.
+
+#include <metal_stdlib>
+using namespace metal;
+
+// One gate, as the host writes it: 32-bit words only, so that both sides lay
+// it out the same. `m` are the gate's words: row 0 is (m0 + i·m1)·a0 +
+// (m2 + i·m3)·a1, row 1 is (m4 + i·m5)·a0 + (m6 + i·m7)·a1. `tbit` is
+// 1 << target; `ctrl` the OR of the controls' bits; `index` the gate's
+// position in the op log. The host checks the words before a dispatch.
+struct Gate {
+    uint pairs_lo;
+    uint pairs_hi;
+    uint tbit;
+    uint ctrl;
+    int m[8];
+    uint index;
+};
+
+// The exact fixed-point product.
+static inline long fx(int a, int b) {
+    return ((long)a * (long)b + (1L << 29)) >> 30;
+}
+
+// The low 32 bits of `v`, as the lane keeps them.
+static inline int low32(long v) {
+    return (int)(uint)(ulong)v;
+}
+
+// Whether `v` fits a 32-bit lane.
+static inline bool fits32(long v) {
+    return v >= -2147483648L && v <= 2147483647L;
+}
+
+// One gate on a state of `2 * pairs` lanes, one thread per pair, striding.
+kernel void apply_gate(
+    device int2* amps [[buffer(0)]],
+    constant Gate& g [[buffer(1)]],
+    device atomic_uint* report [[buffer(2)]],
+    uint gid [[thread_position_in_grid]],
+    uint grid [[threads_per_grid]])
+{
+    ulong pairs = (ulong)g.pairs_lo | ((ulong)g.pairs_hi << 32);
+    for (ulong p = gid; p < pairs; p += grid) {
+        // The p-th index whose target bit is clear: p with a zero inserted at
+        // the target's position.
+        ulong low = p & ((ulong)g.tbit - 1);
+        ulong i = ((p - low) << 1) | low;
+        if ((i & (ulong)g.ctrl) != (ulong)g.ctrl) {
+            continue;
+        }
+        ulong j = i | (ulong)g.tbit;
+        int2 a0 = amps[i];
+        int2 a1 = amps[j];
+        long o0r = fx(g.m[0], a0.x) - fx(g.m[1], a0.y) + fx(g.m[2], a1.x) - fx(g.m[3], a1.y);
+        long o0i = fx(g.m[0], a0.y) + fx(g.m[1], a0.x) + fx(g.m[2], a1.y) + fx(g.m[3], a1.x);
+        long o1r = fx(g.m[4], a0.x) - fx(g.m[5], a0.y) + fx(g.m[6], a1.x) - fx(g.m[7], a1.y);
+        long o1i = fx(g.m[4], a0.y) + fx(g.m[5], a0.x) + fx(g.m[6], a1.y) + fx(g.m[7], a1.x);
+        if (!(fits32(o0r) && fits32(o0i) && fits32(o1r) && fits32(o1i))) {
+            atomic_fetch_min_explicit(report, g.index, memory_order_relaxed);
+        }
+        amps[i] = int2(low32(o0r), low32(o0i));
+        amps[j] = int2(low32(o1r), low32(o1i));
+    }
+}
+
+// The arithmetic probe: for each case (a, b), the 64-bit two's-complement
+// product a·b and the exact fx(a, b), each as (high, low) words, the outputs
+// of `lanes::mul_wide` and `lanes::fxmul_wide`. `count` is (low, high) words.
+kernel void probe(
+    device const int2* cases [[buffer(0)]],
+    device uint4* out [[buffer(1)]],
+    constant uint2& count [[buffer(2)]],
+    uint gid [[thread_position_in_grid]],
+    uint grid [[threads_per_grid]])
+{
+    ulong n = (ulong)count.x | ((ulong)count.y << 32);
+    for (ulong k = gid; k < n; k += grid) {
+        int2 c = cases[k];
+        ulong prod = (ulong)((long)c.x * (long)c.y);
+        ulong f = (ulong)fx(c.x, c.y);
+        out[k] = uint4((uint)(prod >> 32), (uint)prod, (uint)(f >> 32), (uint)f);
+    }
+}
diff --git a/wai/wai-quantum/src/quantum_fabric/matrix_limbs.metal b/wai/wai-quantum/src/quantum_fabric/matrix_limbs.metal
new file mode 100644
index 000000000..d74b43d10
--- /dev/null
+++ b/wai/wai-quantum/src/quantum_fabric/matrix_limbs.metal
@@ -0,0 +1,227 @@
+// wai-quantum fabric: the gate kernel for the platform's GPU whose every
+// product is formed by its SIMD-group matrix multiply, bf16 inputs
+// accumulated in f32.
+//
+// The same entry points as lanes64.metal, with the same parameters, so the
+// member runs either one. Each 32-bit value is split into four limbs of 8
+// bits (the lower three unsigned, the top one signed), which bf16 holds
+// 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, every limb product at most
+// 2^16 in size, every coefficient below 2^18. So every value the multiply
+// forms is an integer f32 holds exactly, in whatever order it adds them
+// (quantum_fabric::tensor_engine has the argument and the layout). The limb
+// weights put each of a gate's sixteen products per pair in its own seven
+// columns of a 16 × 112 matrix; one SIMD group multiplies an 8-pair tile of
+// limbs (8 × 16) by it in fourteen 8 × 8 column blocks of two steps each, and
+// each thread joins its four products' coefficients, shifted into place, in
+// 64-bit integers. The rest is lanes64.metal's arithmetic on those exact
+// products, one row per thread: fx = (product + 2^29) >> 30, the row the sum
+// of four, the range reported and the low 32 bits kept.
+//
+// The four threads of a pair each read one component of its amplitudes and
+// write the row that replaces it, after the SIMD group's barrier; a dispatch
+// is one gate. So no thread reads what another writes within a dispatch.
+
+#include <metal_stdlib>
+#include <metal_simdgroup_matrix>
+using namespace metal;
+
+// SIMD groups per threadgroup: the member dispatches 128 threads a
+// threadgroup, in SIMD groups of 32.
+#define SIMDGROUPS 4
+
+// Bits per limb. Eight is what bf16 holds exactly; the radix is a mutant's
+// one edit, and nine-bit limbs are not exact.
+#define LIMB 8
+
+typedef bfloat limb_t;
+
+// As lanes64.metal's.
+struct Gate {
+    uint pairs_lo;
+    uint pairs_hi;
+    uint tbit;
+    uint ctrl;
+    int m[8];
+    uint index;
+};
+
+// The sixteen products, as (entry, component): rows 0r, 0i, 1r and 1i, four
+// terms each, in the order lanes::exact_row adds them. Component c of a pair
+// is a0.x, a0.y, a1.x, a1.y for c = 0, 1, 2, 3.
+constant int TERM_ENTRY[16] = {0, 1, 2, 3, 0, 1, 2, 3, 4, 5, 6, 7, 4, 5, 6, 7};
+constant int TERM_COMP[16] = {0, 1, 2, 3, 1, 0, 3, 2, 0, 1, 2, 3, 1, 0, 3, 2};
+
+// Limb j of x: the lower three unsigned, the top one the signed rest.
+static inline int limb(int x, int j) {
+    return j < 3 ? (int)(((uint)x >> (LIMB * j)) & ((1u << LIMB) - 1)) : (x >> (3 * LIMB));
+}
+
+// v as a matrix input. An integer of at most 256 in size has at most 8
+// significant bits, so bf16 holds it exactly.
+static inline limb_t as_limb(int v) {
+    return limb_t(float(v));
+}
+
+// The exact fixed-point rounding of a product.
+static inline long fx(long v) {
+    return (v + (1L << 29)) >> 30;
+}
+
+static inline int low32(long v) {
+    return (int)(uint)(ulong)v;
+}
+
+static inline bool fits32(long v) {
+    return v >= -2147483648L && v <= 2147483647L;
+}
+
+// D (8 × 112, f32) = A (8 × 16 limbs) · W (16 × 112), row-major, in fourteen
+// 8 × 8 column blocks. Called by all 32 threads of a SIMD group.
+static void products(threadgroup const limb_t* A, threadgroup const limb_t* W, threadgroup float* D) {
+    simdgroup_matrix<limb_t, 8, 8> a0, a1, w0, w1;
+    simdgroup_load(a0, A, 16);
+    simdgroup_load(a1, A + 8, 16);
+    for (int nb = 0; nb < 14; nb++) {
+        simdgroup_matrix<float, 8, 8> d = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
+        simdgroup_load(w0, W + nb * 8, 112);
+        simdgroup_load(w1, W + 8 * 112 + nb * 8, 112);
+        simdgroup_multiply_accumulate(d, a0, w0, d);
+        simdgroup_multiply_accumulate(d, a1, w1, d);
+        simdgroup_store(d, D + nb * 8, 112);
+    }
+}
+
+// Product q of row r of D: its seven coefficients, each an integer below
+// 2^24 in size (so the conversion is exact), shifted into place and added in
+// 64-bit two's complement.
+static long assemble(threadgroup const float* D, int r, int q) {
+    ulong v = 0;
+    for (int s = 0; s < 7; s++) {
+        v += (ulong)(long)D[r * 112 + q * 7 + s] << (LIMB * s);
+    }
+    return (long)v;
+}
+
+// One gate on a state of `2 * pairs` lanes, eight pairs a SIMD group, four
+// threads a pair. The parameters are lanes64.metal's, with `amps` as words.
+kernel void apply_gate(
+    device int* amps [[buffer(0)]],
+    constant Gate& g [[buffer(1)]],
+    device atomic_uint* report [[buffer(2)]],
+    uint tid [[thread_index_in_threadgroup]],
+    uint group [[threadgroup_position_in_grid]],
+    uint groups [[threadgroups_per_grid]],
+    uint sg [[simdgroup_index_in_threadgroup]],
+    uint lane [[thread_index_in_simdgroup]],
+    uint width [[threads_per_threadgroup]],
+    uint simd_width [[threads_per_simdgroup]])
+{
+    threadgroup limb_t W[16 * 112];
+    threadgroup limb_t A[SIMDGROUPS][8 * 16];
+    threadgroup float D[SIMDGROUPS][8 * 112];
+    if (width != SIMDGROUPS * 32 || simd_width != 32) {
+        return;
+    }
+    // The gate's limb weights, once a threadgroup: row c·4 + j is limb j of
+    // component c; column p·7 + s takes limb s − j of product p's entry.
+    for (uint x = tid; x < 16 * 112; x += width) {
+        int k = x / 112, col = x % 112;
+        int c = k / 4, j = k % 4, p = col / 7, i = col % 7 - j;
+        W[x] = (c == TERM_COMP[p] && i >= 0 && i <= 3) ? as_limb(limb(g.m[TERM_ENTRY[p]], i)) : limb_t(0.0f);
+    }
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    ulong pairs = (ulong)g.pairs_lo | ((ulong)g.pairs_hi << 32);
+    ulong tiles = (pairs + 7) / 8;
+    ulong stride = (ulong)groups * SIMDGROUPS;
+    int r = lane / 4, c = lane % 4;
+    for (ulong tile = (ulong)group * SIMDGROUPS + sg; tile < tiles; tile += stride) {
+        ulong p = tile * 8 + r;
+        bool active = false;
+        ulong at = 0;
+        int v = 0;
+        if (p < pairs) {
+            // The p-th index whose target bit is clear.
+            ulong low = p & ((ulong)g.tbit - 1);
+            ulong i = ((p - low) << 1) | low;
+            if ((i & (ulong)g.ctrl) == (ulong)g.ctrl) {
+                active = true;
+                ulong j = i | (ulong)g.tbit;
+                // Component c: a0.x, a0.y, a1.x, a1.y.
+                at = c < 2 ? 2 * i + c : 2 * j + (c - 2);
+                v = amps[at];
+            }
+        }
+        for (int jj = 0; jj < 4; jj++) {
+            A[sg][r * 16 + c * 4 + jj] = as_limb(limb(v, jj));
+        }
+        simdgroup_barrier(mem_flags::mem_threadgroup);
+        products(A[sg], W, D[sg]);
+        simdgroup_barrier(mem_flags::mem_threadgroup);
+        if (active) {
+            // Row c of the pair: products 4c … 4c + 3.
+            long f0 = fx(assemble(D[sg], r, c * 4));
+            long f1 = fx(assemble(D[sg], r, c * 4 + 1));
+            long f2 = fx(assemble(D[sg], r, c * 4 + 2));
+            long f3 = fx(assemble(D[sg], r, c * 4 + 3));
+            long o = (c & 1) ? f0 + f1 + f2 + f3 : f0 - f1 + f2 - f3;
+            if (!fits32(o)) {
+                atomic_fetch_min_explicit(report, g.index, memory_order_relaxed);
+            }
+            amps[at] = low32(o);
+        }
+        simdgroup_barrier(mem_flags::mem_threadgroup);
+    }
+}
+
+// The arithmetic probe: for each case (a, b), the 64-bit product a·b from the
+// matrix multiply and its fx, as (high, low) words. Sixteen cases a tile:
+// product q takes b_q's limbs against component q % 4, and row q / 4 carries
+// a_{4r}, …, a_{4r+3} (quantum_fabric::tensor_engine::probe_weights).
+kernel void probe(
+    device const int2* cases [[buffer(0)]],
+    device uint4* out [[buffer(1)]],
+    constant uint2& count [[buffer(2)]],
+    uint group [[threadgroup_position_in_grid]],
+    uint groups [[threadgroups_per_grid]],
+    uint sg [[simdgroup_index_in_threadgroup]],
+    uint lane [[thread_index_in_simdgroup]],
+    uint width [[threads_per_threadgroup]],
+    uint simd_width [[threads_per_simdgroup]])
+{
+    threadgroup limb_t W[SIMDGROUPS][16 * 112];
+    threadgroup limb_t A[SIMDGROUPS][8 * 16];
+    threadgroup float D[SIMDGROUPS][8 * 112];
+    if (width != SIMDGROUPS * 32 || simd_width != 32) {
+        return;
+    }
+    ulong n = (ulong)count.x | ((ulong)count.y << 32);
+    ulong tiles = (n + 15) / 16;
+    ulong stride = (ulong)groups * SIMDGROUPS;
+    for (ulong tile = (ulong)group * SIMDGROUPS + sg; tile < tiles; tile += stride) {
+        ulong k0 = tile * 16;
+        for (uint x = lane; x < 16 * 112; x += 32) {
+            int k = x / 112, col = x % 112;
+            int c = k / 4, j = k % 4, p = col / 7, i = col % 7 - j;
+            int b = k0 + p < n ? cases[k0 + p].y : 0;
+            W[sg][x] = (c == p % 4 && i >= 0 && i <= 3) ? as_limb(limb(b, i)) : limb_t(0.0f);
+        }
+        // Lanes 0–15 write rows 0–3; lanes 16–31 clear rows 4–7.
+        int r = lane < 16 ? lane / 4 : 4 + (lane - 16) / 4;
+        int c = lane % 4;
+        int a = lane < 16 && k0 + lane < n ? cases[k0 + lane].x : 0;
+        for (int jj = 0; jj < 4; jj++) {
+            A[sg][r * 16 + c * 4 + jj] = as_limb(limb(a, jj));
+        }
+        simdgroup_barrier(mem_flags::mem_threadgroup);
+        products(A[sg], W[sg], D[sg]);
+        simdgroup_barrier(mem_flags::mem_threadgroup);
+        if (lane < 16 && k0 + lane < n) {
+            long v = assemble(D[sg], lane / 4, lane);
+            ulong prod = (ulong)v;
+            ulong f = (ulong)fx(v);
+            out[k0 + lane] = uint4((uint)(prod >> 32), (uint)prod, (uint)(f >> 32), (uint)f);
+        }
+        simdgroup_barrier(mem_flags::mem_threadgroup);
+    }
+}
diff --git a/wai/wai-quantum/src/quantum_fabric/metal.rs b/wai/wai-quantum/src/quantum_fabric/metal.rs
new file mode 100644
index 000000000..fead29e44
--- /dev/null
+++ b/wai/wai-quantum/src/quantum_fabric/metal.rs
@@ -0,0 +1,364 @@
+//! The gate kernels for the platform's GPU in the Metal shading language, and
+//! what names a Metal member: the kernels' text, their implementation ids, 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 Metal member itself is `quantum_fabric_metal`.
+//!
+//! **The integer kernel** ([`SOURCE`], `lanes64.metal`) is
+//! [`cuda::SOURCE`](super::cuda::SOURCE)'s arithmetic in the Metal shading
+//! language: 64-bit integers throughout, so it is bit-exact for the reasons
+//! that module gives. A run's gates are serial dispatches of one compute
+//! encoder, so each sees the last one's writes.
+//!
+//! **The matrix kernel** ([`MATRIX_SOURCE`], `matrix_limbs.metal`) forms every
+//! product with the GPU's SIMD-group matrix multiply, bf16 inputs accumulated
+//! in f32. Each 32-bit value is split into four 8-bit limbs, the lower three
+//! unsigned and the top one signed, so:
+//!
+//! - every limb is an integer of at most `2^8` in size, which bf16 holds
+//!   exactly (8 significant bits);
+//! - every limb product is at most `2^16` in size, and every coefficient of a
+//!   product, `c_s = Σ_{i+j=s} e_i·a_j`, sums at most four of them, so it is
+//!   an integer below `2^18`: f32 holds every partial sum exactly, in whatever
+//!   order the multiply adds them;
+//! - each thread adds its products' coefficients, shifted by `8s` bits, in
+//!   64-bit two's complement, which gives the exact product `e·a`; the rest is
+//!   [`SOURCE`]'s arithmetic on exact products, one row per thread.
+//!
+//! One SIMD group multiplies eight pairs' limbs (8 × 16) by the gate's limb
+//! weights (16 × 112: sixteen products, seven coefficients each) in fourteen
+//! 8 × 8 column blocks; the layout is the tensor engine's
+//! ([`tensor_engine`](super::tensor_engine)). It has its own implementation id
+//! ([`matrix_implementation_id`]) and its own mutants ([`MATRIX_MUTANTS`]),
+//! one of which, nine-bit limbs, bf16 cannot hold.
+//!
+//! Devices are named by number, not by vendor: [`DeviceIdentity::family`] is
+//! the highest of the platform's own GPU families (`MTLGPUFamily` values from
+//! 1001) the device reports.
+
+use super::FabricClass;
+use super::wgsl::{Detector, Mutant};
+
+/// The integer kernel, `lanes64.metal`, exactly as it is compiled.
+pub const SOURCE: &str = include_str!("lanes64.metal");
+
+/// The integer kernel's implementation name.
+pub const IMPLEMENTATION: &str = "metal-lanes64";
+
+/// The matrix kernel, `matrix_limbs.metal`, exactly as it is compiled.
+pub const MATRIX_SOURCE: &str = include_str!("matrix_limbs.metal");
+
+/// The matrix kernel's implementation name.
+pub const MATRIX_IMPLEMENTATION: &str = "metal-matrix-limbs8";
+
+/// Version of the host-side layout: interleaved `int2` lanes in one shared
+/// buffer; one serial dispatch per gate, the gate's words as 13 32-bit words
+/// at buffer index 1 (the pair count's low and high words, `tbit`, `ctrl`,
+/// the eight entries, the op index); the report word (the lowest
+/// out-of-range op index, `u32::MAX` for none); the probe's `int2` cases,
+/// `uint4` outputs and the case count as two words; and the launch shapes
+/// ([`Shape`]). Bumped on any change to them; it is part of both
+/// implementation ids.
+pub const HOST_LAYOUT: u32 = 1;
+
+/// Blocks per dispatch, at most; the kernels stride over the rest.
+pub const GRID_MAX: u64 = 4096;
+
+/// How a kernel is dispatched: threads per threadgroup, and the threads each
+/// pair (or probe case) takes. The kernels stride, so any grid covers the
+/// work; this one covers it at once up to [`GRID_MAX`] threadgroups.
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+pub struct Shape {
+    pub threads: u64,
+    /// Threads per pair in a gate dispatch.
+    pub per_pair: u64,
+    /// Threads per case in a probe dispatch.
+    pub per_case: u64,
+}
+
+impl Shape {
+    /// The integer kernel's: 256 threads a threadgroup, one thread an item.
+    pub const LANES: Shape = Shape { threads: 256, per_pair: 1, per_case: 1 };
+    /// The matrix kernel's: four SIMD groups of 32 a threadgroup; a SIMD
+    /// group takes eight pairs (four threads a pair) or sixteen cases (two
+    /// threads a case).
+    pub const MATRIX: Shape = Shape { threads: 128, per_pair: 4, per_case: 2 };
+
+    /// Threadgroups for `items` at `per_item` threads each: at least one, at
+    /// most [`GRID_MAX`].
+    pub fn groups(&self, items: u64, per_item: u64) -> u64 {
+        items.saturating_mul(per_item).div_ceil(self.threads).clamp(1, GRID_MAX)
+    }
+}
+
+/// What the runtime compiler is given: the language version the matrix
+/// kernel's `bfloat` needs, and safe math (no fast-math reassociation or
+/// approximation), though every float either kernel forms is an exact
+/// integer. Part of both implementation ids.
+pub const COMPILE_OPTIONS: [&str; 2] = ["language=3.1", "math=safe"];
+
+/// The GPU family the matrix kernel needs at least: SIMD-group matrices.
+pub const MATRIX_MIN_FAMILY: u32 = 1007;
+
+const DOMAIN_IMPLEMENTATION: &[u8] = b"wai:quantum-fabric-metal\x01";
+const DOMAIN_MATRIX_IMPLEMENTATION: &[u8] = b"wai:quantum-fabric-metal-matrix\x01";
+const DOMAIN_MEMBER: &[u8] = b"wai:quantum-fabric-metal-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);
+}
+
+fn id_of(domain: &[u8], name: &str, source: &str, shape: Shape) -> String {
+    let mut o = domain.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());
+    for v in [shape.threads, shape.per_pair, shape.per_case] {
+        o.extend_from_slice(&v.to_be_bytes());
+    }
+    format!("{name}/{}", hex(&blake3::hash(&o).as_bytes()[..8]))
+}
+
+/// `"metal-lanes64/<16 hex>"`: the first 8 bytes of
+/// BLAKE3(`"wai:quantum-fabric-metal\x01"` ‖ put([`SOURCE`]) ‖ put of each of
+/// [`COMPILE_OPTIONS`] ‖ [`HOST_LAYOUT`] ‖ [`Shape::LANES`]'s three numbers,
+/// big-endian), in hex. Any edit to the kernel's text, its options or the host
+/// layout changes it.
+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 {
+    id_of(DOMAIN_IMPLEMENTATION, IMPLEMENTATION, source, Shape::LANES)
+}
+
+/// `"metal-matrix-limbs8/<16 hex>"`: as [`implementation_id`], with the
+/// domain `"wai:quantum-fabric-metal-matrix\x01"`, [`MATRIX_SOURCE`] and
+/// [`Shape::MATRIX`].
+pub fn matrix_implementation_id() -> String {
+    matrix_implementation_id_of(MATRIX_SOURCE)
+}
+
+/// [`matrix_implementation_id`] for another kernel text, such as a mutant's.
+pub fn matrix_implementation_id_of(source: &str) -> String {
+    id_of(DOMAIN_MATRIX_IMPLEMENTATION, MATRIX_IMPLEMENTATION, source, Shape::MATRIX)
+}
+
+/// What a Metal member is, as its id hashes it. The member fills it from the
+/// device; the fields are plain values here so that the id is testable
+/// without one.
+#[derive(Clone, Debug, PartialEq, Eq)]
+#[non_exhaustive]
+pub struct DeviceIdentity {
+    pub name: String,
+    /// The device's registry id, unique on the machine.
+    pub registry_id: u64,
+    /// The highest of the platform's own GPU families (`MTLGPUFamily` values
+    /// 1001 and up) the device reports, or 0 for none.
+    pub family: u32,
+    /// Whether the device shares the host's memory.
+    pub unified_memory: bool,
+    /// The memory the device recommends a process keep resident, bytes.
+    pub working_set: u64,
+    /// The longest buffer the device allocates, bytes.
+    pub max_buffer: u64,
+    /// The operating system's version, whose runtime compiler builds the
+    /// kernels: `"<major>.<minor>.<patch>"`.
+    pub os: String,
+}
+
+impl DeviceIdentity {
+    /// Always [`FabricClass::Gpu`].
+    pub fn class(&self) -> FabricClass {
+        FabricClass::Gpu
+    }
+
+    /// `"local/gpu-msl-<d12>"`. `d12` is the first 12 hex digits of
+    /// BLAKE3(`"wai:quantum-fabric-metal-member\x01"` ‖ put(name) ‖ registry
+    /// id ‖ family ‖ unified memory as one byte ‖ working set ‖ max buffer ‖
+    /// put(os)), 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());
+        o.extend_from_slice(&self.registry_id.to_be_bytes());
+        o.extend_from_slice(&self.family.to_be_bytes());
+        o.push(u8::from(self.unified_memory));
+        o.extend_from_slice(&self.working_set.to_be_bytes());
+        o.extend_from_slice(&self.max_buffer.to_be_bytes());
+        put(&mut o, self.os.as_bytes());
+        format!("local/gpu-msl-{}", hex(&blake3::hash(&o).as_bytes()[..6]))
+    }
+
+    /// The key a runner declares in CI: `"gpu-msl-<family>"`, for instance
+    /// `"gpu-msl-1009"`.
+    pub fn key(&self) -> String {
+        format!("gpu-msl-{}", self.family)
+    }
+
+    /// Whether the device can run the matrix kernel: family
+    /// [`MATRIX_MIN_FAMILY`] or later.
+    pub fn has_simdgroup_matrices(&self) -> bool {
+        self.family >= MATRIX_MIN_FAMILY
+    }
+}
+
+/// The integer kernel's mutants: textual edits of [`SOURCE`] that a
+/// conforming device must be caught running, each by the detector named. Each
+/// `find` occurs exactly once in [`SOURCE`].
+pub const MUTANTS: [Mutant; 6] = [
+    // The rounding constant is half what it should be.
+    Mutant { name: "G1", find: "((long)a * (long)b + (1L << 29)) >> 30;", replace: "((long)a * (long)b + (1L << 28)) >> 30;", detector: Detector::Probe },
+    // The product is shifted by one bit too few.
+    Mutant { name: "G2", find: "((long)a * (long)b + (1L << 29)) >> 30;", replace: "((long)a * (long)b + (1L << 29)) >> 29;", detector: Detector::Probe },
+    // The probe's product is taken in 32 bits.
+    Mutant { name: "G3", find: "ulong prod = (ulong)((long)c.x * (long)c.y);", replace: "ulong prod = (ulong)(uint)(c.x * c.y);", detector: Detector::Probe },
+    // The control test is inverted.
+    Mutant { name: "G4", find: "if ((i & (ulong)g.ctrl) != (ulong)g.ctrl) {", replace: "if ((i & (ulong)g.ctrl) == (ulong)g.ctrl) {", detector: Detector::Certificate },
+    // A sign of row 0's real part is wrong, on the imaginary part of the
+    // top-right entry, which `Y` has.
+    Mutant { name: "G5", find: "fx(g.m[2], a1.x) - fx(g.m[3], a1.y);", replace: "fx(g.m[2], a1.x) + fx(g.m[3], a1.y);", detector: Detector::Certificate },
+    // The range guard never fires.
+    Mutant { name: "G6", find: "if (!(fits32(o0r) && fits32(o0i) && fits32(o1r) && fits32(o1i))) {", replace: "if (false) {", detector: Detector::Range },
+];
+
+/// The matrix kernel's mutants, as [`MUTANTS`] are the other's. Each `find`
+/// occurs exactly once in [`MATRIX_SOURCE`].
+pub const MATRIX_MUTANTS: [Mutant; 7] = [
+    // The rounding constant is half what it should be.
+    Mutant { name: "X1", find: "return (v + (1L << 29)) >> 30;", replace: "return (v + (1L << 28)) >> 30;", detector: Detector::Probe },
+    // The product is shifted by one bit too few.
+    Mutant { name: "X2", find: "return (v + (1L << 29)) >> 30;", replace: "return (v + (1L << 29)) >> 29;", detector: Detector::Probe },
+    // Nine-bit limbs: a limb above 2^8 has nine significant bits, which bf16
+    // does not hold, so the products are not exact.
+    Mutant { name: "X3", find: "#define LIMB 8", replace: "#define LIMB 9", detector: Detector::Probe },
+    // The second limb block multiplies the first block's weights.
+    Mutant { name: "X4", find: "simdgroup_multiply_accumulate(d, a1, w1, d);", replace: "simdgroup_multiply_accumulate(d, a1, w0, d);", detector: Detector::Probe },
+    // The control test is inverted.
+    Mutant { name: "X5", find: "if ((i & (ulong)g.ctrl) == (ulong)g.ctrl) {", replace: "if ((i & (ulong)g.ctrl) != (ulong)g.ctrl) {", detector: Detector::Certificate },
+    // A sign of each real row's last term is wrong (row 0's is on the
+    // imaginary part of the top-right entry, which `Y` has).
+    Mutant { name: "X6", find: "f0 - f1 + f2 - f3;", replace: "f0 - f1 + f2 + f3;", detector: Detector::Certificate },
+    // The range guard never fires.
+    Mutant { name: "X7", find: "if (!fits32(o)) {", replace: "if (false) {", detector: Detector::Range },
+];
+
+fn apply(source: &str, m: &Mutant, what: &str) -> String {
+    let n = source.matches(m.find).count();
+    assert_eq!(n, 1, "mutant {}: `{}` occurs {n} times in the {what}", m.name, m.find);
+    source.replacen(m.find, m.replace, 1)
+}
+
+/// [`SOURCE`] with `m` applied. Panics unless `m.find` occurs exactly once.
+pub fn mutate(m: &Mutant) -> String {
+    apply(SOURCE, m, "integer kernel")
+}
+
+/// [`MATRIX_SOURCE`] with `m` applied. Panics unless `m.find` occurs exactly
+/// once.
+pub fn matrix_mutate(m: &Mutant) -> String {
+    apply(MATRIX_SOURCE, m, "matrix kernel")
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    fn identity() -> DeviceIdentity {
+        DeviceIdentity {
+            name: "a device".into(),
+            registry_id: 0x1000_0000_0abc,
+            family: 1009,
+            unified_memory: true,
+            working_set: 96 << 30,
+            max_buffer: 64 << 30,
+            os: "27.0.1".into(),
+        }
+    }
+
+    #[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-msl-") && id.len() == "local/gpu-msl-".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 { registry_id: 1, ..identity() },
+            DeviceIdentity { family: 1010, ..identity() },
+            DeviceIdentity { unified_memory: false, ..identity() },
+            DeviceIdentity { working_set: 1 << 30, ..identity() },
+            DeviceIdentity { max_buffer: 1 << 30, ..identity() },
+            DeviceIdentity { os: "27.0.2".into(), ..identity() },
+        ];
+        for v in variants {
+            assert_ne!(v.member_id(), id, "{v:?}");
+        }
+        assert_eq!(base.key(), "gpu-msl-1009");
+        assert_eq!(base.class(), FabricClass::Gpu);
+    }
+
+    #[test]
+    fn the_implementation_ids_follow_the_kernel_texts() {
+        let id = implementation_id();
+        assert!(id.starts_with("metal-lanes64/") && id.len() == "metal-lanes64/".len() + 16, "{id}");
+        let mid = matrix_implementation_id();
+        assert!(mid.starts_with("metal-matrix-limbs8/") && mid.len() == "metal-matrix-limbs8/".len() + 16, "{mid}");
+        assert_ne!(matrix_implementation_id_of(SOURCE).split('/').nth(1), implementation_id_of(SOURCE).split('/').nth(1), "the domains differ");
+        for m in &MUTANTS {
+            assert_ne!(implementation_id_of(&mutate(m)), id, "mutant {}", m.name);
+        }
+        for m in &MATRIX_MUTANTS {
+            assert_ne!(matrix_implementation_id_of(&matrix_mutate(m)), mid, "mutant {}", m.name);
+        }
+    }
+
+    #[test]
+    fn every_mutant_edits_its_kernel_once_and_names_distinct_edits() {
+        for (source, mutants, edit) in [(SOURCE, &MUTANTS[..], mutate as fn(&Mutant) -> String), (MATRIX_SOURCE, &MATRIX_MUTANTS[..], matrix_mutate)] {
+            let mut seen = std::collections::BTreeSet::new();
+            for m in mutants {
+                let edited = edit(m);
+                assert_ne!(edited, source, "mutant {} changes nothing", m.name);
+                assert!(seen.insert(edited), "mutant {} duplicates another", m.name);
+            }
+            for d in [Detector::Probe, Detector::Certificate, Detector::Range] {
+                assert!(mutants.iter().any(|m| m.detector == d), "no mutant for {d:?}");
+            }
+        }
+        // The kernels' entry points are the same two.
+        for entry in ["kernel void apply_gate(", "kernel void probe("] {
+            assert_eq!((SOURCE.matches(entry).count(), MATRIX_SOURCE.matches(entry).count()), (1, 1), "{entry}");
+        }
+    }
+
+    #[test]
+    fn simdgroup_matrices_start_at_family_1007() {
+        for (family, has) in [(0, false), (1006, false), (1007, true), (1009, true), (1010, true)] {
+            assert_eq!(DeviceIdentity { family, ..identity() }.has_simdgroup_matrices(), has, "{family}");
+        }
+    }
+
+    #[test]
+    fn the_launch_shapes_cover_the_work_and_cap_the_groups() {
+        assert_eq!(Shape::LANES.groups(0, 1), 1);
+        assert_eq!(Shape::LANES.groups(256, 1), 1);
+        assert_eq!(Shape::LANES.groups(257, 1), 2);
+        assert_eq!(Shape::LANES.groups(u64::MAX, 1), GRID_MAX);
+        // A matrix threadgroup takes 32 pairs (4 SIMD groups of 8) or 64 cases.
+        assert_eq!(Shape::MATRIX.groups(32, Shape::MATRIX.per_pair), 1);
+        assert_eq!(Shape::MATRIX.groups(33, Shape::MATRIX.per_pair), 2);
+        assert_eq!(Shape::MATRIX.groups(64, Shape::MATRIX.per_case), 1);
+        assert_eq!(Shape::MATRIX.groups(65, Shape::MATRIX.per_case), 2);
+    }
+}
diff --git a/wai/wai-quantum/src/quantum_fabric/tensor_limbs.cu b/wai/wai-quantum/src/quantum_fabric/tensor_limbs.cu
new file mode 100644
index 000000000..238e273ba
--- /dev/null
+++ b/wai/wai-quantum/src/quantum_fabric/tensor_limbs.cu
@@ -0,0 +1,233 @@
+// wai-quantum fabric: the gate kernel for CUDA devices whose tensor cores
+// multiply bf16 and accumulate in f32 (compute capability 8.0 and later).
+//
+// The same entry points as lanes64.cu, with the same parameters, so the
+// member runs either one; here every product the gate needs is formed on the
+// tensor cores. Each 32-bit value is split into four limbs of 8 bits (the
+// lower three unsigned, the top one signed), which bf16 holds 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, every limb product at most 2^16 in size,
+// every coefficient below 2^18. So every value a tensor core forms is an
+// integer fp32 holds exactly, in whatever order it adds them
+// (quantum_fabric::tensor_engine has the argument and the layout). The limb
+// weights put each of a gate's sixteen products per pair in its own seven
+// columns of a 16 × 112 matrix; one warp multiplies a 16-pair tile of limbs by
+// it in fourteen m16n8k16 steps, and adds each coefficient, shifted into
+// place, into its product in 64-bit integers. The rest is lanes64.cu's
+// arithmetic on those exact products: fx = (product + 2^29) >> 30, each row
+// the sum of four, the range reported and the low 32 bits kept.
+//
+// Each thread owns its pair's two amplitudes for the launch, and a launch is
+// one gate, so no thread reads what another writes within a launch.
+
+typedef long long i64;
+typedef unsigned long long u64;
+typedef unsigned int u32;
+typedef unsigned short u16;
+
+// Warps per block: the member launches 256 threads a block.
+#define WARPS 8
+
+// Bits per limb. Eight is what bf16 holds exactly; the radix is a mutant's
+// one edit, and nine-bit limbs are not exact.
+#define LIMB 8
+
+// The sixteen products, as (entry, component): rows 0r, 0i, 1r and 1i, four
+// terms each, in the order lanes::exact_row adds them.
+__constant__ int TERM_ENTRY[16] = {0, 1, 2, 3, 0, 1, 2, 3, 4, 5, 6, 7, 4, 5, 6, 7};
+__constant__ int TERM_COMP[16] = {0, 1, 2, 3, 1, 0, 3, 2, 0, 1, 2, 3, 1, 0, 3, 2};
+
+// Limb j of x: the lower three unsigned, the top one the signed rest.
+__device__ __forceinline__ int limb(int x, int j) {
+    return j < 3 ? (int)(((u32)x >> (LIMB * j)) & ((1u << LIMB) - 1)) : (x >> (3 * LIMB));
+}
+
+// v as bf16 bits. An integer of at most 256 in size has at most 8 significant
+// bits, so its f32 has zero low 16 bits and the truncation is exact.
+__device__ __forceinline__ u16 bf16(int v) {
+    return (u16)(__float_as_uint((float)v) >> 16);
+}
+
+__device__ __forceinline__ u32 pack(u16 lo, u16 hi) {
+    return (u32)lo | ((u32)hi << 16);
+}
+
+// The exact fixed-point rounding of a product.
+__device__ __forceinline__ i64 fx(i64 v) {
+    return (v + ((i64)1 << 29)) >> 30;
+}
+
+__device__ __forceinline__ int low32(i64 v) {
+    return (int)(u32)(u64)v;
+}
+
+__device__ __forceinline__ bool fits32(i64 v) {
+    return v >= -(i64)2147483648LL && v <= (i64)2147483647LL;
+}
+
+// d = a·b for one m16n8k16 step, bf16 inputs, f32 accumulation from zero.
+__device__ __forceinline__ void mma(const u32 a[4], u32 b0, u32 b1, float d[4]) {
+    asm volatile(
+        "mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%0,%1,%2,%3}, {%4,%5,%6,%7}, {%8,%9}, {%10,%11,%12,%13};\n"
+        : "=f"(d[0]), "=f"(d[1]), "=f"(d[2]), "=f"(d[3])
+        : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), "r"(b0), "r"(b1), "f"(0.f), "f"(0.f), "f"(0.f), "f"(0.f));
+}
+
+// P[row][p] += product p of row `row`, for the first `rows` rows of the tile:
+// A (16 × 16 limbs) times W (16 × 112), coefficient s of product p at column
+// p·7 + s, each shifted into place. Called by all 32 lanes of a warp.
+__device__ void products(const u16 (*A)[16], const u16 (*W)[112], i64 (*P)[16], int rows, int lane) {
+    int g = lane >> 2, t = lane & 3;
+    u32 a[4];
+    a[0] = pack(A[g][2 * t], A[g][2 * t + 1]);
+    a[1] = pack(A[g + 8][2 * t], A[g + 8][2 * t + 1]);
+    a[2] = pack(A[g][2 * t + 8], A[g][2 * t + 9]);
+    a[3] = pack(A[g + 8][2 * t + 8], A[g + 8][2 * t + 9]);
+    for (int nt = 0; nt < 14; nt++) {
+        int col = nt * 8 + g;
+        u32 b0 = pack(W[2 * t][col], W[2 * t + 1][col]);
+        u32 b1 = pack(W[2 * t + 8][col], W[2 * t + 9][col]);
+        float d[4];
+        mma(a, b0, b1, d);
+        for (int e = 0; e < 4; e++) {
+            int row = g + (e >= 2 ? 8 : 0);
+            if (row >= rows) {
+                continue;
+            }
+            int c = nt * 8 + 2 * t + (e & 1);
+            int p = c / 7, s = c % 7;
+            // An integer below 2^24 in size, so the conversion is exact; the
+            // shift and the add are in two's complement.
+            u64 v = (u64)(i64)d[e] << (LIMB * s);
+            atomicAdd((u64*)&P[row][p], v);
+        }
+    }
+}
+
+// One gate on a state of `2 * pairs` lanes, one pair per lane of a warp's
+// half, sixteen pairs a tile. The parameters are lanes64.cu's.
+extern "C" __global__ void apply_gate(
+    int2* __restrict__ amps,
+    u64 pairs,
+    unsigned int tbit,
+    unsigned int ctrl,
+    int m0, int m1, int m2, int m3, int m4, int m5, int m6, int m7,
+    unsigned int index,
+    unsigned int* __restrict__ report)
+{
+    __shared__ u16 W[16][112];
+    __shared__ u16 A[WARPS][16][16];
+    __shared__ i64 P[WARPS][16][16];
+    if (blockDim.x != WARPS * 32) {
+        return;
+    }
+    int m[8] = {m0, m1, m2, m3, m4, m5, m6, m7};
+    // The gate's limb weights, once a block: row c·4 + j is limb j of
+    // component c; column p·7 + s takes limb s − j of product p's entry.
+    for (int x = threadIdx.x; x < 16 * 112; x += blockDim.x) {
+        int k = x / 112, col = x % 112;
+        int c = k / 4, j = k % 4, p = col / 7, i = col % 7 - j;
+        W[k][col] = (c == TERM_COMP[p] && i >= 0 && i <= 3) ? bf16(limb(m[TERM_ENTRY[p]], i)) : (u16)0;
+    }
+    __syncthreads();
+    int w = threadIdx.x >> 5, lane = threadIdx.x & 31;
+    u64 tiles = (pairs + 15) / 16;
+    u64 stride = (u64)gridDim.x * WARPS;
+    for (u64 tile = (u64)blockIdx.x * WARPS + w; tile < tiles; tile += stride) {
+        bool active = false;
+        u64 i = 0, j = 0;
+        if (lane < 16) {
+            u64 p = tile * 16 + lane;
+            int2 a0 = make_int2(0, 0), a1 = make_int2(0, 0);
+            if (p < pairs) {
+                // The p-th index whose target bit is clear.
+                u64 low = p & ((u64)tbit - 1);
+                i = ((p - low) << 1) | low;
+                if ((i & (u64)ctrl) == (u64)ctrl) {
+                    active = true;
+                    j = i | (u64)tbit;
+                    a0 = amps[i];
+                    a1 = amps[j];
+                }
+            }
+            int comp[4] = {a0.x, a0.y, a1.x, a1.y};
+            for (int c = 0; c < 4; c++) {
+                for (int jj = 0; jj < 4; jj++) {
+                    A[w][lane][c * 4 + jj] = bf16(limb(comp[c], jj));
+                }
+            }
+            for (int q = 0; q < 16; q++) {
+                P[w][lane][q] = 0;
+            }
+        }
+        __syncwarp();
+        products(A[w], W, P[w], 16, lane);
+        __syncwarp();
+        if (active) {
+            const i64* pr = P[w][lane];
+            i64 o0r = fx(pr[0]) - fx(pr[1]) + fx(pr[2]) - fx(pr[3]);
+            i64 o0i = fx(pr[4]) + fx(pr[5]) + fx(pr[6]) + fx(pr[7]);
+            i64 o1r = fx(pr[8]) - fx(pr[9]) + fx(pr[10]) - fx(pr[11]);
+            i64 o1i = fx(pr[12]) + fx(pr[13]) + fx(pr[14]) + fx(pr[15]);
+            if (!(fits32(o0r) && fits32(o0i) && fits32(o1r) && fits32(o1i))) {
+                atomicMin(&report[0], index);
+            }
+            amps[i] = make_int2(low32(o0r), low32(o0i));
+            amps[j] = make_int2(low32(o1r), low32(o1i));
+        }
+        __syncwarp();
+    }
+}
+
+// The arithmetic probe: for each case (a, b), the 64-bit product a·b from the
+// tensor cores and its fx, as (high, low) words. Sixteen cases a tile: product
+// q takes b_q's limbs against component q % 4, and row q / 4 carries
+// a_{4r}, …, a_{4r+3} (quantum_fabric::tensor_engine::probe_weights).
+extern "C" __global__ void probe(
+    const int2* __restrict__ cases,
+    uint4* __restrict__ out,
+    u64 count)
+{
+    __shared__ u16 W[WARPS][16][112];
+    __shared__ u16 A[WARPS][16][16];
+    __shared__ i64 P[WARPS][4][16];
+    if (blockDim.x != WARPS * 32) {
+        return;
+    }
+    int w = threadIdx.x >> 5, lane = threadIdx.x & 31;
+    u64 tiles = (count + 15) / 16;
+    u64 stride = (u64)gridDim.x * WARPS;
+    for (u64 tile = (u64)blockIdx.x * WARPS + w; tile < tiles; tile += stride) {
+        u64 k0 = tile * 16;
+        for (int x = lane; x < 16 * 112; x += 32) {
+            int k = x / 112, col = x % 112;
+            int c = k / 4, j = k % 4, p = col / 7, i = col % 7 - j;
+            int b = k0 + p < count ? cases[k0 + p].y : 0;
+            W[w][k][col] = (c == p % 4 && i >= 0 && i <= 3) ? bf16(limb(b, i)) : (u16)0;
+        }
+        if (lane < 16) {
+            int r = lane / 4, c = lane % 4;
+            int a = k0 + lane < count ? cases[k0 + lane].x : 0;
+            for (int jj = 0; jj < 4; jj++) {
+                A[w][r][c * 4 + jj] = bf16(limb(a, jj));
+            }
+            for (int rr = 4; rr < 16; rr++) {
+                A[w][rr][lane] = 0;
+            }
+            P[w][r][lane] = 0;
+            P[w][r][(lane + 4) % 16] = 0;
+            P[w][r][(lane + 8) % 16] = 0;
+            P[w][r][(lane + 12) % 16] = 0;
+        }
+        __syncwarp();
+        products(A[w], W[w], P[w], 4, lane);
+        __syncwarp();
+        if (lane < 16 && k0 + lane < count) {
+            i64 v = P[w][lane / 4][lane];
+            u64 prod = (u64)v;
+            u64 f = (u64)fx(v);
+            out[k0 + lane] = make_uint4((u32)(prod >> 32), (u32)prod, (u32)(f >> 32), (u32)f);
+        }
+        __syncwarp();
+    }
+}
diff --git a/wai/wai-quantum/src/quantum_fabric_cuda.rs b/wai/wai-quantum/src/quantum_fabric_cuda.rs
index b110076d7..3837e0480 100644
--- a/wai/wai-quantum/src/quantum_fabric_cuda.rs
+++ b/wai/wai-quantum/src/quantum_fabric_cuda.rs
@@ -21,6 +21,12 @@
 //!   [`Certification`]. The corpus digest covers inputs and outputs only, so a
 //!   bit-exact member reproduces the CPU's pinned digests; a device that runs
 //!   nothing, skips a case or gets one bit wrong cannot.
+//! - **Tensor cores** ([`CudaMember::tensor`]): the same device running the
+//!   second kernel, [`cuda::TENSOR_SOURCE`], whose every product is formed on
+//!   the tensor cores from 8-bit limbs; probed before it is returned, with its
+//!   own implementation id, so its certificates are certificates of that
+//!   silicon path. A device below [`cuda::TENSOR_MIN_COMPUTE`] is
+//!   [`CudaError::Unsupported`].
 //! - **Placement, receipts and metering** are as for the WebGPU member:
 //!   [`place`](crate::quantum_fabric::place),
 //!   [`FabricReceipt`](crate::quantum_fabric::FabricReceipt) and a caller's
@@ -85,6 +91,8 @@ pub enum CudaError {
     Internal(&'static str),
     /// The device's arithmetic disagreed with the CPU's on this pair.
     Probe { a: i32, b: i32 },
+    /// The device lacks what the kernel needs, and what.
+    Unsupported(String),
 }
 
 impl std::fmt::Display for CudaError {
@@ -99,6 +107,7 @@ impl std::fmt::Display for CudaError {
             CudaError::OutOfRange { gate } => write!(f, "the result of op {gate} leaves 32 bits"),
             CudaError::Internal(why) => write!(f, "host defect: {why}"),
             CudaError::Probe { a, b } => write!(f, "the device's arithmetic differs from the CPU's at ({a}, {b})"),
+            CudaError::Unsupported(why) => write!(f, "the device cannot run this kernel: {why}"),
         }
     }
 }
@@ -250,7 +259,7 @@ fn open(ordinal: usize) -> Result<CudaMember, Refusal> {
     };
     let key = identity.key();
     let refuse = |error: CudaError| Refusal { device: name.clone(), key: key.clone(), error };
-    let kernel = Kernel::compile(&ctx, &identity, cuda::SOURCE).map_err(refuse)?;
+    let kernel = Kernel::compile(&ctx, &identity, cuda::SOURCE, 1).map_err(refuse)?;
     let (free, _) = ctx.mem_get_info().map_err(|e| refuse(driver(e)))?;
     let max_qubits = qubits_for(free as u64);
     let detail = format!(
@@ -264,29 +273,25 @@ fn open(ordinal: usize) -> Result<CudaMember, Refusal> {
         id: identity.member_id(),
         key: key.clone(),
         implementation: cuda::implementation_id(),
+        discovered: detail.clone(),
         detail,
         max_qubits,
         identity,
     };
-    let pairs = lanes::probe_pairs(4096, 0x5eed);
-    match member.probe(&pairs) {
-        Ok(out) => match lanes::first_mismatch(&pairs, &out) {
-            None => Ok(member),
-            Some((a, b)) => Err(refuse(CudaError::Probe { a, b })),
-        },
-        Err(e) => Err(refuse(e)),
-    }
+    member.probed().map_err(refuse)
 }
 
-/// The kernel compiled for one device: its two entry points.
+/// The kernel compiled for one device: its two entry points, and the threads a
+/// launch gives each item (pair or probe case).
 #[derive(Clone)]
 struct Kernel {
     apply: CudaFunction,
     probe: CudaFunction,
+    spread: u64,
 }
 
 impl Kernel {
-    fn compile(ctx: &Arc<CudaContext>, identity: &DeviceIdentity, source: &str) -> Result<Kernel, CudaError> {
+    fn compile(ctx: &Arc<CudaContext>, identity: &DeviceIdentity, source: &str, spread: u64) -> Result<Kernel, CudaError> {
         let opts = CompileOptions {
             arch: Some(arch_option(identity.arch())),
             options: cuda::COMPILE_OPTIONS.iter().map(|s| s.to_string()).collect(),
@@ -294,7 +299,12 @@ impl Kernel {
         };
         let ptx = compile_ptx_with_opts(source, opts).map_err(|e| CudaError::Compile(e.to_string()))?;
         let module = ctx.load_module(ptx).map_err(driver)?;
-        Ok(Kernel { apply: module.load_function("apply_gate").map_err(driver)?, probe: module.load_function("probe").map_err(driver)? })
+        Ok(Kernel { apply: module.load_function("apply_gate").map_err(driver)?, probe: module.load_function("probe").map_err(driver)?, spread })
+    }
+
+    /// The launch shape for `work` items, at `spread` threads each.
+    fn launch(&self, work: u64) -> LaunchConfig {
+        launch(work.saturating_mul(self.spread))
     }
 }
 
@@ -316,6 +326,9 @@ pub struct CudaMember {
     id: String,
     key: String,
     implementation: String,
+    /// What discovery found, which [`Self::detail`] extends for a kernel that
+    /// says more.
+    discovered: String,
     detail: String,
     max_qubits: u8,
 }
@@ -377,7 +390,7 @@ impl CudaMember {
         args.arg(&cases).arg(&mut out).arg(&count);
         // SAFETY: `cases` holds `count` int2 cases and `out` `count` uint4
         // outputs, the kernel's parameters in order.
-        unsafe { args.launch(launch(count)) }.map_err(driver)?;
+        unsafe { args.launch(self.kernel.launch(count)) }.map_err(driver)?;
         let words = self.stream.clone_dtoh(&out).map_err(driver)?;
         Ok(words.as_chunks::<4>().0.iter().map(|w| ((w[0], w[1]), (w[2], w[3]))).collect())
     }
@@ -470,7 +483,7 @@ impl CudaMember {
             args.arg(&index).arg(&mut report);
             // SAFETY: `amps` holds `2·pairs` int2 lanes and `report` one word,
             // and the words were checked against the state above.
-            unsafe { args.launch(launch(pairs)) }.map_err(driver)?;
+            unsafe { args.launch(self.kernel.launch(pairs)) }.map_err(driver)?;
         }
         let bad = self.stream.clone_dtoh(&report).map_err(driver)?;
         let words = self.stream.clone_dtoh(&amps).map_err(driver)?;
@@ -576,8 +589,43 @@ impl CudaMember {
     /// require every mutant to be caught.
     #[doc(hidden)]
     pub fn with_source(&self, source: &str) -> Result<CudaMember, CudaError> {
-        let kernel = Kernel::compile(&self.ctx, &self.identity, source)?;
-        Ok(CudaMember { kernel, implementation: cuda::implementation_id_of(source), ..self.clone() })
+        let kernel = Kernel::compile(&self.ctx, &self.identity, source, 1)?;
+        Ok(CudaMember { kernel, implementation: cuda::implementation_id_of(source), detail: self.discovered.clone(), ..self.clone() })
+    }
+
+    /// The same device running the tensor-core kernel
+    /// ([`cuda::TENSOR_SOURCE`]), compiled for it and probed in this call:
+    /// the member id is the device's, the implementation id the tensor-core
+    /// kernel's. A device below [`cuda::TENSOR_MIN_COMPUTE`] is
+    /// [`CudaError::Unsupported`]; one whose products disagree with the CPU's
+    /// is [`CudaError::Probe`].
+    pub fn tensor(&self) -> Result<CudaMember, CudaError> {
+        self.with_tensor_source(cuda::TENSOR_SOURCE)?.probed()
+    }
+
+    /// The same device running `source` as the tensor-core kernel, unprobed,
+    /// under the implementation id that text earns: for the mutation tests.
+    #[doc(hidden)]
+    pub fn with_tensor_source(&self, source: &str) -> Result<CudaMember, CudaError> {
+        if !self.identity.has_tensor_cores() {
+            let (major, minor) = self.identity.compute;
+            let (need_major, need_minor) = cuda::TENSOR_MIN_COMPUTE;
+            return Err(CudaError::Unsupported(format!("compute capability {major}.{minor}; the tensor-core kernel needs {need_major}.{need_minor} or later")));
+        }
+        let kernel = Kernel::compile(&self.ctx, &self.identity, source, cuda::TENSOR_THREADS_PER_ITEM)?;
+        let detail = format!("{}; every product on the tensor cores (bf16 limbs of 8 bits, f32 accumulation)", self.discovered);
+        Ok(CudaMember { kernel, implementation: cuda::tensor_implementation_id_of(source), detail, ..self.clone() })
+    }
+
+    /// This member, once its kernel's arithmetic agrees with the CPU's on the
+    /// probe's pairs.
+    fn probed(self) -> Result<CudaMember, CudaError> {
+        let pairs = lanes::probe_pairs(4096, 0x5eed);
+        let out = self.probe(&pairs)?;
+        match lanes::first_mismatch(&pairs, &out) {
+            None => Ok(self),
+            Some((a, b)) => Err(CudaError::Probe { a, b }),
+        }
     }
 }
 
@@ -766,6 +814,12 @@ mod tests {
         assert_eq!(launch(u64::from(cuda::BLOCK) + 1).grid_dim.0, 2);
         assert_eq!(launch(1 << 40).grid_dim.0, cuda::GRID_MAX);
         assert_eq!(launch(7).block_dim, (cuda::BLOCK, 1, 1));
+        // A tensor-core launch: a warp takes sixteen items, so a block of 256
+        // threads takes 128 and the grid covers every item at once.
+        let threads = |work: u64| launch(work.saturating_mul(cuda::TENSOR_THREADS_PER_ITEM)).grid_dim.0;
+        assert_eq!(threads(128), 1);
+        assert_eq!(threads(129), 2);
+        assert_eq!(threads(u64::MAX), cuda::GRID_MAX);
     }
 
     #[test]
diff --git a/wai/wai-quantum/src/quantum_fabric_metal.rs b/wai/wai-quantum/src/quantum_fabric_metal.rs
new file mode 100644
index 000000000..512083f39
--- /dev/null
+++ b/wai/wai-quantum/src/quantum_fabric_metal.rs
@@ -0,0 +1,781 @@
+//! The fabric's Metal member: the gate kernels of
+//! [`quantum_fabric::metal`](crate::quantum_fabric::metal) on the platform's
+//! GPU, compiled at run time by the operating system's own compiler. See that
+//! module for the kernels and the argument that each is bit-exact.
+//!
+//! - **Discovery** ([`discover`]): on macOS, every device the platform lists
+//!   is opened, the integer kernel is compiled for it, and its arithmetic is
+//!   probed against [`lanes`] in this process. A device that cannot host the
+//!   kernel, or whose arithmetic disagrees, is listed with the reason
+//!   ([`Refusal`]); one that passes is a [`MetalMember`], present as
+//!   [`Presence::Detected`]. Elsewhere nothing is found, and
+//!   [`Discovery::unavailable`] says why.
+//! - **The matrix kernel** ([`MetalMember::matrix`]): the same device running
+//!   [`metal::MATRIX_SOURCE`], whose every product is formed by the GPU's
+//!   SIMD-group matrix multiply from 8-bit limbs; probed before it is
+//!   returned, with its own implementation id, so its certificates are
+//!   certificates of that path. A device below
+//!   [`metal::MATRIX_MIN_FAMILY`] is [`MetalError::Unsupported`].
+//! - **Running** ([`MetalMember::run_program`]): the state stays in one
+//!   buffer the device and the host share; one serial dispatch per gate, in
+//!   op-log order, in one command buffer; then the report word and the state
+//!   are read. The kernel reports exactly when the reference's result leaves
+//!   32 bits; the member then returns [`MetalError::OutOfRange`], never a
+//!   state.
+//! - **Certification, placement, receipts and metering** are as for the CUDA
+//!   member: [`MetalMember::certify`], [`MetalMember::member`],
+//!   [`place`](crate::quantum_fabric::place),
+//!   [`FabricReceipt`](crate::quantum_fabric::FabricReceipt) and a caller's
+//!   [`RunMeter`], whose window runs from the first upload to the last
+//!   readback.
+//!
+//! Every call is synchronous; the device is waited on before a call returns.
+//! No type of the platform's bindings appears in this module's public API.
+//!
+//! # Honest boundaries
+//!
+//! - **A member is what this process measured.** The probe and the
+//!   certificates are earned per process; the member id binds the device, its
+//!   registry id, its family, its memory and the operating system's version,
+//!   whose compiler built the kernels.
+//! - **Plans are per gate.** [`Config::PER_GATE`].
+//! - **Unmetered by default**, as for every member.
+
+use crate::quantum::{Amp, Circuit, Gateset, MAX_QUBITS, ONE, QuantumError, StateVector};
+use crate::quantum_energy::EnergyClass;
+use crate::quantum_fabric::lanes::{self, GateWords, Lane};
+use crate::quantum_fabric::metal::{self, DeviceIdentity, Shape};
+use crate::quantum_fabric::plan::{self, Config, LaneRun, Program, Step};
+use crate::quantum_fabric::{
+    Certificate, Certification, Corpus, EnergyCalibration, FabricClass, KernelEntry, Member, MemberError, Presence, RunFuture, RunMeter, Runner, Unmetered, apply_kernel,
+    pair_updates,
+};
+use std::sync::Arc;
+
+/// One probe case's device results: [`lanes::mul_wide`] and
+/// [`lanes::fxmul_wide`], each as `(high, low)` words.
+pub type ProbeOutput = ((u32, u32), (u32, u32));
+
+/// A run's lanes and report, with the meter's accepted reading (zero and
+/// `None` when unmetered).
+type MeteredRun = (LaneRun, (u64, Option<EnergyClass>));
+
+/// Why a device is not a member, or a run has no state.
+#[derive(Debug, PartialEq, Eq)]
+#[non_exhaustive]
+pub enum MetalError {
+    /// The platform has no Metal device this member reaches, and why.
+    Unavailable(String),
+    /// The runtime compiler refused a kernel, or a pipeline could not be made.
+    Compile(String),
+    /// The device reported an error, or a command buffer did not complete.
+    Device(String),
+    /// Out of memory on the device or the host.
+    Memory(String),
+    /// Exactly what the CPU reference refuses.
+    Circuit(QuantumError),
+    /// More qubits than this member's memory holds.
+    TooWide { n_qubits: u8, max: u8 },
+    /// The reference's result leaves 32 bits at this op; the member returns
+    /// no state.
+    OutOfRange { gate: u32 },
+    /// A program the host cannot lay out: 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 },
+    /// The device lacks what the kernel needs, and what.
+    Unsupported(String),
+}
+
+impl std::fmt::Display for MetalError {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        match self {
+            MetalError::Unavailable(why) => write!(f, "Metal is unavailable: {why}"),
+            MetalError::Compile(e) => write!(f, "the kernel did not compile: {e}"),
+            MetalError::Device(e) => write!(f, "device: {e}"),
+            MetalError::Memory(e) => write!(f, "out of memory: {e}"),
+            MetalError::Circuit(e) => write!(f, "{e}"),
+            MetalError::TooWide { n_qubits, max } => write!(f, "{n_qubits} qubits; this member holds at most {max}"),
+            MetalError::OutOfRange { gate } => write!(f, "the result of op {gate} leaves 32 bits"),
+            MetalError::Internal(why) => write!(f, "host defect: {why}"),
+            MetalError::Probe { a, b } => write!(f, "the device's arithmetic differs from the CPU's at ({a}, {b})"),
+            MetalError::Unsupported(why) => write!(f, "the device cannot run this kernel: {why}"),
+        }
+    }
+}
+
+impl std::error::Error for MetalError {}
+
+impl From<MetalError> for MemberError {
+    fn from(e: MetalError) -> MemberError {
+        match e {
+            MetalError::TooWide { n_qubits, max } => MemberError::TooWide { n_qubits, max },
+            MetalError::OutOfRange { gate } => MemberError::OutOfRange { gate },
+            MetalError::Memory(why) => MemberError::Memory(why),
+            MetalError::Device(why) => MemberError::Device(why),
+            MetalError::Circuit(e) => MemberError::Circuit(e),
+            other => MemberError::Other(other.to_string()),
+        }
+    }
+}
+
+/// A device that is not a member, and why.
+#[derive(Debug)]
+pub struct Refusal {
+    /// The device's name, or its position in the platform's list.
+    pub device: String,
+    /// Its CI key ([`DeviceIdentity::key`]), when the device could be read.
+    pub key: String,
+    pub error: MetalError,
+}
+
+/// What [`discover`] found.
+#[derive(Debug)]
+pub struct Discovery {
+    /// Devices that compiled the kernel and passed the probe.
+    pub found: Vec<MetalMember>,
+    /// Devices that did not, with why.
+    pub refused: Vec<Refusal>,
+    /// Why no device was even listed. `None` when the platform listed its
+    /// devices.
+    pub unavailable: Option<String>,
+}
+
+/// Every device the platform lists, each opened, compiled for and probed in
+/// this process before it is listed. Off macOS, nothing is found and
+/// [`Discovery::unavailable`] says why.
+pub fn discover() -> Discovery {
+    sys::discover()
+}
+
+/// The widest state of `2^n` interleaved 8-byte lanes that half of the
+/// device's recommended working set holds, in one buffer no longer than
+/// `max_buffer`, at most [`MAX_QUBITS`].
+#[cfg_attr(not(target_os = "macos"), allow(dead_code))]
+fn qubits_for(working_set: u64, max_buffer: u64) -> u8 {
+    let lanes = (working_set / 2 / 8).min(max_buffer / 8);
+    if lanes < 2 { 0 } else { (lanes.ilog2() as u8).min(MAX_QUBITS) }
+}
+
+/// One gate as the kernels read it: 13 32-bit words ([`metal::HOST_LAYOUT`]).
+fn gate_words(pairs: u64, w: &GateWords, index: u32) -> [u32; 13] {
+    let mut o = [0u32; 13];
+    o[0] = pairs as u32;
+    o[1] = (pairs >> 32) as u32;
+    o[2] = w.tbit;
+    o[3] = w.ctrl;
+    o[4..12].copy_from_slice(&w.m);
+    o[12] = index;
+    o
+}
+
+/// One Metal device as a fabric member: the device and a queue on it, a
+/// kernel compiled for it, and the arithmetic probe passed in this process.
+#[derive(Clone)]
+pub struct MetalMember {
+    gpu: Arc<sys::Gpu>,
+    kernel: Arc<sys::Kernel>,
+    identity: DeviceIdentity,
+    id: String,
+    key: String,
+    implementation: String,
+    /// What discovery found, which [`Self::detail`] extends for a kernel that
+    /// says more.
+    discovered: String,
+    detail: String,
+    max_qubits: u8,
+}
+
+impl std::fmt::Debug for MetalMember {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        f.debug_struct("MetalMember").field("id", &self.id).field("key", &self.key).field("implementation", &self.implementation).field("max_qubits", &self.max_qubits).finish()
+    }
+}
+
+impl MetalMember {
+    /// The member id ([`DeviceIdentity::member_id`]).
+    pub fn id(&self) -> &str {
+        &self.id
+    }
+
+    /// The key a runner declares for it ([`DeviceIdentity::key`]).
+    pub fn key(&self) -> &str {
+        &self.key
+    }
+
+    /// Always [`FabricClass::Gpu`].
+    pub fn class(&self) -> FabricClass {
+        self.identity.class()
+    }
+
+    /// What discovery found: device, family, memory, the operating system,
+    /// and what a metered run's window covers. Not signed.
+    pub fn detail(&self) -> &str {
+        &self.detail
+    }
+
+    /// The most qubits a run holds on this member.
+    pub fn max_qubits(&self) -> u8 {
+        self.max_qubits
+    }
+
+    /// The implementation id ([`metal::implementation_id`],
+    /// [`metal::matrix_implementation_id`], or a mutant's).
+    pub fn implementation(&self) -> &str {
+        &self.implementation
+    }
+
+    /// What the member id hashes.
+    pub fn identity(&self) -> &DeviceIdentity {
+        &self.identity
+    }
+
+    /// The device's [`lanes::mul_wide`] and [`lanes::fxmul_wide`] of each
+    /// pair, in order.
+    pub fn probe(&self, pairs: &[(i32, i32)]) -> Result<Vec<ProbeOutput>, MetalError> {
+        if pairs.is_empty() {
+            return Ok(Vec::new());
+        }
+        let flat: Vec<i32> = pairs.iter().flat_map(|&(a, b)| [a, b]).collect();
+        let words = sys::probe(&self.gpu, &self.kernel, &flat)?;
+        Ok(words.as_chunks::<4>().0.iter().map(|w| ((w[0], w[1]), (w[2], w[3]))).collect())
+    }
+
+    /// The planning this member runs: [`Config::PER_GATE`].
+    pub fn config(&self) -> Config {
+        Config::PER_GATE
+    }
+
+    /// Runs `p` from its basis state: the lanes, and the first op whose exact
+    /// result left 32 bits. A program another config planned (one with tile
+    /// steps) is [`MetalError::Internal`]: this member runs gate steps only.
+    pub fn run_program(&self, p: &Program) -> Result<LaneRun, MetalError> {
+        Ok(self.run_program_metered(p, &mut Unmetered)?.0)
+    }
+
+    fn run_program_metered<M: RunMeter + ?Sized>(&self, p: &Program, meter: &mut M) -> Result<MeteredRun, MetalError> {
+        let n = p.n_qubits;
+        if n == 0 {
+            return Ok((LaneRun { amps: vec![[ONE as i32, 0]], out_of_range: None }, (0, None)));
+        }
+        if n > self.max_qubits {
+            return Err(MetalError::TooWide { n_qubits: n, max: self.max_qubits });
+        }
+        if n > MAX_QUBITS || p.start >= 1usize << n {
+            return Err(MetalError::Internal("the start is outside the state"));
+        }
+        let len = 1usize << n;
+        let mut gates = Vec::with_capacity(p.steps.len());
+        for step in &p.steps {
+            let Step::Gate { gate } = *step else { return Err(MetalError::Internal("a tile step; this member runs gate steps only")) };
+            let g = p.gates.get(gate as usize).ok_or(MetalError::Internal("a step names a gate the program does not hold"))?;
+            if !g.is_well_formed() {
+                return Err(MetalError::Internal("malformed gate words for this state"));
+            }
+            gates.push((&g.words, g.index));
+        }
+        let mut init: Vec<i32> = Vec::new();
+        init.try_reserve_exact(2 * len).map_err(|e| MetalError::Memory(e.to_string()))?;
+        init.resize(2 * len, 0);
+        init[2 * p.start] = ONE as i32;
+        self.execute(init, &gates, meter)
+    }
+
+    /// Gates in order from `state`, as [`lanes::run_words`] runs them: the
+    /// lanes, and the lowest position in `gates` of a gate whose exact result
+    /// left 32 bits. A state that is not `2^n` lanes (`n ≥ 1`) or words
+    /// [`lanes::apply_words`] refuses for it are [`MetalError::Internal`].
+    pub fn run_lanes(&self, state: &[Lane], gates: &[GateWords]) -> Result<LaneRun, MetalError> {
+        let len = state.len();
+        if len < 2 || !len.is_power_of_two() {
+            return Err(MetalError::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(MetalError::TooWide { n_qubits: n, max: self.max_qubits });
+        }
+        let planned: Vec<(&GateWords, u32)> = gates.iter().zip(0u32..).collect();
+        let mut init: Vec<i32> = Vec::new();
+        init.try_reserve_exact(2 * len).map_err(|e| MetalError::Memory(e.to_string()))?;
+        init.extend(state.iter().flat_map(|l| [l[0], l[1]]));
+        Ok(self.execute(init, &planned, &mut Unmetered)?.0)
+    }
+
+    /// Uploads `init` (interleaved lanes, `2·2^n` words), runs `gates` in order
+    /// inside a window of `meter`, and reads the report and the state back.
+    /// Every gate's words are checked against the state first: the words
+    /// [`lanes::apply_words`] refuses, the host refuses.
+    fn execute<M: RunMeter + ?Sized>(&self, init: Vec<i32>, gates: &[(&GateWords, u32)], meter: &mut M) -> Result<MeteredRun, MetalError> {
+        let len = init.len() / 2;
+        for (w, _) in gates {
+            let (tbit, ctrl) = (w.tbit as usize, w.ctrl as usize);
+            if tbit == 0 || !tbit.is_power_of_two() || tbit > len / 2 || ctrl & tbit != 0 || ctrl & !(len - 1) != 0 {
+                return Err(MetalError::Internal("malformed gate words for this state"));
+            }
+        }
+        let pairs = (len / 2) as u64;
+        let words: Vec<[u32; 13]> = gates.iter().map(|&(w, index)| gate_words(pairs, w, index)).collect();
+        let (state, bad, reading) = sys::execute(&self.gpu, &self.kernel, init, pairs, &words, meter)?;
+        let out_of_range = Some(bad).filter(|&k| k != u32::MAX);
+        let mut lanes_out: Vec<Lane> = Vec::new();
+        lanes_out.try_reserve_exact(len).map_err(|e| MetalError::Memory(e.to_string()))?;
+        lanes_out.extend(state.as_chunks::<2>().0.iter().copied());
+        Ok((LaneRun { amps: lanes_out, out_of_range }, reading))
+    }
+
+    /// Equal to `c.simulate_from(start)`, or a typed refusal: what the CPU
+    /// reference refuses ([`MetalError::Circuit`]), more qubits than the
+    /// member holds, or a result that leaves 32 bits.
+    pub fn simulate_from(&self, c: &Circuit, start: usize) -> Result<StateVector, MetalError> {
+        Ok(self.simulate_with(c, start, &mut Unmetered)?.0)
+    }
+
+    /// [`Self::simulate_from`] inside a window of `meter`: the state, and the
+    /// microjoules and class the meter read, which are zero and `None` unless
+    /// the crate's label check accepts them. The window opens just before the
+    /// first upload and closes after the last readback.
+    pub fn simulate_metered(&self, c: &Circuit, start: usize, meter: &mut dyn RunMeter) -> Result<(StateVector, u64, Option<EnergyClass>), MetalError> {
+        self.simulate_with(c, start, meter)
+    }
+
+    fn simulate_with<M: RunMeter + ?Sized>(&self, c: &Circuit, start: usize, meter: &mut M) -> Result<(StateVector, u64, Option<EnergyClass>), MetalError> {
+        let p = plan::compile_with(c, start, &self.config()).map_err(MetalError::Circuit)?;
+        if p.n_qubits > self.max_qubits {
+            return Err(MetalError::TooWide { n_qubits: p.n_qubits, max: self.max_qubits });
+        }
+        let (run, (joules_micro, class)) = self.run_program_metered(&p, meter)?;
+        if let Some(gate) = run.out_of_range {
+            return Err(MetalError::OutOfRange { gate });
+        }
+        let mut amps = Vec::new();
+        amps.try_reserve_exact(run.amps.len()).map_err(|e| MetalError::Memory(e.to_string()))?;
+        amps.extend(run.amps.iter().map(|l| Amp { re: i64::from(l[0]), im: i64::from(l[1]) }));
+        Ok((StateVector { n_qubits: p.n_qubits, amps }, joules_micro, class))
+    }
+
+    /// [`Self::simulate_from`] the all-zeros state.
+    pub fn simulate(&self, c: &Circuit) -> Result<StateVector, MetalError> {
+        self.simulate_from(c, 0)
+    }
+
+    /// The certificate of [`apply_kernel`]`(gs)` as [`Self::implementation`]
+    /// on this member, over `corpus`: each circuit run from `|0…0⟩`, against
+    /// the native simulator. A range report inside a corpus is an error: its
+    /// circuits are valid and within the horizon.
+    pub fn certify(&self, gs: Gateset, corpus: Corpus) -> Result<Certificate, MetalError> {
+        Ok(self.certify_metered(gs, corpus, &mut Unmetered)?.0)
+    }
+
+    fn certify_metered<M: RunMeter + ?Sized>(&self, gs: Gateset, corpus: Corpus, meter: &mut M) -> Result<(Certificate, Option<(u64, EnergyClass)>), MetalError> {
+        let mut cert = Certification::new();
+        let mut calibration = EnergyCalibration::new();
+        for c in corpus.circuits(gs) {
+            let reference = lanes::reference_output(&c).map_err(MetalError::Circuit)?;
+            let (run, (joules_micro, class)) = self.run_program_metered(&plan::compile_with(&c, 0, &self.config()).map_err(MetalError::Circuit)?, &mut *meter)?;
+            if let Some(gate) = run.out_of_range {
+                return Err(MetalError::OutOfRange { gate });
+            }
+            cert.add(&lanes::circuit_bytes(&c), &reference, &lanes::output(&run.amps));
+            calibration.add(pair_updates(&c), joules_micro, class.as_ref());
+        }
+        Ok((cert.finish(apply_kernel(gs), &self.implementation, &self.id, &corpus.generator(gs)), calibration.finish()))
+    }
+
+    /// A fabric [`Member`]: present as [`Presence::Detected`], with one entry
+    /// per gateset, each certified in this call on `corpus`, and no energy
+    /// figure.
+    pub fn member(&self, corpus: Corpus) -> Result<Member, MetalError> {
+        self.member_with(corpus, &mut Unmetered)
+    }
+
+    /// [`Self::member`], each corpus run inside its own window of `meter`:
+    /// every entry also carries the energy figure its runs earn
+    /// ([`EnergyCalibration`]), or none if any run went unmetered.
+    pub fn member_metered(&self, corpus: Corpus, meter: &mut dyn RunMeter) -> Result<Member, MetalError> {
+        self.member_with(corpus, meter)
+    }
+
+    fn member_with<M: RunMeter + ?Sized>(&self, corpus: Corpus, meter: &mut M) -> Result<Member, MetalError> {
+        let mut kernels = Vec::new();
+        for gs in Gateset::ALL {
+            let (certificate, energy) = self.certify_metered(gs, corpus, &mut *meter)?;
+            kernels.push(KernelEntry { kernel: apply_kernel(gs).into(), implementation: self.implementation.clone(), certificate, energy });
+        }
+        Ok(Member { id: self.id.clone(), class: self.class(), presence: Presence::Detected, kernels, detail: self.detail.clone() })
+    }
+
+    /// The same member holding at most `max_qubits` qubits (no more than it
+    /// holds), so that tests can make it refuse a circuit as too wide.
+    #[doc(hidden)]
+    pub fn with_max_qubits(&self, max_qubits: u8) -> MetalMember {
+        MetalMember { max_qubits: max_qubits.min(self.max_qubits), ..self.clone() }
+    }
+
+    /// The same device running `source` as the integer kernel, under the
+    /// implementation id that text earns: for the mutation tests, which
+    /// require every mutant to be caught.
+    #[doc(hidden)]
+    pub fn with_source(&self, source: &str) -> Result<MetalMember, MetalError> {
+        let kernel = sys::compile(&self.gpu, source, Shape::LANES)?;
+        Ok(MetalMember { kernel: Arc::new(kernel), implementation: metal::implementation_id_of(source), detail: self.discovered.clone(), ..self.clone() })
+    }
+
+    /// The same device running the matrix kernel ([`metal::MATRIX_SOURCE`]),
+    /// compiled for it and probed in this call: the member id is the
+    /// device's, the implementation id the matrix kernel's. A device below
+    /// [`metal::MATRIX_MIN_FAMILY`] is [`MetalError::Unsupported`]; one whose
+    /// products disagree with the CPU's is [`MetalError::Probe`].
+    pub fn matrix(&self) -> Result<MetalMember, MetalError> {
+        self.with_matrix_source(metal::MATRIX_SOURCE)?.probed()
+    }
+
+    /// The same device running `source` as the matrix kernel, unprobed, under
+    /// the implementation id that text earns: for the mutation tests.
+    #[doc(hidden)]
+    pub fn with_matrix_source(&self, source: &str) -> Result<MetalMember, MetalError> {
+        if !self.identity.has_simdgroup_matrices() {
+            return Err(MetalError::Unsupported(format!(
+                "GPU family {}; the matrix kernel needs SIMD-group matrices, family {} or later",
+                self.identity.family,
+                metal::MATRIX_MIN_FAMILY
+            )));
+        }
+        let kernel = sys::compile(&self.gpu, source, Shape::MATRIX)?;
+        let detail = format!("{}; every product by the SIMD-group matrix multiply (bf16 limbs of 8 bits, f32 accumulation)", self.discovered);
+        Ok(MetalMember { kernel: Arc::new(kernel), implementation: metal::matrix_implementation_id_of(source), detail, ..self.clone() })
+    }
+
+    /// This member, once its kernel's arithmetic agrees with the CPU's on the
+    /// probe's pairs.
+    fn probed(self) -> Result<MetalMember, MetalError> {
+        let pairs = lanes::probe_pairs(4096, 0x5eed);
+        let out = self.probe(&pairs)?;
+        match lanes::first_mismatch(&pairs, &out) {
+            None => Ok(self),
+            Some((a, b)) => Err(MetalError::Probe { a, b }),
+        }
+    }
+}
+
+/// The member on the fabric ([`Fabric`](crate::quantum_fabric::Fabric)): its
+/// own implementation, through [`MetalMember::simulate_metered`]. Pair it with
+/// a [`Member`] whose entries name that implementation, such as
+/// [`MetalMember::member`]'s; another implementation is
+/// [`MemberError::Other`].
+impl Runner for MetalMember {
+    fn member_id(&self) -> String {
+        self.id.clone()
+    }
+
+    fn run<'r>(&'r self, implementation: &'r str, c: &'r Circuit, start: usize, meter: &'r mut dyn RunMeter) -> RunFuture<'r> {
+        Box::pin(async move {
+            if implementation != self.implementation {
+                return Err(MemberError::Other(format!("member {} runs implementation {}, not {implementation:?}", self.id, self.implementation)));
+            }
+            Ok(self.simulate_with(c, start, meter)?)
+        })
+    }
+}
+
+/// The platform's side: on macOS, the device, its queue, the compiled
+/// kernels, and the calls on them.
+#[cfg(target_os = "macos")]
+mod sys {
+    use super::{Discovery, MetalError, MetalMember, Refusal, qubits_for};
+    use crate::quantum_energy::EnergyClass;
+    use crate::quantum_fabric::metal::{self, DeviceIdentity, Shape};
+    use crate::quantum_fabric::{RunMeter, Window};
+    use objc2::rc::{Retained, autoreleasepool};
+    use objc2::runtime::ProtocolObject;
+    use objc2_foundation::{NSError, NSProcessInfo, NSString};
+    use objc2_metal::{
+        MTLBuffer, MTLCommandBuffer, MTLCommandBufferStatus, MTLCommandEncoder, MTLCommandQueue, MTLCompileOptions, MTLComputeCommandEncoder, MTLComputePipelineState,
+        MTLCopyAllDevices, MTLDevice, MTLGPUFamily, MTLLanguageVersion, MTLLibrary, MTLMathMode, MTLResourceOptions, MTLSize,
+    };
+    use std::ptr::NonNull;
+    use std::sync::Arc;
+
+    /// A device and the queue the member submits on.
+    pub struct Gpu {
+        device: Retained<ProtocolObject<dyn MTLDevice>>,
+        queue: Retained<ProtocolObject<dyn MTLCommandQueue>>,
+    }
+
+    /// A kernel compiled for one device: its two entry points and its shape.
+    pub struct Kernel {
+        apply: Retained<ProtocolObject<dyn MTLComputePipelineState>>,
+        probe: Retained<ProtocolObject<dyn MTLComputePipelineState>>,
+        shape: Shape,
+    }
+
+    type Buffer = Retained<ProtocolObject<dyn MTLBuffer>>;
+
+    fn ns_error(e: &NSError) -> String {
+        e.localizedDescription().to_string()
+    }
+
+    /// `"<major>.<minor>.<patch>"`.
+    fn os_version() -> String {
+        let v = NSProcessInfo::processInfo().operatingSystemVersion();
+        format!("{}.{}.{}", v.majorVersion, v.minorVersion, v.patchVersion)
+    }
+
+    /// The highest of the platform's own GPU families the device reports.
+    fn family(device: &ProtocolObject<dyn MTLDevice>) -> u32 {
+        (1001..=1012u32).rev().find(|&f| device.supportsFamily(MTLGPUFamily(f as isize))).unwrap_or(0)
+    }
+
+    pub fn discover() -> Discovery {
+        autoreleasepool(|_| {
+            let mut d = Discovery { found: Vec::new(), refused: Vec::new(), unavailable: None };
+            let devices = MTLCopyAllDevices();
+            if devices.count() == 0 {
+                d.unavailable = Some("the platform lists no Metal device".into());
+                return d;
+            }
+            for k in 0..devices.count() {
+                match open(devices.objectAtIndex(k)) {
+                    Ok(m) => d.found.push(m),
+                    Err(r) => d.refused.push(r),
+                }
+            }
+            d
+        })
+    }
+
+    /// Reads the device, compiles the integer kernel for it and probes it.
+    fn open(device: Retained<ProtocolObject<dyn MTLDevice>>) -> Result<MetalMember, Refusal> {
+        let identity = DeviceIdentity {
+            name: device.name().to_string(),
+            registry_id: device.registryID(),
+            family: family(&device),
+            unified_memory: device.hasUnifiedMemory(),
+            working_set: device.recommendedMaxWorkingSetSize(),
+            max_buffer: device.maxBufferLength() as u64,
+            os: os_version(),
+        };
+        let (name, key) = (identity.name.clone(), identity.key());
+        let refuse = |error: MetalError| Refusal { device: name.clone(), key: key.clone(), error };
+        let queue = device.newCommandQueue().ok_or_else(|| refuse(MetalError::Device("the device gave no command queue".into())))?;
+        let gpu = Arc::new(Gpu { device, queue });
+        let kernel = compile(&gpu, metal::SOURCE, Shape::LANES).map_err(refuse)?;
+        let max_qubits = qubits_for(identity.working_set, identity.max_buffer);
+        let detail = format!(
+            "metal; device {name:?}; GPU family {}; registry id {:x}; unified memory {}; working set {} bytes; buffers up to {} bytes; operating system {}; up to {max_qubits} qubits; a metered run's window runs from the first upload to the last readback",
+            identity.family, identity.registry_id, identity.unified_memory, identity.working_set, identity.max_buffer, identity.os,
+        );
+        let member = MetalMember {
+            gpu,
+            kernel: Arc::new(kernel),
+            id: identity.member_id(),
+            key: key.clone(),
+            implementation: metal::implementation_id(),
+            discovered: detail.clone(),
+            detail,
+            max_qubits,
+            identity,
+        };
+        member.probed().map_err(refuse)
+    }
+
+    /// `source` compiled for the device ([`metal::COMPILE_OPTIONS`]), its two
+    /// entry points as pipelines that take `shape.threads` a threadgroup.
+    pub fn compile(gpu: &Gpu, source: &str, shape: Shape) -> Result<Kernel, MetalError> {
+        autoreleasepool(|_| {
+            let options = MTLCompileOptions::new();
+            options.setLanguageVersion(MTLLanguageVersion::Version3_1);
+            options.setMathMode(MTLMathMode::Safe);
+            let library = gpu.device.newLibraryWithSource_options_error(&NSString::from_str(source), Some(&options)).map_err(|e| MetalError::Compile(ns_error(&e)))?;
+            let pipeline = |entry: &str| -> Result<Retained<ProtocolObject<dyn MTLComputePipelineState>>, MetalError> {
+                let function = library.newFunctionWithName(&NSString::from_str(entry)).ok_or_else(|| MetalError::Compile(format!("no entry point {entry}")))?;
+                let p = gpu.device.newComputePipelineStateWithFunction_error(&function).map_err(|e| MetalError::Compile(ns_error(&e)))?;
+                let most = p.maxTotalThreadsPerThreadgroup() as u64;
+                if most < shape.threads {
+                    return Err(MetalError::Unsupported(format!("{entry} takes at most {most} threads a threadgroup; the kernel needs {}", shape.threads)));
+                }
+                Ok(p)
+            };
+            Ok(Kernel { apply: pipeline("apply_gate")?, probe: pipeline("probe")?, shape })
+        })
+    }
+
+    /// A buffer of `bytes` (at least 4) the device and the host share.
+    fn buffer(gpu: &Gpu, bytes: usize) -> Result<Buffer, MetalError> {
+        gpu.device.newBufferWithLength_options(bytes.max(4), MTLResourceOptions::StorageModeShared).ok_or_else(|| MetalError::Memory(format!("a buffer of {bytes} bytes")))
+    }
+
+    /// A shared buffer holding `data`.
+    fn upload<T: Copy>(gpu: &Gpu, data: &[T]) -> Result<Buffer, MetalError> {
+        let bytes = std::mem::size_of_val(data);
+        let b = buffer(gpu, bytes)?;
+        // SAFETY: the buffer holds at least `bytes` bytes, and no command
+        // using it has been encoded.
+        unsafe { std::ptr::copy_nonoverlapping(data.as_ptr().cast::<u8>(), b.contents().as_ptr().cast::<u8>(), bytes) };
+        Ok(b)
+    }
+
+    /// The first `n` values of type `T` in `b`.
+    fn read<T: Copy>(b: &Buffer, n: usize) -> Result<Vec<T>, MetalError> {
+        let mut out: Vec<T> = Vec::new();
+        out.try_reserve_exact(n).map_err(|e| MetalError::Memory(e.to_string()))?;
+        // SAFETY: the buffer was made for at least `n` values of `T`, and the
+        // command buffer that wrote it has completed.
+        unsafe {
+            std::ptr::copy_nonoverlapping(b.contents().as_ptr().cast::<T>(), out.as_mut_ptr(), n);
+            out.set_len(n);
+        }
+        Ok(out)
+    }
+
+    /// Encodes with `encode` into one compute encoder (serial dispatches),
+    /// commits, and waits; an error unless the command buffer completed.
+    fn submit(gpu: &Gpu, encode: impl FnOnce(&ProtocolObject<dyn MTLComputeCommandEncoder>)) -> Result<(), MetalError> {
+        let cb = gpu.queue.commandBuffer().ok_or_else(|| MetalError::Device("the queue gave no command buffer".into()))?;
+        let encoder = cb.computeCommandEncoder().ok_or_else(|| MetalError::Device("the command buffer gave no compute encoder".into()))?;
+        encode(&encoder);
+        encoder.endEncoding();
+        cb.commit();
+        cb.waitUntilCompleted();
+        let status = cb.status();
+        if status != MTLCommandBufferStatus::Completed {
+            let why = cb.error().map(|e| ns_error(&e)).unwrap_or_else(|| format!("status {status:?}"));
+            return Err(MetalError::Device(format!("the command buffer did not complete: {why}")));
+        }
+        Ok(())
+    }
+
+    fn dispatch(encoder: &ProtocolObject<dyn MTLComputeCommandEncoder>, groups: u64, threads: u64) {
+        let size = |width: u64| MTLSize { width: width as usize, height: 1, depth: 1 };
+        encoder.dispatchThreadgroups_threadsPerThreadgroup(size(groups), size(threads));
+    }
+
+    /// Sets `value`'s bytes at buffer index `index`.
+    fn bytes<T: Copy>(encoder: &ProtocolObject<dyn MTLComputeCommandEncoder>, value: &T, index: usize) {
+        // SAFETY: the encoder copies `size_of::<T>()` bytes from `value`, which
+        // is live for the call, and every kernel reads the layout written.
+        unsafe { encoder.setBytes_length_atIndex(NonNull::from(value).cast(), std::mem::size_of::<T>(), index) };
+    }
+
+    fn bind(encoder: &ProtocolObject<dyn MTLComputeCommandEncoder>, b: &Buffer, index: usize) {
+        // SAFETY: each buffer is bound at the index its kernel parameter
+        // names, and is sized for what the kernel reads and writes there.
+        unsafe { encoder.setBuffer_offset_atIndex(Some(b), 0, index) };
+    }
+
+    /// The probe of `cases` (interleaved (a, b) pairs): four words a case.
+    pub fn probe(gpu: &Gpu, k: &Kernel, cases: &[i32]) -> Result<Vec<u32>, MetalError> {
+        autoreleasepool(|_| {
+            let n = cases.len() / 2;
+            let input = upload(gpu, cases)?;
+            let out = buffer(gpu, n * 16)?;
+            let count = [n as u32, ((n as u64) >> 32) as u32];
+            submit(gpu, |e| {
+                e.setComputePipelineState(&k.probe);
+                bind(e, &input, 0);
+                bind(e, &out, 1);
+                bytes(e, &count, 2);
+                dispatch(e, k.shape.groups(n as u64, k.shape.per_case), k.shape.threads);
+            })?;
+            read(&out, n * 4)
+        })
+    }
+
+    /// The gates `words` on `init` (interleaved lanes), each a serial dispatch
+    /// of one encoder, inside a window of `meter` from the upload to the
+    /// readback: the lanes, the report word and the reading.
+    #[allow(clippy::type_complexity)]
+    pub fn execute<M: RunMeter + ?Sized>(gpu: &Gpu, k: &Kernel, init: Vec<i32>, pairs: u64, words: &[[u32; 13]], meter: &mut M) -> Result<(Vec<i32>, u32, (u64, Option<EnergyClass>)), MetalError> {
+        let mut window = Window::new(meter);
+        window.open();
+        let (state, bad) = autoreleasepool(|_| -> Result<(Vec<i32>, u32), MetalError> {
+            let amps = upload(gpu, &init)?;
+            let report = upload(gpu, &[u32::MAX])?;
+            let groups = k.shape.groups(pairs, k.shape.per_pair);
+            submit(gpu, |e| {
+                e.setComputePipelineState(&k.apply);
+                bind(e, &amps, 0);
+                bind(e, &report, 2);
+                for w in words {
+                    bytes(e, w, 1);
+                    dispatch(e, groups, k.shape.threads);
+                }
+            })?;
+            let bad = read::<u32>(&report, 1)?[0];
+            Ok((read(&amps, init.len())?, bad))
+        })?;
+        Ok((state, bad, window.close()))
+    }
+}
+
+/// Off macOS there is no device, so no member can be made: these types have
+/// no values.
+#[cfg(not(target_os = "macos"))]
+mod sys {
+    use super::{Discovery, MetalError};
+    use crate::quantum_energy::EnergyClass;
+    use crate::quantum_fabric::RunMeter;
+    use crate::quantum_fabric::metal::Shape;
+
+    pub enum Gpu {}
+    pub enum Kernel {}
+
+    pub fn discover() -> Discovery {
+        Discovery { found: Vec::new(), refused: Vec::new(), unavailable: Some("this platform has no Metal; the member runs on macOS".into()) }
+    }
+
+    pub fn compile(gpu: &Gpu, _: &str, _: Shape) -> Result<Kernel, MetalError> {
+        match *gpu {}
+    }
+
+    pub fn probe(gpu: &Gpu, _: &Kernel, _: &[i32]) -> Result<Vec<u32>, MetalError> {
+        match *gpu {}
+    }
+
+    #[allow(clippy::type_complexity)]
+    pub fn execute<M: RunMeter + ?Sized>(gpu: &Gpu, _: &Kernel, _: Vec<i32>, _: u64, _: &[[u32; 13]], _: &mut M) -> Result<(Vec<i32>, u32, (u64, Option<EnergyClass>)), MetalError> {
+        match *gpu {}
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[test]
+    fn refusals_map_to_member_errors() {
+        assert!(MemberError::from(MetalError::Memory("x".into())).falls_back());
+        assert!(MemberError::from(MetalError::OutOfRange { gate: 3 }).falls_back());
+        assert!(MemberError::from(MetalError::TooWide { n_qubits: 5, max: 4 }).falls_back());
+        assert!(!MemberError::from(MetalError::Probe { a: 1, b: 2 }).falls_back());
+        assert!(!MemberError::from(MetalError::Internal("x")).falls_back());
+        assert!(!MemberError::from(MetalError::Unsupported("x".into())).falls_back());
+    }
+
+    #[test]
+    fn the_state_a_member_holds_is_half_its_working_set_in_one_buffer() {
+        assert_eq!(qubits_for(0, u64::MAX), 0);
+        assert_eq!(qubits_for(31, u64::MAX), 0);
+        assert_eq!(qubits_for(32, u64::MAX), 1);
+        // A 96 GiB working set is 6·2^30 lanes in half of it, so 32 qubits
+        // before the cap; a 1 GiB buffer limit holds 2^27 lanes.
+        assert_eq!(qubits_for(96 << 30, u64::MAX), MAX_QUBITS.min(32));
+        assert_eq!(qubits_for(96 << 30, 1 << 30), MAX_QUBITS.min(27));
+        assert_eq!(qubits_for(u64::MAX, u64::MAX), MAX_QUBITS);
+    }
+
+    #[test]
+    fn a_gate_is_thirteen_words_in_the_kernels_order() {
+        let w = GateWords { m: [1, 2, 3, 4, 5, 6, 7, 8], tbit: 4, ctrl: 3 };
+        assert_eq!(gate_words((7 << 32) | 9, &w, 11), [9, 7, 4, 3, 1, 2, 3, 4, 5, 6, 7, 8, 11]);
+        assert_eq!(std::mem::size_of::<[u32; 13]>(), 52, "the kernels' Gate struct: 13 four-byte fields");
+    }
+
+    #[cfg(not(target_os = "macos"))]
+    #[test]
+    fn off_macos_nothing_is_found_and_discovery_says_why() {
+        let d = discover();
+        assert!(d.found.is_empty() && d.refused.is_empty());
+        assert!(d.unavailable.as_deref().is_some_and(|w| w.contains("macOS")), "{:?}", d.unavailable);
+    }
+}
diff --git a/wai/wai-quantum/tests/cuda_tensor.rs b/wai/wai-quantum/tests/cuda_tensor.rs
new file mode 100644
index 000000000..b71bbff6b
--- /dev/null
+++ b/wai/wai-quantum/tests/cuda_tensor.rs
@@ -0,0 +1,445 @@
+//! The CUDA member's tensor-core kernel on the devices of the machine that runs
+//! this: every product the gate needs formed on the tensor cores, from 8-bit
+//! limbs in bf16 accumulated in f32, held to the same tests as the other
+//! kernel (`cuda_member`).
+//!
+//! The machine declares its devices in `WAI_QUANTUM_CUDA_EXPECT`, as for
+//! `cuda_member`: their keys, comma-separated and sorted, or `none`. Every
+//! declared device with tensor cores (compute capability 8.0 or later) must
+//! run the kernel and pass its probe; one without must be refused as
+//! unsupported. Each test prints what it established and, when
+//! `WAI_QUANTUM_CUDA_REPORT` names a file, appends the same lines to it.
+//!
+//! ```text
+//! WAI_QUANTUM_CUDA_EXPECT=gpu-cuda-sm89 \
+//!   cargo test --features quantum_fabric_cuda --test cuda_tensor -- --test-threads=1
+//! ```
+#![cfg(not(target_arch = "wasm32"))]
+
+use std::io::Write;
+use std::sync::OnceLock;
+use wai_quantum::ed25519_dalek::SigningKey;
+use wai_quantum::quantum::{BaseGate, Circuit, Gate, Gateset, MAX_QUBITS, QuantumError};
+use wai_quantum::quantum_fabric::cuda::{self, TENSOR_MUTANTS};
+use wai_quantum::quantum_fabric::lanes::{self, GateWords, Lane};
+use wai_quantum::quantum_fabric::plan::{Program, Step};
+use wai_quantum::quantum_fabric::wgsl::{Detector, Mutant};
+use wai_quantum::quantum_fabric::{
+    Corpus, CpuReference, FABRIC_DEFAULT, Fabric, FabricClass, FabricReceipt, MemberError, Parity, Presence, REFERENCE, apply_kernel, detect_local, place, ready,
+};
+use wai_quantum::quantum_fabric_cuda::{CudaError, CudaMember, Discovery, discover};
+
+/// The CPU's corpus digests: random-v1 (seed 7, 200 cases) and fabric-v1
+/// (seed 7, 200 + 24 cases), per gateset. A bit-exact member reproduces them.
+const PINS: [(Gateset, &str, &str); 2] = [
+    (Gateset::V1, "d18454c639acfd9c7b66450f4ded1aac4ab46cc10a7bedbcf4958a96e204fb36", "41c726a0ac80f7defef92e15ecf893bdb7a3076b5e7f670fa0b7142e5c033df7"),
+    (Gateset::V2, "3dbde276c1095bf827faec71d5980e050f8a55a5ba7d97cd3bbdaec3f43a13b2", "9a84b55af68f1a8b07390b234cf562dd8d85cef3185a697ab85690c8790a7809"),
+];
+
+const RANDOM_V1: Corpus = Corpus::Random { seed: 7, cases: 200 };
+
+static DISCOVERY: OnceLock<Discovery> = OnceLock::new();
+static TENSOR: OnceLock<Vec<CudaMember>> = OnceLock::new();
+
+fn discovery() -> &'static Discovery {
+    DISCOVERY.get_or_init(discover)
+}
+
+/// Prints `line` and appends it to the file `WAI_QUANTUM_CUDA_REPORT` names.
+fn report(line: &str) {
+    println!("{line}");
+    if let Some(path) = std::env::var_os("WAI_QUANTUM_CUDA_REPORT") {
+        let mut f = std::fs::OpenOptions::new().create(true).append(true).open(&path).expect("the report file opens");
+        writeln!(f, "{line}").expect("the report line is written");
+    }
+}
+
+/// The declared device set, as written, if it is exactly what was found;
+/// otherwise why not, with what was found and refused and how to declare it.
+fn declared(d: &Discovery) -> Result<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_CUDA_EXPECT={found_list}");
+    let expect = match std::env::var("WAI_QUANTUM_CUDA_EXPECT") {
+        Ok(v) if !v.is_empty() => v,
+        _ => return Err(format!("WAI_QUANTUM_CUDA_EXPECT is not set. {how}")),
+    };
+    let keys: Vec<&str> = if expect == "none" { Vec::new() } else { expect.split(',').collect() };
+    let well_formed = keys.windows(2).all(|w| w[0] <= w[1]) && keys.iter().all(|k| !k.is_empty() && *k != "none" && k.trim() == *k);
+    if !well_formed {
+        return Err(format!("WAI_QUANTUM_CUDA_EXPECT={expect:?} is not a sorted, comma-separated list of keys, nor none. {how}"));
+    }
+    for k in &keys {
+        if !found.contains(k)
+            && let Some(r) = d.refused.iter().find(|r| r.key == *k)
+        {
+            return Err(format!("the declared device {k} was refused: {}. {how}", r.error));
+        }
+    }
+    if found != keys {
+        return Err(format!("the devices found are not the declared ones ({expect}). {how}"));
+    }
+    Ok(expect)
+}
+
+/// The tensor-core members: one per declared device with tensor cores, once
+/// the declared devices are exactly the ones found.
+fn members() -> &'static [CudaMember] {
+    let d = discovery();
+    if let Err(why) = declared(d) {
+        panic!("{why}");
+    }
+    TENSOR.get_or_init(|| {
+        d.found.iter().filter(|m| m.identity().has_tensor_cores()).map(|m| m.tensor().unwrap_or_else(|e| panic!("{}: the tensor-core kernel: {e}", m.key()))).collect()
+    })
+}
+
+fn hex(b: &[u8]) -> String {
+    b.iter().map(|x| format!("{x:02x}")).collect()
+}
+
+fn next(s: &mut u64) -> u64 {
+    *s = s.wrapping_add(0x9E37_79B9_7F4A_7C15);
+    let mut z = *s;
+    z = (z ^ (z >> 30)).wrapping_mul(0xBF58_476D_1CE4_E5B9);
+    z = (z ^ (z >> 27)).wrapping_mul(0x94D0_49BB_1331_11EB);
+    z ^ (z >> 31)
+}
+
+/// 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!("cuda tensor devices: as declared ({expect})")),
+        Err(why) => panic!("{why}"),
+    }
+    for m in &d.found {
+        if !m.identity().has_tensor_cores() {
+            assert!(matches!(m.tensor(), Err(CudaError::Unsupported(_))), "{}", m.key());
+            report(&format!("cuda tensor unsupported: key={} compute {:?}", m.key(), m.identity().compute));
+        }
+    }
+    for t in members() {
+        assert_eq!(t.class(), FabricClass::Gpu);
+        assert!(t.id().starts_with("local/gpu-cuda-"), "{}", t.id());
+        assert_eq!(t.implementation(), cuda::tensor_implementation_id());
+        assert_ne!(t.implementation(), cuda::implementation_id());
+        assert_eq!(t.identity().member_id(), t.id());
+        let lanes64 = d.found.iter().find(|m| m.id() == t.id()).expect("the device's member");
+        assert_eq!((lanes64.implementation(), lanes64.max_qubits()), (cuda::implementation_id().as_str(), t.max_qubits()));
+        assert!(t.max_qubits() >= 20, "{}: holds only {} qubits", t.key(), t.max_qubits());
+        assert!(t.detail().starts_with(lanes64.detail()) && t.detail().contains("tensor cores"), "{}", t.detail());
+        report(&format!("cuda tensor member: key={} id={} implementation={} {}", t.key(), t.id(), t.implementation(), t.detail()));
+    }
+}
+
+#[test]
+fn c1_probe_matches_lanes() {
+    for m in members() {
+        let pairs = lanes::probe_pairs(1 << 16, 0x0cda);
+        let out = m.probe(&pairs).unwrap_or_else(|e| panic!("{}: {e}", m.key()));
+        assert_eq!(lanes::first_mismatch(&pairs, &out), None, "{}", m.key());
+        report(&format!("cuda tensor 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!("cuda tensor kernel: key={} ok", m.key()));
+    }
+}
+
+#[test]
+fn c3_certificates_reproduce_the_pins() {
+    let ms = members();
+    for (gs, random_pin, fabric_pin) in PINS {
+        let cpu = lanes::certify_corpus(gs, "local/test", RANDOM_V1);
+        assert_eq!(hex(&cpu.corpus), random_pin, "{gs:?}");
+        let other = Corpus::Random { seed: 8, cases: 300 };
+        let cpu_other = lanes::certify_corpus(gs, "local/test", other);
+        for m in ms {
+            for (corpus, cases, pin) in [(RANDOM_V1, 200, random_pin), (FABRIC_DEFAULT, 224, fabric_pin)] {
+                let c = m.certify(gs, corpus).unwrap_or_else(|e| panic!("{} {gs:?} {}: {e}", m.key(), corpus.name()));
+                assert_eq!(c.parity, Parity::BitExact, "{} {gs:?} {}", m.key(), corpus.name());
+                assert_eq!(c.cases, cases);
+                assert_eq!(hex(&c.corpus), pin, "{} {gs:?} {}", m.key(), corpus.name());
+                assert_eq!((c.kernel.as_str(), c.implementation.as_str(), c.member.as_str()), (apply_kernel(gs), m.implementation(), m.id()));
+                assert_eq!(c.generator, corpus.generator(gs));
+                report(&format!(
+                    "cuda tensor certified: key={} member={} kernel={} corpus={} implementation={} parity=bit-exact cases={} digest={}",
+                    m.key(),
+                    m.id(),
+                    c.kernel,
+                    corpus.name(),
+                    c.implementation,
+                    c.cases,
+                    hex(&c.corpus)
+                ));
+            }
+            // A corpus with no pin: the member and the CPU agree on it.
+            let c = m.certify(gs, other).unwrap_or_else(|e| panic!("{} {gs:?} seed 8: {e}", m.key()));
+            assert_eq!((c.parity, c.cases), (Parity::BitExact, 300), "{} {gs:?}", m.key());
+            assert_eq!(c.corpus, cpu_other.corpus, "{} {gs:?}", m.key());
+        }
+    }
+}
+
+#[test]
+fn c4_wide_circuits_match_native() {
+    for m in members() {
+        let widest = m.max_qubits().min(22);
+        let mut s = 0x0077_6964_6500_u64;
+        for (k, n) in [12u8, 16, 18, 20, widest].into_iter().enumerate() {
+            for gs in Gateset::ALL {
+                let c = random_circuit(n, 40, 0x1000 + k as u64, gs);
+                let start = (next(&mut s) as usize) % (1usize << n);
+                let want = c.simulate_from(start).expect("the native run");
+                let got = m.simulate_from(&c, start).unwrap_or_else(|e| panic!("{}: n = {n} {gs:?}: {e}", m.key()));
+                assert!(got == want, "{}: n = {n} {gs:?}: the states differ", m.key());
+            }
+        }
+        // The transforms the conformance corpus carries.
+        for n in [3u8, 5, 8, 12] {
+            let qft = Circuit::qft(n);
+            assert!(m.simulate(&qft).unwrap() == qft.simulate().unwrap(), "{}: QFT-{n}", m.key());
+        }
+        report(&format!("cuda tensor wide: key={} up to {widest} qubits match native", m.key()));
+    }
+}
+
+#[test]
+fn c5_refusals_are_typed() {
+    for m in members() {
+        for c in refused_circuits() {
+            let native = c.simulate_from(0).unwrap_err();
+            assert_eq!(m.simulate_from(&c, 0).unwrap_err(), CudaError::Circuit(native), "{c:?}");
+        }
+        assert_eq!(m.simulate(&Circuit::new(MAX_QUBITS + 1)).unwrap_err(), CudaError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1)));
+        let small = m.with_max_qubits(4);
+        assert_eq!(small.simulate(&Circuit::new(5)).unwrap_err(), CudaError::TooWide { n_qubits: 5, max: 4 });
+        let wide = Program { n_qubits: 5, start: 0, gates: Vec::new(), steps: Vec::new() };
+        assert_eq!(small.run_program(&wide).unwrap_err(), CudaError::TooWide { n_qubits: 5, max: 4 });
+        // A result that leaves 32 bits is reported at its op, and the lane
+        // holds its low 32 bits.
+        let r = m.run_lanes(&Z_STATE, &[z_words()]).unwrap();
+        assert_eq!(r.out_of_range, Some(0), "{}", m.key());
+        assert_eq!(r.amps[1][0] as u32, 0x8000_0000);
+        // Malformed words, a state the host cannot lay out and a tile step are
+        // host defects, never results.
+        let bad = GateWords { tbit: 3, ..z_words() };
+        assert!(matches!(m.run_lanes(&[[0, 0]; 4], &[bad]), Err(CudaError::Internal(_))));
+        assert!(matches!(m.run_lanes(&[[0, 0]; 3], &[z_words()]), Err(CudaError::Internal(_))));
+        let tiled = Program { n_qubits: 2, start: 0, gates: Vec::new(), steps: vec![Step::Tile { first: 0, count: 1, size: 2 }] };
+        assert!(matches!(m.run_program(&tiled), Err(CudaError::Internal(_))));
+        // The refusals that may still fall back to another member, do.
+        assert!(MemberError::from(CudaError::TooWide { n_qubits: 5, max: 4 }).falls_back());
+        report(&format!("cuda tensor refusals: key={} typed", m.key()));
+    }
+}
+
+/// Whether a member running `mutant`'s text is caught by the mutant's
+/// detector.
+fn caught(mm: &CudaMember, mutant: &Mutant, key: &str) -> bool {
+    match mutant.detector {
+        Detector::Probe => {
+            let edges = lanes::probe_pairs(0, 0);
+            let out = mm.probe(&edges).expect("the mutant's probe runs");
+            let differing = edges.iter().zip(&out).filter(|&(&(a, b), o)| *o != (lanes::mul_wide(a, b), lanes::fxmul_wide(a, b))).count();
+            println!("cuda tensor mutant {} on {key}: {differing} of 484 edge pairs differ", mutant.name);
+            differing >= 1
+        }
+        Detector::Certificate => Gateset::ALL.into_iter().all(|gs| match mm.certify(gs, RANDOM_V1) {
+            Err(e) => {
+                println!("cuda tensor 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 tensor-core 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 &TENSOR_MUTANTS {
+            let mm = m.with_tensor_source(&cuda::tensor_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, TENSOR_MUTANTS.len());
+        report(&format!("cuda tensor mutants: key={} killed {killed} of {}", m.key(), TENSOR_MUTANTS.len()));
+    }
+}
+
+#[test]
+fn c7_receipt_placement_and_the_fabric() {
+    for m in members() {
+        let member = m.member(FABRIC_DEFAULT).unwrap_or_else(|e| panic!("{}: {e}", m.key()));
+        assert_eq!((member.id.as_str(), member.class, member.presence), (m.id(), FabricClass::Gpu, Presence::Detected));
+        assert_eq!(member.kernels.len(), 2);
+        for (entry, gs) in member.kernels.iter().zip(Gateset::ALL) {
+            assert_eq!((entry.kernel.as_str(), entry.implementation.as_str()), (apply_kernel(gs), m.implementation()));
+            assert!(entry.energy.is_none());
+            assert_eq!(entry.certificate.parity, Parity::BitExact, "{gs:?}");
+        }
+        // A receipt over a member's run.
+        let qft = Circuit::qft(5);
+        let state = m.simulate(&qft).unwrap();
+        assert_eq!(state, qft.simulate().unwrap());
+        let entry = member.kernels.iter().find(|k| k.kernel == apply_kernel(qft.gateset)).unwrap();
+        let signer = SigningKey::from_bytes(&[7; 32]);
+        let receipt = FabricReceipt::seal(&signer, "local/test", state.statevector_hash(), &member, entry, 0, None);
+        assert!(receipt.verify());
+        assert!(receipt.certificate_matches(&entry.certificate));
+        // Placement: beside the CPU reference with no energy figures, the
+        // reference; alone, the member.
+        let local = detect_local();
+        for gs in Gateset::ALL {
+            let both = [local.clone(), member.clone()];
+            let (placed, k) = place(&both, apply_kernel(gs), &Parity::BitExact).unwrap();
+            assert_eq!((placed.id.as_str(), k.implementation.as_str()), (local.id.as_str(), REFERENCE));
+            let (placed, k) = place(std::slice::from_ref(&member), apply_kernel(gs), &Parity::BitExact).unwrap();
+            assert_eq!((placed.id.as_str(), k.implementation.as_str()), (m.id(), m.implementation()));
+        }
+        // On the fabric: the member alone runs the circuit; a pairing with
+        // another member's entries is refused.
+        let mut fabric = Fabric::new();
+        fabric.add(member.clone(), m);
+        let run = ready(fabric.run(&qft, 0, &Parity::BitExact, &signer, "local/test")).expect("resolves").expect("runs");
+        assert_eq!(run.state, state);
+        assert_eq!(run.receipt.member_id, m.id());
+        assert!(run.receipt.verify());
+        let mut mismatched = Fabric::new();
+        mismatched.add(local.clone(), m);
+        assert!(ready(mismatched.run(&qft, 0, &Parity::BitExact, &signer, "local/test")).expect("resolves").is_err());
+        let mut both = Fabric::new();
+        both.add(local.clone(), CpuReference).add(member.clone(), m);
+        let run = ready(both.run(&qft, 0, &Parity::BitExact, &signer, "local/test")).expect("resolves").expect("runs");
+        assert_eq!(run.state, state);
+        report(&format!("cuda tensor receipt: key={} ok", m.key()));
+    }
+}
+
+#[test]
+fn c8_repeated_runs_agree() {
+    for m in members() {
+        let c = random_circuit(16, 60, 0xabc, Gateset::V2);
+        let first = m.simulate(&c).unwrap();
+        for _ in 0..5 {
+            assert!(m.simulate(&c).unwrap() == first, "{}", m.key());
+        }
+        report(&format!("cuda tensor repeat: key={} ok", m.key()));
+    }
+}
diff --git a/wai/wai-quantum/tests/metal_matrix.rs b/wai/wai-quantum/tests/metal_matrix.rs
new file mode 100644
index 000000000..27519a7ca
--- /dev/null
+++ b/wai/wai-quantum/tests/metal_matrix.rs
@@ -0,0 +1,445 @@
+//! The Metal member's matrix kernel on the devices of the machine that runs
+//! this: every product the gate needs formed by the GPU's SIMD-group matrix
+//! multiply, from 8-bit limbs in bf16 accumulated in f32, held to the same
+//! tests as the other kernel (`metal_member`).
+//!
+//! The machine declares its devices in `WAI_QUANTUM_METAL_EXPECT`, as for
+//! `metal_member`: their keys, comma-separated and sorted, or `none`. Every
+//! declared device with SIMD-group matrices (GPU family 1007 or later) must
+//! run the kernel and pass its probe; one without must be refused as
+//! unsupported. Each test prints what it established and, when
+//! `WAI_QUANTUM_METAL_REPORT` names a file, appends the same lines to it.
+//!
+//! ```text
+//! WAI_QUANTUM_METAL_EXPECT=gpu-msl-1009 \
+//!   cargo test --features quantum_fabric_metal --test metal_matrix -- --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::metal::{self, MATRIX_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_metal::{Discovery, MetalError, MetalMember, 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();
+static MATRIX: OnceLock<Vec<MetalMember>> = OnceLock::new();
+
+fn discovery() -> &'static Discovery {
+    DISCOVERY.get_or_init(discover)
+}
+
+/// Prints `line` and appends it to the file `WAI_QUANTUM_METAL_REPORT` names.
+fn report(line: &str) {
+    println!("{line}");
+    if let Some(path) = std::env::var_os("WAI_QUANTUM_METAL_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_METAL_EXPECT={found_list}");
+    let expect = match std::env::var("WAI_QUANTUM_METAL_EXPECT") {
+        Ok(v) if !v.is_empty() => v,
+        _ => return Err(format!("WAI_QUANTUM_METAL_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_METAL_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 matrix members: one per declared device with SIMD-group matrices, once
+/// the declared devices are exactly the ones found.
+fn members() -> &'static [MetalMember] {
+    let d = discovery();
+    if let Err(why) = declared(d) {
+        panic!("{why}");
+    }
+    MATRIX.get_or_init(|| {
+        d.found.iter().filter(|m| m.identity().has_simdgroup_matrices()).map(|m| m.matrix().unwrap_or_else(|e| panic!("{}: the matrix kernel: {e}", m.key()))).collect()
+    })
+}
+
+fn hex(b: &[u8]) -> String {
+    b.iter().map(|x| format!("{x:02x}")).collect()
+}
+
+fn next(s: &mut u64) -> u64 {
+    *s = s.wrapping_add(0x9E37_79B9_7F4A_7C15);
+    let mut z = *s;
+    z = (z ^ (z >> 30)).wrapping_mul(0xBF58_476D_1CE4_E5B9);
+    z = (z ^ (z >> 27)).wrapping_mul(0x94D0_49BB_1331_11EB);
+    z ^ (z >> 31)
+}
+
+/// 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!("metal matrix devices: as declared ({expect})")),
+        Err(why) => panic!("{why}"),
+    }
+    for m in &d.found {
+        if !m.identity().has_simdgroup_matrices() {
+            assert!(matches!(m.matrix(), Err(MetalError::Unsupported(_))), "{}", m.key());
+            report(&format!("metal matrix unsupported: key={} family {}", m.key(), m.identity().family));
+        }
+    }
+    for t in members() {
+        assert_eq!(t.class(), FabricClass::Gpu);
+        assert!(t.id().starts_with("local/gpu-msl-"), "{}", t.id());
+        assert_eq!(t.implementation(), metal::matrix_implementation_id());
+        assert_ne!(t.implementation(), metal::implementation_id());
+        assert_eq!(t.identity().member_id(), t.id());
+        let lanes64 = d.found.iter().find(|m| m.id() == t.id()).expect("the device's member");
+        assert_eq!((lanes64.implementation(), lanes64.max_qubits()), (metal::implementation_id().as_str(), t.max_qubits()));
+        assert!(t.max_qubits() >= 20, "{}: holds only {} qubits", t.key(), t.max_qubits());
+        assert!(t.detail().starts_with(lanes64.detail()) && t.detail().contains("SIMD-group matrix"), "{}", t.detail());
+        report(&format!("metal matrix member: key={} id={} implementation={} {}", t.key(), t.id(), t.implementation(), t.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!("metal matrix 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!("metal matrix 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!(
+                    "metal matrix 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!("metal matrix 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(), MetalError::Circuit(native), "{c:?}");
+        }
+        assert_eq!(m.simulate(&Circuit::new(MAX_QUBITS + 1)).unwrap_err(), MetalError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1)));
+        let small = m.with_max_qubits(4);
+        assert_eq!(small.simulate(&Circuit::new(5)).unwrap_err(), MetalError::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(), MetalError::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(MetalError::Internal(_))));
+        assert!(matches!(m.run_lanes(&[[0, 0]; 3], &[z_words()]), Err(MetalError::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(MetalError::Internal(_))));
+        // The refusals that may still fall back to another member, do.
+        assert!(MemberError::from(MetalError::TooWide { n_qubits: 5, max: 4 }).falls_back());
+        report(&format!("metal matrix refusals: key={} typed", m.key()));
+    }
+}
+
+/// Whether a member running `mutant`'s text is caught by the mutant's
+/// detector.
+fn caught(mm: &MetalMember, 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!("metal matrix 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!("metal matrix 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 matrix 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 &MATRIX_MUTANTS {
+            let mm = m.with_matrix_source(&metal::matrix_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, MATRIX_MUTANTS.len());
+        report(&format!("metal matrix mutants: key={} killed {killed} of {}", m.key(), MATRIX_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!("metal matrix 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!("metal matrix repeat: key={} ok", m.key()));
+    }
+}
diff --git a/wai/wai-quantum/tests/metal_member.rs b/wai/wai-quantum/tests/metal_member.rs
new file mode 100644
index 000000000..3a9768e1c
--- /dev/null
+++ b/wai/wai-quantum/tests/metal_member.rs
@@ -0,0 +1,429 @@
+//! The fabric's Metal member on the devices of the machine that runs this.
+//!
+//! The machine declares the devices it must have in `WAI_QUANTUM_METAL_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_METAL_REPORT` names a file, appends the
+//! same lines to it, which a CI step can count per device.
+//!
+//! ```text
+//! WAI_QUANTUM_METAL_EXPECT=gpu-msl-1009 \
+//!   cargo test --features quantum_fabric_metal --test metal_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::metal::{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_metal::{Discovery, MetalError, MetalMember, 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_METAL_REPORT` names.
+fn report(line: &str) {
+    println!("{line}");
+    if let Some(path) = std::env::var_os("WAI_QUANTUM_METAL_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_METAL_EXPECT={found_list}");
+    let expect = match std::env::var("WAI_QUANTUM_METAL_EXPECT") {
+        Ok(v) if !v.is_empty() => v,
+        _ => return Err(format!("WAI_QUANTUM_METAL_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_METAL_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 [MetalMember] {
+    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!("metal 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-msl-"), "{}", m.id());
+        assert_eq!(m.implementation(), metal::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!("metal 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!("metal 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!("metal 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!(
+                    "metal 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!("metal 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(), MetalError::Circuit(native), "{c:?}");
+        }
+        assert_eq!(m.simulate(&Circuit::new(MAX_QUBITS + 1)).unwrap_err(), MetalError::Circuit(QuantumError::TooManyQubits(MAX_QUBITS + 1)));
+        let small = m.with_max_qubits(4);
+        assert_eq!(small.simulate(&Circuit::new(5)).unwrap_err(), MetalError::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(), MetalError::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(MetalError::Internal(_))));
+        assert!(matches!(m.run_lanes(&[[0, 0]; 3], &[z_words()]), Err(MetalError::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(MetalError::Internal(_))));
+        // The refusals that may still fall back to another member, do.
+        assert!(MemberError::from(MetalError::TooWide { n_qubits: 5, max: 4 }).falls_back());
+        report(&format!("metal refusals: key={} typed", m.key()));
+    }
+}
+
+/// Whether a member running `mutant`'s text is caught by the mutant's
+/// detector.
+fn caught(mm: &MetalMember, 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!("metal 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!("metal 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 Metal 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(&metal::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!("metal 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!("metal 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!("metal repeat: key={} ok", m.key()));
+    }
+}
diff --git a/wai/wai-rs/Cargo.lock b/wai/wai-rs/Cargo.lock
index 4ef93f9de..537bb75e9 100644
--- a/wai/wai-rs/Cargo.lock
+++ b/wai/wai-rs/Cargo.lock
@@ -4310,7 +4310,7 @@ version = "0.1.0"
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "ed25519-dalek",
diff --git a/wai/wai-web/Cargo.lock b/wai/wai-web/Cargo.lock
index 8295b5692..67cf542a0 100644
--- a/wai/wai-web/Cargo.lock
+++ b/wai/wai-web/Cargo.lock
@@ -1241,7 +1241,7 @@ version = "0.1.0"
 
 [[package]]
 name = "wai-quantum"
-version = "0.6.0"
+version = "0.6.1"
 dependencies = [
  "blake3",
  "ed25519-dalek",