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

Merged
dcharlot merged 24 commits from wai/quantum-dataflow-vector into main 2026-10-15 20:58:49 -04:00
Owner

wai-quantum 0.6.0: four more members (HIP, npu-cores, tensor-engine, avx512), the dataflow device's vector engine, the gate as a portable tensor program, and CI on a second hardware host.

Members

  • HIP (quantum_fabric_hip, Linux): the CUDA member's kernel text, unchanged, compiled at run time by HIP's runtime compiler for each device's own architecture. Both libraries load when first used. The device's properties are read in the runtime's frozen layout (hipDeviceProp_tR0600); properties that don't read as that layout are refused. With quantum_sv, the f64 state vector runs the CUDA kernels' text too.
  • npu-cores (quantum_fabric_npu_cores, Linux): the gate on one compute core of a neural processor's tile array, programmed directly. The core forms each pair's four rows in exact 64-bit integers; the host checks range and keeps the low 32 bits.
    • The core's program is built at run time with the tile array's open toolchain. The Python it runs under is named by WAI_QUANTUM_NPU_CORES_PYTHON. Programs are cached per variant, generation and toolchain.
    • The host side is compiled against the device runtime and loaded.
    • Each mutant is its own program.
  • tensor-engine (quantum_fabric_tensor_engine): the gate's products as exact limb convolutions on a tensor processor's matrix engine, using bf16 multiplies with fp32 accumulation.
    • Four 8-bit limbs per value; every partial sum is an integer below 2^24, so it's exact in any order on any such engine.
    • The host assembles the products.
    • The target here is the kernel toolchain's simulator (Apache-2.0), and the certificates say so.
    • One mutant is nine-bit limbs, which the engine can't hold.
  • avx512 (quantum_fabric_vector): eight 64-bit lanes, with the unit's own 64-bit arithmetic shift. Discovery opens one member per instruction set the processor reports.
  • Dataflow vector engine (dataflow-1e52-b140-vector): the gate on the simulated device's vector engine, 32 lanes per tile.
    • The tile's instructions are transcribed from the simulator's own table, and CI checks the transcription against it.
    • The listings are reproduced by an independent assembler.

The gate as a portable tensor program

quantum_fabric::stablehlo: StableHLO, with every value a 64-bit integer that never overflows.

  • Two independent compilers run it, both pinned: IREE 3.12.0 and XLA (through JAX 0.11.2).
  • ci/tensor-program/check.sh holds every run to the outcome it must have. A run that ends without stating its outcome fails, so a dead device can't read as a mutant caught.

Measured

All bit-exact, with every mutant killed:

Member Device Result
HIP gfx1151, silicon 9/9 member tests; 3/3 state-vector tests (the pin b7d129d0…)
npu-cores XDNA 2 tile array, silicon 9/9
avx2 and avx512 Zen 5, silicon 9/9 each
WebGPU, unchanged same GPU (gpu-vulkan-1002) bit-exact
OpenCL, unchanged same GPU (opencl-gpu-1002) bit-exact, the state vector included
tensor-engine the kernel toolchain's simulator 9/9
dataflow vector engine simulated b140 9/9

The tensor program is exact under both compilers on:

  • the CPU;
  • every Vulkan device on both GPU hosts;
  • CUDA;
  • HIP;
  • Metal on two GPU generations;
  • the arm64 CPU.

CI

  • A new runner linux-amd (framework-128gb), for hardware jobs only.
  • New legs: HIP; npu-cores; avx2 and avx512 on silicon; WebGPU; OpenCL; the tensor program.
  • The tensor-engine member's simulator runs on the container runner.
  • The new "absent" legs run on linux-arm64.
  • Each new leg was dry-run as the runner's own user before this PR. That dry run caught the tensor program's HIP compile flag.
  • New workflows: wai-quantum-hip, wai-quantum-npu-cores, wai-quantum-tensor-engine, wai-quantum-tensor-program.

Breaking (hence 0.6.0)

  • OpenCL keys are opencl-<kind>-<vendor id>. DeviceIdentity gains vendor_id, and vendor_slug is removed.
    • The hardware host reads the new declaration from WAI_QUANTUM_HW_OPENCL_KEYS, already set beside the old variable.
  • Isa gains Avx512.
  • Chip gains vector.
  • The device identities, Chip and Isa are now #[non_exhaustive], so the next addition isn't a break.
  • wai-quantum-qpu becomes 0.2.0, against 0.6.0.

🤖 Generated with Claude Code

wai-quantum 0.6.0: four more members (HIP, npu-cores, tensor-engine, avx512), the dataflow device's vector engine, the gate as a portable tensor program, and CI on a second hardware host. ## Members - **HIP** (`quantum_fabric_hip`, Linux): the CUDA member's kernel text, unchanged, compiled at run time by HIP's runtime compiler for each device's own architecture. Both libraries load when first used. The device's properties are read in the runtime's frozen layout (`hipDeviceProp_tR0600`); properties that don't read as that layout are refused. With `quantum_sv`, the f64 state vector runs the CUDA kernels' text too. - **npu-cores** (`quantum_fabric_npu_cores`, Linux): the gate on one compute core of a neural processor's tile array, programmed directly. The core forms each pair's four rows in exact 64-bit integers; the host checks range and keeps the low 32 bits. - The core's program is built at run time with the tile array's open toolchain. The Python it runs under is named by `WAI_QUANTUM_NPU_CORES_PYTHON`. Programs are cached per variant, generation and toolchain. - The host side is compiled against the device runtime and loaded. - Each mutant is its own program. - **tensor-engine** (`quantum_fabric_tensor_engine`): the gate's products as exact limb convolutions on a tensor processor's matrix engine, using bf16 multiplies with fp32 accumulation. - Four 8-bit limbs per value; every partial sum is an integer below 2^24, so it's exact in any order on any such engine. - The host assembles the products. - The target here is the kernel toolchain's simulator (Apache-2.0), and the certificates say so. - One mutant is nine-bit limbs, which the engine can't hold. - **avx512** (`quantum_fabric_vector`): eight 64-bit lanes, with the unit's own 64-bit arithmetic shift. Discovery opens one member per instruction set the processor reports. - **Dataflow vector engine** (`dataflow-1e52-b140-vector`): the gate on the simulated device's vector engine, 32 lanes per tile. - The tile's instructions are transcribed from the simulator's own table, and CI checks the transcription against it. - The listings are reproduced by an independent assembler. ## The gate as a portable tensor program `quantum_fabric::stablehlo`: StableHLO, with every value a 64-bit integer that never overflows. - Two independent compilers run it, both pinned: IREE 3.12.0 and XLA (through JAX 0.11.2). - `ci/tensor-program/check.sh` holds every run to the outcome it must have. A run that ends without stating its outcome fails, so a dead device can't read as a mutant caught. ## Measured All bit-exact, with every mutant killed: | Member | Device | Result | |---|---|---| | HIP | gfx1151, silicon | 9/9 member tests; 3/3 state-vector tests (the pin `b7d129d0…`) | | npu-cores | XDNA 2 tile array, silicon | 9/9 | | avx2 and avx512 | Zen 5, silicon | 9/9 each | | WebGPU, unchanged | same GPU (`gpu-vulkan-1002`) | bit-exact | | OpenCL, unchanged | same GPU (`opencl-gpu-1002`) | bit-exact, the state vector included | | tensor-engine | the kernel toolchain's simulator | 9/9 | | dataflow vector engine | simulated b140 | 9/9 | The tensor program is exact under both compilers on: - the CPU; - every Vulkan device on both GPU hosts; - CUDA; - HIP; - Metal on two GPU generations; - the arm64 CPU. ## CI - A new runner `linux-amd` (framework-128gb), for hardware jobs only. - New legs: HIP; npu-cores; avx2 and avx512 on silicon; WebGPU; OpenCL; the tensor program. - The tensor-engine member's simulator runs on the container runner. - The new "absent" legs run on linux-arm64. - Each new leg was dry-run as the runner's own user before this PR. That dry run caught the tensor program's HIP compile flag. - New workflows: `wai-quantum-hip`, `wai-quantum-npu-cores`, `wai-quantum-tensor-engine`, `wai-quantum-tensor-program`. ## Breaking (hence 0.6.0) - OpenCL keys are `opencl-<kind>-<vendor id>`. `DeviceIdentity` gains `vendor_id`, and `vendor_slug` is removed. - The hardware host reads the new declaration from `WAI_QUANTUM_HW_OPENCL_KEYS`, already set beside the old variable. - `Isa` gains `Avx512`. - `Chip` gains `vector`. - The device identities, `Chip` and `Isa` are now `#[non_exhaustive]`, so the next addition isn't a break. - `wai-quantum-qpu` becomes 0.2.0, against 0.6.0. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
On the device generation whose vector engine multiplies 23-bit numbers
(PCI 1e52:b140), a second dataflow member, dataflow-1e52-b140-vector,
runs the gate's arithmetic on that engine, 32 lanes at a time, on 32
compute tiles.

- Each tile's first core writes the unpacker's and the packer's
  configuration, then hands the tile its instructions inline (a word with
  low bits other than 0b11 in its own instruction stream): the unpacker
  copies the pairs' columns into the destination registers as 32-bit
  integers, the vector engine forms each pair's four rows, the packer
  copies them back, and the core signals when the packer is done.
- Exact arithmetic: each operand split into 23 and 9 bits; four partial
  products give the unsigned high word, corrected for sign to P_hi, and
  the bits of the low word the rounding reaches; fx = P_hi * 2^(32-s) + t,
  held as 30 low bits and the rest; each row adds or subtracts its terms
  with a carry after each. The three mutants are the other members'.
- quantum_fabric::tensix: the instructions used, transcribed from the
  device simulator's own table; CI checks the transcription against it.
  quantum_fabric::vector_engine builds the programs (rows and probe); the
  assembler check covers them (14 programs, identical words).
- A pass is sized to its share of the pairs in steps of 64 lanes: one
  iteration covers the even or the odd columns of four destination rows,
  so two cover the next 64 pairs in order.

Measured on the simulator, built from source: 66020 probe pairs agree,
2000 kernel cases match run_words, all four pinned certificates, circuits
up to 20 qubits match native, 3 of 3 mutants killed; the instruction table
matches for all 20 instructions; the other generation's cores member is
unchanged (9 of 9).

Co-Authored-By: Claude Opus 5.5 <[email protected]>
quantum_fabric::stablehlo writes the gate's rows and the arithmetic
probe as a StableHLO module: every value a 64-bit integer that never
overflows (|e·a| < 2^62; a row of four fx below 2^35), so any tensor
compiler must give the exact rows. The stablehlo_gate example writes the
module, its three mutants, a batch of 4096 pairs reaching the 32-bit
edges, and the exact outputs (raw little-endian i64, compared under the
module's own signless type).

The wai-quantum-tensor-program workflow builds it with a pinned
open-source tensor compiler (IREE 3.12.0) and has its runtime compare:
the CPU build and the three mutants in a container, plus builds for
three GPU generations the fleet lacks (compiled only); every Vulkan
device of the hardware host and CUDA; Metal on the macOS runner. Each
GPU leg also runs one mutant, so its comparison is shown to catch a
difference.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
XLA (through JAX 0.11.2, pinned) now builds and runs the gate's tensor
program beside IREE: on the x86-64 CPU (the program and all three
mutants), on the discrete GPU through CUDA, and on the macOS runner's
arm64 CPU. ci/tensor-program/xla_run.py renames the chosen function to
main, runs it and compares every output with the exact one.

ci/tensor-program/check.sh runs one build on one device and holds it to
the outcome it must have (rows, then probe: exact or differs). A run
counts only when it states its outcome: the runtime's SUCCESS or FAILED
comparison line, or the driver's exact or differs line. Any other ending
(a device that will not open, an exception) fails the check, so it can
never read as a mutant caught. The three workflow legs call it instead
of repeating the argument lists.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
quantum_fabric_hip (opt-in, Linux): the gate kernel the CUDA member runs,
unchanged, compiled at run time by HIP's runtime compiler for each
device's own architecture (--offload-arch from the device), with
-ffp-contract=off. Both libraries are loaded when first used, so a build
with the feature runs anywhere and finds no member where there is no
HIP. The device's properties are read in the runtime's frozen layout
(hipDeviceProp_tR0600); properties that do not read as that layout are
refused, never guessed at.

quantum_fabric::hip holds what names a member (the kernel text, the
implementation id hip-lanes64/..., distinct from the CUDA one over the
same text, the device identity and the key gpu-hip-<arch>); the mutants
are the CUDA kernel's, the same edits of the same text. With quantum_sv
the dense f64 state vector runs the CUDA kernels' text too
(hip-sv64/..., contraction off).

Measured on gfx1151: hip_member 9/9 (66020 probe pairs, all four pinned
certificates, circuits to 22 qubits, 6 of 6 mutants killed) and hip_sv
3/3 (the state vector's pin b7d129d0..., 16 and 20 qubits bit for bit,
4 of 4 mutants killed).

Co-Authored-By: Claude Opus 5.5 <[email protected]>
Co-Authored-By: Claude Opus 5.5 <[email protected]>
A key was the vendor's first word, so the GPU vendor whose name begins
"Advanced" was keyed opencl-gpu-advanced. A key is now
opencl-<kind>-<vendor id, hex> from CL_DEVICE_VENDOR_ID, the PCI vendor
id for a device on a PCI bus, as the WebGPU member's keys already are;
DeviceIdentity gains vendor_id, which the member id binds, and
vendor_slug goes. Measured: opencl-cpu-8086, opencl-gpu-10de and
opencl-gpu-8086 on the hardware host, opencl-gpu-1027f00 on the macOS
runner, opencl-gpu-1002 on the AMD host.

The workflow reads the hardware host's declaration from
WAI_QUANTUM_HW_OPENCL_KEYS (set beside the old variable, which main's
workflow still reads), adds the AMD host's leg (opencl-amd: 9/9 and the
state vector, bit-exact there by hand) and names the macOS job by its
runner.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
The WebGPU member on its integrated GPU and the software Vulkan driver
(gpu-vulkan-1002, swgpu-vulkan-10005) with the conformance corpus through
it; the tensor program through HIP (built for gfx1151) and on both Vulkan
devices, the third mutant caught on each. Job names carry the host's name.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
quantum_fabric_vector gains avx512 (x86-64, the foundation instructions):
512-bit vectors, eight 64-bit lanes, each fx an exact 32x32 multiply
(vpmuldq), an add and the unit's own 64-bit arithmetic shift (vpsraq),
where avx2 needs a logical shift with the sign flipped back. Discovery
now opens one member per instruction set the processor reports
(native_isas), so a processor with both units has both members.

Measured on framework-128gb's processor (silicon): avx2 and avx512 each
9/9 (66020 probe pairs, all four pinned certificates, circuits to 20
qubits, 3 of 3 mutants killed). The workflow's new leg there holds each
unit to its vector length, silicon provenance, four certificates and
three mutants.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
quantum_fabric_npu_cores (opt-in, Linux): the gate's fixed-point
arithmetic on one compute core of a neural processor's tile array,
programmed directly rather than through a model runtime. The kernel
(npu_cores.cc) is the same 64-bit transcription of lanes::exact_row as
the other members'; the core forms each gathered pair's four rows, the
host checks their range and keeps the low 32 bits. A line carries a
16-word header (the mode and the gate's words), so one program serves
every gate and the arithmetic probe; each kernel variant (the conforming
one and the three mutants) is its own program.

At run time the member finds the device in the accelerator class by its
driver, builds the program with the tile array's open toolchain
(npu_cores_design.py, run by the Python WAI_QUANTUM_NPU_CORES_PYTHON
names; cached per variant, generation and toolchain), compiles the host
side (npu_cores_host.cpp) with the system C++ compiler against the
device runtime and loads it. The key is npu-cores-<pci vendor>-<device>;
the member id binds the PCI ids, revision, runtime name, firmware and
toolchain versions.

Measured on the XDNA 2 processor (silicon), npu_cores_member 9/9 in
100 s: 66020 probe pairs, all four pinned certificates, circuits to 20
qubits, 3 of 3 mutants killed.

Also: the tensor program's HIP build is --iree-hal-target-backends=rocm
--iree-rocm-target (the compiler has no --iree-hip-target), found by the
leg's dry run on the CI user.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
Co-Authored-By: Claude Opus 5.5 <[email protected]>
Co-Authored-By: Claude Opus 5.5 <[email protected]>
The branch carries three API breaks since 0.5.1 (OpenCL keys and
DeviceIdentity::vendor_id, Isa::Avx512, Chip::vector), so it releases as
0.6.0. The device identities (cuda, hip, opencl, npu_cores), the
dataflow Chip and the vector Isa are #[non_exhaustive] now, so adding a
field or an instruction set is not a break again. wai-quantum-qpu
becomes 0.2.0 against wai-quantum 0.6.0, and wai-rs requires 0.6.0.
README rows for the HIP, npu-cores and avx512 members and the OpenCL
key; the CHANGELOG entry.

Verified on the build host: wai-quantum --features full --lib 597
passed; every member feature checks with its tests and examples;
wai-quantum-qpu's tests pass; wai-rs builds with its quantum features.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
# Conflicts:
#	wai/CHANGELOG.md
quantum_fabric_tensor_engine (opt-in): the gate's sixteen products per
pair as exact limb convolutions in one bf16 matrix multiply per tile of
512 pairs, accumulated in fp32, on a tensor processor's matrix engine.
Each 32-bit value is four 8-bit limbs (bf16 holds each exactly), every
limb product is at most 2^16 and every coefficient a sum of at most four,
so fp32 holds every partial sum exactly in any order: the argument holds
for any engine that multiplies bf16 and accumulates in fp32. The host
checks every coefficient is an integer below 2^24, assembles the products
in 64 bits and applies the fixed-point arithmetic, as the neural-processor
member's host does. The probe uses its own weight layout, sixteen
independent pairs a multiply.

The kernel (tensor_engine_helper.py) runs as a long-lived helper under the
Python WAI_QUANTUM_TENSOR_ENGINE_PYTHON names; today its target is the
kernel toolchain's simulator (Apache-2.0), so every certificate's
provenance is simulator:<versions>. Four mutants: the three host
constants, and Radix9, nine-bit limbs the engine's bf16 inputs cannot
hold, which the engine gets wrong.

Measured on the simulator: tensor_engine_member 9/9 in 340 s (66020
probe pairs, all four pinned certificates, circuits to 20 qubits, 4 of 4
mutants killed). Workflow: the simulator on the container runner, and
the absent case on linux-arm64, where the HIP and npu-cores absent legs
move too.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
Co-Authored-By: Claude Opus 5.5 <[email protected]>
stablehlo::limb_module: the tensor-engine member's limb layout as one
StableHLO dot_general of bf16 operands, accumulated and returned in f32,
the program a compiler maps to a matrix unit. stablehlo_gate writes it
with the batch's limb weights and limbs and the exact coefficients, and a
nine-bit-limb data set whose exact coefficients bf16 cannot reproduce (a
run that matches them is not comparing). check.sh holds the limb module
to "=!" (exact on eight-bit limbs, different on nine); both drivers read
typed raw files, and floats are compared exactly (IREE's
--expected_f32_threshold=0).

Measured: exact under both compilers on the x86-64 and arm64 CPUs and on
the discrete GPU through CUDA; the nine-bit limbs differ everywhere. Not
exact: the IREE 3.12 HIP build for gfx1151, whose bf16 multiply returns
non-integers (-14.0002 for -14), so the workflow shows that step and
does not require it. Not built: the IREE Vulkan and Metal builds do not
lower bf16 in this version.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
Co-Authored-By: Claude Opus 5.5 <[email protected]>
Merge remote-tracking branch 'origin/main' into wai/quantum-dataflow-vector
Some checks failed
wai-quantum tensor-engine member / member · tensor engine absent (linux-arm64) (pull_request) Has been cancelled
wai-quantum portable tensor program / does this change touch it (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · CPU, mutants, GPU builds (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · GPUs (jetson-hub) (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · integrated GPU (framework-128gb) (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · Metal and the arm64 CPU (macos) (pull_request) Has been cancelled
wai-quantum CPU vector units / does this change touch it (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, rvv 1024-bit (riscv64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 128-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, rvv 128-bit (riscv64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 2048-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, avx2 (x86_64 silicon) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 512-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector units, avx2 and avx512 (framework-128gb silicon) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit absent (macos) (pull_request) Has been cancelled
wai-quantum WebGPU member / does this change touch it (pull_request) Has been cancelled
wai-quantum WebGPU member / shader checks · browser build · wasip2 inertness · MSRV (pull_request) Has been cancelled
wai-quantum WebGPU member / member · software Vulkan (linux) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · Metal (macos) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · discrete GPU, Vulkan (jetson-hub) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · integrated GPU, Vulkan (framework-128gb) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · browser WebGPU (macos) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · DX12 and Vulkan (cbio-windows) (pull_request) Has been cancelled
WAI mbt2018_mean integer decode (regenerated vectors) / does this change touch it (pull_request) Has been cancelled
WAI mbt2018_mean integer decode (regenerated vectors) / regenerate mbt2018_mean integer vectors → Rust decode, required (pull_request) Has been cancelled
wai-web reach builds / wai-web reach · does this change touch it (pull_request) Has been cancelled
wai-web reach builds / wai-web reach builds · the package of record, every output as recorded (pull_request) Has been cancelled
wai-web reach builds / wai-web reach builds · this tree's wai-web, every rung on its engine line (pull_request) Has been cancelled
wai-web wasm32 / does this change touch it (pull_request) Has been cancelled
wai-web wasm32 / wai-web wasm32 build · corpus, menu and writers (pull_request) Has been cancelled
bebb6a46cf
# Conflicts:
#	wai/CHANGELOG.md
wai-quantum: the tensor-engine helper runs on a named device platform too
All checks were successful
WAI conformance / codec wrappers · fuzz smoke, text decode limits, C entry points, provenance, the semantic path (pull_request) Successful in 49s
WAI energy receipts / energy receipts · macos (pull_request) Successful in 1m53s
wai-pack / wai-pack · macos (pull_request) Successful in 1m17s
wai-quantum dataflow member / member · dataflow device absent (macos) (pull_request) Successful in 14s
wai-web wasm32 / wai-web wasm32 build · corpus, menu and writers (pull_request) Successful in 4m46s
WAI mbt2018_mean integer decode (regenerated vectors) / regenerate mbt2018_mean integer vectors → Rust decode, required (pull_request) Successful in 5m13s
wai-quantum FPGA member / member · FPGA datapath, simulated (macos) (pull_request) Successful in 1m18s
wai-quantum neural-processor member / member · neural processor (macos) (pull_request) Successful in 6m4s
wai-quantum OpenCL member / member · OpenCL GPU (macos) (pull_request) Successful in 50s
wai-quantum matrix-extension member / member · matrix extension absent (macos) (pull_request) Successful in 9s
wai-quantum portable tensor program / tensor program · Metal and the arm64 CPU (macos) (pull_request) Successful in 15s
wai-quantum CPU vector units / member · vector unit absent (macos) (pull_request) Successful in 13s
WAI HTTP delivery / HTTP corpus · falsify mutations 0/3 (pull_request) Successful in 59m57s
WAI HTTP delivery / HTTP corpus · falsify mutations 1/3 (pull_request) Successful in 57m40s
wai-quantum WebGPU member / member · Metal (macos) (pull_request) Successful in 5m46s
WAI HTTP delivery / HTTP corpus · falsify mutations 2/3 (pull_request) Successful in 57m5s
wai-quantum WebGPU member / member · browser WebGPU (macos) (pull_request) Successful in 2m6s
WAI HTTP delivery / HTTP corpus · falsify patches 0/3 (pull_request) Successful in 6m1s
WAI HTTP delivery / HTTP corpus · every patch and mutation in exactly one shard (pull_request) Successful in 4s
WAI HTTP delivery / HTTP corpus · verified both ways, and the differential (pull_request) Successful in 2m53s
WAI HTTP delivery / HTTP corpus · falsify patches 1/3 (pull_request) Successful in 7m21s
WAI HTTP delivery / cache rules · controls and labelled caches (pull_request) Successful in 1m13s
WAI HTTP delivery / HTTP corpus · falsify patches 2/3 (pull_request) Successful in 6m52s
wai-pack / wai-pack · linux-arm64 (pull_request) Successful in 1m21s
wai-pack / wai-pack · the floor through a media prober (pull_request) Successful in 16s
wai-quantum HIP member / member · HIP absent (linux-arm64) (pull_request) Successful in 21s
wai-quantum npu-cores member / member · npu-cores absent (linux-arm64) (pull_request) Successful in 21s
wai-quantum tensor-engine member / member · tensor engine absent (linux-arm64) (pull_request) Successful in 17s
wai-pack / wai-pack · dependency licences (pull_request) Successful in 1m13s
wai-pack / wai-pack · independent verifier (Python) (pull_request) Successful in 4m21s
cec78c1f37
WAI_QUANTUM_TENSOR_ENGINE_TARGET names a device platform (trn1, trn2, ...):
the helper then runs the same kernel compiled for that device (nki.jit on
NumPy inputs), reports that target, and the member's key is
tensor-engine-<target> with silicon provenance. Unset, the target is the
toolchain's simulator, as before.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
ci: the dataflow workflow checks the instruction table on the vector generation only, and keeps it in the cache
Some checks failed
wai-quantum tensor-engine member / member · tensor engine absent (linux-arm64) (pull_request) Has been cancelled
wai-quantum portable tensor program / does this change touch it (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · CPU, mutants, GPU builds (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · GPUs (jetson-hub) (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · integrated GPU (framework-128gb) (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · Metal and the arm64 CPU (macos) (pull_request) Has been cancelled
wai-quantum CPU vector units / does this change touch it (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, rvv 1024-bit (riscv64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 128-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, rvv 128-bit (riscv64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 2048-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, avx2 (x86_64 silicon) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 512-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector units, avx2 and avx512 (framework-128gb silicon) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit absent (macos) (pull_request) Has been cancelled
wai-quantum WebGPU member / does this change touch it (pull_request) Has been cancelled
wai-quantum WebGPU member / shader checks · browser build · wasip2 inertness · MSRV (pull_request) Has been cancelled
wai-quantum WebGPU member / member · software Vulkan (linux) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · Metal (macos) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · discrete GPU, Vulkan (jetson-hub) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · integrated GPU, Vulkan (framework-128gb) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · browser WebGPU (macos) (pull_request) Has been cancelled
WAI mbt2018_mean integer decode (regenerated vectors) / does this change touch it (pull_request) Has been cancelled
WAI mbt2018_mean integer decode (regenerated vectors) / regenerate mbt2018_mean integer vectors → Rust decode, required (pull_request) Has been cancelled
wai-web reach builds / wai-web reach · does this change touch it (pull_request) Has been cancelled
wai-web reach builds / wai-web reach builds · the package of record, every output as recorded (pull_request) Has been cancelled
wai-web reach builds / wai-web reach builds · this tree's wai-web, every rung on its engine line (pull_request) Has been cancelled
wai-web wasm32 / does this change touch it (pull_request) Has been cancelled
wai-web wasm32 / wai-web wasm32 build · corpus, menu and writers (pull_request) Has been cancelled
wai-quantum WebGPU member / member · DX12 and Vulkan (cbio-windows) (pull_request) Has been cancelled
05b8cca575
The table test failed on both legs of #266: the cached simulator had been
built before the table was copied beside it, so the table was missing.
A cache entry without its table is rebuilt now. And the encoder is a
transcription of the b140 generation's table, which the 401e table does
not match (its PACR has other fields; measured), so only the b140 leg
names the table and requires its check.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
ci: the dataflow workflow checks each declared member, not the list as one key
Some checks failed
wai-quantum tensor-engine member / member · tensor engine absent (linux-arm64) (pull_request) Has been cancelled
wai-quantum portable tensor program / does this change touch it (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · CPU, mutants, GPU builds (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · GPUs (jetson-hub) (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · integrated GPU (framework-128gb) (pull_request) Has been cancelled
wai-quantum portable tensor program / tensor program · Metal and the arm64 CPU (macos) (pull_request) Has been cancelled
wai-quantum CPU vector units / does this change touch it (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, rvv 1024-bit (riscv64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 128-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, rvv 128-bit (riscv64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 2048-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, avx2 (x86_64 silicon) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit, sve 512-bit (aarch64, emulated) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector units, avx2 and avx512 (framework-128gb silicon) (pull_request) Has been cancelled
wai-quantum CPU vector units / member · vector unit absent (macos) (pull_request) Has been cancelled
wai-quantum WebGPU member / does this change touch it (pull_request) Has been cancelled
wai-quantum WebGPU member / shader checks · browser build · wasip2 inertness · MSRV (pull_request) Has been cancelled
wai-quantum WebGPU member / member · software Vulkan (linux) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · Metal (macos) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · discrete GPU, Vulkan (jetson-hub) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · integrated GPU, Vulkan (framework-128gb) (pull_request) Has been cancelled
wai-quantum WebGPU member / member · browser WebGPU (macos) (pull_request) Has been cancelled
WAI mbt2018_mean integer decode (regenerated vectors) / does this change touch it (pull_request) Has been cancelled
WAI mbt2018_mean integer decode (regenerated vectors) / regenerate mbt2018_mean integer vectors → Rust decode, required (pull_request) Has been cancelled
wai-web reach builds / wai-web reach · does this change touch it (pull_request) Has been cancelled
wai-web reach builds / wai-web reach builds · the package of record, every output as recorded (pull_request) Has been cancelled
wai-web reach builds / wai-web reach builds · this tree's wai-web, every rung on its engine line (pull_request) Has been cancelled
wai-web wasm32 / does this change touch it (pull_request) Has been cancelled
wai-web wasm32 / wai-web wasm32 build · corpus, menu and writers (pull_request) Has been cancelled
wai-quantum WebGPU member / member · DX12 and Vulkan (cbio-windows) (pull_request) Has been cancelled
9225edce71
The b140 leg declares two members (the cores and the vector engine), and
the checks used the comma-joined list as one key, so no report line
matched. They loop over the declared keys now. Found by running the leg
by hand in its CI image; with this change all three members pass there
(the simulator named, four certificates, three of three mutants each).

Co-Authored-By: Claude Opus 5.5 <[email protected]>
Merge remote-tracking branch 'origin/main' into wai/quantum-dataflow-vector
All checks were successful
wai-pack / wai-pack · the floor through a media prober (pull_request) Successful in 41s
wai-pack / wai-pack · linux (pull_request) Successful in 1m23s
wai-pack / wai-pack · linux-arm64 (pull_request) Successful in 1m42s
wai-quantum HIP member / member · HIP absent (linux-arm64) (pull_request) Successful in 24s
wai-pack / wai-pack · dependency licences (pull_request) Successful in 1m31s
wai-quantum npu-cores member / member · npu-cores absent (linux-arm64) (pull_request) Successful in 28s
wai-quantum tensor-engine member / member · tensor engine absent (linux-arm64) (pull_request) Successful in 37s
wai-pack / wai-pack · independent verifier (Python) (pull_request) Successful in 4m58s
wai-quantum dataflow member / member · dataflow device, simulated (pci 1e52:b140) (pull_request) Successful in 3m52s
wai-quantum dataflow member / member · dataflow device, simulated (pci 1e52:401e) (pull_request) Successful in 1m22s
wai-quantum FPGA member / member · FPGA simulation absent (linux) (pull_request) Successful in 1m10s
wai-quantum neural-processor member / member · neural processor absent (linux) (pull_request) Successful in 21s
wai-quantum OpenCL member / member · OpenCL absent (linux) (pull_request) Successful in 23s
wai-quantum matrix-extension member / member · matrix extension, emulated (aarch64 Linux) (pull_request) Successful in 17s
wai-quantum-qpu / clients against their contracts (linux) (pull_request) Successful in 23s
wai-quantum signal processor / processor and member · signal processor, emulated (pull_request) Successful in 13m16s
wai-quantum portable tensor program / tensor program · CPU, mutants, GPU builds (pull_request) Successful in 27s
wai-quantum CPU vector units / member · vector unit, rvv 1024-bit (riscv64, emulated) (pull_request) Successful in 17s
wai-quantum CPU vector units / member · vector unit, sve 128-bit (aarch64, emulated) (pull_request) Successful in 26s
wai-quantum CPU vector units / member · vector unit, rvv 128-bit (riscv64, emulated) (pull_request) Successful in 13s
wai-quantum CPU vector units / member · vector unit, sve 2048-bit (aarch64, emulated) (pull_request) Successful in 10s
wai-quantum CPU vector units / member · vector unit, avx2 (x86_64 silicon) (pull_request) Successful in 13s
wai-quantum CPU vector units / member · vector unit, sve 512-bit (aarch64, emulated) (pull_request) Successful in 12s
wai-quantum tensor-engine member / member · tensor engine, the kernel toolchain's simulator (linux) (pull_request) Successful in 7m59s
wai-quantum WebGPU member / member · software Vulkan (linux) (pull_request) Successful in 2m8s
wai-quantum WebGPU member / shader checks · browser build · wasip2 inertness · MSRV (pull_request) Successful in 3m20s
wai-web wasm32 / wai-web wasm32 build · corpus, menu and writers (pull_request) Successful in 3m19s
WAI mbt2018_mean integer decode (regenerated vectors) / regenerate mbt2018_mean integer vectors → Rust decode, required (pull_request) Successful in 6m7s
wai-web reach builds / wai-web reach builds · the package of record, every output as recorded (pull_request) Successful in 2m54s
wai-web reach builds / wai-web reach builds · this tree's wai-web, every rung on its engine line (pull_request) Successful in 5m3s
d1f1bd6ed7
# Conflicts:
#	wai/CHANGELOG.md
Sign in to join this conversation.
No reviewers
No labels
No milestone
No project
No assignees
1 participant
Notifications
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!266
No description provided.