wai-quantum 0.6.1: the CUDA tensor-core kernel and the Metal member #275

Merged

dcharlot merged 3 commits from wai/quantum-tensor-cores into main 2026-10-16 09:45:04 -04:00

Conversation 0 Commits 3 Files changed 25 +3486 -30

dcharlot commented 2026-10-16 04:21:30 -04:00

Owner

Copy link

Two additions to wai-quantum, released together as 0.6.1 (additive; every pinned digest unchanged):

1. The CUDA member's tensor-core kernel: every product on the tensor cores.
2. The Metal member (quantum_fabric_metal, macOS), with an exact integer kernel and a second kernel whose every product uses the GPU's SIMD-group matrix multiply.

Both matrix paths use the same argument (below), and each kernel has its own implementation id and mutants.

1. The CUDA tensor-core kernel

tensor_limbs.cu: a second CUDA kernel for the gate, whose every product is formed on the tensor cores. It runs on devices of compute capability 8.0 or later.

How it stays exact

• Each 32-bit value is four 8-bit limbs in bf16: the lower three unsigned, the top one signed. bf16 holds every one exactly.
• One warp multiplies sixteen pairs' limbs (16 × 16) by the gate's limb weights (16 × 112: sixteen products per pair, seven coefficients each) in fourteen mma.sync m16n8k16 steps, accumulated in f32.
• Every limb product is at most 2^16, and every coefficient sums at most four of them. So every value the tensor core forms is an integer below 2^18, which f32 holds exactly in any order of addition.
• The kernel adds the coefficients, shifted by 8s bits, in 64-bit integers. That gives the exact product; the rest of the gate is the other kernel's arithmetic.

API (additive, 0.6.1)

• cuda::TENSOR_SOURCE, TENSOR_IMPLEMENTATION, tensor_implementation_id{,_of}, TENSOR_MUTANTS, tensor_mutate, TENSOR_MIN_COMPUTE, TENSOR_THREADS_PER_ITEM, DeviceIdentity::has_tensor_cores.
• CudaMember::tensor(): the same device, member id unchanged, under the kernel's own implementation id (cuda-tensor-limbs8/…). It is probed before it is returned.
• CudaError::Unsupported (new) for a device below 8.0.
• The compiled kernel carries its launch spread (two threads per pair), so the probe and gate launches cover every item.

Mutants (7)

	Mutant	Caught by
T1	rounding constant 2^28	probe
T2	shift by 29	probe
T3	nine-bit limbs, which bf16 does not hold	probe
T4	the B fragment's second half loaded with its rows swapped	probe
T5	control test inverted	certificate
T6	the sign of row 0's last term	certificate
T7	range guard never fires	range

Measured

On jetson-hub's discrete GPU (sm_89, silicon), tests/cuda_tensor.rs passed 9/9:

• 66020 probe pairs agree.
• All four pinned certificates reproduce (d18454c6…, 41c726a0…, 3dbde276…, 9a84b55a…).
• Circuits up to 22 qubits match native.
• 7 of 7 mutants are killed. On the 484 edge pairs, T1 differs on 80, T2 on 436, T3 on 297 and T4 on 210.
• cuda_member (9/9, 6 of 6 mutants) and cuda_sv (3/3, pin b7d129d0…) pass unchanged on the same tree.

On the build host, which has no CUDA library, cuda_tensor and cuda_member each pass 9/9 over no device (none).

I checked the workflow's report checks against the real report: the device set, the member line, 4 certificates and 7 of 7 mutants all match.

CI

• wai-quantum-cuda, cuda-hw job: a new tensor-core step that requires 9 passed, the member line, 4 cuda-tensor-limbs8 certificates and 7 of 7 mutants.
• cuda-absent job: runs the new tests over no device.

2. The Metal member

quantum_fabric_metal (opt-in, macOS) runs the gate on the platform's GPU through its own interface. The operating system's compiler builds the kernels at run time. The state lives in one buffer the GPU and the host share, and each gate is one serial dispatch in a single command buffer.

• lanes64.metal (metal-lanes64/…): the CUDA kernel's exact 64-bit integer arithmetic in the Metal shading language.
• matrix_limbs.metal (metal-matrix-limbs8/…, MetalMember::matrix()): every product formed by the SIMD-group matrix multiply, from 8-bit limbs in bf16 accumulated in f32.
  • One SIMD group multiplies eight pairs' limbs (8 × 16) by the gate's limb weights (16 × 112) in fourteen 8 × 8 column blocks.
  • Each of a pair's four threads reads one amplitude component, joins its four products, and writes the row that replaces that component.
  • A device below GPU family 1007 gets MetalError::Unsupported.
• Keys are the GPU family number (gpu-msl-<family>), not a vendor name. The member id binds the device, its registry id, family and memory, and the operating system's version.
• Off macOS the member still builds (its platform types have no values there), finds nothing, and says why.

Mutants

• Integer kernel (G1–G6): the CUDA kernel's six, in Metal text.
• Matrix kernel (X1–X7): rounding, shift, nine-bit limbs, the second limb block multiplying the first block's weights, control inverted, a real row's sign, and the range guard.

Measured

On the M3 Ultra build host, the same machine as the macos runner (GPU family 1009, silicon), metal_member and metal_matrix each passed 9/9:

• 66020 probe pairs agree.
• All four pinned certificates reproduce.
• Circuits up to 22 qubits match native.
• 6 of 6 and 7 of 7 mutants are killed. On the 484 edge pairs, X3 (nine-bit limbs) breaks 297, the same count as on the CUDA tensor cores, which shows the multiply really takes bf16 inputs. X4 breaks 443.

On Linux (jetson-hub), the unit tests pass, both suites pass 9/9 over no device, and clippy is clean. Clippy is also clean on macOS, and I checked the workflow's report greps against the real report.

Not done: Metal 4's tensor-op multiply, the path to M5's GPU neural accelerators. The only M5 in the fleet is the dev Mac, where builds and tests are off-limits.

CI

wai-quantum-metal:

• metal-hw, on the macos runner: unit tests; 9 passed, 4 certificates and 6 of 6 mutants for the integer kernel; then the member line, 4 matrix certificates and 7 of 7 mutants for the matrix kernel.
• metal-absent, on linux: both suites over no device, plus the off-platform discovery test.

The CHANGELOG also records two hand measurements made outside CI: the tensor program on a cloud tensor processor, and the QPU crate against its service's simulator.

All repo checks are clean: prose, vanity metrics, tests-in-CI, concurrency, lockfiles, orphans and public names.

🤖 Generated with Claude Code

Two additions to wai-quantum, released together as 0.6.1 (additive; every pinned digest unchanged): 1. **The CUDA member's tensor-core kernel**: every product on the tensor cores. 2. **The Metal member** (`quantum_fabric_metal`, macOS), with an exact integer kernel and a second kernel whose every product uses the GPU's SIMD-group matrix multiply. Both matrix paths use the same argument (below), and each kernel has its own implementation id and mutants. # 1. The CUDA tensor-core kernel `tensor_limbs.cu`: a second CUDA kernel for the gate, whose every product is formed on the tensor cores. It runs on devices of compute capability 8.0 or later. ## How it stays exact - Each 32-bit value is four 8-bit limbs in bf16: the lower three unsigned, the top one signed. bf16 holds every one exactly. - One warp multiplies sixteen pairs' limbs (16 × 16) by the gate's limb weights (16 × 112: sixteen products per pair, seven coefficients each) in fourteen `mma.sync` `m16n8k16` steps, accumulated in f32. - Every limb product is at most `2^16`, and every coefficient sums at most four of them. So every value the tensor core forms is an integer below `2^18`, which f32 holds exactly in any order of addition. - The kernel adds the coefficients, shifted by `8s` bits, in 64-bit integers. That gives the exact product; the rest of the gate is the other kernel's arithmetic. ## API (additive, 0.6.1) - `cuda::TENSOR_SOURCE`, `TENSOR_IMPLEMENTATION`, `tensor_implementation_id{,_of}`, `TENSOR_MUTANTS`, `tensor_mutate`, `TENSOR_MIN_COMPUTE`, `TENSOR_THREADS_PER_ITEM`, `DeviceIdentity::has_tensor_cores`. - `CudaMember::tensor()`: the same device, member id unchanged, under the kernel's own implementation id (`cuda-tensor-limbs8/…`). It is probed before it is returned. - `CudaError::Unsupported` (new) for a device below 8.0. - The compiled kernel carries its launch spread (two threads per pair), so the probe and gate launches cover every item. ## Mutants (7) | | Mutant | Caught by | |---|---|---| | T1 | rounding constant `2^28` | probe | | T2 | shift by 29 | probe | | T3 | nine-bit limbs, which bf16 does not hold | probe | | T4 | the B fragment's second half loaded with its rows swapped | probe | | T5 | control test inverted | certificate | | T6 | the sign of row 0's last term | certificate | | T7 | range guard never fires | range | ## Measured On jetson-hub's discrete GPU (sm_89, silicon), `tests/cuda_tensor.rs` passed 9/9: - 66020 probe pairs agree. - All four pinned certificates reproduce (`d18454c6…`, `41c726a0…`, `3dbde276…`, `9a84b55a…`). - Circuits up to 22 qubits match native. - 7 of 7 mutants are killed. On the 484 edge pairs, T1 differs on 80, T2 on 436, T3 on 297 and T4 on 210. - `cuda_member` (9/9, 6 of 6 mutants) and `cuda_sv` (3/3, pin `b7d129d0…`) pass unchanged on the same tree. On the build host, which has no CUDA library, `cuda_tensor` and `cuda_member` each pass 9/9 over no device (`none`). I checked the workflow's report checks against the real report: the device set, the member line, 4 certificates and 7 of 7 mutants all match. ## CI - `wai-quantum-cuda`, cuda-hw job: a new tensor-core step that requires 9 passed, the member line, 4 `cuda-tensor-limbs8` certificates and 7 of 7 mutants. - cuda-absent job: runs the new tests over no device. # 2. The Metal member `quantum_fabric_metal` (opt-in, macOS) runs the gate on the platform's GPU through its own interface. The operating system's compiler builds the kernels at run time. The state lives in one buffer the GPU and the host share, and each gate is one serial dispatch in a single command buffer. - **`lanes64.metal`** (`metal-lanes64/…`): the CUDA kernel's exact 64-bit integer arithmetic in the Metal shading language. - **`matrix_limbs.metal`** (`metal-matrix-limbs8/…`, `MetalMember::matrix()`): every product formed by the SIMD-group matrix multiply, from 8-bit limbs in bf16 accumulated in f32. - One SIMD group multiplies eight pairs' limbs (8 × 16) by the gate's limb weights (16 × 112) in fourteen 8 × 8 column blocks. - Each of a pair's four threads reads one amplitude component, joins its four products, and writes the row that replaces that component. - A device below GPU family 1007 gets `MetalError::Unsupported`. - **Keys** are the GPU family number (`gpu-msl-<family>`), not a vendor name. The member id binds the device, its registry id, family and memory, and the operating system's version. - **Off macOS** the member still builds (its platform types have no values there), finds nothing, and says why. ## Mutants - **Integer kernel (G1–G6):** the CUDA kernel's six, in Metal text. - **Matrix kernel (X1–X7):** rounding, shift, nine-bit limbs, the second limb block multiplying the first block's weights, control inverted, a real row's sign, and the range guard. ## Measured On the M3 Ultra build host, the same machine as the `macos` runner (GPU family 1009, silicon), `metal_member` and `metal_matrix` each passed 9/9: - 66020 probe pairs agree. - All four pinned certificates reproduce. - Circuits up to 22 qubits match native. - 6 of 6 and 7 of 7 mutants are killed. On the 484 edge pairs, X3 (nine-bit limbs) breaks 297, the same count as on the CUDA tensor cores, which shows the multiply really takes bf16 inputs. X4 breaks 443. On Linux (jetson-hub), the unit tests pass, both suites pass 9/9 over no device, and clippy is clean. Clippy is also clean on macOS, and I checked the workflow's report greps against the real report. **Not done: Metal 4's tensor-op multiply,** the path to M5's GPU neural accelerators. The only M5 in the fleet is the dev Mac, where builds and tests are off-limits. ## CI `wai-quantum-metal`: - metal-hw, on the `macos` runner: unit tests; 9 passed, 4 certificates and 6 of 6 mutants for the integer kernel; then the member line, 4 matrix certificates and 7 of 7 mutants for the matrix kernel. - metal-absent, on `linux`: both suites over no device, plus the off-platform discovery test. The CHANGELOG also records two hand measurements made outside CI: the tensor program on a cloud tensor processor, and the QPU crate against its service's simulator. All repo checks are clean: prose, vanity metrics, tests-in-CI, concurrency, lockfiles, orphans and public names. 🤖 Generated with [Claude Code](https://claude.com/claude-code)

dcharlot added 1 commit 2026-10-16 04:21:31 -04:00

wai-quantum 0.6.1: the CUDA member's tensor-core kernel ... 246e953e20

A second CUDA kernel for the same gate (tensor_limbs.cu), for devices of
compute capability 8.0 or later, whose every product is formed on the
tensor cores: each 32-bit value as four 8-bit limbs in bf16, one warp
multiplying sixteen pairs' limbs by the gate's limb weights in fourteen
m16n8k16 steps, accumulated in f32. Every coefficient is an integer below
2^18, so f32 holds every partial sum exactly in any order; the kernel adds
the coefficients, shifted into place, in 64-bit integers and does the rest
of the gate as lanes64.cu does.

CudaMember::tensor gives the same device under the kernel's own
implementation id (cuda-tensor-limbs8/...), probed before it is returned;
a device below 8.0 is CudaError::Unsupported. Seven mutants, one of them
nine-bit limbs, which bf16 does not hold.

Measured on the discrete GPU (sm_89): tests/cuda_tensor.rs 9/9, 66020
probe pairs, all four pinned certificates, circuits to 22 qubits, 7 of 7
mutants; cuda_member 9/9 and cuda_sv 3/3 unchanged. With no CUDA, both
suites pass over no device. CI: a step of wai-quantum-cuda and the absent
leg.

Co-Authored-By: Claude Opus 5.5 <[email protected]>

dcharlot added 1 commit 2026-10-16 04:59:07 -04:00

wai-quantum: the Metal member, with a SIMD-group matrix kernel ... 9af060620d

quantum_fabric_metal (opt-in, macOS): the gate on the platform's GPU
through its own interface, kernels compiled at run time by the operating
system's compiler, the state in a buffer the GPU and the host share, one
serial dispatch per gate in one command buffer.

- lanes64.metal (metal-lanes64/...): the CUDA kernel's exact 64-bit
  integer arithmetic in the Metal shading language.
- matrix_limbs.metal (metal-matrix-limbs8/..., MetalMember::matrix):
  every product formed by the SIMD-group matrix multiply from 8-bit limbs
  in bf16 accumulated in f32; one SIMD group takes eight pairs in fourteen
  8x8 column blocks, each thread joins four products and forms one row.
  GPU family 1007 or later, else MetalError::Unsupported.

Keyed by GPU family number (gpu-msl-<family>). Off macOS the member
builds, finds nothing and says why (uninhabited platform types).

Measured on the M3 Ultra runner's GPU (family 1009): metal_member and
metal_matrix 9/9 each, 66020 probe pairs, all four pinned certificates,
22 qubits, 6 of 6 and 7 of 7 mutants (nine-bit limbs break 297 of 484
edge pairs). On Linux both suites pass over no device. CI:
wai-quantum-metal (macos and linux legs).

Co-Authored-By: Claude Opus 5.5 <[email protected]>

dcharlot changed title from wai-quantum 0.6.1: the CUDA member's tensor-core kernel to wai-quantum 0.6.1: the CUDA tensor-core kernel and the Metal member 2026-10-16 04:59:42 -04:00

dcharlot added 1 commit 2026-10-16 09:44:38 -04:00

Merge remote-tracking branch 'origin/main' into wai/quantum-metal ... 6b9b7b98d1

# Conflicts:
#	wai/CHANGELOG.md

dcharlot merged commit 29e9b95bb5 into main 2026-10-16 09:45:04 -04:00

dcharlot referenced this pull request from a commit 2026-10-16 09:45:05 -04:00

Merge pull request 'wai-quantum 0.6.1: the CUDA tensor-core kernel and the Metal member' (#275) from wai/quantum-tensor-cores into main

dcharlot referenced this pull request from a commit 2026-10-18 07:19:04 -04:00

Merge remote-tracking branch 'origin/main' into wai/d2c-m7a

Sign in to join this conversation.

Reviewers

No reviewers

Labels

Clear labels

No items

No labels

Milestone

Clear milestone

No items

No milestone

Projects

Clear projects

No items

No project

Assignees

Clear assignees

No assignees

1 participant

Notifications

Subscribe

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set

Reference

Transaction-Science/open-standards!275

No description provided.