### wai-quantum 0.6.2: the HIP member's matrix-instruction kernel, and the Metal member's tensor-operation kernel

Additive. Every pinned digest is unchanged.

- **Every product by the shading language's tensor operation**
  (`quantum_fabric::metal::TENSOR_OPS_SOURCE`, `tensor_ops.metal`; `MetalMember::tensor_ops`):
  a third Metal kernel for the gate, through the language's version 4.0 tensor operation (a
  bf16 matrix multiply accumulated in f32), which the platform maps to the GPU's matrix
  hardware where it has it. The operation reads and writes device memory, so each threadgroup
  owns a slice of scratch (`TENSOR_OPS_SCRATCH_BYTES`): a 64-pair tile is one multiply of
  64 × 32 limbs by 32 × 128 weights, in four 32-column runs. Its operands are offset-binary, as
  for the HIP kernel below, so the result does not depend on how an accumulator treats a
  negative term. A device or operating system without the operation
  (`DeviceIdentity::has_tensor_ops`: GPU family 1007 or later, operating system 26 or later)
  is `MetalError::Unsupported`. Eight mutants (`TENSOR_OPS_MUTANTS`), among them nine-bit
  limbs, every run reading the first columns of the weights, and a wrong offset correction.
- **Measured** on two GPU generations, silicon: family 1009 (the macOS runner's) and family
  1010, the generation whose GPU cores carry matrix accelerators. On each: 9/9, 66020 probe
  pairs, all four pinned certificates, circuits to 22 qubits, 8 of 8 mutants killed. On family
  1010 the integer and SIMD-group matrix kernels are bit-exact too (9/9 each, 6 of 6 and 7 of 7
  mutants). With signed limbs instead, the tensor operation still agrees on all 66020 pairs on
  both generations: the offset is for the instruction below, not needed here. Where the
  platform runs the operation is its own choice; it is not observed here. CI: a step of
  `wai-quantum-metal` on family 1009; family 1010 was run by hand, the binaries built on the
  build host.

- **Every product by the bf16 matrix instruction** (`quantum_fabric::hip::WMMA_SOURCE`,
  `wmma_limbs.hip`; `HipMember::wmma`): a second HIP kernel for the same gate, for the gfx11
  architectures (`DeviceIdentity::has_wmma`), whose wavefronts of 32 multiply bf16 at
  16 × 16 × 16 and accumulate in f32. One wavefront multiplies sixteen pairs' limbs by the
  gate's limb weights in seven 16-column steps, in the instruction's own operand layout. A
  device without it is `HipError::Unsupported` (new). Eight mutants (`WMMA_MUTANTS`): nine-bit
  limbs, the result's rows in another instruction's layout, and a wrong offset correction
  among them.
- **Offset-binary limbs, because of what the instruction does.** Its f32 accumulation is not
  exact once a negative term enters a sum. Measured on gfx1151, on the instruction alone:
  `1·1 + (−1)·0` returns `1 − 2^-24`, `1 − 1` returns `−2^-24`, `128 − 64` returns
  `64 − 2^-17`; sums of non-negative terms come back exact. The CUDA kernel's signed top limb
  therefore fails the probe on this GPU, off by one in a coefficient. So each value is taken
  as `x + 2^31`, all four limbs unsigned, the instruction forms `(e + 2^31)·(a + 2^31)`
  exactly, and the kernel removes the offset in 64-bit integers:
  `e·a = (e + 2^31)·(a + 2^31) − 2^31·(e + a) − 2^62`, modulo `2^64`. The non-integer
  results the limb program gave through one compiler's HIP build for this GPU (0.6.0) are
  consistent with the same behaviour, though which instruction that build chose was not
  checked.
- **Measured** on the integrated GPU (gfx1151, silicon): 9/9, 66020 probe pairs, all four
  pinned certificates, circuits to 22 qubits, 8 of 8 mutants killed; the other kernel and the
  state vector unchanged there. CI: a step of `wai-quantum-hip`; its absent leg runs the new
  tests over no device.

