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

wai-quantum: energy-landscape jobs on a GPU, through WebGPU, CUDA and Metal, each member certified by the host #672

Merged
dcharlot merged 5 commits from wai/quantum-landscape-gpu into main 2026-11-07 01:28:30 -05:00
Conversation 0 Commits 5 Files changed 31 +47297 -3104
31 changed files with 47297 additions and 3104 deletions
Download patch file Download diff file Expand all files Collapse all files

160
.github/workflows/wai-quantum-landscape-devices.yml vendored Normal file
View file

@ -0,0 +1,160 @@
name: wai-quantum energy-landscape GPU members
# The energy-landscape kernels on GPUs (wai_quantum::quantum_landscape::{webgpu, cuda, metal};
# extensions/quantum-landscape.md §5), each device held to the reference as a host holds a member:
# certified through member::certify on the conformance set, bit-exact for the settle kernel and
# the sampling kernel; every worked example at three temperatures over 50 seeds, and jobs past
# every window and workgroup the devices split them into, giving the reference's bits; each
# fault a port of the kernels can have (a dropped carry, a high-word compare, fields held in 16
# or 32 bits, a draw counter that starts again each window, cells last to first, a settle that
# stops at one workgroup), put into the kernels on the device, costing them their grade; the
# coloured WebGPU kernel earning statistical and no more; a WebGPU wait covering one window, and
# a job a member cannot finish going on to the reference; a workload a device declines placed
# on the reference, the answer naming the refusal. Each host declares the members it must have
# (WAI_LANDSCAPE_EXPECT): one missing, or whose kernels no longer compile, fails its leg. The
# timings that follow are wall-clock on that shared host, with its load printed beside them;
# their bits are asserted and their times are not.
#
# - jetson-hub (linux-hw): its laptop GPU through CUDA and through Vulkan, its integrated GPU and
# the software Vulkan driver through WebGPU. This is the only leg with a CUDA device: while its
# laptop GPU is off its bus, the leg fails, naming the members it is missing, and nothing here
# runs the CUDA kernels. The CUDA kernels' compile test still runs there.
# - framework-128gb (linux-amd): its integrated GPU and the software Vulkan driver through WebGPU.
# - the Mac Studio (macos-m3-ultra): its GPU through WebGPU on Metal, and through Metal itself.
# No leg runs WebGPU on DX12, and no parity is claimed for it.
concurrency:
group: ${{ github.workflow }}-${{ github.ref }}
cancel-in-progress: true
on:
pull_request:
paths:
- 'wai/wai-quantum/src/quantum_landscape/**'
- 'wai/wai-quantum/src/quantum_landscape.rs'
- 'wai/wai-quantum/src/quantum_fabric.rs'
- 'wai/wai-quantum/src/quantum_fabric/**'
- 'wai/wai-quantum/src/quantum_fabric_webgpu.rs'
- 'wai/wai-quantum/src/quantum_fabric_cuda.rs'
- 'wai/wai-quantum/src/quantum_fabric_metal.rs'
- 'wai/wai-quantum/src/quantum_energy.rs'
- 'wai/wai-quantum/src/lib.rs'
- 'wai/wai-quantum/tests/landscape_devices.rs'
- 'wai/wai-quantum/Cargo.toml'
- 'wai/wai-quantum/Cargo.lock'
- 'wai/quantum-conformance/landscape/vectors.json'
- 'rust-toolchain.toml'
- '.github/workflows/wai-quantum-landscape-devices.yml'
schedule:
- cron: '17 8 * * *'
jobs:
relevant:
name: does this change touch it
runs-on: linux
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-landscape-devices.yml \
"${{ github.event.pull_request.base.sha }}" \
"${{ github.event.pull_request.head.sha }}")
echo "this change touches the landscape GPU members: $r"
echo "run=$r" >> "$GITHUB_OUTPUT"
devices:
name: members · ${{ matrix.host }}
needs: relevant
if: needs.relevant.outputs.run == 'true'
runs-on: ${{ matrix.runner }}
timeout-minutes: 120
strategy:
fail-fast: false
matrix:
include:
# Each host's members, declared: a prefix of each member id it must have. A member
# missing, for any reason, fails the leg with why.
- host: jetson-hub
runner: linux-hw
features: quantum_landscape_webgpu,quantum_landscape_cuda
expect: gpu:cuda.nvidia,gpu:wgpu-vulkan.nvidia,gpu:wgpu-vulkan.intel,cpu:wgpu-vulkan.llvmpipe
units: 9
target: hw-target
- host: framework-128gb
runner: linux-amd
features: quantum_landscape_webgpu
expect: gpu:wgpu-vulkan.radeon,cpu:wgpu-vulkan.llvmpipe
units: 8
target: hw-target
- host: mac-studio
runner: macos-m3-ultra
features: quantum_landscape_webgpu,quantum_landscape_metal
expect: gpu:metal.apple-m3-ultra,gpu:wgpu-metal.apple-m3-ultra
units: 8
# The job's own workspace, on the Studio's external volume: never its internal disk.
target: ''
env:
WAI_REQUIRE_GPU: '1'
WAI_LANDSCAPE_EXPECT: ${{ matrix.expect }}
WAI_LANDSCAPE_NVRTC: ${{ contains(matrix.features, 'cuda') && '1' || '' }}
CARGO_PROFILE_DEV_DEBUG: '0'
CARGO_BUILD_JOBS: '8'
steps:
- name: Fetch the source
shell: bash
run: |
set -euo pipefail
want="${{ github.event.pull_request.head.sha || github.sha }}"
rm -rf src && mkdir src && cd src
git init -q
# Fetch the commit itself, not the ref. The host names its own route to the forge.
git fetch -q --depth 1 "${WAI_FORGE_URL:-http://100.64.224.0:3400}/Transaction-Science/open-standards.git" "$want"
git checkout -q FETCH_HEAD
got="$(git rev-parse HEAD)"
[ "$got" = "$want" ] || { echo "::error::fetched $got, not $want"; exit 1; }
- name: The members' own unit tests (ids, labels, the meter's slot, colouring, bindings, the CUDA kernels compiling)
shell: bash
working-directory: src/wai/wai-quantum
run: |
set -euo pipefail
export PATH="$HOME/.cargo/bin:$PATH"
if [ -n "${{ matrix.target }}" ]; then export CARGO_TARGET_DIR="$HOME/${{ matrix.target }}/wai-quantum-landscape"; fi
nice -n 10 cargo test --locked --release --features "${{ matrix.features }}" --lib -- \
quantum_landscape::device:: quantum_landscape::webgpu:: quantum_landscape::cuda:: quantum_landscape::metal:: 2>&1 | tee "$RUNNER_TEMP/units.log"
grep -qE '^test result: ok\. ${{ matrix.units }} passed; 0 failed' "$RUNNER_TEMP/units.log" || { echo "::error::not every unit test of the members ran and passed (${{ matrix.units }} expected)"; exit 1; }
# Each test that certifies a member writes its certificates to a report, which is read
# here rather than the test's output, so that no test's printing can hide a line.
- name: Every member, certified, held to the reference past every window, and every fault losing its grade
shell: bash
working-directory: src/wai/wai-quantum
run: |
set -euo pipefail
export PATH="$HOME/.cargo/bin:$PATH"
if [ -n "${{ matrix.target }}" ]; then export CARGO_TARGET_DIR="$HOME/${{ matrix.target }}/wai-quantum-landscape"; fi
report="$RUNNER_TEMP/landscape-report.txt"
rm -f "$report"
WAI_LANDSCAPE_REPORT="$report" nice -n 10 \
cargo test --locked --release --features "${{ matrix.features }}" --test landscape_devices -- --test-threads=1 2>&1 | tee "$RUNNER_TEMP/devices.log"
grep -qE '^test result: ok\. 7 passed; 0 failed' "$RUNNER_TEMP/devices.log" || { echo "::error::not every device test ran and passed (7 expected)"; exit 1; }
cat "$report"
IFS=',' read -ra want <<< "$WAI_LANDSCAPE_EXPECT"
for p in "${want[@]}"; do
grep -qE "^certified ${p}[^ ]*: settle bit-exact, sample bit-exact$" "$report" || { echo "::error::$p was not certified bit-exact for both kernels"; exit 1; }
n=$(grep -cE "^fault .* in ${p}[^ ]*: " "$report" || true)
[ "$n" -ge 3 ] || { echo "::error::$p lost its grade under $n faults, and 3 or more are put into every backend's kernels"; exit 1; }
done
grep -qE "^waits [^ ]+: 40 windows" "$report" || { echo "::error::no WebGPU member showed a wait covering one window"; exit 1; }
- name: Timings, wall-clock on this host, with its load (the bits asserted, the times printed)
shell: bash
working-directory: src/wai/wai-quantum
run: |
set -euo pipefail
export PATH="$HOME/.cargo/bin:$PATH"
if [ -n "${{ matrix.target }}" ]; then export CARGO_TARGET_DIR="$HOME/${{ matrix.target }}/wai-quantum-landscape"; fi
nice -n 10 cargo test --locked --release --features "${{ matrix.features }}" --test landscape_devices -- --ignored --nocapture --test-threads=1 timings 2>&1 | tee "$RUNNER_TEMP/timings.log"
grep -qE '^test result: ok\. 1 passed; 0 failed' "$RUNNER_TEMP/timings.log" || { echo "::error::the timings' bits differ from the reference's"; exit 1; }

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

@ -36,6 +36,13 @@ on:
- 'wai/wai-quantum-cli/Cargo.toml'
- 'wai/wai-quantum-cli/Cargo.lock'
- 'wai/wai-web/src/quantum_engine.rs'
- 'wai/wai-rs/src/sink_meter.rs'
- 'wai/wai-rs/src/sink_energy.rs'
- 'wai/wai-rs/src/energy_report.rs'
- 'wai/wai-rs/Cargo.toml'
- 'wai/wai-rs/Cargo.lock'
- 'wai/wai-quantum/src/quantum_fabric.rs'
- 'wai/wai-quantum/src/quantum_energy.rs'
- 'wai/wai-web/Cargo.toml'
- 'wai/wai-web/Cargo.lock'
- 'wai/wai-web/demo/verify_landscape_wasm.mjs'
@ -61,6 +68,13 @@ on:
- 'wai/wai-quantum-cli/Cargo.toml'
- 'wai/wai-quantum-cli/Cargo.lock'
- 'wai/wai-web/src/quantum_engine.rs'
- 'wai/wai-rs/src/sink_meter.rs'
- 'wai/wai-rs/src/sink_energy.rs'
- 'wai/wai-rs/src/energy_report.rs'
- 'wai/wai-rs/Cargo.toml'
- 'wai/wai-rs/Cargo.lock'
- 'wai/wai-quantum/src/quantum_fabric.rs'
- 'wai/wai-quantum/src/quantum_energy.rs'
- 'wai/wai-web/Cargo.toml'
- 'wai/wai-web/Cargo.lock'
- 'wai/wai-web/demo/verify_landscape_wasm.mjs'
@ -111,7 +125,7 @@ jobs:
run: |
set -euo pipefail
cargo test --locked --release --features full --lib -- quantum_landscape:: 2>&1 | tee "$RUNNER_TEMP/landscape.log"
grep -qE '^test result: ok\. 83 passed; 0 failed' "$RUNNER_TEMP/landscape.log" || { echo "::error::the landscape tests did not all pass (83 expected)"; exit 1; }
grep -qE '^test result: ok\. 95 passed; 0 failed' "$RUNNER_TEMP/landscape.log" || { echo "::error::the landscape tests did not all pass (95 expected)"; exit 1; }
for t in quantum_landscape::classical::tests::a_solver_keeps_to_its_budget_and_repeats_itself \
quantum_landscape::classical::tests::each_solver_finds_the_exact_lowest_on_small_instances_with_a_generous_budget \
quantum_landscape::classical::tests::with_no_budget_a_solver_misses \
@ -124,6 +138,12 @@ jobs:
quantum_landscape::energy::tests::the_lowered_energy_is_the_declared_one \
quantum_landscape::energy::tests::the_regime_reads_the_cost_against_the_couplings \
quantum_landscape::job::tests::a_bad_request_is_refused_plainly \
quantum_landscape::job::tests::a_job_goes_to_the_first_member_that_takes_it \
quantum_landscape::job::tests::a_member_never_runs_what_it_declines \
quantum_landscape::job::tests::a_member_that_fails_a_job_passes_it_on \
quantum_landscape::job::tests::an_uncertified_kernel_is_never_placed \
quantum_landscape::job::tests::every_answer_says_who_ran_it_at_what_parity \
quantum_landscape::job::tests::a_member_may_decline_only_what_does_not_fit_32_bits \
quantum_landscape::job::tests::a_sealed_answer_verifies_and_follows_its_schema \
quantum_landscape::job::tests::a_seed_keeps_all_64_bits \
quantum_landscape::job::tests::a_temperature_in_text_keeps_every_bit \
@ -158,11 +178,17 @@ jobs:
quantum_landscape::program::tests::temperatures \
quantum_landscape::program::tests::warnings_come_when_a_program_is_read \
quantum_landscape::program::tests::weights_go_through_quantisation_and_say_what_was_lost \
quantum_landscape::reach::tests::a_table_holds_only_what_the_cells_can_reach \
quantum_landscape::reach::tests::a_table_is_bounded_by_the_smaller_of_the_two_forms \
quantum_landscape::reach::tests::a_window_ends_at_its_work_or_at_its_table \
quantum_landscape::reach::tests::no_job_can_outgrow_a_table \
quantum_landscape::reach::tests::the_table_is_the_reference_s_cut_points \
quantum_landscape::run::tests::a_comparison_gives_every_solver_the_same_budget \
quantum_landscape::run::tests::a_job_runs_on_another_member_and_its_answer_names_it \
quantum_landscape::run::tests::a_member_is_held_to_its_ending \
quantum_landscape::run::tests::a_member_s_bad_answer_is_refused \
quantum_landscape::run::tests::a_member_s_claim_of_parity_is_not_taken_on_its_word \
quantum_landscape::run::tests::a_run_s_reading_is_its_calls_added_up \
quantum_landscape::run::tests::a_run_too_big_is_refused_before_it_starts \
quantum_landscape::run::tests::a_sweep_keeps_the_budget_and_the_cold_row_of_a_minimize_runs_once \
quantum_landscape::run::tests::the_same_seed_gives_the_same_answer_and_another_seed_another_run \
@ -203,7 +229,7 @@ jobs:
run: |
set -euo pipefail
cargo test --locked --release --features quantum_landscape --test landscape_conformance 2>&1 | tee "$RUNNER_TEMP/conformance.log"
grep -qE '^test result: ok\. 8 passed; 0 failed' "$RUNNER_TEMP/conformance.log" || { echo "::error::the conformance tests did not all pass (8 expected)"; exit 1; }
grep -qE '^test result: ok\. 12 passed; 0 failed' "$RUNNER_TEMP/conformance.log" || { echo "::error::the conformance tests did not all pass (12 expected)"; exit 1; }
# With the module's feature alone: a temperature given as a JSON number is read to the nearest double.
- name: Numbers read to the nearest double
shell: bash
@ -212,6 +238,14 @@ jobs:
set -euo pipefail
cargo test --locked --release --no-default-features --features quantum_landscape --test landscape_float_read 2>&1 | tee "$RUNNER_TEMP/float.log"
grep -qE '^test result: ok\. 1 passed; 0 failed' "$RUNNER_TEMP/float.log" || { echo "::error::the float reading test did not pass"; exit 1; }
# The bridge that makes a run meter of an on-chip counter, which a GPU member reads.
- name: The meter bridge (wai-rs sink_meter::RunCounter)
shell: bash
working-directory: wai/wai-rs
run: |
set -euo pipefail
cargo test --locked --release --no-default-features --features sink_energy,quantum_fabric --lib -- sink_meter::tests::a_counter_meters_a_run_only_where_it_was_live_and_says_what_it_covers 2>&1 | tee "$RUNNER_TEMP/meter.log"
grep -qE '^test result: ok\. 1 passed; 0 failed' "$RUNNER_TEMP/meter.log" || { echo "::error::the meter bridge's test did not run and pass"; exit 1; }
- name: The command line's verb
shell: bash
working-directory: wai/wai-quantum-cli
@ -271,7 +305,7 @@ jobs:
quantum_landscape::settle::tests::four_recall_vectors_are_pinned 2>&1 | tee "$RUNNER_TEMP/bits.log"
grep -qE '^test result: ok\. 2 passed; 0 failed' "$RUNNER_TEMP/bits.log" || { echo "::error::the pinned bits differ here ($(uname -sm), ${profile:-debug})"; exit 1; }
nice -n 10 cargo test --locked $profile --no-default-features --features quantum_landscape --test landscape_conformance 2>&1 | tee "$RUNNER_TEMP/vectors.log"
grep -qE '^test result: ok\. 8 passed; 0 failed' "$RUNNER_TEMP/vectors.log" || { echo "::error::the conformance set differs here ($(uname -sm), ${profile:-debug})"; exit 1; }
grep -qE '^test result: ok\. 12 passed; 0 failed' "$RUNNER_TEMP/vectors.log" || { echo "::error::the conformance set differs here ($(uname -sm), ${profile:-debug})"; exit 1; }
done
wasip2:
@ -320,7 +354,7 @@ jobs:
run: |
set -euo pipefail
cargo test --locked --target wasm32-wasip2 --no-default-features --features quantum_engine,quantum_receipt --lib -- quantum_landscape:: 2>&1 | tee "$RUNNER_TEMP/wasip2.log"
grep -qE '^test result: ok\. 82 passed; 0 failed' "$RUNNER_TEMP/wasip2.log" || { echo "::error::the landscape tests did not all pass under wasip2 (82 expected)"; exit 1; }
grep -qE '^test result: ok\. 94 passed; 0 failed' "$RUNNER_TEMP/wasip2.log" || { echo "::error::the landscape tests did not all pass under wasip2 (94 expected)"; exit 1; }
# The browser build's answers against the command line's, byte for byte, in Node.
parity:

7
wai/changes/2026-11-06-landscape-gpu-members.md Normal file
View file

@ -0,0 +1,7 @@
### wai-quantum: energy-landscape jobs on a GPU, through WebGPU, CUDA and Metal
- **`quantum_landscape_webgpu`, `quantum_landscape_cuda` and `quantum_landscape_metal`** run the settle kernel (one thread a cell, one dispatch a sweep) and the sampling kernel (one thread a chain, its cells in the reference's order, the reference's draws and the host's cut points for the fields the cells can reach) on a GPU, a window of sweeps a dispatch, each a member the host certifies through `member::certify`. Each certifies bit-exact for both kernels on the devices it ran on here, and each fault a port of the kernels can have, put into them on those devices, costs them the grade. The WebGPU member's coloured kernel, a colour class of cells at a time, earns `statistical` and no more.
- **The conformance set reaches what large jobs reach**: settles past one workgroup of cells and past one batch of sweeps; chains that cross a device's windows of sweeps, more chains than a workgroup, chains that stop hot; one-cell chains at the seeds where a port of the generator in 32-bit words first goes wrong; fields that only their full width tells apart; and field bounds at the decline line and just past it. The battery adds workloads past four cells, one whose schedule matters to its law, and fields past 16 and 32 bits.
- **Declining and placing.** A workload's field bound is the largest field any state can give a cell. A member may decline a workload only where that is above 2³¹ − 1, and the host never runs it on one it declines. `job::call_placed` passes over a member whose kernel for the job is not certified, one that declines it, and one that fails it, naming each under `refusals`. Every answer says who ran it at what parity: a replay its member, a sweep and a comparison each kernel they used.
- **A job's budget holds a device's work**: a sweep's table of cut points holds the fields its cells can reach, at most `4P + n` entries, counted in the job's operations and memory.
- **`wai::sink_meter::RunCounter`** makes a run meter of an on-chip counter (IOReport's GPU channel; NVML once it has shown itself plausible) which reads the counter through each call, gives a figure only from a live window of the profile's minimum or more, and says what it covers: the whole rail or board, its idle draw and any other work on the device in the window. An answer carries that beside the figure; a run's calls' figures and windows are added up.

160
wai/extensions/quantum-landscape.md
View file

@ -84,12 +84,15 @@ With `n` variables, `P` coupled pairs, `C` items in the checks and a plan of `R`
starts of `S` sweeps, the operations of a job are `n²` (reading the couplings into
rows), plus, where `n ≤ 16`, `3ⁿ·(n + 1 + C)` (the exhaustive truth) and, for a
program with `weight`s, `3ⁿ·n²` more; plus, for each run, `(2P + n)·R·S` (the
sweeps) `+ (2P + n)·R` (each start's energy) `+ n²`, or `2n²` with weights (the
answer's energy) `+ C` (its checks). A `run` is one run; a `replay` and each row of
sweeps) `+ (4P + n)·S·⌈R / D⌉` for a plan above temperature 0, with `D` the smaller of
4096 and `⌊1,048,576 / n⌋` (one sweep's cut points as a device works them out, §4.2,
once for each batch of starts, §5) `+ (2P + n)·R` (each start's energy) `+ n²`, or
`2n²` with weights (the answer's energy) `+ C` (its checks). A `run` is one run; a `replay` and each row of
a `sweep` are one run a seed, a settle row one run; a `compare` is five runs a seed
of its sampling plan. The memory is the program's bytes `+ n²`, `+ 8n²` with
weights, `+ 32P + 64n`, `+ 2,359,296` (a batch of chains, §5), `+ 256` a line of a
`replay`, `+ 1,048,576`.
weights, `+ 32P + 64n`, `+ 64P + 16n` (one sweep's cut points as a device holds
them, §4.2), `+ 2,359,296` (a batch of chains, §5), `+ 256` a line of a `replay`,
`+ 1,048,576`.
A host does not hold more than one batch of chains at once: at most 1,048,576 cells
and at most 4096 chains a call to a member.
@ -191,6 +194,13 @@ A `run` answers with these fields (the schema has them all):
| `work` | `sweeps`, `accumulates`, `cell_updates`, `draws`, `ended`, `rests` |
| `joules` | §4.4 |
| `state_hash`, `replay` | §3.2 |
| `refusals` | each member that a placing host passed over, with why (§5); absent where there was none |
A `replay` answers with `protocol`, `label`, `member` (`id`, `class`, `kernel`,
`parity`), `seeds` and `lines`. A `sweep` and a `compare` give `member` as `id`,
`class` and `kernels`, each kernel the job used with the parity the host found for it;
each row of a `sweep` gives its own `kernel` and `parity`, and each of a `compare` the
member's `parity`, null for a classical method. Each carries `refusals` as a `run` does.
An accumulate is one read of one coupling into one cell's field, and is the unit
of work. The accumulates of a sweep are the non-zero entries of `W`, counting each
@ -258,6 +268,10 @@ the state reported is the one the last sweep computed) or the bound. Everything
on the order a field is added in, on how many lanes a device has or on how it
splits the cells. A settle kernel MUST give the reference's bits.
A workload's **field bound** is the largest `|h_i|` any state can give a cell:
`max_i (|b_i| + deg_i)`, with `deg_i` the number of cells coupled to `i`. It, and not
the largest `|b_i|`, says how wide a lane must be to hold every field.
A `quench` runs chains through the kernel of §4.2 at temperature 0, whose every
sweep is this one.
@ -298,6 +312,9 @@ and `exp` the function of Appendix A.
A 53-bit number `u` gives −1 if `u < a`, 0 if `u < b`, else +1. A device needs a table
of cut points for the fields it can see at each sweep's temperature and two comparisons
a cell. It needs no exponential and no floating point. The host may give it the table.
Cell `i` only ever sees `b_i − deg_i` to `b_i + deg_i`, so a table of those fields, where
two overlap taken once, holds at most `4P + n` entries a sweep however wide the fields
are, and at most `2B + 1` for a field bound `B`.
**The reference's draws.** They come from a counter, not from a state carried along:
@ -345,10 +362,16 @@ multiplied by a rate measured elsewhere, with that rate's source named. A host:
- MUST NOT print an uncertainty it does not have.
A member with a meter puts its reading in `joules.metered`: `{ "joules_micro": …,
"class": … }`, with the class and the evidence energy-measurement asks of that class.
The reading stands beside the priced figures and does not replace them, so the two can
be compared. Where there is no meter, `metered` is null. No figure is `0` for being
unmeasured.
"window_us": …, "class": …, "scope": … }`, with the class and the evidence
energy-measurement asks of that class. Over a run of several calls the figure and the
window are each the calls' added up, and so is the window of the class's uncertainty;
its relative part is the largest a call declared. `scope` says what the figure covers
where that is more than the run: a counter that reads a whole rail or board holds its
idle draw and any other work on the device in the window, and says so; null is the run
alone. The reading stands beside the priced figures and does not replace them, so the
two can be compared, and the answer's `label` says it is there and what it covers only
where it is. Where there is no meter or no reading, `metered` is null. No figure is `0`
for being unmeasured.
The reference's rates:
@ -376,6 +399,29 @@ batch of chains at a time (§2.0).
What a member gives back, for a settle or each chain: the state, the sweeps, how it
ended, the accumulates and the draws.
**Declining.** Before anything runs, a member MAY decline a workload whose field bound
(§4.1) is above 2³¹ − 1, which a device with 32-bit lanes cannot hold, saying why in
words that are true of the workload. It MAY decline no other workload a job can give;
§2.0's limits keep every other need of a job, such as a sweep's table (§4.2), within
what any member can hold. A member may decline for one kernel and not the other: the
rule is the same for each. A host runs a member on no workload it declines, however the
member was reached. The reference declines nothing.
**Placing a job.** A host offered several members gives the job to the first that can
take it: one whose kernel is certified for every kernel the job runs (§5.1), that does
not decline the workload, and that gives an answer the host's checks keep. Each member
passed over, for any of those, is named under `refusals`, in the order offered, with
why; with none, `refusals` is absent. A member that took the job and failed it, or
whose answer the checks refused, is passed over the same way and the next is asked.
Where no member answers, the job is refused, naming each with why. Offering the
reference last gives every job an answer.
**Batches.** A host MAY give a member the starts of several seeds' runs in one call,
up to one batch (§2.0). Each run is the one it would be alone: the same chains, its
best state chosen in the same order. A meter's reading (§4.4) stands for a run only
where every call of that run carried that run alone, and a settle that answers several
seeds from one call carried none of them alone; elsewhere `metered` is null.
A host MUST check what comes back, and refuse the answer with one plain sentence if
any of these does not hold:
@ -392,6 +438,16 @@ only by returning a worse one, and it cannot price itself low. How many sweeps a
took before it came to rest can be checked only by running it again: a `bit-exact`
certificate is what vouches for it.
The reference implementation has three GPU members besides the CPU reference, each
certified by the host on its own run of the conformance set (informative): through WebGPU
(32-bit lanes; it declines a workload whose field bound is above 2³¹ − 1, holds a cost
above every field at its lanes' top, where it holds every cell at 0 as it does at its
own size, and its coloured kernel can be `statistical` at most), through CUDA and through
Metal (64-bit fields, so they take every workload). All three run the sampling kernel one
chain a thread, the cells in the reference's order, with the reference's draws and the
host's cut points for the fields the cells can reach, a window of sweeps a dispatch, each
window waited for before the next is sent; and give its bits.
A member names itself `class:name`, with `class` one of the fabric's (`cpu`, `gpu`,
`fpga`) and the class it gives, and `name` lower-case letters, digits, `.`, `_` and `-`,
at most 64 characters in all. `cpu:reference` is the reference's, and no other member's.
@ -405,16 +461,22 @@ run used (`member.kernel`, `member.parity`).
| parity | granted when |
|---|---|
| `reference` | to the host's own reference kernels, and to nothing else |
| `bit-exact` | settle: the member gives every `settle` vector exactly. Sample: it gives every `chains` vector exactly, each case's chains in one call |
| `bit-exact` | settle: the member gives every `settle` vector it takes exactly. Sample: it gives every `chains` vector it takes exactly, each case's chains in one call. Either way it takes every vector §5 does not let it decline |
| `statistical` | sample: not `bit-exact`, and every workload of the `statistical` battery passes (§7.2) |
| `uncertified` | anything else, or a member the host was asked to run by name without certifying it |
A member is held to every vector it takes, and declining a vector §5 does not let it
decline is failing that vector. The set holds vectors on both sides of the line, at a
field bound of exactly 2³¹ − 1 and just above it by the couplings alone, so that a
member is held to the edge of what it takes.
A settle kernel can earn `bit-exact` on any device. A sampling kernel can, if it uses the
reference's draws and order and the cut points of §4.2.
The fabric's placement never places an uncertified kernel. A host MAY still run one when
it is asked for that member by name, which is how a board gets its first answers; the
answer says `uncertified`, and nothing should be concluded from its state but its energy.
A host placing a job never places an uncertified kernel (§5). A host MAY still run one
when it is asked for that member by name, without placing it, which is how a board gets
its first answers; the answer says `uncertified`, and nothing should be concluded from
its state but its energy.
### 5.2. The receipt
@ -435,7 +497,10 @@ reference's, word for word:
> scale. No photonic hardware, custom chip or quantum effect is involved in this run.
A member states its own, and MUST say what device ran the job, whether its joules are
priced or read from a meter, and that no quantum effect is involved where none is.
priced or read from a meter, and that no quantum effect is involved where none is. It
calls a device what it is: a software driver that runs on the processor is not a GPU.
An answer adds that a reading stands beside the priced figures, and what it covers, only
where that answer carries one (§4.4).
## 6. The comparison
@ -458,19 +523,35 @@ way, is caught.
### 7.1. Bits
- `settle`: whole workloads, starts and bounds, and the state, sweeps, ending,
accumulates and state hash the settle kernel gives: the four recall vectors, every
ending, a two-cycle with fields, a bound of one sweep and fields at 2⁴⁰. A settle
kernel is `bit-exact` when it gives every one exactly.
- `settle`: whole workloads, starts and bounds, each with its field bound, and the state,
sweeps, ending, accumulates and state hash the settle kernel gives: the four recall
vectors, every ending, a two-cycle with fields, a bound of one sweep, fields at 2⁴⁰;
waves down paths of 65 to 300 cells, which run past one workgroup of cells and past
64 sweeps, one stopped at its bound in a later batch, a two-cycle of 70 cells and 1000
cells of mixed couplings; fields that only their full width tells apart past 16 and
32 bits; and field bounds at 2³¹ − 1 and just above it (§5.1). A settle kernel is
`bit-exact` when it gives every one it takes exactly and takes every one §5 does not
let it decline.
- `exp`: the bits of Appendix A's function at the arguments the cut points use and at its
edges, as `[x, exp(x)]` pairs of 16 hexadecimal digits.
- `schedule`: the bits of every sweep's temperature for starting temperatures and chain
lengths that tell `T₀ − T₀·q` from other ways of writing it.
- `cells`: the two cut points for fields from −40 to 40, costs from 0 to 8 and temperatures
from 0.25 to 25, the temperature given by its bits too.
- `draws`: the generator's outputs `out(s, j)` at several seeds, and at each seed of the
`a-draw` chains below.
- `chains`: whole chains from their starts and chain seeds, with every start of some runs,
cooling from starting temperatures that include one that is not a short decimal. A
sampling kernel is `bit-exact` when it gives every case's chains exactly, from one call.
cooling from starting temperatures that include one that is not a short decimal, or
held at temperatures given sweep by sweep (`held`), which need not end at 0, so that
the last sweep's draws are in the final state. Among them: one cell at seeds where a
port of the generator in 32-bit words first goes wrong (a carry out of the middle of a
product, a cut point compared by its high word alone); 16,384 draws in one sweep;
more chains than a workgroup; chains that stop hot, in which the order of the cells
shows; chains long enough to cross a device's windows of sweeps, by table and by
work, hot in every window; and the same fields past 16 and 32 bits and at the line as
the settle vectors, at 0 and at a temperature of their own size. A sampling kernel is
`bit-exact` when it gives every case it takes exactly, each from one call, and takes
every one §5 does not let it decline.
- `jobs`: whole requests and what their answers hold, refusals included: every worked
example at seed 1, and requests that pin a rule each (a `minimize` above temperature 0,
a temperature given as text, a pin's size, quantised weights, the state a two-cycle
@ -481,24 +562,35 @@ way, is caught.
### 7.2. The battery
A sampling kernel with draws of its own is held to `statistical`: workloads of two to four
cells, each with a schedule. Workload `k` runs `statistical.chains` (20,000) chains, chain
`r` with the chain seed `out(statistical.seed + k, r + 1)` and its start from that seed
(§4.2), through the workload's temperatures. The final states are counted over the `3ⁿ`
states, in the order `Σ (s_i + 1)·3ⁱ`, and held to `exact`: the Boltzmann distribution at
the schedule's last temperature above 0, carried through one deterministic sweep where the
schedule ends at 0. The statistic is Pearson's, every state expected fewer than
`statistical.pool_below` (5) times pooled into one bin; with `df` the bins less one, a
workload passes when it is at most `df + 5·√(2·df)`.
A sampling kernel with draws of its own is held to `statistical`: workloads of two to six
cells, and one of eight blocks of three side by side, none coupled to another, each with
a schedule. A member may decline a workload of the battery only where §5 lets it decline
a vector, and it then passes that workload without running it. Workload `k` runs `statistical.chains` (20,000) chains, chain `r` with the
chain seed `out(statistical.seed + k, r + 1)` and its start from that seed (§4.2),
through the workload's temperatures. The final states of each `block` of cells are
counted over its `3ⁿ` states, in the order `Σ (s_i + 1)·3ⁱ`, and held to `exact`. Where
`law` is `settled`, that is the Boltzmann distribution at the schedule's last temperature
above 0, reached on a plateau, carried through one deterministic sweep where the schedule
ends at 0. Where it is `walked`, it is the law of the walk itself from the start's
uniform law, every sweep of the schedule cell by cell in order, on a block whose cells
are all alike; there the counts are read over how many cells of the block hold each
state, which any order of the cells leaves the same. The statistic is Pearson's, every
bin expected fewer than `statistical.pool_below` (5) times pooled into one; with `df` the
bins less one, a block passes when it is at most `df + 5·√(2·df)`, and a workload when
every block does.
The workloads hold pairs of coupled cells next to each other, two apart and three apart,
a triangle no two colours can split, fields to 40 and a cost of 8 at a high temperature, a
small field at a low one, a schedule that ends at 0, and one that cools from 6 to 0.7. A
kernel that moves two coupled cells at once in any order of cells, that ignores the
schedule, that samples the last sweep rather than running it at 0, or whose table is wrong
for large fields fails at least one of them; the reference's own kernel, and one that
colours its cells properly, pass every one. A test holds both statements, with kernels
that are wrong in each of those ways.
small field at a low one, a schedule that ends at 0, one that cools from 6 to 0.7, a ring
of six cells, eight triangles side by side, a triangle pulled one way that is warm
and then cold, whose end alone gives another law, and fields past 16 bits and past 32 at
a temperature of their own size. A kernel that moves two coupled cells
at once in any order of cells, that ignores the schedule or keeps only its end, that
samples the last sweep rather than running it at 0, whose table is wrong for large
fields, that holds its fields in 16 or 32 bits, or that leaves cells from the fifth on
where they started fails at least one of them; the reference's own kernel, one that colours its cells properly, and one that
takes its cells last to first pass every one. A test holds both statements, with
kernels that are wrong in each of those ways.
## 8. What does not run on this model

14
wai/quantum-conformance/landscape/README.md
View file

@ -132,8 +132,12 @@ of board power the priced line leaves out. A run's joules are not measured here,
## A member's conformance
A member is held to [`vectors.json`](vectors.json) as §7 of the extension says, and the host finds its parity by running its kernels there: `settle` bit for bit
for a `bit-exact` settle kernel; `chains` bit for bit for a `bit-exact` sampling kernel; otherwise the `statistical` battery, eight small workloads of 20,000 chains each,
their final states against the exact distribution. `wai-quantum/tests/landscape_conformance.rs` reads the file as a member's checker would, from the file alone,
and runs the battery on the reference's own kernel, which passes; on five kernels each wrong in one way (cells moved all at once, even then odd, the schedule
ignored, a table clamped at 5, the last sweep sampled), each of which fails at least one workload; and on one that colours its cells properly, which passes
and is granted `statistical`.
for a `bit-exact` settle kernel, every vector it takes, and it takes every vector whose field bound is at most 2³¹ − 1; `chains` bit for bit for a `bit-exact`
sampling kernel, the same way; otherwise the `statistical` battery, thirteen workloads of 20,000 chains each, each block's final states against its exact
distribution. `wai-quantum/tests/landscape_conformance.rs` reads the file as a member's checker would, from the file alone, and holds both parts to kernels
that are wrong in one way each. The vectors refuse a dropped carry in a 32 by 32 product, a cut point compared by its high word, fields held in 16 or 32
bits, a check of the fields that forgets the couplings, cells last to first, a draw counter that starts again at each of a device's windows of sweeps, a
run that stops after one window, chains past the first 64 left alone, a settle whose cells past 64 or 256 never move or that reports its end a sweep late,
and a member that declines below the line. The battery passes the reference's own kernel, one that colours its cells properly and one that takes its
cells last to first, and refuses kernels that move cells all at once, even then odd, ignore the schedule or keep only its end, clamp the table at 5,
sample the last sweep, hold their fields in 16 or 32 bits, or leave cells from the fifth on where they started.

39754
wai/quantum-conformance/landscape/vectors.json
View file

File diff suppressed because it is too large Load diff

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

@ -14,7 +14,7 @@ readme = "README.md"
# quantum-conformance/qasm3/, quantum-conformance/qec-text/, the answer
# layer's conformance set and quantum-conformance/formats/) to the crate;
# they run in the repository only.
exclude = ["tests/landscape_conformance.rs", "tests/fabric_expected.rs", "tests/qasm3_conformance.rs", "tests/qec_text_conformance.rs", "tests/answer_conformance.rs", "tests/formats_conformance.rs", "tests/formats_mutation.rs"]
exclude = ["tests/landscape_conformance.rs", "tests/landscape_devices.rs", "tests/fabric_expected.rs", "tests/qasm3_conformance.rs", "tests/qec_text_conformance.rs", "tests/answer_conformance.rs", "tests/formats_conformance.rs", "tests/formats_mutation.rs"]
# docs.rs builds with DEFAULT features only, which would document 3 of 21
# modules and hide the entire operational + ML layers. Document everything.
@ -143,6 +143,13 @@ quantum_qio = []
# as the reference, every arrangement checked where a program is small, the joules
# priced, and standard classical methods on the same instance (extensions/quantum-landscape.md).
quantum_landscape = ["quantum_qio", "serde_json/float_roundtrip"]
# The landscape kernels on a GPU through WebGPU, natively (quantum_landscape::webgpu): the
# settle kernel and the sampling kernel bit-exact in 32-bit lanes, and a coloured kernel.
quantum_landscape_webgpu = ["quantum_landscape", "quantum_fabric_webgpu"]
# The landscape kernels on an NVIDIA GPU through CUDA (quantum_landscape::cuda), and on an Apple
# GPU through Metal (quantum_landscape::metal): 64-bit fields, both kernels bit-exact.
quantum_landscape_cuda = ["quantum_landscape", "quantum_fabric_cuda"]
quantum_landscape_metal = ["quantum_landscape", "quantum_fabric_metal"]
# Union-find decoding in near-linear time over the matching graph.
quantum_uf = ["quantum_frame"]
# Noisy circuits: quantum trajectories on the state vector, an exact density-
@ -252,6 +259,10 @@ required-features = ["quantum_landscape"]
name = "landscape_float_read"
required-features = ["quantum_landscape"]
[[test]]
name = "landscape_devices"
required-features = ["quantum_landscape"]
[[test]]
name = "webgpu_member"
required-features = ["quantum_fabric_webgpu"]

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

@ -80,6 +80,7 @@ never creates information.
| `quantum_resource_receipt` | a signed, energy-accounted receipt over an estimate, re-checkable bit for bit on any machine |
| `quantum_qio` | **quantum-inspired optimisation**: Ising and Max-Cut problems (G-set reader, spin-glass generator), simulated bifurcation in its ballistic, discrete and heated forms (batched eight trials per pass, bit-identical to one at a time), simulated annealing, single-flip descent, an exact Gray-code referee up to 34 spins, and step-to-solution. Reaches the best-known cut on seven G-set instances |
| `quantum_landscape` | **energy-landscape jobs**: a program in epuML states an energy over variables of three states (spin-1, with a single-ion cost), run on a classical emulator of the settling model: a deterministic kernel in whole numbers and a sampling kernel of one draw between two cut points, both the same bits on every platform; every arrangement checked up to 16 variables, each lowest one graded against the problem's own check; the joules priced at rates whose sources are named, never measured here; the same instance on simulated annealing, tabu search and greedy descent at the same budget, with the published table; and the seam a GPU or an FPGA board runs behind (`extensions/quantum-landscape.md`). `quantum_engine` and `full` both turn it on |
| `quantum_landscape_webgpu`, `quantum_landscape_cuda`, `quantum_landscape_metal` | **energy-landscape jobs on a GPU**: the settle and sampling kernels through WebGPU (32-bit lanes, and a coloured kernel that can be statistical at most), CUDA and Metal (64-bit fields), each member certified by the host on the conformance set and given only the workloads it takes, a job passing over an uncertified kernel or a member that fails it (`job::call_placed`) |
| `quantum_uf` | **union-find decoding in near-linear time** over the same matching graph: event-driven cluster growth on integer weights (a radix heap orders the fill events), merges, and peeling. Measured against exact matching on the same shots: within 0.92–1.18× the failures, up to 77× faster at `d = 11` |
| `quantum_dmrg` | **ground states of spin chains by DMRG**: matrix-product operators for nearest-neighbour chains (XXZ Heisenberg in a field, transverse-field Ising), two-site sweeps with Lanczos and density-matrix truncation, entanglement entropies, a symmetric eigensolver, and two exact referees: full diagonalisation up to 22 sites, and the free-fermion solution of the Ising chain at any length. The critical Ising chain is reproduced to about `10⁻¹¹` per site at 200 sites |
| `quantum_tdvp` | **time evolution of spin chains by TDVP**: two-site sweeps forward and back, Lanczos exponentials that halve a step they cannot converge, bonds split by singular value decomposition, and every basis completed with null vectors so a product state's small bonds do not hold the evolution back. Measures site and bond observables, entanglement at every bond, energy and the Loschmidt echo. Held to two exact referees: Krylov evolution of all amplitudes up to 20 sites, and the free-fermion solution of an Ising quench at any length. On a 100-site critical quench at bond 64 it stays within `10⁻⁹` of the exact answer until the entanglement fills the bond, and past that its error follows the discarded weight it reports |

1
wai/wai-quantum/README.zh-CN.md
View file

@ -63,6 +63,7 @@ wai-quantum lab qutrit # d = 3，不只是量子比特
| `quantum_resource_receipt` | 针对一次估算的、带能耗核算的签名回执，可在任何机器上逐位重新校验 |
| `quantum_qio` | **量子启发式优化**：Ising 与 Max-Cut 问题（G-set 读取器、自旋玻璃生成器），弹道式、离散式和加热式的模拟分岔（每遍批量运行八次试验，与逐次运行的结果逐位一致），模拟退火，单翻转下降，最多 34 个自旋的精确 Gray 码参照，以及求解所需步数（step-to-solution）。在七个 G-set 实例上达到已知最优割 |
| `quantum_landscape` | **能量景观作业**：epuML 程序对三态变量（自旋 1，带单离子代价）给出能量，在沉降模型的经典模拟器上运行：整数运算的确定性核，以及在两个切点之间做一次抽样的采样核，两者在每个平台上比特一致；最多 16 个变量时检查每一种排列，并按问题自身的检查为每个最低能量状态评分；焦耳按注明来源的费率定价，这里从不测量；同一实例上以相同预算运行模拟退火、禁忌搜索和贪心下降，并发布结果表；以及 GPU 或 FPGA 板运行所经由的接缝（`extensions/quantum-landscape.md`）。`quantum_engine` 和 `full` 都会启用它 |
| `quantum_landscape_webgpu`、`quantum_landscape_cuda`、`quantum_landscape_metal` | **GPU 上的能量景观作业**：沉降核与采样核分别通过 WebGPU（32 位通道，另有至多只能达到统计一致的着色核）、CUDA 与 Metal（64 位场）运行，每个成员都由主机在一致性集合上认证，且只接收它能承接的工作负载；作业会跳过未经认证的核或运行失败的成员（`job::call_placed`） |
| `quantum_uf` | **近线性时间的并查集（union-find）解码**，基于同一个匹配图：在整数权重上由事件驱动的簇生长（用基数堆为填充事件排序）、合并以及剥离（peeling）。在相同的采样上与精确匹配对比：失败次数为其 0.92–1.18 倍，在 `d = 11` 时最多快 77 倍 |
| `quantum_dmrg` | **用 DMRG 求自旋链的基态**：最近邻链的矩阵乘积算符（外场中的 XXZ Heisenberg 链、横场 Ising 链），带 Lanczos 和密度矩阵截断的双格点扫描，纠缠熵，一个对称矩阵本征求解器，以及两个精确参照：最多 22 个格点的完全对角化，以及任意长度 Ising 链的自由费米子解。在 200 个格点上，临界 Ising 链的复现精度约为每格点 `10⁻¹¹` |
| `quantum_tdvp` | **用 TDVP 计算自旋链的时间演化**：前后往返的双格点扫描，遇到无法收敛的步长就将其减半的 Lanczos 指数运算，用奇异值分解拆分键，并用零空间向量补全每个基，使乘积态的小键不会阻碍演化。测量格点和键上的可观测量、每个键上的纠缠、能量以及 Loschmidt 回波。以两个精确参照为准：最多 20 个格点的全振幅 Krylov 演化，以及任意长度 Ising 淬火的自由费米子解。在键维数 64、100 个格点的临界淬火中，直到纠缠填满键之前，它与精确解的偏差都保持在 `10⁻⁹` 以内；超过之后，其误差随它所报告的舍弃权重变化 |

49
wai/wai-quantum/docs/landscape.md
View file

@ -170,6 +170,55 @@ In a page, the browser build answers the same request: `QuantumEngine.landscape(
answer with the page's key, and `QuantumEngine.landscape_answer(op, request)` gives the bytes without a
seal. `wai-web/demo/landscape.html` is the page.
## On a GPU
The kernels run on a GPU too, as members the host certifies and checks like any other:
| feature | member | lanes | parity it can earn |
|---|---|---|---|
| `quantum_landscape_webgpu` | `quantum_landscape::webgpu`: Metal, Vulkan or DX12 through WebGPU | 32-bit: it declines a workload whose field bound is above 2³¹ − 1 | bit-exact for both kernels; its coloured kernel, a colour class of cells at once, statistical at most |
| `quantum_landscape_cuda` | `quantum_landscape::cuda`: an NVIDIA GPU, compiled at run time | 64-bit: it takes every workload | bit-exact for both kernels |
| `quantum_landscape_metal` | `quantum_landscape::metal`: an Apple GPU, macOS | 64-bit: it takes every workload | bit-exact for both kernels |
A parity in this table is what a member can earn, not what any one device has: the host finds it
on each device, every time it certifies one, and states it in every answer.
The settle kernel runs one thread a cell and one dispatch a sweep. The sampling kernel runs one
thread a chain, its cells in the reference's order, with the reference's draws and the cut points the
host works out for the fields the cells can reach (at most `4P + n` a sweep, counted in the job's
operations and memory): so its bits are the reference's, and many chains run side by side, the
starts of a run and the seeds of a sweep, a comparison or a replay. A member runs a window of sweeps
a dispatch and waits for each before it sends the next.
A workload's field bound is the largest field any state can give a cell, `max_i (|b_i| + deg_i)`.
A member may decline a workload only where that is above 2³¹ − 1, and the host never runs it on one
it declines. Placing a job, the host passes over a member whose kernel for the job is not certified,
one that declines the workload, and one that fails the job or whose answer its checks refuse, and
names each under `refusals`; offered last, the reference answers what the others did not.
```rust norun=account
use wai_quantum::quantum_landscape::{job, member, webgpu, MAX_JOB_BYTES};
use serde_json::json;
let gpu = webgpu::discover().into_iter().find_map(Result::ok).expect("a GPU");
let placed = [member::certify(&gpu).unwrap(), member::Certified::reference()];
let answer = job::call_placed(&placed, "run", &json!({ "program": "vars 2\noppose 0 1\nsolve\n" }), MAX_JOB_BYTES).unwrap();
// answer["member"] names what ran the job and the parity the host found for the kernel it used;
// answer["refusals"], if any, names each member passed over, with why.
```
A member can carry a meter: `with_meter` takes the engine's run meter, and `wai::sink_meter::RunCounter`
makes one of a machine's own on-chip counter: the graphics rail through IOReport on macOS, and the
board through NVML on Linux once NVML has shown itself plausible (`RunCounter::nvml`). It reads the
counter through each call and gives a figure only from a window in which the counter was live and
which lasted the profile's minimum, labelled `OnChipCounter` with the uncertainty the operator declares
and the meter's own terms beside it. **The figure is the whole rail or board over the call's window:
the run's work, the device's idle draw and any other work on the device then.** No share is
attributed to the run, and the answer says so: `joules.metered.scope` states it, and the `label`
says a reading is there, and what it covers, only in an answer that has one. Over a run of several
calls the figures and the windows are added up. A reading stands for a run only where every call
carried that run alone; elsewhere `metered` is null.
## A member of your own
A member is what runs the kernels: the trait

269
wai/wai-quantum/schemas/landscape.v1.json
View file

@ -19,7 +19,7 @@
},
"$defs": {
"result": {
"description": "A run, a temperature sweep or a comparison, with its receipt.",
"description": "A run, a temperature sweep, a comparison or a replay, with its receipt where it has one.",
"anyOf": [
{
"$ref": "#/$defs/run"
@ -29,6 +29,9 @@
},
{
"$ref": "#/$defs/compare"
},
{
"$ref": "#/$defs/replay"
}
]
},
@ -598,19 +601,39 @@
}
},
"metered": {
"description": "A reading from a meter on the member (extensions/energy-measurement.md); null where there is none.",
"description": "A reading from a meter on the member (extensions/energy-measurement.md; extensions/quantum-landscape.md §4.4); null where there is none or the run has none.",
"anyOf": [
{
"type": "object",
"required": [
"joules_micro",
"class"
"window_us",
"class",
"scope"
],
"properties": {
"joules_micro": {
"type": "integer"
"type": "integer",
"description": "The run's calls' figures added up."
},
"class": {}
"window_us": {
"type": "integer",
"description": "The meter's windows over the run's calls, added up."
},
"class": {
"description": "The acquisition class and the evidence energy-measurement asks of it."
},
"scope": {
"anyOf": [
{
"type": "string"
},
{
"type": "null"
}
],
"description": "What the figure covers where that is more than this run: a whole rail or board, its idle draw and any other work on it in the windows. Null is this run alone."
}
}
},
{
@ -627,6 +650,25 @@
"replay": {
"type": "string",
"description": "BLAKE3 of the protocol, the program, the seed, the plan, the state and its energy."
},
"refusals": {
"type": "array",
"description": "Members a placing host passed over before one answered, each with why (an uncertified kernel, a declined workload, a failed run or an answer the host's checks refused), in the order they were offered; absent where there was none (extensions/quantum-landscape.md §5).",
"items": {
"type": "object",
"required": [
"member",
"why"
],
"properties": {
"member": {
"type": "string"
},
"why": {
"type": "string"
}
}
}
}
}
},
@ -675,7 +717,8 @@
"type": "object",
"required": [
"id",
"class"
"class",
"kernels"
],
"properties": {
"id": {
@ -683,6 +726,34 @@
},
"class": {
"type": "string"
},
"kernels": {
"type": "array",
"description": "Each kernel the job used, in the order first used, with the parity the host found for it.",
"items": {
"type": "object",
"required": [
"kernel",
"parity"
],
"properties": {
"kernel": {
"enum": [
"wai.quantum.landscape/settle",
"wai.quantum.landscape/sample"
]
},
"parity": {
"enum": [
"reference",
"bit-exact",
"statistical",
"uncertified"
],
"description": "What the member has earned for this kernel."
}
}
}
}
}
},
@ -714,7 +785,9 @@
"mean_energy",
"accumulates_per_run",
"priced_per_run",
"left_out_per_run"
"left_out_per_run",
"kernel",
"parity"
],
"properties": {
"temperature": {
@ -765,6 +838,40 @@
"left_out_per_run": {
"type": "string",
"description": "The board's own draw over the fewest clocks one run's sweeps take, which priced_per_run leaves out."
},
"kernel": {
"enum": [
"wai.quantum.landscape/settle",
"wai.quantum.landscape/sample"
]
},
"parity": {
"enum": [
"reference",
"bit-exact",
"statistical",
"uncertified"
],
"description": "What the member has earned for this kernel."
}
}
}
},
"refusals": {
"type": "array",
"description": "Members a placing host passed over before one answered, each with why (an uncertified kernel, a declined workload, a failed run or an answer the host's checks refused), in the order they were offered; absent where there was none (extensions/quantum-landscape.md §5).",
"items": {
"type": "object",
"required": [
"member",
"why"
],
"properties": {
"member": {
"type": "string"
},
"why": {
"type": "string"
}
}
}
@ -813,7 +920,8 @@
"type": "object",
"required": [
"id",
"class"
"class",
"kernels"
],
"properties": {
"id": {
@ -821,6 +929,34 @@
},
"class": {
"type": "string"
},
"kernels": {
"type": "array",
"description": "Each kernel the job used, in the order first used, with the parity the host found for it.",
"items": {
"type": "object",
"required": [
"kernel",
"parity"
],
"properties": {
"kernel": {
"enum": [
"wai.quantum.landscape/settle",
"wai.quantum.landscape/sample"
]
},
"parity": {
"enum": [
"reference",
"bit-exact",
"statistical",
"uncertified"
],
"description": "What the member has earned for this kernel."
}
}
}
}
}
},
@ -856,7 +992,8 @@
"passed_checks",
"best_energy",
"operations_per_run",
"accumulates_per_run"
"accumulates_per_run",
"parity"
],
"properties": {
"solver": {
@ -889,6 +1026,120 @@
},
"accumulates_per_run": {
"type": "integer"
},
"parity": {
"anyOf": [
{
"enum": [
"reference",
"bit-exact",
"statistical",
"uncertified"
],
"description": "What the member has earned for this kernel."
},
{
"type": "null"
}
],
"description": "The member's parity for the sampling kernel; null for a classical method, which ran on no member."
}
}
}
},
"refusals": {
"type": "array",
"description": "Members a placing host passed over before one answered, each with why (an uncertified kernel, a declined workload, a failed run or an answer the host's checks refused), in the order they were offered; absent where there was none (extensions/quantum-landscape.md §5).",
"items": {
"type": "object",
"required": [
"member",
"why"
],
"properties": {
"member": {
"type": "string"
},
"why": {
"type": "string"
}
}
}
}
}
},
"replay": {
"type": "object",
"description": "Many seeds, one line each: the seed, a digest of that run's whole answer, the state, its energy, the verdict, the accumulates and the replay digest.",
"required": [
"protocol",
"label",
"member",
"seeds",
"lines"
],
"properties": {
"protocol": {
"const": "wai.quantum.landscape/1"
},
"label": {
"type": "string"
},
"member": {
"type": "object",
"required": [
"id",
"class",
"kernel",
"parity"
],
"properties": {
"id": {
"type": "string",
"description": "The member that ran it: cpu:reference, gpu:..., fpga:..."
},
"class": {
"type": "string",
"description": "A class of the fabric: cpu, gpu, fpga."
},
"kernel": {
"enum": [
"wai.quantum.landscape/settle",
"wai.quantum.landscape/sample"
]
},
"parity": {
"enum": [
"reference",
"bit-exact",
"statistical",
"uncertified"
],
"description": "What the member has earned for this kernel."
}
}
},
"seeds": {
"type": "integer"
},
"lines": {
"type": "string"
},
"refusals": {
"type": "array",
"description": "Members a placing host passed over before one answered, each with why (an uncertified kernel, a declined workload, a failed run or an answer the host's checks refused), in the order they were offered; absent where there was none (extensions/quantum-landscape.md §5).",
"items": {
"type": "object",
"required": [
"member",
"why"
],
"properties": {
"member": {
"type": "string"
},
"why": {
"type": "string"
}
}
}

1
wai/wai-quantum/src/engine/job.rs
View file

@ -791,6 +791,7 @@ pub(crate) mod tests {
/// What a host seals, its own verifier accepts: a request holding a number
/// no double holds is refused where the request is read, on every path to
/// a seal, with the sentence a run gives.
#[cfg(feature = "quantum_qec")]
#[test]
fn a_host_never_seals_what_its_verifier_refuses() {
const BIG: u64 = 1793907034945123456;

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

@ -624,8 +624,15 @@ pub trait RunMeter {
/// After the last readback: the microjoules in the window and how they
/// were acquired, or `None` when the window cannot be attributed to this
/// run alone.
/// run alone, unless [`RunMeter::scope`] says what else the figure holds.
fn stop(&mut self) -> Result<Option<(u64, EnergyClass)>, String>;
/// What a figure covers, where that is more than this run: an answer says
/// it beside the figure. A meter that attributes its window to the run
/// alone has nothing to add.
fn scope(&self) -> Option<String> {
None
}
}
impl<M: RunMeter + ?Sized> RunMeter for &mut M {
@ -636,6 +643,10 @@ impl<M: RunMeter + ?Sized> RunMeter for &mut M {
fn stop(&mut self) -> Result<Option<(u64, EnergyClass)>, String> {
(**self).stop()
}
fn scope(&self) -> Option<String> {
(**self).scope()
}
}
/// The meter of an unmetered run: it reads nothing.

9
wai/wai-quantum/src/quantum_landscape.rs
View file

@ -89,12 +89,21 @@ pub mod job;
pub mod member;
pub mod price;
pub mod program;
pub mod reach;
pub mod run;
pub mod sample;
pub mod settle;
pub mod table;
pub mod trial;
pub mod truth;
#[cfg(all(any(feature = "quantum_landscape_webgpu", feature = "quantum_landscape_cuda", feature = "quantum_landscape_metal"), not(target_arch = "wasm32")))]
pub mod device;
#[cfg(all(feature = "quantum_landscape_webgpu", not(target_arch = "wasm32")))]
pub mod webgpu;
#[cfg(all(feature = "quantum_landscape_cuda", not(target_arch = "wasm32")))]
pub mod cuda;
#[cfg(all(feature = "quantum_landscape_metal", not(target_arch = "wasm32")))]
pub mod metal;
/// The protocol of a job and its answer (`extensions/quantum-landscape.md`).
pub const PROTOCOL: &str = "wai.quantum.landscape/1";

383
wai/wai-quantum/src/quantum_landscape/conformance.rs
View file

@ -3,18 +3,30 @@
//!
//! - `settle`: whole workloads, the state, the sweeps, the ending and the
//! accumulates a settle kernel gives from a start. A member's settle kernel
//! is `bit-exact` when it gives exactly these.
//! is `bit-exact` when it gives exactly every one it takes, and it takes
//! every one it must ([`MUST_TAKE_FIELD_BOUND`]). They reach past one
//! workgroup and past one batch of sweeps, pin the width of a field at 16, 32
//! and 40 bits, and sit on both sides of the line a member may decline at.
//! - `exp`, `schedule`, `cells`: the reference's exponential, its cooling
//! schedule and its cut points, bit for bit, so that a text-only
//! implementation that computes any of them another way is caught here.
//! - `draws`: the generator itself, at the seeds where a port of it in 32-bit
//! words first goes wrong.
//! - `chains`: whole chains, the reference's bits. A sampling kernel is
//! `bit-exact` when it gives exactly these.
//! - `statistical`: a battery of small workloads, each with a schedule and
//! the exact distribution of the final states; a sampling kernel with draws
//! of its own is `statistical` when it passes every one. They are built to
//! catch a kernel that moves two coupled cells at once in any order of
//! cells, one that ignores its schedule or mishandles the last sweep at 0,
//! and one whose table is wrong for large fields or costs.
//! `bit-exact` when it gives exactly these, as the settle kernel does. Some
//! stop hot, so that the last sweep's draws are in the answer; some are long
//! enough to cross a device's windows, by work and by table, with every
//! window hot; some hold more chains than a workgroup; some have fields that
//! only their full width tells apart, above temperature 0 and at it.
//! - `statistical`: a battery of workloads, each with a schedule and the exact
//! distribution of the final states; a sampling kernel with draws of its own
//! is `statistical` when it passes every one. They are built to catch a kernel
//! that moves two coupled cells at once in any order of cells, one that
//! ignores its schedule, keeps only the end of it or mishandles the last sweep
//! at 0, one whose table is wrong for large fields or costs, one that is
//! wrong past the first four cells of a workload, and one whose fields are
//! held in 16 or 32 bits. A member may decline a workload of it only where
//! it may decline a vector.
//! - `jobs`: whole requests, what their answers hold, refusals included: the
//! worked examples and the rules a reader must follow.
@ -23,7 +35,7 @@ use super::energy::Landscape;
use super::job;
use super::member::{Chain, Member, check_chain, check_settle};
use super::run::state_text;
use super::sample::{anneal_with, chain_seed, cuts, schedule, start};
use super::sample::{anneal_with, chain_seed, cuts, out, schedule, start};
use super::settle::{Outcome, state_hash};
use serde_json::{Value, json};
@ -116,13 +128,90 @@ pub fn settle_cases() -> Vec<SettleCase> {
case("fields-at-the-bound", &wide, &[1, 0, -1, 1], 20);
let cycle = landscape(2, &[(0, 1, 1)], &[(0, 1), (1, -1)], 0);
case("two-cycle-with-fields", &cycle, &[0, 0], 60);
// Past one workgroup of cells and past one batch of sweeps: a wave down a path moves one
// cell a sweep, so every cell's move depends on the one before it.
for n in [65usize, 130, 300] {
case(&format!("path/n={n}"), &path(n), &vec![0; n], 2 * n as u32);
}
case("path-stopped-in-a-later-batch", &path(300), &[0; 300], 130);
case("random/n=1000", &random(1000, 6, 0x5E77), &pattern(0x51DE, 1000), 300);
let mut ring = Landscape::new(70);
for i in 0..70 {
ring.couple(i, (i + 1) % 70, -1);
}
case("a-two-cycle-past-a-workgroup", &ring, &vec![1; 70], 200);
for (name, l) in width_cases() {
let start = vec![0; l.len()];
case(&name, &l, &start, 20);
}
out
}
/// Whether a member's settle kernel gives every `settle` vector exactly.
/// A wave along a path: a field on the first cell, every cell agreeing with the next.
fn path(n: usize) -> Landscape {
let mut l = Landscape::new(n);
for i in 0..n - 1 {
l.couple(i, i + 1, 1);
}
l.set_field(0, 3);
l
}
/// `n` cells with about `degree` couplings each, of both signs.
fn random(n: usize, degree: usize, seed: u64) -> Landscape {
let mut l = Landscape::new(n);
for k in 0..(n * degree / 2) as u64 {
let (i, j) = ((out(seed, 2 * k + 1) % n as u64) as usize, (out(seed, 2 * k + 2) % n as u64) as usize);
if i != j {
l.couple(i, j, if out(seed ^ 1, k + 1) & 1 == 0 { 1 } else { -1 });
}
}
l
}
/// Workloads whose fields only their full width tells apart, on both sides of
/// the line a member may decline at: `(name, landscape)`.
fn width_cases() -> Vec<(String, Landscape)> {
const LANE: i64 = i32::MAX as i64;
let star = |b: i64, cost: i64| landscape(3, &[(0, 1, 1), (0, 2, 1)], &[(0, b), (1, b.signum()), (2, b.signum())], cost);
vec![
// 65,536 is 0 in 16 bits, 40,000 negative, and −70,000 is inside the cost.
("fields-past-16-bits".into(), landscape(4, &[], &[(0, 65_536), (1, 40_000), (2, -70_000), (3, 1)], 5_000)),
// 2³³ is 0 in 32 bits, 3·2³¹ negative, and −(2³³ + 5) is inside the cost.
("fields-past-32-bits".into(), landscape(4, &[], &[(0, 1 << 33), (1, 3 << 31), (2, -((1 << 33) + 5)), (3, 1)], 7)),
// At the line: a field bound of exactly 2³¹ − 1 is one a member must take.
("field-bound-at-the-line".into(), star(LANE - 2, 0)),
("field-bound-at-the-line-below-zero".into(), star(-(LANE - 2), 0)),
// Its largest field is 2³¹ − 1 and its field bound 2³¹ + 1: a member that checks
// the fields and forgets the couplings overflows here.
("field-bound-over-the-line-by-its-couplings".into(), star(LANE, 0)),
("field-bound-over-the-line-below-zero".into(), star(-LANE, 0)),
// A cost past any lane on small fields: the cost holds every cell, at any width.
("a-cost-past-the-lanes".into(), landscape(3, &[(0, 1, 1), (1, 2, -1)], &[(0, 3)], 1 << 33)),
("a-cost-just-past-the-lanes".into(), landscape(3, &[(0, 1, 1), (1, 2, -1)], &[(0, LANE - 1)], LANE + 1)),
]
}
/// The widest field bound a member must take: every vector at or below it. The
/// bound is [`super::settle::Workload::field_bound`], the largest field any
/// state can give a cell, couplings included.
pub const MUST_TAKE_FIELD_BOUND: i64 = i32::MAX as i64;
/// Whether a member declines a vector it may decline, before running it: one
/// whose field bound is above [`MUST_TAKE_FIELD_BOUND`]. `None` when it takes
/// the vector, `Some(true)` when it may skip it, `Some(false)` when it must not.
fn declines(m: &dyn Member, w: &super::settle::Workload, sampling: bool) -> Option<bool> {
m.takes(w, sampling).err().map(|_| w.field_bound() > MUST_TAKE_FIELD_BOUND)
}
/// Whether a member's settle kernel gives every `settle` vector it takes
/// exactly, and takes every one it must.
pub fn settle_vectors_hold(m: &dyn Member) -> bool {
settle_cases().iter().all(|c| {
let w = c.landscape.lower();
if let Some(may) = declines(m, &w, false) {
return may;
}
let want = w.settle(&c.start, c.bound);
m.settle(&w, &c.start, c.bound).is_ok_and(|o| check_settle(&o, &w, c.bound).is_ok() && o == want)
})
@ -136,6 +225,7 @@ fn settle_json() -> Vec<Value> {
json!({
"name": c.name,
"landscape": landscape_json(&c.landscape),
"field_bound": c.landscape.lower().field_bound(),
"start": state_text(&c.start),
"bound": c.bound,
"expect": { "state": state_text(&o.state), "sweeps": o.sweeps, "term": o.term.label(), "accumulates": o.accumulates, "state_hash": format!("{:016x}", state_hash(&o.state)) },
@ -189,12 +279,31 @@ fn schedule_json() -> Vec<Value> {
// ── chains ────────────────────────────────────────────────────────────────────────────────
/// The temperatures of a chain's sweeps.
#[derive(Clone, Debug, PartialEq)]
pub enum Temperatures {
/// The schedule of a run: cooling in a line from `t0` to 0.
Cooling { t0: f64, sweeps: u32 },
/// Runs of sweeps at one temperature, in order, which need not end at 0:
/// a member is given every sweep's temperature, so a vector can hold a
/// chain hot to its last sweep, where its draws are in the final state.
Held(Vec<(f64, u32)>),
}
impl Temperatures {
pub fn each(&self) -> Vec<f64> {
match self {
Temperatures::Cooling { t0, sweeps } => schedule(*t0, *sweeps),
Temperatures::Held(runs) => runs.iter().flat_map(|&(t, k)| std::iter::repeat_n(t, k as usize)).collect(),
}
}
}
/// Chains the reference runs: every start of a run with this seed.
pub struct ChainCase {
pub name: String,
pub landscape: Landscape,
pub t0: f64,
pub sweeps: u32,
pub temperatures: Temperatures,
pub seed: u64,
pub restarts: u64,
}
@ -210,6 +319,13 @@ impl ChainCase {
}
}
/// Seeds whose first draw, for one cell at field 0 and cost 0, a port of the
/// generator in 32-bit words gets wrong: found by search, and held by a test.
/// The first needs the carry out of the middle of a 32 by 32 product; the other
/// two fall in the same high word as a cut point, so that only a compare of all
/// 53 bits puts them on the right side of it.
pub const DRAW_EDGES: [(&str, u64); 3] = [("the-carry-in-a-product", 1982), ("under-the-first-cut-by-its-low-word", 1_174_725), ("under-the-second-cut-by-its-low-word", 6_472_161)];
pub fn chain_cases() -> Vec<ChainCase> {
let mixed = landscape(8, &[(0, 1, -1), (1, 2, -1), (0, 2, -1), (3, 4, 1), (4, 5, 1), (3, 5, 1), (5, 6, -1), (3, 6, -1), (6, 7, 1)], &[], 1);
let dense = stored(&pattern(0xBEEF, 16));
@ -217,17 +333,49 @@ pub fn chain_cases() -> Vec<ChainCase> {
let mut out = Vec::new();
for (name, l) in [("sparse-with-a-cost", &mixed), ("dense", &dense), ("wide-fields", &wide)] {
for (t0, sweeps, seed, restarts) in [(2.0f64, 30u32, 1u64, 1u64), (0.5, 12, 2, 1), (0.0, 7, 3, 1), (1.8708286933869707, 41, 4, 8), (13.7, 9, 5, 3)] {
out.push(ChainCase { name: format!("{name}/t0={t0}/sweeps={sweeps}"), landscape: l.clone(), t0, sweeps, seed, restarts });
out.push(ChainCase { name: format!("{name}/t0={t0}/sweeps={sweeps}"), landscape: l.clone(), temperatures: Temperatures::Cooling { t0, sweeps }, seed, restarts });
}
}
let held = |runs: &[(f64, u32)]| Temperatures::Held(runs.to_vec());
// The draws, held to the last sweep: at temperature 1 with no field a cell takes each
// state a third of the time, so a wrong draw shows in the state at once.
for (name, seed) in DRAW_EDGES {
out.push(ChainCase { name: format!("a-draw/{name}"), landscape: Landscape::new(1), temperatures: held(&[(1.0, 1)]), seed, restarts: 1 });
}
out.push(ChainCase { name: "one-sweep-of-16384-draws".into(), landscape: Landscape::new(128), temperatures: held(&[(1.0, 1)]), seed: 6, restarts: 128 });
// More chains than a workgroup holds, and hot to the end.
out.push(ChainCase { name: "more-chains-than-a-workgroup".into(), landscape: wide.clone(), temperatures: held(&[(13.7, 3)]), seed: 7, restarts: 130 });
// In the reference's order or not: a chain that stops hot shows the order of its cells.
out.push(ChainCase { name: "the-order-of-the-cells".into(), landscape: mixed.clone(), temperatures: held(&[(1.5, 2)]), seed: 8, restarts: 16 });
// Windows of sweeps, by table: every cell's reach is its own, so 320 entries a sweep
// and about fifty sweeps a window. Hot from start to end, with sweeps at 0 inside.
let mut spread = Landscape::new(64);
for i in 0..64 {
spread.couple(i, (i + 1) % 64, 1);
spread.set_field(i, 7 * i as i64 - 224);
}
out.push(ChainCase { name: "windows-by-their-table".into(), landscape: spread, temperatures: held(&[(400.0, 70), (0.0, 2), (250.0, 70), (150.0, 60)]), seed: 9, restarts: 4 });
// Windows of sweeps, by work: about 170 sweeps a window, couplings of both signs.
out.push(ChainCase { name: "windows-by-their-work".into(), landscape: random(1024, 4, 0x71D0), temperatures: held(&[(3.0, 200), (1.5, 200)]), seed: 10, restarts: 2 });
// Fields only their full width tells apart: at temperature 0, and at a temperature
// of their own size, so that every cell's field is in its draw.
for (name, l) in width_cases() {
let t = l.field().iter().map(|b| b.unsigned_abs()).max().unwrap_or(1).max(1) as f64;
out.push(ChainCase { name: format!("{name}/at-0"), landscape: l.clone(), temperatures: Temperatures::Cooling { t0: 0.0, sweeps: 5 }, seed: 11, restarts: 4 });
out.push(ChainCase { name: format!("{name}/hot"), landscape: l, temperatures: held(&[(t, 4)]), seed: 12, restarts: 8 });
}
out
}
/// Whether a member's sampling kernel gives every `chains` vector exactly, every case in one call.
/// Whether a member's sampling kernel gives every `chains` vector it takes
/// exactly, every case in one call, and takes every one it must.
pub fn chain_vectors_hold(m: &dyn Member) -> bool {
chain_cases().iter().all(|c| {
let w = c.landscape.lower();
let temperatures = schedule(c.t0, c.sweeps);
let temperatures = c.temperatures.each();
if let Some(may) = declines(m, &w, temperatures.iter().any(|&t| t > 0.0)) {
return may;
}
let chains = c.chains();
let want: Vec<Outcome> = chains.iter().map(|ch| anneal_with(&w, &ch.start, &temperatures, ch.seed)).collect();
m.chains(&w, &chains, &temperatures).is_ok_and(|got| got.len() == want.len() && got.iter().all(|o| check_chain(o, &w, &temperatures).is_ok()) && got == want)
@ -239,23 +387,38 @@ fn chains_json() -> Vec<Value> {
.iter()
.map(|c| {
let w = c.landscape.lower();
let temperatures = schedule(c.t0, c.sweeps);
json!({
let temperatures = c.temperatures.each();
let mut v = json!({
"name": c.name,
"landscape": landscape_json(&c.landscape),
"field_bound": w.field_bound(),
"seed": c.seed.to_string(),
"t0": c.t0,
"t0_bits": bits(c.t0),
"sweeps": c.sweeps,
"chains": c.chains().iter().map(|ch| {
let o = anneal_with(&w, &ch.start, &temperatures, ch.seed);
json!({ "chain": ch.seed.to_string(), "start": state_text(&ch.start), "expect": { "state": state_text(&o.state), "accumulates": o.accumulates, "draws": o.draws, "term": o.term.label(), "state_hash": format!("{:016x}", state_hash(&o.state)) } })
}).collect::<Vec<_>>(),
})
});
match &c.temperatures {
Temperatures::Cooling { t0, sweeps } => {
v["t0"] = json!(t0);
v["t0_bits"] = json!(bits(*t0));
v["sweeps"] = json!(sweeps);
}
Temperatures::Held(runs) => v["held"] = runs.iter().map(|&(t, k)| json!({ "temperature": t, "temperature_bits": bits(t), "sweeps": k })).collect(),
}
v
})
.collect()
}
/// The generator's outputs: the first few of several seeds, and every seed of
/// [`DRAW_EDGES`] with the draw its case turns on.
fn draws_json() -> Vec<Value> {
let mut at: Vec<(u64, u64)> = [0u64, 1, 2, u64::MAX, 0x9E37_79B9_7F4A_7C15].iter().flat_map(|&s| (1..=4).map(move |j| (s, j))).collect();
at.extend(DRAW_EDGES.iter().map(|&(_, seed)| (chain_seed(seed, 0), 1)));
at.iter().map(|&(s, j)| json!({ "seed": s.to_string(), "j": j, "out": out(s, j).to_string() })).collect()
}
// ── the statistical battery ───────────────────────────────────────────────────────────────
/// How many chains each battery workload runs, and the test its final states are held to.
@ -265,25 +428,66 @@ pub const BATTERY_POOL_BELOW: f64 = 5.0;
/// The seed battery workload `k` draws its chains from: `chain_seed(BATTERY_SEED + k, r)`.
pub const BATTERY_SEED: u64 = 0xBA77_E124;
/// One workload of the battery: a landscape, a schedule, and the temperature
/// the schedule holds before its end, at which the chains have mixed.
/// What a battery workload's final states are held to.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub enum Law {
/// Boltzmann at the schedule's last temperature above 0, reached on a
/// plateau; carried through one deterministic sweep where the schedule ends at 0.
Settled,
/// The law of the walk itself, from the start's uniform law through every
/// sweep of the schedule, on a block whose cells are all alike. Which order a
/// sweep takes the cells in only renames them, so the law is read over how
/// many cells hold each state, and any order of cells passes.
Walked,
}
/// One workload of the battery: a landscape, a schedule, what its final states
/// are held to, and how many cells a block of it has. A landscape of several
/// blocks holds copies of one block, none coupled to another, each held to the
/// block's law by itself.
pub struct BatteryCase {
pub name: &'static str,
pub landscape: Landscape,
pub temperatures: Vec<f64>,
pub law: Law,
pub block: usize,
}
impl BatteryCase {
fn of(name: &'static str, landscape: Landscape, temperatures: Vec<f64>) -> BatteryCase {
let block = landscape.len();
BatteryCase { name, landscape, temperatures, law: Law::Settled, block }
}
/// The temperature the schedule settles at: its last sweep above 0.
pub fn equilibrium(&self) -> f64 {
self.temperatures.iter().rev().copied().find(|&t| t > 0.0).unwrap_or(0.0)
}
/// The exact distribution of the final state, over the `3^n` states in the
/// order `Σ (s_i + 1) · 3^i`: Boltzmann at [`Self::equilibrium`], and if the
/// schedule ends at 0, that distribution carried through one deterministic sweep.
/// The first block of the landscape by itself.
fn unit(&self) -> Landscape {
let mut l = Landscape::new(self.block);
for i in 0..self.block {
for j in (i + 1)..self.block {
l.couple(i, j, self.landscape.coupling(i, j));
}
l.set_field(i, self.landscape.field()[i]);
}
l.set_cost(self.landscape.cost());
l
}
/// The exact distribution of a block's final state, over its `3^block`
/// states in the order `Σ (s_i + 1) · 3^i`.
pub fn exact(&self) -> Vec<f64> {
let l = &self.landscape;
match self.law {
Law::Settled => self.settled(),
Law::Walked => self.walked(),
}
}
fn settled(&self) -> Vec<f64> {
let l = self.unit();
let n = l.len();
let all: Vec<Vec<i8>> = (0..3usize.pow(n as u32)).map(|k| state_of(k, n)).collect();
let energies: Vec<i64> = all.iter().map(|s| l.energy(s)).collect();
@ -304,6 +508,34 @@ impl BatteryCase {
}
}
/// Every sweep of the schedule, cell by cell in order, on the whole law at once.
fn walked(&self) -> Vec<f64> {
let l = self.unit();
let (n, w) = (l.len(), l.lower());
let states = 3usize.pow(n as u32);
let mut p = vec![1.0 / states as f64; states];
for &t in &self.temperatures {
for i in 0..n {
let mut q = vec![0.0; states];
for (k, &pk) in p.iter().enumerate() {
let mut s = state_of(k, n);
let h = w.field(&s, i) as f64;
let c = w.cost() as f64;
let weights = [(-h - c) / t, 0.0, (h - c) / t];
let top = weights.iter().copied().fold(f64::MIN, f64::max);
let e = weights.map(|x| crate::repro::exp(x - top));
let z = e[0] + e[1] + e[2];
for (v, ev) in [-1i8, 0, 1].into_iter().zip(e) {
s[i] = v;
q[index_of(&s)] += pk * ev / z;
}
}
p = q;
}
}
p
}
pub fn chains(&self, k: usize) -> Vec<Chain> {
(0..u64::from(BATTERY_CHAINS))
.map(|r| {
@ -334,19 +566,46 @@ pub fn battery() -> Vec<BatteryCase> {
cooling.extend(vec![0.7; 24]);
let mut quench = vec![1.3; 31];
quench.push(0.0);
// A triangle pulled one way: warm, it lies mostly that way; cold, a chain cannot leave
// the side it is on in the sweeps it has. So the end of the schedule alone gives another law.
let alike = landscape(3, &[(0, 1, 1), (1, 2, 1), (0, 2, 1)], &[(0, 1), (1, 1), (2, 1)], 0);
let mut two_stages = vec![1.2; 16];
two_stages.extend(vec![0.2; 6]);
// Past the first cells of a workload: eight triangles side by side, each held to its law.
let (blocks, side) = (8usize, 3usize);
let mut triangles = Landscape::new(blocks * side);
for b in 0..blocks {
let at = b * side;
triangles.couple(at, at + 1, 1).couple(at + 1, at + 2, 1).couple(at, at + 2, 1);
triangles.set_field(at, 1);
}
triangles.set_cost(1);
let mut ring = Landscape::new(6);
for i in 0..6 {
ring.couple(i, (i + 1) % 6, if i % 3 == 0 { -1 } else { 1 });
}
ring.set_field(0, 1).set_field(3, -1).set_cost(1);
vec![
BatteryCase { name: "one-pair", landscape: landscape(2, &[(0, 1, 1)], &[], 0), temperatures: constant(1.0, 32) },
BatteryCase { name: "a-pair-apart", landscape: landscape(3, &[(0, 2, 1)], &[(1, 1)], 0), temperatures: constant(1.0, 32) },
BatteryCase { name: "a-pair-far-apart", landscape: landscape(4, &[(0, 3, -1), (1, 2, 1)], &[(1, -1)], 0), temperatures: constant(0.8, 32) },
BatteryCase { name: "a-triangle", landscape: landscape(3, &[(0, 1, 1), (1, 2, 1), (0, 2, 1)], &[], 1), temperatures: constant(1.2, 32) },
BatteryCase { name: "wide-fields-and-cost", landscape: landscape(3, &[(0, 1, 1), (1, 2, -1)], &[(0, 17), (1, -9), (2, 40)], 8), temperatures: constant(13.7, 32) },
BatteryCase { name: "a-small-field-cold", landscape: landscape(2, &[(0, 1, 1)], &[(0, 5), (1, -3)], 2), temperatures: constant(0.45, 32) },
BatteryCase { name: "ending-at-zero", landscape: landscape(4, &[(0, 1, 1), (1, 2, 1), (2, 3, -1)], &[(0, 1)], 1), temperatures: quench },
BatteryCase { name: "cooling", landscape: landscape(3, &[(0, 1, -1), (1, 2, -1), (0, 2, -1)], &[(2, -2)], 0), temperatures: cooling },
BatteryCase::of("one-pair", landscape(2, &[(0, 1, 1)], &[], 0), constant(1.0, 32)),
BatteryCase::of("a-pair-apart", landscape(3, &[(0, 2, 1)], &[(1, 1)], 0), constant(1.0, 32)),
BatteryCase::of("a-pair-far-apart", landscape(4, &[(0, 3, -1), (1, 2, 1)], &[(1, -1)], 0), constant(0.8, 32)),
BatteryCase::of("a-triangle", landscape(3, &[(0, 1, 1), (1, 2, 1), (0, 2, 1)], &[], 1), constant(1.2, 32)),
BatteryCase::of("wide-fields-and-cost", landscape(3, &[(0, 1, 1), (1, 2, -1)], &[(0, 17), (1, -9), (2, 40)], 8), constant(13.7, 32)),
BatteryCase::of("a-small-field-cold", landscape(2, &[(0, 1, 1)], &[(0, 5), (1, -3)], 2), constant(0.45, 32)),
BatteryCase::of("ending-at-zero", landscape(4, &[(0, 1, 1), (1, 2, 1), (2, 3, -1)], &[(0, 1)], 1), quench),
BatteryCase::of("cooling", landscape(3, &[(0, 1, -1), (1, 2, -1), (0, 2, -1)], &[(2, -2)], 0), cooling),
BatteryCase::of("a-ring-of-six", ring, constant(1.1, 32)),
BatteryCase { name: "triangles-side-by-side", landscape: triangles, temperatures: constant(1.2, 32), law: Law::Settled, block: side },
BatteryCase { name: "a-schedule-that-matters", landscape: alike, temperatures: two_stages, law: Law::Walked, block: 3 },
// Fields that only their full width tells apart, at a temperature of their own size:
// past 16 bits, which every member takes, and past 32, which a member may decline.
BatteryCase::of("fields-past-16-bits", landscape(3, &[(0, 1, 1), (1, 2, -1)], &[(0, 70_000), (1, -65_536), (2, 40_000)], 30_000), constant(40_000.0, 32)),
BatteryCase::of("fields-past-32-bits", landscape(3, &[(0, 1, 1), (1, 2, -1)], &[(0, 3 << 31), (1, -(1 << 33)), (2, (1 << 32) + 5)], 1 << 32), constant(4_294_967_296.0, 32)),
]
}
/// What one workload of the battery found.
/// What one workload of the battery found: over its blocks, the one furthest
/// from its law.
#[derive(Clone, Debug, PartialEq)]
pub struct BatteryResult {
pub name: &'static str,
@ -354,6 +613,9 @@ pub struct BatteryResult {
pub degrees: usize,
pub critical: f64,
pub passed: bool,
/// The member declined the workload, which it may where the field bound is above
/// [`MUST_TAKE_FIELD_BOUND`], as for a vector; it then passes without running.
pub declined: bool,
}
/// Pearson's statistic of `counts` against `expected` (probabilities) over
@ -382,30 +644,60 @@ pub fn chi_square(counts: &[u64], expected: &[f64], total: u64) -> (f64, usize,
(stat, df, df as f64 + 5.0 * (2.0 * df as f64).sqrt())
}
/// Run the battery on a member: every workload's chains, their final states
/// against the exact distribution. A member that fails to answer, or whose
/// answers fail [`check_chain`], fails that workload.
/// Bins as a block's law reads them: every state, or for [`Law::Walked`] each
/// count of cells at −1, 0 and +1.
fn binned(law: Law, n: usize, by_state: &[f64]) -> Vec<f64> {
if law == Law::Settled {
return by_state.to_vec();
}
let mut bins = std::collections::BTreeMap::new();
for (k, &v) in by_state.iter().enumerate() {
let mut s = state_of(k, n);
s.sort_unstable();
*bins.entry(s).or_insert(0.0) += v;
}
bins.into_values().collect()
}
/// Run the battery on a member: every workload's chains, the final states of
/// each block against the exact distribution. A member that fails to answer, or
/// whose answers fail [`check_chain`], fails that workload.
pub fn run_battery(m: &dyn Member) -> Vec<BatteryResult> {
battery()
.iter()
.enumerate()
.map(|(k, case)| {
let w = case.landscape.lower();
let expected = case.exact();
let mut counts = vec![0u64; expected.len()];
if let Some(may) = declines(m, &w, true) {
return BatteryResult { name: case.name, chi_square: 0.0, degrees: 0, critical: 0.0, passed: may, declined: true };
}
let exact = case.exact();
let blocks = case.landscape.len() / case.block;
let mut counts = vec![vec![0u64; exact.len()]; blocks];
let mut ok = true;
for chunk in case.chains(k).chunks(4096) {
match m.chains(&w, chunk, &case.temperatures) {
Ok(got) if got.len() == chunk.len() && got.iter().all(|o| check_chain(o, &w, &case.temperatures).is_ok()) => {
for o in got {
counts[index_of(&o.state)] += 1;
for (b, c) in counts.iter_mut().enumerate() {
c[index_of(&o.state[b * case.block..(b + 1) * case.block])] += 1;
}
}
}
_ => ok = false,
}
}
let (chi, df, critical) = chi_square(&counts, &expected, u64::from(BATTERY_CHAINS));
BatteryResult { name: case.name, chi_square: chi, degrees: df, critical, passed: ok && chi <= critical }
let expected = binned(case.law, case.block, &exact);
let worst = counts
.iter()
.map(|c| {
let as_real: Vec<f64> = c.iter().map(|&v| v as f64).collect();
let got: Vec<u64> = binned(case.law, case.block, &as_real).iter().map(|&v| v as u64).collect();
chi_square(&got, &expected, u64::from(BATTERY_CHAINS))
})
.fold((0.0f64, 1usize, 0.0f64), |w, r| if r.0 - r.2 > w.0 - w.2 || w.2 == 0.0 { r } else { w });
let (chi, df, critical) = worst;
BatteryResult { name: case.name, chi_square: chi, degrees: df, critical, passed: ok && chi <= critical, declined: false }
})
.collect()
}
@ -420,6 +712,8 @@ fn battery_json() -> Value {
"name": c.name,
"landscape": landscape_json(&c.landscape),
"temperatures": c.temperatures,
"law": if c.law == Law::Walked { "walked" } else { "settled" },
"block": c.block,
"state_order": "the sum of (s_i + 1) · 3^i",
"exact": c.exact(),
})).collect::<Vec<_>>(),
@ -476,6 +770,7 @@ pub fn vectors() -> Value {
"exp": exp_arguments().iter().map(|&x| json!([bits(x), bits(crate::repro::exp(x))])).collect::<Vec<_>>(),
"schedule": schedule_json(),
"cells": cells_json(),
"draws": draws_json(),
"chains": chains_json(),
"statistical": battery_json(),
"jobs": jobs_json(),

311
wai/wai-quantum/src/quantum_landscape/cuda.rs Normal file
View file

@ -0,0 +1,311 @@
//! A CUDA member of the landscape fabric: the kernels of [`SOURCE`] on an
//! NVIDIA GPU, compiled at run time for the device's architecture.
//!
//! The kernels are those of `landscape.wgsl` with native 64-bit words: the
//! settle kernel one thread a cell, one launch a sweep and one to advance; the
//! sampling kernel one thread a chain, its cells in order, the reference's
//! draws and the host's cut points. Fields are 64-bit, so it takes every
//! workload a job can give. Every launch is on one stream, so each sees what
//! the one before it wrote.
//!
//! The driver and the runtime compiler are loaded when first asked for, so a
//! build with this feature runs on a host with no CUDA and finds no member.
//! The member says nothing of its own parity: the host certifies it.
use super::device::{self, MeterSlot, SETTLE_BATCH};
use super::member::{Chain, Member, Metered};
use super::settle::{Outcome, Workload};
use crate::quantum_fabric::RunMeter;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use cudarc::nvrtc::{CompileOptions, compile_ptx_with_opts};
use std::sync::Arc;
/// The kernels, in CUDA C.
pub const SOURCE: &str = r#"
typedef unsigned long long u64;
typedef long long i64;
__device__ __forceinline__ int threshold(i64 h, i64 c) {
return h > c ? 1 : (h < -c ? -1 : 0);
}
__device__ __forceinline__ i64 field_of(const unsigned* rows, const unsigned* cols, const int* vals, const i64* fields, const int* s, unsigned i) {
i64 h = fields[i];
for (unsigned k = rows[i]; k < rows[i + 1]; ++k) {
h += (i64)vals[k] * (i64)s[cols[k]];
}
return h;
}
extern "C" __global__ void settle_sweep(const unsigned* rows, const unsigned* cols, const int* vals, const i64* fields, i64 cost, unsigned n, int* ring, unsigned* status) {
unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
volatile unsigned* st = status;
if (i >= n || st[1] != 0u) return;
unsigned done = st[0];
unsigned now = (done % 3u) * n, next = ((done + 1u) % 3u) * n, back = ((done + 2u) % 3u) * n;
int s = threshold(field_of(rows, cols, vals, fields, ring + now, i), cost);
ring[next + i] = s;
if (s != ring[now + i]) atomicOr(&status[2], 1u);
if (s != ring[back + i]) atomicOr(&status[3], 1u);
}
extern "C" __global__ void settle_advance(unsigned bound, unsigned* status) {
if (blockIdx.x != 0 || threadIdx.x != 0 || status[1] != 0u) return;
unsigned k = status[0] + 1u;
status[0] = k;
if (status[2] == 0u) status[1] = 1u;
else if (k >= 2u && status[3] == 0u) status[1] = 2u;
else if (k >= bound) status[1] = 3u;
status[2] = 0u;
status[3] = 0u;
}
__device__ __forceinline__ u64 splitmix(u64 s, u64 j) {
u64 x = s + j * 0x9E3779B97F4A7C15ull;
x ^= x >> 30; x *= 0xBF58476D1CE4E5B9ull;
x ^= x >> 27; x *= 0x94D049BB133111EBull;
return x ^ (x >> 31);
}
extern "C" __global__ void chains_in_order(const unsigned* rows, const unsigned* cols, const int* vals, const i64* fields, i64 cost, unsigned n, unsigned chains, const i64* offsets, unsigned k0, unsigned k1, const u64* seeds, const unsigned* sweeps, const u64* table, int* cells) {
unsigned c = blockIdx.x * blockDim.x + threadIdx.x;
if (c >= chains) return;
int* s = cells + (u64)c * n;
int* spare = cells + (u64)chains * n + (u64)c * n;
u64 seed = seeds[c];
for (unsigned k = k0; k < k1; ++k) {
unsigned kind = sweeps[k - k0];
if (kind == 0u) {
for (unsigned i = 0; i < n; ++i) spare[i] = threshold(field_of(rows, cols, vals, fields, s, i), cost);
for (unsigned i = 0; i < n; ++i) s[i] = spare[i];
} else {
for (unsigned i = 0; i < n; ++i) {
i64 h = field_of(rows, cols, vals, fields, s, i);
u64 u = splitmix(seed, (u64)k * n + i + 1) >> 11;
u64 e = (u64)(kind - 1u) + (u64)h + (u64)offsets[i];
s[i] = u < table[2 * e] ? -1 : (u < table[2 * e + 1] ? 0 : 1);
}
}
}
}
"#;
/// Threads a block.
const BLOCK: u32 = 256;
fn launch(work: u64) -> LaunchConfig {
LaunchConfig { grid_dim: (work.div_ceil(u64::from(BLOCK)).max(1) as u32, 1, 1), block_dim: (BLOCK, 1, 1), shared_mem_bytes: 0 }
}
struct Kernels {
settle_sweep: CudaFunction,
settle_advance: CudaFunction,
chains: CudaFunction,
}
/// One CUDA device as a landscape member.
pub struct CudaLandscape {
_ctx: Arc<CudaContext>,
stream: Arc<CudaStream>,
kernels: Kernels,
id: String,
name: String,
meter: MeterSlot,
}
impl std::fmt::Debug for CudaLandscape {
fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
f.debug_struct("CudaLandscape").field("id", &self.id).field("name", &self.name).finish()
}
}
fn driver(e: cudarc::driver::DriverError) -> String {
format!("the CUDA driver: {e}")
}
/// Every device the driver lists, opened and compiled for, each a member or
/// the reason it is not one. Nothing, with the reason, where there is no CUDA.
pub fn discover() -> Vec<Result<CudaLandscape, String>> {
discover_with(SOURCE)
}
/// [`discover`] with kernels other than [`SOURCE`]: how a test shows that a
/// fault in the kernels loses the grade it would otherwise earn on this device.
#[doc(hidden)]
pub fn discover_with(source: &str) -> Vec<Result<CudaLandscape, String>> {
if let Err(why) = crate::quantum_fabric_cuda::available() {
return vec![Err(why)];
}
match CudaContext::device_count() {
Ok(count) => (0..usize::try_from(count).unwrap_or(0)).map(|k| open(k, source)).collect(),
Err(e) => vec![Err(driver(e))],
}
}
fn open(ordinal: usize, source: &str) -> Result<CudaLandscape, String> {
let ctx = CudaContext::new(ordinal).map_err(driver)?;
let name = ctx.name().map_err(driver)?;
let (major, minor) = ctx.compute_capability().map_err(driver)?;
let uuid: String = ctx.uuid().map_err(driver)?.bytes.iter().fold(String::new(), |mut s, b| {
s.push_str(&format!("{:02x}", *b as u8));
s
});
let arch: &'static str = Box::leak(format!("compute_{major}{minor}").into_boxed_str());
let opts = CompileOptions { arch: Some(arch), options: vec!["--std=c++17".into(), "--fmad=false".into()], ..Default::default() };
let ptx = compile_ptx_with_opts(source, opts).map_err(|e| format!("{name}: the kernels did not compile: {e}"))?;
let module = ctx.load_module(ptx).map_err(driver)?;
let kernels = Kernels {
settle_sweep: module.load_function("settle_sweep").map_err(driver)?,
settle_advance: module.load_function("settle_advance").map_err(driver)?,
chains: module.load_function("chains_in_order").map_err(driver)?,
};
Ok(CudaLandscape { stream: ctx.default_stream(), _ctx: ctx, kernels, id: device::id_of("gpu", "cuda", &name, &format!("{uuid}/sm_{major}{minor}")), name, meter: MeterSlot::default() })
}
impl CudaLandscape {
/// Read a meter around each call: its reading stands beside the priced figures.
pub fn with_meter(mut self, meter: Box<dyn RunMeter + Send>) -> CudaLandscape {
self.meter = MeterSlot::new(meter);
self
}
/// The device's name, as the driver reports it.
pub fn name(&self) -> &str {
&self.name
}
fn settle_on_device(&self, w: &Workload, start: &[i8], bound: u32) -> Result<Outcome, String> {
let n = w.n();
let r = device::rows(w);
let s = &self.stream;
let (rows, cols, vals, fields) = (s.clone_htod(&r.rows).map_err(driver)?, s.clone_htod(&nonempty(&r.cols)).map_err(driver)?, s.clone_htod(&nonempty(&r.vals)).map_err(driver)?, s.clone_htod(&r.fields).map_err(driver)?);
let mut init = vec![0i32; 3 * n];
for (i, &v) in start.iter().enumerate() {
init[i] = i32::from(v);
}
let mut ring = s.clone_htod(&init).map_err(driver)?;
let mut status = s.alloc_zeros::<u32>(4).map_err(driver)?;
let (cost, n32) = (w.cost(), n as u32);
let (mut done, mut ended) = (0u32, 0u32);
while ended == 0 && done < bound {
for _ in 0..SETTLE_BATCH.min(bound - done) {
let mut a = s.launch_builder(&self.kernels.settle_sweep);
a.arg(&rows).arg(&cols).arg(&vals).arg(&fields).arg(&cost).arg(&n32).arg(&mut ring).arg(&mut status);
// SAFETY: the buffers are sized as the kernel reads them: n + 1 rows, the couplings,
// n fields, 3n cells of ring and four status words.
unsafe { a.launch(launch(n as u64)) }.map_err(driver)?;
let mut b = s.launch_builder(&self.kernels.settle_advance);
b.arg(&bound).arg(&mut status);
// SAFETY: four status words.
unsafe { b.launch(launch(1)) }.map_err(driver)?;
}
let st = s.clone_dtoh(&status).map_err(driver)?;
(done, ended) = (st[0], st[1]);
}
let all = s.clone_dtoh(&ring).map_err(driver)?;
let third = (done % 3) as usize * n;
Ok(device::settled(w, done, ended, all[third..third + n].iter().map(|&v| v as i8).collect()))
}
fn chains_on_device(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String> {
let sampling = temperatures.iter().any(|&t| t > 0.0);
let reach = device::reach(w);
if sampling {
device::table_fits(&reach)?;
}
let (n, count) = (w.n(), chains.len());
if count == 0 {
return Ok(Vec::new());
}
let r = device::rows(w);
let s = &self.stream;
let (rows, cols, vals, fields) = (s.clone_htod(&r.rows).map_err(driver)?, s.clone_htod(&nonempty(&r.cols)).map_err(driver)?, s.clone_htod(&nonempty(&r.vals)).map_err(driver)?, s.clone_htod(&r.fields).map_err(driver)?);
let mut init = vec![0i32; 2 * count * n];
for (c, chain) in chains.iter().enumerate() {
if chain.start.len() != n {
return Err("a chain's start is not one cell a variable".into());
}
for (i, &v) in chain.start.iter().enumerate() {
init[c * n + i] = i32::from(v);
}
}
let mut cells = s.clone_htod(&init).map_err(driver)?;
drop(init);
let seeds = s.clone_htod(&chains.iter().map(|c| c.seed).collect::<Vec<u64>>()).map_err(driver)?;
let offsets = s.clone_htod(&reach.offsets).map_err(driver)?;
let (cost, n32, count32) = (w.cost(), n as u32, count as u32);
let window = device::window(w, &reach, sampling);
let mut k0 = 0usize;
while k0 < temperatures.len() {
let k1 = (k0 + window).min(temperatures.len());
let (sweeps, table) = device::window_sweeps(w, &reach, &temperatures[k0..k1]);
let (sweeps, table) = (s.clone_htod(&sweeps).map_err(driver)?, s.clone_htod(&nonempty(&table)).map_err(driver)?);
let (a0, a1) = (k0 as u32, k1 as u32);
let mut a = s.launch_builder(&self.kernels.chains);
a.arg(&rows).arg(&cols).arg(&vals).arg(&fields).arg(&cost).arg(&n32).arg(&count32).arg(&offsets).arg(&a0).arg(&a1).arg(&seeds).arg(&sweeps).arg(&table).arg(&mut cells);
// SAFETY: 2·count·n cells, an offset a cell, a seed a chain, a word a sweep of the
// window, and the window's table of (a, b) pairs for every field a cell can reach.
unsafe { a.launch(launch(count as u64)) }.map_err(driver)?;
// One window at a time: only its table is held, and a failure leaves no more queued.
s.synchronize().map_err(driver)?;
k0 = k1;
}
let got = s.clone_dtoh(&cells).map_err(driver)?;
Ok(device::outcomes(w, &got, count, temperatures))
}
}
/// A buffer the driver can make: at least one element.
fn nonempty<T: Copy + Default>(v: &[T]) -> Vec<T> {
if v.is_empty() { vec![T::default()] } else { v.to_vec() }
}
impl Member for CudaLandscape {
fn id(&self) -> String {
self.id.clone()
}
fn class(&self) -> &'static str {
"gpu"
}
fn label(&self) -> String {
device::label("gpu", &self.name, "CUDA", "its cells in the reference's order")
}
fn takes(&self, w: &Workload, sampling: bool) -> Result<(), String> {
if sampling { device::table_fits(&device::reach(w)) } else { Ok(()) }
}
fn settle(&self, w: &Workload, start: &[i8], bound: u32) -> Result<Outcome, String> {
self.meter.around(|| self.settle_on_device(w, start, bound))
}
fn chains(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String> {
self.meter.around(|| self.chains_on_device(w, chains, temperatures))
}
fn meter(&self) -> Option<Metered> {
self.meter.last()
}
}
#[cfg(test)]
mod tests {
use super::*;
/// The kernels compile, where the runtime compiler is, with no device to run them on: a
/// build host whose card is off its bus still catches a kernel that no longer compiles.
/// With `WAI_LANDSCAPE_NVRTC` set, a host with no runtime compiler fails.
#[test]
fn the_kernels_compile_where_the_runtime_compiler_is() {
// SAFETY: only tries to open the runtime compiler's library by its known names.
if !unsafe { cudarc::nvrtc::sys::is_culib_present() } {
assert!(std::env::var_os("WAI_LANDSCAPE_NVRTC").is_none(), "WAI_LANDSCAPE_NVRTC is set, and this host has no runtime compiler");
eprintln!("no runtime compiler on this host: nothing compiled");
return;
}
let opts = CompileOptions { arch: Some("compute_80"), options: vec!["--std=c++17".into(), "--fmad=false".into()], ..Default::default() };
compile_ptx_with_opts(SOURCE, opts).unwrap();
}
}

216
wai/wai-quantum/src/quantum_landscape/device.rs Normal file
View file

@ -0,0 +1,216 @@
//! What the GPU members of the landscape fabric share: the workload in rows,
//! the cut-point tables the host works out and the windows of sweeps a
//! dispatch runs ([`super::reach`]), the outcomes a member returns, its id, its
//! label and its meter.
//! The kernels themselves are each member's ([`super::webgpu`], [`super::cuda`],
//! [`super::metal`]).
use super::member::Metered;
pub use super::reach::{MAX_TABLE_ENTRIES, Reach, WINDOW_ENTRIES, WINDOW_WORK, cut_pairs, reach, table_fits, window, window_sweeps};
use super::settle::{Outcome, Term, Workload};
use crate::quantum_energy::check_label;
use crate::quantum_fabric::RunMeter;
use std::collections::HashMap;
use std::sync::Mutex;
use std::thread::{self, ThreadId};
use std::time::Instant;
/// Sweeps a settle submits before it reads where it got to.
pub const SETTLE_BATCH: u32 = 64;
/// The workload as a device reads it: each cell's first coupling (and one
/// past the last), the couplings' columns and values, and the fields.
pub struct Rows {
pub rows: Vec<u32>,
pub cols: Vec<u32>,
pub vals: Vec<i32>,
pub fields: Vec<i64>,
}
pub fn rows(w: &Workload) -> Rows {
let n = w.n();
let zero = vec![0i8; n];
let mut r = Rows { rows: Vec::with_capacity(n + 1), cols: Vec::new(), vals: Vec::new(), fields: (0..n).map(|i| w.field(&zero, i)).collect() };
r.rows.push(0);
for i in 0..n {
for (j, c) in w.row(i) {
r.cols.push(j as u32);
r.vals.push(c as i32);
}
r.rows.push(r.cols.len() as u32);
}
r
}
/// The outcomes of `count` chains from their cells: the state, every sweep
/// run, the work and the draws the schedule comes to, and whether it rests.
pub fn outcomes(w: &Workload, cells: &[i32], count: usize, temperatures: &[f64]) -> Vec<Outcome> {
let n = w.n();
let draws = (n as u64) * temperatures.iter().filter(|&&t| t > 0.0).count() as u64;
(0..count)
.map(|c| {
let state: Vec<i8> = cells[c * n..(c + 1) * n].iter().map(|&v| v as i8).collect();
let term = if w.rests(&state) { Term::FixedPoint } else { Term::Bound };
Outcome { state, sweeps: temperatures.len() as u32, term, accumulates: w.nonzeros() * temperatures.len() as u64, draws }
})
.collect()
}
/// A settle's outcome from where it ended: the sweeps run, how it ended
/// (1 rest, 2 two-cycle, else the bound) and the state the last sweep computed.
pub fn settled(w: &Workload, done: u32, ended: u32, state: Vec<i8>) -> Outcome {
let term = match ended {
1 => Term::FixedPoint,
2 => Term::TwoCycle,
_ => Term::Bound,
};
Outcome { state, sweeps: done, term, accumulates: w.nonzeros() * u64::from(done), draws: 0 }
}
/// A member id: `<class>:<backend>.<name>.<4 hex>`, at most 64 characters, the
/// name as the platform reports it in lower case and the hex from what tells
/// two of the same name apart.
pub fn id_of(class: &str, backend: &str, name: &str, hardware: &str) -> String {
let mut slug = String::new();
for c in name.to_lowercase().chars() {
if c.is_ascii_lowercase() || c.is_ascii_digit() {
slug.push(c);
} else if !slug.ends_with('-') && !slug.is_empty() {
slug.push('-');
}
}
let mut h = 0xcbf2_9ce4_8422_2325u64;
for b in hardware.bytes() {
h = (h ^ u64::from(b)).wrapping_mul(0x0100_0000_01b3);
}
let room = 64usize.saturating_sub(class.len() + 1 + backend.len() + 1 + 5);
let slug: String = slug.trim_end_matches('-').chars().take(room).collect();
format!("{class}:{backend}.{}.{:04x}", slug.trim_end_matches('-'), h & 0xffff)
}
/// What a device member's answers say of themselves. `class` is the member's:
/// a software driver is `cpu`, and is not called a GPU.
pub fn label(class: &str, device: &str, through: &str, order: &str) -> String {
let on = if class == "gpu" { format!("this machine's GPU, {device}") } else { format!("a software driver on this machine's CPU, {device}") };
format!(
"A classical emulator of the settling model of a planned energy-based processor, running on {on}, through {through}, {order}. Energies are computed exactly. The joule line is this run's operation count priced at rates measured on an FPGA, with a published figure for one integer add beside it for scale. No photonic hardware, custom chip or quantum effect is involved in this run."
)
}
/// A member's meter, read around each call. A reading goes back to the thread
/// that made the call, so that runs sharing a member never carry each other's.
#[derive(Default)]
pub struct MeterSlot {
meter: Option<Mutex<Box<dyn RunMeter + Send>>>,
readings: Mutex<HashMap<ThreadId, Metered>>,
}
impl MeterSlot {
pub fn new(meter: Box<dyn RunMeter + Send>) -> MeterSlot {
MeterSlot { meter: Some(Mutex::new(meter)), readings: Mutex::new(HashMap::new()) }
}
/// Run `call` inside the meter's window, keeping its reading for this
/// thread. A meter that fails leaves the call unmetered and the call whole,
/// and a figure its label does not hold for gives no reading.
pub fn around<T>(&self, call: impl FnOnce() -> Result<T, String>) -> Result<T, String> {
let Some(meter) = &self.meter else { return call() };
let mut meter = meter.lock().unwrap_or_else(|p| p.into_inner());
let me = thread::current().id();
let started = meter.start().is_ok();
let from = Instant::now();
let out = call();
let window_us = from.elapsed().as_micros() as u64;
let reading = if started { meter.stop().ok().flatten() } else { None };
let scope = meter.scope();
drop(meter);
let mut readings = self.readings.lock().unwrap_or_else(|p| p.into_inner());
readings.remove(&me);
if let Some((joules_micro, class)) = reading.filter(|(j, c)| check_label(*j, c).is_ok()) {
let class = serde_json::Value::Object(class.json_fields().into_iter().map(|(k, v)| (k.to_string(), v)).collect());
readings.insert(me, Metered { joules_micro, window_us, class, scope });
}
out
}
/// The reading of this thread's last call, once.
pub fn last(&self) -> Option<Metered> {
self.readings.lock().unwrap_or_else(|p| p.into_inner()).remove(&thread::current().id())
}
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn an_id_is_a_member_id() {
let id = id_of("gpu", "metal", "Apple M3 Ultra (some driver build)", "4203/1/drv/1.2");
assert!(id.starts_with("gpu:metal.apple-m3-ultra-some-driver-build."), "{id}");
assert!(id.len() <= 64);
assert!(id_of("gpu", "vulkan", &"x".repeat(200), "").len() <= 64);
}
/// Two threads on one meter each get the reading of their own call.
#[test]
fn a_reading_goes_back_to_the_thread_that_made_the_call() {
use crate::quantum_energy::EnergyClass;
struct Ticks(u64);
impl RunMeter for Ticks {
fn start(&mut self) -> Result<(), String> {
Ok(())
}
fn stop(&mut self) -> Result<Option<(u64, EnergyClass)>, String> {
self.0 += 1;
Ok(Some((self.0, EnergyClass::Estimator)))
}
}
let slot = MeterSlot::new(Box::new(Ticks(0)));
let (first, second) = (std::sync::Barrier::new(2), std::sync::Barrier::new(2));
thread::scope(|s| {
s.spawn(|| {
slot.around(|| Ok(())).unwrap();
first.wait();
// The other thread's whole call falls here, before this one reads its own.
second.wait();
assert_eq!(slot.last().unwrap().joules_micro, 1);
assert!(slot.last().is_none(), "a reading is read once");
});
s.spawn(|| {
first.wait();
slot.around(|| Ok(())).unwrap();
second.wait();
assert_eq!(slot.last().unwrap().joules_micro, 2);
});
});
}
/// A figure its class does not hold for gives no reading; a meter that fails
/// leaves the call whole and unmetered.
#[test]
fn a_meter_never_costs_a_call_its_answer() {
use crate::quantum_energy::EnergyClass;
struct Broken(bool);
impl RunMeter for Broken {
fn start(&mut self) -> Result<(), String> {
if self.0 { Err("no counter".into()) } else { Ok(()) }
}
fn stop(&mut self) -> Result<Option<(u64, EnergyClass)>, String> {
Ok(Some((0, EnergyClass::Estimator)))
}
}
for broken in [true, false] {
let slot = MeterSlot::new(Box::new(Broken(broken)));
assert_eq!(slot.around(|| Ok(7)), Ok(7));
assert!(slot.last().is_none());
}
}
/// A software driver is not called a GPU.
#[test]
fn the_label_says_what_ran() {
assert!(label("gpu", "Radeon", "Vulkan", "in order").contains("this machine's GPU, Radeon"));
let soft = label("cpu", "llvmpipe", "WebGPU (vulkan)", "in order");
assert!(soft.contains("a software driver on this machine's CPU, llvmpipe") && !soft.contains("GPU, llvmpipe"), "{soft}");
}
}

335
wai/wai-quantum/src/quantum_landscape/job.rs
View file

@ -20,7 +20,7 @@
//! [`call`] answers on the CPU reference, and [`call_certified`] on a member the
//! host has certified ([`super::member::certify`]).
use super::member::Certified;
use super::member::{Certified, Parity};
use super::program::{Program, Refusal, decimal, read};
use super::run::{self, DEFAULT_SEED, Job, Plan};
use super::{EXAMPLES, MAX_JOB_BYTES, MAX_OPERATIONS, PROTOCOL};
@ -93,17 +93,30 @@ pub fn run_text(program: &str, seed: u64) -> Result<Value, Refusal> {
call("run", &json!({ "program": program, "seed": seed.to_string() }))
}
/// Answer a request on a certified member, holding no more than `max_bytes`
/// (and never more than [`MAX_JOB_BYTES`]). `op` is `run`, `sweep`, `compare`,
/// `replay` or `examples`.
pub fn call_certified(c: &Certified<'_>, op: &str, request: &Value, max_bytes: u64) -> Result<Value, Refusal> {
/// A request read and its budget checked: the job, ready to run.
pub struct Prepared {
pub job: Job,
op: String,
seed: u64,
temperature: Option<f64>,
seeds: u32,
temperatures: Vec<f64>,
/// Whether any sweep of the job is above temperature 0.
pub sampling: bool,
/// The kernels the job runs.
pub kernels: Vec<&'static str>,
}
/// Read a request and check every bound of its job before any of it runs:
/// `None` for `examples`, which runs nothing.
pub fn prepare(op: &str, request: &Value, max_bytes: u64) -> Result<Option<Prepared>, Refusal> {
if let Some(p) = request.get("protocol")
&& p != PROTOCOL
{
return Err(Refusal::general(format!("This reads {PROTOCOL}, not {p}.")));
}
if op == "examples" {
return Ok(examples());
return Ok(None);
}
if !matches!(op, "run" | "sweep" | "compare" | "replay") {
return Err(Refusal::general(format!("`{op}` is not an operation here: run, sweep, compare, replay and examples are.")));
@ -112,22 +125,24 @@ pub fn call_certified(c: &Certified<'_>, op: &str, request: &Value, max_bytes: u
let p = read(source)?;
let seed = seed_of(request)?;
let temperature = temperature_of(request)?;
let (mut seeds, mut temperatures) = (1u32, Vec::new());
// What the job will run, and how many times, before any of it runs.
let (runs, lines): (Vec<(Plan, u64)>, u64) = match op {
"run" => (vec![(Plan::of(&p, temperature)?, 1)], 0),
"sweep" => {
let temperatures = temperatures_of(request, &p)?;
(run::sweep_plans(&p, &temperatures, seeds_of(request, 20)?)?.into_iter().map(|(_, plan, k)| (plan, u64::from(k))).collect(), 0)
temperatures = temperatures_of(request, &p)?;
seeds = seeds_of(request, 20)?;
(run::sweep_plans(&p, &temperatures, seeds)?.into_iter().map(|(_, plan, k)| (plan, u64::from(k))).collect(), 0)
}
"compare" => {
let seeds = u64::from(seeds_of(request, 20)?);
seeds = seeds_of(request, 20)?;
let plan = run::compare_plan(&p, temperature)?;
// The emulator twice, and three classical methods at the sampling run's budget.
(vec![(plan, seeds * 5)], 0)
(vec![(plan, u64::from(seeds) * 5)], 0)
}
_ => {
let seeds = u64::from(seeds_of(request, 20)?);
(vec![(Plan::of(&p, temperature)?, seeds)], seeds)
seeds = seeds_of(request, 20)?;
(vec![(Plan::of(&p, temperature)?, u64::from(seeds))], u64::from(seeds))
}
};
let operations = run::job_operations(&p, &runs);
@ -146,35 +161,122 @@ pub fn call_certified(c: &Certified<'_>, op: &str, request: &Value, max_bytes: u
if bytes > most {
return Err(Refusal::general(format!("This job could hold up to {} bytes, and this host allows {}.", super::price::group(bytes), super::price::group(most))));
}
let job = Job::new(p);
match op {
"run" => run::run_on(c, &job, seed, temperature),
"sweep" => run::sweep_on(c, &job, &temperatures_of(request, &job.program)?, seeds_of(request, 20)?, seed),
"compare" => run::compare_on(c, &job, seeds_of(request, 20)?, seed, temperature),
let sampling = op == "compare" || runs.iter().any(|(plan, _)| plan.mode() == super::price::Mode::Sample);
let mut kernels: Vec<&'static str> = runs.iter().map(|(plan, _)| plan.kernel()).collect();
kernels.sort_unstable();
kernels.dedup();
Ok(Some(Prepared { job: Job::new(p), op: op.into(), seed, temperature, seeds, temperatures, sampling, kernels }))
}
/// Answer a prepared job on a certified member.
pub fn answer(c: &Certified<'_>, prepared: &Prepared) -> Result<Value, Refusal> {
let Prepared { job, op, seed, temperature, seeds, temperatures, .. } = prepared;
let (seed, temperature, seeds) = (*seed, *temperature, *seeds);
match op.as_str() {
"run" => run::run_on(c, job, seed, temperature),
"sweep" => run::sweep_on(c, job, temperatures, seeds, seed),
"compare" => run::compare_on(c, job, seeds, seed, temperature),
_ => {
// Many seeds, one line each: the seed, a digest of that run's whole answer, the
// state, its energy, the verdict, the accumulates and the replay digest.
let seeds = seeds_of(request, 20)?;
// state, its energy, the verdict, the accumulates and the replay digest. The
// starts of several seeds go to the member together.
let plan = Plan::of(&job.program, temperature)?;
let all: Vec<u64> = (0..seeds).map(|k| seed.wrapping_add(u64::from(k))).collect();
let mut lines = String::new();
for k in 0..seeds {
let s = seed.wrapping_add(u64::from(k));
let a = run::run_on(c, &job, s, temperature)?;
let text = a.to_string();
lines.push_str(&format!(
"{s} {} {} {} {} {} {}\n",
&blake3::hash(text.as_bytes()).to_hex()[..16],
a["state"].as_str().unwrap_or(""),
a["energy"],
a["verdict"].as_str().unwrap_or(""),
a["work"]["accumulates"],
&a["replay"].as_str().unwrap_or("")[..16],
));
for group in all.chunks(run::seeds_a_batch(job, plan)) {
for (s, found) in group.iter().zip(run::execute_seeds(job, plan, group, c)?) {
let a = run::answer_of(c, job, plan, *s, &found);
let text = a.to_string();
lines.push_str(&format!(
"{s} {} {} {} {} {} {}\n",
&blake3::hash(text.as_bytes()).to_hex()[..16],
a["state"].as_str().unwrap_or(""),
a["energy"],
a["verdict"].as_str().unwrap_or(""),
a["work"]["accumulates"],
&a["replay"].as_str().unwrap_or("")[..16],
));
}
}
Ok(json!({ "protocol": PROTOCOL, "seeds": seeds, "lines": lines }))
Ok(json!({ "protocol": PROTOCOL, "label": c.member().label(), "member": super::member::describe(c, plan.kernel()), "seeds": seeds, "lines": lines }))
}
}
}
/// Answer a request on a certified member, holding no more than `max_bytes`
/// (and never more than [`MAX_JOB_BYTES`]). `op` is `run`, `sweep`, `compare`,
/// `replay` or `examples`.
pub fn call_certified(c: &Certified<'_>, op: &str, request: &Value, max_bytes: u64) -> Result<Value, Refusal> {
match prepare(op, request, max_bytes)? {
None => Ok(examples()),
Some(p) => answer(c, &p),
}
}
/// Answer a request on the first of `members` that can: one certified for
/// every kernel the job runs (an uncertified kernel is never placed; a member
/// asked for by name goes through [`call_certified`]), that takes the workload
/// ([`super::member::Member::takes`]), and that gives an answer. Each member
/// passed over, for any of those, is named in the answer's `refusals` with why,
/// and the next is asked. Put the CPU reference last, and a job that a member
/// could not finish is still answered.
pub fn call_placed(members: &[Certified<'_>], op: &str, request: &Value, max_bytes: u64) -> Result<Value, Refusal> {
let Some(p) = prepare(op, request, max_bytes)? else { return Ok(examples()) };
let mut refusals = Vec::new();
for c in members {
let m = c.member();
let why = if let Some(k) = p.kernels.iter().find(|k| c.parity(k) == Parity::Uncertified) {
format!("its kernel {k} is not certified")
} else if let Err(why) = m.takes(&p.job.workload, p.sampling) {
why
} else {
match answer(c, &p) {
Ok(mut a) => {
if !refusals.is_empty() {
a["refusals"] = Value::Array(refusals);
}
return Ok(a);
}
Err(r) => r.plain(),
}
};
refusals.push(json!({ "member": m.id(), "why": why }));
}
Err(Refusal::general(format!("No member offered answered this job: {}.", refusals.iter().map(|r| format!("{}: {}", r["member"].as_str().unwrap_or(""), r["why"].as_str().unwrap_or(""))).collect::<Vec<_>>().join("; "))))
}
/// A digest of what a member gives for every worked example at its own
/// temperature, at 0 and at 0.5, over `seeds` seeds, that does not depend on
/// which member it was: each run's state, energy, verdict, work and replay
/// digest. Two members that give the same bits give the same digest.
pub fn corpus_bits_on(c: &Certified<'_>, seeds: u32) -> String {
let mut h = blake3::Hasher::new();
for (name, text) in EXAMPLES {
for temperature in [Value::Null, json!(0.0), json!(0.5)] {
h.update(name.as_bytes());
h.update(&[0]);
let req = json!({ "program": text, "seed": 0, "seeds": seeds, "temperature": temperature });
match prepare("replay", &req, MAX_JOB_BYTES).and_then(|p| {
let p = p.expect("a replay runs");
let plan = Plan::of(&p.job.program, p.temperature)?;
let all: Vec<u64> = (0..seeds).map(u64::from).collect();
let mut out = String::new();
for group in all.chunks(run::seeds_a_batch(&p.job, plan)) {
for (s, found) in group.iter().zip(run::execute_seeds(&p.job, plan, group, c)?) {
let a = run::answer_of(c, &p.job, plan, *s, &found);
out.push_str(&format!("{s} {} {} {} {} {} {}\n", a["state"], a["energy"], a["verdict"], a["work"]["accumulates"], a["work"]["draws"], a["replay"]));
}
}
Ok(out)
}) {
Ok(lines) => h.update(lines.as_bytes()),
Err(r) => h.update(r.plain().as_bytes()),
};
}
}
h.finalize().to_hex().to_string()
}
/// The worked examples: each with what it is there to show.
pub fn examples() -> Value {
let note = |text: &str, key: &str| -> Value {
@ -441,6 +543,173 @@ mod tests {
assert_eq!(a["rows"][0]["mean_energy"].as_f64(), Some(want), "{}", a["rows"][0]);
}
/// A stand-in member with narrow lanes: it declines a workload whose field
/// bound is above `widest`, and runs the reference's kernels otherwise.
struct Narrow(i64);
impl super::super::member::Member for Narrow {
fn id(&self) -> String {
format!("fpga:narrow-{}", self.0)
}
fn class(&self) -> &'static str {
"fpga"
}
fn label(&self) -> String {
"A board with narrow lanes.".into()
}
fn settle(&self, w: &super::super::settle::Workload, start: &[i8], bound: u32) -> Result<super::super::settle::Outcome, String> {
Ok(w.settle(start, bound))
}
fn chains(&self, w: &super::super::settle::Workload, chains: &[super::super::member::Chain], temperatures: &[f64]) -> Result<Vec<super::super::settle::Outcome>, String> {
super::super::member::CpuReference.chains(w, chains, temperatures)
}
fn takes(&self, w: &super::super::settle::Workload, _: bool) -> Result<(), String> {
if w.field_bound() > self.0 { Err(format!("its fields reach {}, beyond {}", w.field_bound(), self.0)) } else { Ok(()) }
}
}
/// A job goes to the first member that takes it, and the answer names every member that declined.
#[test]
fn a_job_goes_to_the_first_member_that_takes_it() {
use super::super::member::certify;
let lanes32 = Narrow(i64::from(i32::MAX));
let placed = [certify(&lanes32).unwrap(), Certified::reference()];
let wide = call_placed(&placed, "run", &json!({ "program": "vars 2\nfield 0 1099511627776\nagree 0 1\nminimize\n" }), MAX_JOB_BYTES).unwrap();
assert_eq!(wide["member"]["id"], "cpu:reference");
assert_eq!(wide["refusals"], json!([{ "member": "fpga:narrow-2147483647", "why": "its fields reach 1099511627777, beyond 2147483647" }]));
let narrow = call_placed(&placed, "run", &json!({ "program": example("talks"), "seed": "3" }), MAX_JOB_BYTES).unwrap();
assert_eq!(narrow["member"]["id"], "fpga:narrow-2147483647");
assert!(narrow.get("refusals").is_none());
assert_eq!(narrow["replay"], call("run", &json!({ "program": example("talks"), "seed": "3" })).unwrap()["replay"]);
let none = call_placed(&placed[..1], "run", &json!({ "program": "vars 2\nfield 0 1099511627776\nminimize\n" }), MAX_JOB_BYTES).unwrap_err();
assert!(none.message.starts_with("No member offered answered this job"), "{}", none.message);
}
/// A member may decline only the vectors whose fields do not fit 32 bits:
/// declining one that does is failing it.
#[test]
fn a_member_may_decline_only_what_does_not_fit_32_bits() {
use super::super::member::{KERNEL_SAMPLE, KERNEL_SETTLE, Parity, certify};
let lanes32 = Narrow(i64::from(i32::MAX));
let c = certify(&lanes32).unwrap();
assert_eq!((c.parity(KERNEL_SETTLE), c.parity(KERNEL_SAMPLE)), (Parity::BitExact, Parity::BitExact));
let lanes8 = Narrow(10);
let c = certify(&lanes8).unwrap();
assert_eq!((c.parity(KERNEL_SETTLE), c.parity(KERNEL_SAMPLE)), (Parity::Uncertified, Parity::Uncertified), "it declines vectors whose fields fit 32 bits");
}
/// A member whose chains come back where they started: its settle kernel is the
/// reference's, its sampling kernel wrong. It fails at run time once told to.
struct Inert(std::cell::Cell<bool>);
impl super::super::member::Member for Inert {
fn id(&self) -> String {
"gpu:inert".into()
}
fn class(&self) -> &'static str {
"gpu"
}
fn label(&self) -> String {
"A board under test.".into()
}
fn settle(&self, w: &super::super::settle::Workload, start: &[i8], bound: u32) -> Result<super::super::settle::Outcome, String> {
if self.0.get() { Err("the device was lost".into()) } else { Ok(w.settle(start, bound)) }
}
fn chains(&self, w: &super::super::settle::Workload, chains: &[super::super::member::Chain], temperatures: &[f64]) -> Result<Vec<super::super::settle::Outcome>, String> {
if self.0.get() {
return Err("the device was lost".into());
}
let draws = w.n() as u64 * temperatures.iter().filter(|&&t| t > 0.0).count() as u64;
Ok(chains
.iter()
.map(|c| {
let term = if w.rests(&c.start) { super::super::settle::Term::FixedPoint } else { super::super::settle::Term::Bound };
super::super::settle::Outcome { state: c.start.clone(), sweeps: temperatures.len() as u32, term, accumulates: w.nonzeros() * temperatures.len() as u64, draws }
})
.collect())
}
}
/// An uncertified kernel is never placed: a job that runs it goes to the next
/// member, and the answer says why. A job that runs only a certified kernel of
/// the same member stays on it. A member asked for by name is not placed at all.
#[test]
fn an_uncertified_kernel_is_never_placed() {
use super::super::member::{KERNEL_SAMPLE, KERNEL_SETTLE, Parity, certify, named};
let board = Inert(std::cell::Cell::new(false));
let c = certify(&board).unwrap();
assert_eq!((c.parity(KERNEL_SETTLE), c.parity(KERNEL_SAMPLE)), (Parity::BitExact, Parity::Uncertified));
let placed = [c, Certified::reference()];
let why = json!([{ "member": "gpu:inert", "why": "its kernel wai.quantum.landscape/sample is not certified" }]);
for program in [example("talks"), "vars 3\nagree 0 1\nagree 1 2\nsolve t0=0 restarts=4\n"] {
for op in ["run", "replay", "sweep", "compare"] {
let a = call_placed(&placed, op, &json!({ "program": program, "seed": "3", "seeds": 2 }), MAX_JOB_BYTES).unwrap();
assert_eq!(a["member"]["id"], "cpu:reference", "{op}");
assert_eq!(a["refusals"], why, "{op}");
}
}
let settle = call_placed(&placed, "run", &json!({ "program": "vars 3\nagree 0 1\nfield 0 1\nminimize\n" }), MAX_JOB_BYTES).unwrap();
assert_eq!((settle["member"]["id"].as_str(), settle["member"]["parity"].as_str()), (Some("gpu:inert"), Some("bit-exact")));
let by_name = [named(&board).unwrap()];
let refused = call_placed(&by_name, "run", &json!({ "program": "vars 3\nagree 0 1\nminimize\n" }), MAX_JOB_BYTES).unwrap_err();
assert!(refused.message.contains("is not certified"), "{}", refused.message);
assert_eq!(call_certified(&by_name[0], "run", &json!({ "program": "vars 3\nagree 0 1\nminimize\n" }), MAX_JOB_BYTES).unwrap()["member"]["parity"], "uncertified");
}
/// A member that takes a job and fails it passes the job on: the next member
/// answers, and every member passed over is named, in order, with why.
#[test]
fn a_member_that_fails_a_job_passes_it_on() {
use super::super::member::certify;
let (lanes32, board) = (Narrow(i64::from(i32::MAX)), Inert(std::cell::Cell::new(false)));
let placed = [certify(&lanes32).unwrap(), certify(&board).unwrap(), Certified::reference()];
board.0.set(true);
let a = call_placed(&placed, "run", &json!({ "program": "vars 3\nagree 0 1\nfield 0 4294967295\nminimize\n" }), MAX_JOB_BYTES).unwrap();
assert_eq!(a["member"]["id"], "cpu:reference");
let refusals = a["refusals"].as_array().unwrap();
assert_eq!(refusals.iter().map(|r| r["member"].as_str().unwrap()).collect::<Vec<_>>(), ["fpga:narrow-2147483647", "gpu:inert"]);
assert!(refusals[1]["why"].as_str().unwrap().contains("the device was lost"), "{a}");
let none = call_placed(&placed[1..2], "run", &json!({ "program": "vars 3\nagree 0 1\nminimize\n" }), MAX_JOB_BYTES).unwrap_err();
assert!(none.message.contains("the device was lost"), "{}", none.message);
}
/// However a member is reached, it is never run on a workload it declines.
#[test]
fn a_member_never_runs_what_it_declines() {
use super::super::member::certify;
let lanes32 = Narrow(i64::from(i32::MAX));
let c = certify(&lanes32).unwrap();
for op in ["run", "replay", "sweep", "compare"] {
for program in ["vars 3\nfield 0 4294967295\nagree 0 1\nagree 0 2\nminimize steps=1\n", "vars 3\nfield 0 4294967295\nagree 0 1\nsolve\n"] {
let r = call_certified(&c, op, &json!({ "program": program, "seeds": 2 }), MAX_JOB_BYTES).unwrap_err();
assert!(r.message.contains("fpga:narrow-2147483647") && r.message.contains("beyond 2147483647"), "{op}: {}", r.message);
}
}
}
/// Every answer says who ran it and at what parity: a replay names its member
/// and kernel, a sweep each kernel it used, and a comparison the emulator's.
#[test]
fn every_answer_says_who_ran_it_at_what_parity() {
let c = Certified::reference();
let replay = call_certified(&c, "replay", &json!({ "program": example("talks"), "seeds": 2 }), MAX_JOB_BYTES).unwrap();
assert_eq!(replay["member"], json!({ "id": "cpu:reference", "class": "cpu", "kernel": "wai.quantum.landscape/sample", "parity": "reference" }));
assert_eq!(replay["label"], super::super::LABEL);
let sweep = call_certified(&c, "sweep", &json!({ "program": example("talks"), "temperatures": [0, 1.5], "seeds": 2 }), MAX_JOB_BYTES).unwrap();
assert_eq!(
sweep["member"]["kernels"],
json!([{ "kernel": "wai.quantum.landscape/sample", "parity": "reference" }])
);
assert_eq!(sweep["rows"][0]["kernel"], "wai.quantum.landscape/sample");
let minimize = call_certified(&c, "sweep", &json!({ "program": "vars 2\nagree 0 1\nminimize\n", "temperatures": [0, 1.5] }), MAX_JOB_BYTES).unwrap();
assert_eq!(minimize["member"]["kernels"].as_array().unwrap().len(), 2, "a settle at 0 and a sampling run above it");
assert_eq!(minimize["rows"][0]["parity"], "reference");
let compare = call_certified(&c, "compare", &json!({ "program": example("talks"), "seeds": 2 }), MAX_JOB_BYTES).unwrap();
assert_eq!(compare["member"]["kernels"][0]["parity"], "reference");
assert_eq!(compare["rows"][0]["parity"], "reference");
assert_eq!(compare["rows"][2]["parity"], Value::Null, "a classical row ran on no member");
}
#[test]
fn a_seed_keeps_all_64_bits() {
let a = call("run", &json!({ "program": example("chain"), "seed": "18446744073709551615" })).unwrap();

239
wai/wai-quantum/src/quantum_landscape/landscape.wgsl Normal file
View file

@ -0,0 +1,239 @@
// The two kernels of wai.quantum.landscape/1 in WGSL (extensions/quantum-landscape.md §4).
//
// Every value a cell's field can take fits a 32-bit lane: the host gives a workload only when
// its field bound does (quantum_landscape::webgpu), and a cost above every field is held at the
// lane's top, where it holds every cell at 0 as it would at its own size. Fields are whole numbers, so a sum
// is the same whatever order a device adds it in, and the settle kernel gives the reference's
// bits. The sampling kernel gives them too: its draws are the reference's (SplitMix64 by
// counter, in pairs of 32-bit words), its cut points come from the host, and each chain runs
// its cells in order. `chains_coloured` moves a colour class of cells at once: its order is
// not the reference's, so its parity can only be statistical.
//
// At most four storage buffers an entry point, the WebGPU floor: the couplings, the fields and
// the offsets in one (`graph`), every chain's cells and scratch in another (`cells`), and the
// seeds, the window's sweeps, the cut points and a colour class's cells in a third (`draws`).
struct Shape {
n: u32,
cost: i32,
chains: u32,
// A settle: the most sweeps.
bound: u32,
// Chains: the first sweep of this window, and one past its last.
k0: u32,
k1: u32,
// Where each part of `graph` starts: rows at 0, then columns, values, fields and, last,
// each cell's table offset (`offsets`, below).
cols: u32,
vals: u32,
fields: u32,
// Where each part of `draws` starts: seeds at 0, then the sweeps and the table.
sweeps: u32,
table: u32,
// A colour class: where its cells start in `draws`, how many it has, its sweep and kind.
class_first: u32,
class_count: u32,
class_k: u32,
class_kind: u32,
offsets: u32,
}
@group(0) @binding(0) var<storage, read> graph: array<u32>;
@group(0) @binding(1) var<uniform> shape: Shape;
fn threshold(h: i32, c: i32) -> i32 {
if (h > c) { return 1; }
if (h < -c) { return -1; }
return 0;
}
// ── settle: one invocation a cell, one dispatch a sweep, then one to advance ────────────────
// Three states in turn: s_k is at k mod 3.
@group(0) @binding(2) var<storage, read_write> ring: array<i32>;
// [sweeps run, how it ended (1 rest, 2 two-cycle, 3 bound), moved, moved from two back]
@group(0) @binding(3) var<storage, read_write> status: array<atomic<u32>, 4>;
fn ring_field(base: u32, i: u32) -> i32 {
var h = bitcast<i32>(graph[shape.fields + i]);
for (var k = graph[i]; k < graph[i + 1u]; k = k + 1u) {
h = h + bitcast<i32>(graph[shape.vals + k]) * ring[base + graph[shape.cols + k]];
}
return h;
}
@compute @workgroup_size(64)
fn settle_sweep(@builtin(global_invocation_id) id: vec3<u32>, @builtin(num_workgroups) groups: vec3<u32>) {
let i = id.x + id.y * groups.x * 64u;
if (i >= shape.n || atomicLoad(&status[1]) != 0u) { return; }
let n = shape.n;
let done = atomicLoad(&status[0]);
let now = (done % 3u) * n;
let next = ((done + 1u) % 3u) * n;
let back = ((done + 2u) % 3u) * n;
let s = threshold(ring_field(now, i), shape.cost);
ring[next + i] = s;
if (s != ring[now + i]) { atomicOr(&status[2], 1u); }
if (s != ring[back + i]) { atomicOr(&status[3], 1u); }
}
@compute @workgroup_size(1)
fn settle_advance() {
if (atomicLoad(&status[1]) != 0u) { return; }
let k = atomicLoad(&status[0]) + 1u;
atomicStore(&status[0], k);
if (atomicLoad(&status[2]) == 0u) {
atomicStore(&status[1], 1u);
} else if (k >= 2u && atomicLoad(&status[3]) == 0u) {
atomicStore(&status[1], 2u);
} else if (k >= shape.bound) {
atomicStore(&status[1], 3u);
}
atomicStore(&status[2], 0u);
atomicStore(&status[3], 0u);
}
// ── chains ──────────────────────────────────────────────────────────────────────────────────
// Every chain's cells, then the same again as scratch for a sweep at 0.
@group(0) @binding(4) var<storage, read_write> cells: array<i32>;
// Each chain's seed (low word, high word); each sweep of the window (0 for a sweep at
// temperature 0, else 1 + its first table entry); the table, four words an entry (a's low and
// high words, then b's) for each field from −B to B; a colour class's cells.
@group(0) @binding(5) var<storage, read> draws: array<u32>;
fn cell_field(base: u32, i: u32) -> i32 {
var h = bitcast<i32>(graph[shape.fields + i]);
for (var k = graph[i]; k < graph[i + 1u]; k = k + 1u) {
h = h + bitcast<i32>(graph[shape.vals + k]) * cells[base + graph[shape.cols + k]];
}
return h;
}
// 64-bit words as (low, high).
fn add64(a: vec2<u32>, b: vec2<u32>) -> vec2<u32> {
let lo = a.x + b.x;
return vec2<u32>(lo, a.y + b.y + select(0u, 1u, lo < a.x));
}
// The full product of two 32-bit words.
fn mul32(a: u32, b: u32) -> vec2<u32> {
let a0 = a & 0xffffu;
let a1 = a >> 16u;
let b0 = b & 0xffffu;
let b1 = b >> 16u;
let p00 = a0 * b0;
let p01 = a0 * b1;
let p10 = a1 * b0;
let p11 = a1 * b1;
let mid = p01 + p10;
let mid_carry = select(0u, 0x10000u, mid < p01);
let lo = p00 + (mid << 16u);
let lo_carry = select(0u, 1u, lo < p00);
return vec2<u32>(lo, p11 + (mid >> 16u) + mid_carry + lo_carry);
}
// The low 64 bits of a product of two 64-bit words.
fn mul64(a: vec2<u32>, b: vec2<u32>) -> vec2<u32> {
let low = mul32(a.x, b.x);
return vec2<u32>(low.x, low.y + a.x * b.y + a.y * b.x);
}
fn shr64(a: vec2<u32>, s: u32) -> vec2<u32> {
return vec2<u32>((a.x >> s) | (a.y << (32u - s)), a.y >> s);
}
fn xor64(a: vec2<u32>, b: vec2<u32>) -> vec2<u32> {
return vec2<u32>(a.x ^ b.x, a.y ^ b.y);
}
// The j-th output of SplitMix64 started at s.
fn splitmix(s: vec2<u32>, j: vec2<u32>) -> vec2<u32> {
var x = add64(s, mul64(j, vec2<u32>(0x7F4A7C15u, 0x9E3779B9u)));
x = xor64(x, shr64(x, 30u));
x = mul64(x, vec2<u32>(0x1CE4E5B9u, 0xBF58476Du));
x = xor64(x, shr64(x, 27u));
x = mul64(x, vec2<u32>(0x133111EBu, 0x94D049BBu));
return xor64(x, shr64(x, 31u));
}
fn below(u: vec2<u32>, a: vec2<u32>) -> bool {
return u.y < a.y || (u.y == a.y && u.x < a.x);
}
fn pick(u: vec2<u32>, entry: u32) -> i32 {
let e = shape.table + entry * 4u;
if (below(u, vec2<u32>(draws[e], draws[e + 1u]))) { return -1; }
if (below(u, vec2<u32>(draws[e + 2u], draws[e + 3u]))) { return 0; }
return 1;
}
// The draw of cell i at sweep k: out(chain, k·n + i + 1) >> 11.
fn draw(seed: vec2<u32>, k: u32, i: u32) -> vec2<u32> {
let j = add64(mul32(k, shape.n), vec2<u32>(i + 1u, 0u));
return shr64(splitmix(seed, j), 11u);
}
fn seed_of(c: u32) -> vec2<u32> {
return vec2<u32>(draws[c * 2u], draws[c * 2u + 1u]);
}
// Cell i's entry for field h in a sweep's table: its offset, modulo 2³², lands inside the table.
fn entry_of(kind: u32, i: u32, h: i32) -> u32 {
return kind - 1u + bitcast<u32>(h) + graph[shape.offsets + i];
}
@compute @workgroup_size(64)
fn chains_in_order(@builtin(global_invocation_id) id: vec3<u32>, @builtin(num_workgroups) groups: vec3<u32>) {
let c = id.x + id.y * groups.x * 64u;
if (c >= shape.chains) { return; }
let n = shape.n;
let base = c * n;
let spare = shape.chains * n + base;
let seed = seed_of(c);
for (var k = shape.k0; k < shape.k1; k = k + 1u) {
let kind = draws[shape.sweeps + k - shape.k0];
if (kind == 0u) {
for (var i = 0u; i < n; i = i + 1u) {
cells[spare + i] = threshold(cell_field(base, i), shape.cost);
}
for (var i = 0u; i < n; i = i + 1u) {
cells[base + i] = cells[spare + i];
}
} else {
for (var i = 0u; i < n; i = i + 1u) {
let h = cell_field(base, i);
cells[base + i] = pick(draw(seed, k, i), entry_of(kind, i, h));
}
}
}
}
// ── chains, a colour class at a time: statistical parity at most ────────────────────────────
// One invocation a (chain, cell of the class). No two cells of a class are coupled, so each
// reads fields no other invocation of this dispatch writes. A sweep at 0 is two dispatches:
// every cell from the old state into scratch, then `chains_copy`.
@compute @workgroup_size(64)
fn chains_coloured(@builtin(global_invocation_id) id: vec3<u32>, @builtin(num_workgroups) groups: vec3<u32>) {
let t = id.x + id.y * groups.x * 64u;
let n = shape.n;
if (shape.class_kind == 0u) {
if (t >= shape.chains * n) { return; }
cells[shape.chains * n + t] = threshold(cell_field((t / n) * n, t % n), shape.cost);
return;
}
if (t >= shape.chains * shape.class_count) { return; }
let c = t / shape.class_count;
let i = draws[shape.class_first + t % shape.class_count];
let base = c * n;
let h = cell_field(base, i);
cells[base + i] = pick(draw(seed_of(c), shape.class_k, i), entry_of(shape.class_kind, i, h));
}
@compute @workgroup_size(64)
fn chains_copy(@builtin(global_invocation_id) id: vec3<u32>, @builtin(num_workgroups) groups: vec3<u32>) {
let t = id.x + id.y * groups.x * 64u;
if (t >= shape.chains * shape.n) { return; }
cells[t] = cells[shape.chains * shape.n + t];
}

66
wai/wai-quantum/src/quantum_landscape/member.rs
View file

@ -17,8 +17,9 @@
//! conformance set and grants what they pass:
//! - [`Parity::Reference`]: only this crate's own CPU kernels, through
//! [`Certified::reference`]; no member can be granted it.
//! - [`Parity::BitExact`]: its settle kernel gives every `settle` vector
//! bit for bit, or its sampling kernel every `chains` vector.
//! - [`Parity::BitExact`]: its settle kernel gives every `settle` vector it
//! takes bit for bit, or its sampling kernel every `chains` vector it takes,
//! and it takes every one it must ([`conformance::MUST_TAKE_FIELD_BOUND`]).
//! - [`Parity::Statistical`]: a sampling kernel with draws of its own whose
//! final states, over the battery of `statistical` workloads, follow the
//! exact distribution by the test that battery states.
@ -58,13 +59,44 @@ impl Parity {
}
}
/// A reading of what a run spent, from a meter on the member. `class` is the
/// acquisition class of `extensions/energy-measurement.md` §2 as a member
/// states it (`OnChipCounter`, `CalibratedInstrument`), with its uncertainty.
/// A reading from a meter on the member, over the calls of a run. `class` is
/// the acquisition class of `extensions/energy-measurement.md` §2 as a member
/// states it, with what that class asks. `window_us` is how long the meter's
/// windows ran, added up over the calls. `scope` is what the figure covers where
/// that is more than the run, as the meter says it; `None` is the run alone.
#[derive(Clone, Debug, PartialEq)]
pub struct Metered {
pub joules_micro: u64,
pub window_us: u64,
pub class: Value,
pub scope: Option<String>,
}
impl Metered {
/// Two calls' readings as one, if they are one reading: the same class and
/// scope. The figure is the two together and so is the window, and an
/// uncertainty's window with it, its relative part the larger of the two.
pub fn and(self, other: Metered) -> Option<Metered> {
if self.scope != other.scope {
return None;
}
let without = |c: &Value| {
let mut c = c.clone();
if let Some(o) = c.as_object_mut() {
o.remove("energy_uncertainty");
}
c
};
if without(&self.class) != without(&other.class) {
return None;
}
let mut class = self.class;
if let (Some(a), Some(b)) = (class.get("energy_uncertainty").cloned(), other.class.get("energy_uncertainty")) {
let of = |v: &Value, k: &str| v[k].as_u64().unwrap_or(0);
class["energy_uncertainty"] = json!({ "relative_ppm": of(&a, "relative_ppm").max(of(b, "relative_ppm")), "window_us": of(&a, "window_us").saturating_add(of(b, "window_us")) });
}
Some(Metered { joules_micro: self.joules_micro.saturating_add(other.joules_micro), window_us: self.window_us.saturating_add(other.window_us), class, scope: self.scope })
}
}
/// One chain to run: where it starts and the seed its draws come from.
@ -90,10 +122,19 @@ pub trait Member {
/// ([`super::sample::anneal_with`]): a sweep at 0 is the deterministic one.
/// A device may run the chains side by side.
fn chains(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String>;
/// A reading of what the run just made spent, if it has a meter.
/// The reading of the last call this thread made, if it has a meter and
/// the meter gave one.
fn meter(&self) -> Option<Metered> {
None
}
/// Whether it takes a workload, said before anything runs: the reason
/// where it does not, and the host then gives the job to the next member it
/// was offered. `sampling` is whether any sweep is above temperature 0.
/// A member declines only what `extensions/quantum-landscape.md` §5 lets it,
/// and the host never runs it on what it declines.
fn takes(&self, _w: &Workload, _sampling: bool) -> Result<(), String> {
Ok(())
}
}
/// The reference: this crate's kernels, on this machine's CPU.
@ -255,6 +296,19 @@ pub(super) fn describe(c: &Certified<'_>, kernel: &str) -> Value {
json!({ "id": m.id(), "class": m.class(), "kernel": kernel, "parity": c.parity(kernel).label() })
}
/// What an answer that may run both kernels says of the member that ran it:
/// each kernel it used, in the order first used, at the parity the host found.
pub(super) fn describe_all(c: &Certified<'_>, kernels: &[&str]) -> Value {
let m = c.member();
let mut used: Vec<&str> = Vec::new();
for k in kernels {
if !used.contains(k) {
used.push(k);
}
}
json!({ "id": m.id(), "class": m.class(), "kernels": used.iter().map(|k| json!({ "kernel": k, "parity": c.parity(k).label() })).collect::<Vec<_>>() })
}
#[cfg(test)]
mod tests {
use super::*;

407
wai/wai-quantum/src/quantum_landscape/metal.rs Normal file
View file

@ -0,0 +1,407 @@
//! A Metal member of the landscape fabric: the kernels of [`SOURCE`] on an
//! Apple GPU, compiled at run time with safe math.
//!
//! The kernels are those of `landscape.wgsl` with native 64-bit words, as the
//! CUDA member's: the settle kernel one thread a cell, one dispatch a sweep and
//! one to advance; the sampling kernel one thread a chain, its cells in order,
//! the reference's draws and the host's cut points. Fields are 64-bit, so it
//! takes every workload a job can give. The dispatches of a command buffer run
//! one after another, each seeing what the one before it wrote.
//!
//! The member says nothing of its own parity: the host certifies it. On a
//! platform other than macOS there is no device, and [`discover`] says so.
use super::device::{self, MeterSlot};
use super::member::{Chain, Member, Metered};
use super::settle::{Outcome, Workload};
use crate::quantum_fabric::RunMeter;
/// The kernels, in the Metal shading language.
pub const SOURCE: &str = r#"
#include <metal_stdlib>
using namespace metal;
struct Shape {
uint n;
uint chains;
uint bound_sweeps;
uint k0;
uint k1;
uint pad;
long cost;
};
inline int threshold(long h, long c) {
return h > c ? 1 : (h < -c ? -1 : 0);
}
inline long field_of(device const uint* rows, device const uint* cols, device const int* vals, device const long* fields, device const int* s, uint i) {
long h = fields[i];
for (uint k = rows[i]; k < rows[i + 1]; ++k) {
h += long(vals[k]) * long(s[cols[k]]);
}
return h;
}
kernel void settle_sweep(device const uint* rows [[buffer(0)]], device const uint* cols [[buffer(1)]], device const int* vals [[buffer(2)]],
device const long* fields [[buffer(3)]], constant Shape& sh [[buffer(4)]], device int* ring [[buffer(5)]],
device atomic_uint* status [[buffer(6)]], uint i [[thread_position_in_grid]]) {
if (i >= sh.n || atomic_load_explicit(&status[1], memory_order_relaxed) != 0u) return;
uint n = sh.n;
uint done = atomic_load_explicit(&status[0], memory_order_relaxed);
uint now = (done % 3u) * n, next = ((done + 1u) % 3u) * n, back = ((done + 2u) % 3u) * n;
int s = threshold(field_of(rows, cols, vals, fields, ring + now, i), sh.cost);
ring[next + i] = s;
if (s != ring[now + i]) atomic_fetch_or_explicit(&status[2], 1u, memory_order_relaxed);
if (s != ring[back + i]) atomic_fetch_or_explicit(&status[3], 1u, memory_order_relaxed);
}
kernel void settle_advance(constant Shape& sh [[buffer(4)]], device atomic_uint* status [[buffer(6)]], uint i [[thread_position_in_grid]]) {
if (i != 0u || atomic_load_explicit(&status[1], memory_order_relaxed) != 0u) return;
uint k = atomic_load_explicit(&status[0], memory_order_relaxed) + 1u;
atomic_store_explicit(&status[0], k, memory_order_relaxed);
if (atomic_load_explicit(&status[2], memory_order_relaxed) == 0u) {
atomic_store_explicit(&status[1], 1u, memory_order_relaxed);
} else if (k >= 2u && atomic_load_explicit(&status[3], memory_order_relaxed) == 0u) {
atomic_store_explicit(&status[1], 2u, memory_order_relaxed);
} else if (k >= sh.bound_sweeps) {
atomic_store_explicit(&status[1], 3u, memory_order_relaxed);
}
atomic_store_explicit(&status[2], 0u, memory_order_relaxed);
atomic_store_explicit(&status[3], 0u, memory_order_relaxed);
}
inline ulong splitmix(ulong s, ulong j) {
ulong x = s + j * 0x9E3779B97F4A7C15ul;
x ^= x >> 30; x *= 0xBF58476D1CE4E5B9ul;
x ^= x >> 27; x *= 0x94D049BB133111EBul;
return x ^ (x >> 31);
}
kernel void chains_in_order(device const uint* rows [[buffer(0)]], device const uint* cols [[buffer(1)]], device const int* vals [[buffer(2)]],
device const long* fields [[buffer(3)]], constant Shape& sh [[buffer(4)]], device const ulong* seeds [[buffer(7)]],
device const uint* sweeps [[buffer(8)]], device const ulong* table [[buffer(9)]], device int* cells [[buffer(10)]],
device const long* offsets [[buffer(11)]], uint c [[thread_position_in_grid]]) {
if (c >= sh.chains) return;
uint n = sh.n;
device int* s = cells + ulong(c) * n;
device int* spare = cells + ulong(sh.chains) * n + ulong(c) * n;
ulong seed = seeds[c];
for (uint k = sh.k0; k < sh.k1; ++k) {
uint kind = sweeps[k - sh.k0];
if (kind == 0u) {
for (uint i = 0; i < n; ++i) spare[i] = threshold(field_of(rows, cols, vals, fields, s, i), sh.cost);
for (uint i = 0; i < n; ++i) s[i] = spare[i];
} else {
for (uint i = 0; i < n; ++i) {
long h = field_of(rows, cols, vals, fields, s, i);
ulong u = splitmix(seed, ulong(k) * n + i + 1) >> 11;
ulong e = ulong(kind - 1u) + ulong(h) + ulong(offsets[i]);
s[i] = u < table[2 * e] ? -1 : (u < table[2 * e + 1] ? 0 : 1);
}
}
}
}
"#;
/// The uniform the kernels read, as the shading language lays it out.
#[cfg_attr(not(target_os = "macos"), allow(dead_code))]
#[repr(C)]
#[derive(Clone, Copy)]
struct Shape {
n: u32,
chains: u32,
bound_sweeps: u32,
k0: u32,
k1: u32,
pad: u32,
cost: i64,
}
/// One Metal device as a landscape member.
pub struct MetalLandscape {
gpu: sys::Gpu,
id: String,
name: String,
meter: MeterSlot,
}
impl std::fmt::Debug for MetalLandscape {
fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
f.debug_struct("MetalLandscape").field("id", &self.id).field("name", &self.name).finish()
}
}
/// Every Metal device, opened and compiled for, each a member or the reason it is not one.
pub fn discover() -> Vec<Result<MetalLandscape, String>> {
sys::discover(SOURCE)
}
/// [`discover`] with kernels other than [`SOURCE`]: how a test shows that a
/// fault in the kernels loses the grade it would otherwise earn on this device.
#[doc(hidden)]
pub fn discover_with(source: &str) -> Vec<Result<MetalLandscape, String>> {
sys::discover(source)
}
impl MetalLandscape {
/// Read a meter around each call: its reading stands beside the priced figures.
pub fn with_meter(mut self, meter: Box<dyn RunMeter + Send>) -> MetalLandscape {
self.meter = MeterSlot::new(meter);
self
}
/// The device's name, as the platform reports it.
pub fn name(&self) -> &str {
&self.name
}
fn chains_on_device(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String> {
let sampling = temperatures.iter().any(|&t| t > 0.0);
let reach = device::reach(w);
if sampling {
device::table_fits(&reach)?;
}
if chains.is_empty() {
return Ok(Vec::new());
}
if chains.iter().any(|c| c.start.len() != w.n()) {
return Err("a chain's start is not one cell a variable".into());
}
sys::chains(&self.gpu, w, &reach, chains, temperatures, device::window(w, &reach, sampling))
}
}
impl Member for MetalLandscape {
fn id(&self) -> String {
self.id.clone()
}
fn class(&self) -> &'static str {
"gpu"
}
fn label(&self) -> String {
device::label("gpu", &self.name, "Metal", "its cells in the reference's order")
}
fn takes(&self, w: &Workload, sampling: bool) -> Result<(), String> {
if sampling { device::table_fits(&device::reach(w)) } else { Ok(()) }
}
fn settle(&self, w: &Workload, start: &[i8], bound: u32) -> Result<Outcome, String> {
self.meter.around(|| sys::settle(&self.gpu, w, start, bound))
}
fn chains(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String> {
self.meter.around(|| self.chains_on_device(w, chains, temperatures))
}
fn meter(&self) -> Option<Metered> {
self.meter.last()
}
}
#[cfg(target_os = "macos")]
mod sys {
use super::{MetalLandscape, Shape};
use crate::quantum_landscape::device::{self, MeterSlot, Reach, SETTLE_BATCH};
use crate::quantum_landscape::member::Chain;
use crate::quantum_landscape::settle::{Outcome, Workload};
use objc2::rc::{Retained, autoreleasepool};
use objc2::runtime::ProtocolObject;
use objc2_foundation::{NSError, NSString};
use objc2_metal::{
MTLBuffer, MTLCommandBuffer, MTLCommandBufferStatus, MTLCommandEncoder, MTLCommandQueue, MTLCompileOptions, MTLComputeCommandEncoder, MTLComputePipelineState,
MTLCopyAllDevices, MTLDevice, MTLLanguageVersion, MTLLibrary, MTLMathMode, MTLResourceOptions, MTLSize,
};
use std::ptr::NonNull;
type Pipeline = Retained<ProtocolObject<dyn MTLComputePipelineState>>;
type Buffer = Retained<ProtocolObject<dyn MTLBuffer>>;
type Encoder = ProtocolObject<dyn MTLComputeCommandEncoder>;
pub struct Gpu {
device: Retained<ProtocolObject<dyn MTLDevice>>,
queue: Retained<ProtocolObject<dyn MTLCommandQueue>>,
settle_sweep: Pipeline,
settle_advance: Pipeline,
chains: Pipeline,
}
fn ns_error(e: &NSError) -> String {
e.localizedDescription().to_string()
}
pub fn discover(source: &str) -> Vec<Result<MetalLandscape, String>> {
autoreleasepool(|_| {
let devices = MTLCopyAllDevices();
if devices.count() == 0 {
return vec![Err("the platform lists no Metal device".into())];
}
(0..devices.count()).map(|k| open(devices.objectAtIndex(k), source)).collect()
})
}
fn open(device: Retained<ProtocolObject<dyn MTLDevice>>, source: &str) -> Result<MetalLandscape, String> {
let name = device.name().to_string();
let queue = device.newCommandQueue().ok_or_else(|| format!("{name}: the device gave no command queue"))?;
let options = MTLCompileOptions::new();
options.setLanguageVersion(MTLLanguageVersion((3 << 16) as usize));
options.setMathMode(MTLMathMode::Safe);
let library = device.newLibraryWithSource_options_error(&NSString::from_str(source), Some(&options)).map_err(|e| format!("{name}: the kernels did not compile: {}", ns_error(&e)))?;
let pipeline = |entry: &str| -> Result<Pipeline, String> {
let f = library.newFunctionWithName(&NSString::from_str(entry)).ok_or_else(|| format!("{name}: no entry point {entry}"))?;
device.newComputePipelineStateWithFunction_error(&f).map_err(|e| format!("{name}: {entry}: {}", ns_error(&e)))
};
let gpu = Gpu { settle_sweep: pipeline("settle_sweep")?, settle_advance: pipeline("settle_advance")?, chains: pipeline("chains_in_order")?, queue, device };
let id = device::id_of("gpu", "metal", &name, &format!("{:x}", gpu.device.registryID()));
Ok(MetalLandscape { gpu, id, name, meter: MeterSlot::default() })
}
fn upload<T: Copy>(gpu: &Gpu, data: &[T]) -> Result<Buffer, String> {
let bytes = std::mem::size_of_val(data);
let b = gpu.device.newBufferWithLength_options(bytes.max(16), MTLResourceOptions::StorageModeShared).ok_or_else(|| format!("a buffer of {bytes} 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)
}
fn read<T: Copy>(b: &Buffer, n: usize) -> Vec<T> {
let mut out: Vec<T> = Vec::with_capacity(n);
// 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);
}
out
}
fn submit(gpu: &Gpu, encode: impl FnOnce(&Encoder)) -> Result<(), String> {
let cb = gpu.queue.commandBuffer().ok_or("the queue gave no command buffer")?;
let encoder = cb.computeCommandEncoder().ok_or("the command buffer gave no compute encoder")?;
encode(&encoder);
encoder.endEncoding();
cb.commit();
cb.waitUntilCompleted();
if cb.status() != MTLCommandBufferStatus::Completed {
return Err(format!("the command buffer did not complete: {}", cb.error().map(|e| ns_error(&e)).unwrap_or_default()));
}
Ok(())
}
fn bind(e: &Encoder, b: &Buffer, index: usize) {
// SAFETY: each buffer is bound at the index its kernel parameter names, sized for what the kernel reads.
unsafe { e.setBuffer_offset_atIndex(Some(b), 0, index) };
}
fn shape(e: &Encoder, s: &Shape) {
// SAFETY: the encoder copies the bytes of `s`, laid out as the kernels' `Shape`.
unsafe { e.setBytes_length_atIndex(NonNull::from(s).cast(), std::mem::size_of::<Shape>(), 4) };
}
fn dispatch(e: &Encoder, p: &Pipeline, threads: u64) {
e.setComputePipelineState(p);
let size = |w: u64| MTLSize { width: w as usize, height: 1, depth: 1 };
e.dispatchThreadgroups_threadsPerThreadgroup(size(threads.div_ceil(64).max(1)), size(64));
}
fn graph(gpu: &Gpu, w: &Workload) -> Result<[Buffer; 4], String> {
let r = device::rows(w);
Ok([upload(gpu, &r.rows)?, upload(gpu, &r.cols)?, upload(gpu, &r.vals)?, upload(gpu, &r.fields)?])
}
pub fn settle(gpu: &Gpu, w: &Workload, start: &[i8], bound: u32) -> Result<Outcome, String> {
autoreleasepool(|_| {
let n = w.n();
let g = graph(gpu, w)?;
let mut init = vec![0i32; 3 * n];
for (i, &v) in start.iter().enumerate() {
init[i] = i32::from(v);
}
let ring = upload(gpu, &init)?;
let status = upload(gpu, &[0u32; 4])?;
let s = Shape { n: n as u32, chains: 0, bound_sweeps: bound, k0: 0, k1: 0, pad: 0, cost: w.cost() };
let (mut done, mut ended) = (0u32, 0u32);
while ended == 0 && done < bound {
submit(gpu, |e| {
for (k, b) in g.iter().enumerate() {
bind(e, b, k);
}
shape(e, &s);
bind(e, &ring, 5);
bind(e, &status, 6);
for _ in 0..SETTLE_BATCH.min(bound - done) {
dispatch(e, &gpu.settle_sweep, n as u64);
dispatch(e, &gpu.settle_advance, 1);
}
})?;
let st: Vec<u32> = read(&status, 4);
(done, ended) = (st[0], st[1]);
}
let all: Vec<i32> = read(&ring, 3 * n);
let third = (done % 3) as usize * n;
Ok(device::settled(w, done, ended, all[third..third + n].iter().map(|&v| v as i8).collect()))
})
}
pub fn chains(gpu: &Gpu, w: &Workload, reach: &Reach, chains: &[Chain], temperatures: &[f64], window: usize) -> Result<Vec<Outcome>, String> {
autoreleasepool(|_| {
let (n, count) = (w.n(), chains.len());
let g = graph(gpu, w)?;
let offsets = upload(gpu, &reach.offsets)?;
let mut init = vec![0i32; 2 * count * n];
for (c, chain) in chains.iter().enumerate() {
for (i, &v) in chain.start.iter().enumerate() {
init[c * n + i] = i32::from(v);
}
}
let cells = upload(gpu, &init)?;
drop(init);
let seeds = upload(gpu, &chains.iter().map(|c| c.seed).collect::<Vec<u64>>())?;
let mut k0 = 0usize;
while k0 < temperatures.len() {
let k1 = (k0 + window).min(temperatures.len());
let (sweeps, table) = device::window_sweeps(w, reach, &temperatures[k0..k1]);
let (sweeps, table) = (upload(gpu, &sweeps)?, upload(gpu, &table)?);
let s = Shape { n: n as u32, chains: count as u32, bound_sweeps: 0, k0: k0 as u32, k1: k1 as u32, pad: 0, cost: w.cost() };
submit(gpu, |e| {
for (k, b) in g.iter().enumerate() {
bind(e, b, k);
}
shape(e, &s);
bind(e, &seeds, 7);
bind(e, &sweeps, 8);
bind(e, &table, 9);
bind(e, &cells, 10);
bind(e, &offsets, 11);
dispatch(e, &gpu.chains, count as u64);
})?;
k0 = k1;
}
Ok(device::outcomes(w, &read::<i32>(&cells, count * n), count, temperatures))
})
}
}
#[cfg(not(target_os = "macos"))]
mod sys {
use super::MetalLandscape;
use crate::quantum_landscape::device::Reach;
use crate::quantum_landscape::member::Chain;
use crate::quantum_landscape::settle::{Outcome, Workload};
pub enum Gpu {}
pub fn discover(_: &str) -> Vec<Result<MetalLandscape, String>> {
vec![Err("this platform has no Metal; the member runs on macOS".into())]
}
pub fn settle(gpu: &Gpu, _: &Workload, _: &[i8], _: u32) -> Result<Outcome, String> {
match *gpu {}
}
pub fn chains(gpu: &Gpu, _: &Workload, _: &Reach, _: &[Chain], _: &[f64], _: usize) -> Result<Vec<Outcome>, String> {
match *gpu {}
}
}

207
wai/wai-quantum/src/quantum_landscape/reach.rs Normal file
View file

@ -0,0 +1,207 @@
//! The cut points a device's sampling kernel reads, and how many sweeps of them
//! one dispatch holds: what any member that takes its table from the host
//! (`extensions/quantum-landscape.md` §4.2) needs, whatever it runs on.
//!
//! A table holds the fields a workload's cells can reach and no others
//! ([`reach`]): at most `2·nonzeros + n` entries a sweep, however wide the
//! fields. A job counts the work of building one for each sampled sweep of
//! each call in its operations ([`super::run::Plan::tables_bound`]), and the
//! memory of holding one in its bytes ([`super::run::job_bytes`]).
use super::sample::cuts;
use super::settle::Workload;
/// The most cut-point entries one sweep's table may hold. No job comes near it:
/// a table holds at most `2·nonzeros + n` entries ([`reach`]), and a program of
/// [`super::MAX_PROGRAM_BYTES`] has room for under 1.7 million. Only a workload
/// made by hand can be declined for this.
pub const MAX_TABLE_ENTRIES: u64 = 1 << 22;
/// The work one chain does in one dispatch, at most: cells × (couplings + 2) × sweeps.
pub const WINDOW_WORK: u64 = 1 << 20;
/// The cut-point entries one dispatch's tables hold, at most, unless one sweep's alone is more.
pub const WINDOW_ENTRIES: u64 = 1 << 14;
/// The fields a workload's cells can reach, which are all a sweep's table needs:
/// cell `i` only ever sees `b_i − degree_i` to `b_i + degree_i`. Where those
/// intervals overlap they are one span, so a table is no larger than `2B + 1`
/// entries for a field bound `B`, nor than `2·nonzeros + n`, however wide the
/// fields are. The work of building it is therefore never more than twice the
/// sweep it serves, and a sampling kernel's table cannot outgrow its job.
pub struct Reach {
/// What added to a cell's field, wrapping, gives its entry in a sweep's table.
pub offsets: Vec<i64>,
/// The spans, lowest first: first field and last field.
spans: Vec<(i64, i64)>,
/// The entries of one sweep's table.
pub entries: u64,
}
/// What the cells of a workload can reach.
pub fn reach(w: &Workload) -> Reach {
let n = w.n();
let zero = vec![0i8; n];
let mut cells: Vec<(i64, i64, usize)> = (0..n)
.map(|i| {
let (b, d) = (w.field(&zero, i), i64::try_from(w.degree(i)).unwrap_or(i64::MAX));
(b.saturating_sub(d), b.saturating_add(d), i)
})
.collect();
cells.sort_unstable();
let (mut offsets, mut spans, mut entries) = (vec![0i64; n], Vec::<(i64, i64)>::new(), 0u64);
let mut first = 0u64;
for (lo, hi, i) in cells {
match spans.last_mut() {
Some(last) if lo <= last.1.saturating_add(1) => {
entries += hi.saturating_sub(last.1).max(0) as u64;
last.1 = last.1.max(hi);
}
_ => {
first = entries;
entries += hi.saturating_sub(lo) as u64 + 1;
spans.push((lo, hi));
}
}
let span = spans[spans.len() - 1];
offsets[i] = (first as i64).wrapping_sub(span.0);
}
Reach { offsets, spans, entries }
}
/// One sweep's cut points at temperature `t`, `(a, b)` for every field the
/// cells can reach: the reference's own ([`cuts`]).
pub fn cut_pairs(cost: i64, t: f64, reach: &Reach, out: &mut Vec<u64>) {
let beta = 1.0 / t;
for &(lo, hi) in &reach.spans {
for h in lo..=hi {
let (a, b) = cuts(h, cost, beta);
out.extend([a, b]);
}
}
}
/// Whether a sampling run's table fits: one sweep's cut points within [`MAX_TABLE_ENTRIES`].
pub fn table_fits(reach: &Reach) -> Result<(), String> {
if reach.entries > MAX_TABLE_ENTRIES {
return Err(format!("its cells reach {} different fields, and a sweep's cut points are at most {MAX_TABLE_ENTRIES} entries here", reach.entries));
}
Ok(())
}
/// How many sweeps one dispatch runs: within [`WINDOW_WORK`] for each chain,
/// and within [`WINDOW_ENTRIES`] of cut points, and at least one.
pub fn window(w: &Workload, reach: &Reach, sampling: bool) -> usize {
let n = w.n() as u64;
let per_sweep = n * (w.nonzeros() / n.max(1) + 2);
let mut window = (WINDOW_WORK / per_sweep.max(1)).max(1);
if sampling {
window = window.min((WINDOW_ENTRIES / reach.entries.max(1)).max(1));
}
window as usize
}
/// The sweeps of a window as a device reads them, and their tables: 0 for a
/// sweep at temperature 0, else 1 + where its table starts, in entries.
pub fn window_sweeps(w: &Workload, reach: &Reach, temperatures: &[f64]) -> (Vec<u32>, Vec<u64>) {
let mut sweeps = Vec::with_capacity(temperatures.len());
let mut table = Vec::new();
for &t in temperatures {
if t > 0.0 {
sweeps.push(1 + (table.len() / 2) as u32);
cut_pairs(w.cost(), t, reach, &mut table);
} else {
sweeps.push(0);
}
}
(sweeps, table)
}
#[cfg(test)]
mod tests {
use super::super::energy::Landscape;
use super::*;
#[test]
fn the_table_is_the_reference_s_cut_points() {
let mut l = Landscape::new(2);
l.couple(0, 1, 1).set_field(0, 2).set_cost(1);
let r = reach(&l.lower());
let mut t = Vec::new();
cut_pairs(1, 0.7, &r, &mut t);
// −1 to 3 are one span: cell 1 reaches −1 to 1, cell 0 reaches 1 to 3.
assert_eq!(r.entries, 5);
assert_eq!(t.len(), 10);
assert_eq!((t[2], t[3]), cuts(0, 1, 1.0 / 0.7));
assert_eq!(r.offsets, vec![1, 1]);
}
/// A table holds the fields the cells can reach, whatever the field bound:
/// the job that once asked for four million entries a sweep has six.
#[test]
fn a_table_holds_only_what_the_cells_can_reach() {
let mut l = Landscape::new(2);
l.couple(0, 1, 1).set_field(0, 2_097_149);
let w = l.lower();
assert_eq!(w.field_bound(), 2_097_150);
let r = reach(&w);
assert_eq!(r.entries, 6);
// Every field a cell can reach has its own entry, and nothing else does.
let mut seen = std::collections::BTreeSet::new();
for s0 in -1i8..=1 {
for s1 in -1i8..=1 {
for i in 0..2 {
let e = w.field(&[s0, s1], i).wrapping_add(r.offsets[i]);
assert!((0..r.entries as i64).contains(&e));
seen.insert((e, w.field(&[s0, s1], i)));
}
}
}
assert_eq!(seen.len(), 6);
assert_eq!(seen.iter().map(|p| p.0).collect::<std::collections::BTreeSet<_>>().len(), 6);
}
/// `2·nonzeros + n` bounds a table, and `2B + 1`: whichever is smaller.
#[test]
fn a_table_is_bounded_by_the_smaller_of_the_two_forms() {
let (n, mut l) = (64, Landscape::new(64));
for i in 0..n {
l.couple(i, (i + 1) % n, 1);
}
let dense = l.clone();
for i in 0..n {
l.set_field(i, 7 * i as i64);
}
let (spread, flat) = (l.lower(), dense.lower());
assert_eq!(reach(&spread).entries, 2 * spread.nonzeros() + n as u64);
assert_eq!(reach(&flat).entries, 2 * flat.field_bound() as u64 + 1);
}
/// The largest program a job can read cannot ask for more than a table holds.
#[test]
fn no_job_can_outgrow_a_table() {
// The shortest line that couples two cells is ten bytes, `agree 0 1` and its end.
let pairs = (crate::quantum_landscape::MAX_PROGRAM_BYTES / 10) as u64;
assert!(4 * pairs + crate::quantum_landscape::MAX_VARS as u64 <= MAX_TABLE_ENTRIES);
}
/// A window ends at the work a chain may do in a dispatch, or at the cut
/// points it may hold, whichever comes first, and a sweep at 0 needs no table.
#[test]
fn a_window_ends_at_its_work_or_at_its_table() {
let (n, mut l) = (64, Landscape::new(64));
for i in 0..n {
l.couple(i, (i + 1) % n, 1);
l.set_field(i, 7 * i as i64);
}
let w = l.lower();
let r = reach(&w);
assert_eq!(r.entries, 320);
assert_eq!(window(&w, &r, true), (WINDOW_ENTRIES / 320) as usize);
assert_eq!(window(&w, &r, false), (WINDOW_WORK / (64 * 4)) as usize);
let (sweeps, table) = window_sweeps(&w, &r, &[0.5, 0.0, 0.25]);
assert_eq!(sweeps, vec![1, 0, 321]);
assert_eq!(table.len(), 2 * 2 * 320);
}
}

218
wai/wai-quantum/src/quantum_landscape/run.rs
View file

@ -103,7 +103,7 @@ impl Plan {
/// The kernel a plan runs: a settle runs the settle kernel, and a quench
/// and a sampling run run chains, a quench's every sweep at 0.
fn kernel(self) -> &'static str {
pub(crate) fn kernel(self) -> &'static str {
match self {
Plan::Settle { .. } => KERNEL_SETTLE,
_ => KERNEL_SAMPLE,
@ -143,15 +143,28 @@ impl Plan {
per_sweep.saturating_mul(self.restarts()).saturating_mul(u64::from(self.sweeps()))
}
/// The most operations one run of this plan can take: its sweeps, the
/// energy of every start's state on the lowered program, the energy of the
/// answer on the program as declared (and on its weights), and its checks.
/// The most operations a device's cut points for this plan can take: for
/// each call of one batch of starts, one sweep's table for every sweep above
/// 0, at most `4P + n` entries ([`super::reach`]). Nothing for a plan with
/// no sweep above 0.
pub fn tables_bound(self, p: &Program) -> u64 {
let Plan::Anneal { restarts, sweeps, .. } = self else { return 0 };
let n = p.landscape.len() as u64;
let batch = (CHUNK_CELLS as u64 / n.max(1)).clamp(1, CHUNK_CHAINS as u64);
let calls = u64::from(restarts).div_ceil(batch);
(4 * p.landscape.pairs() + n).saturating_mul(u64::from(sweeps)).saturating_mul(calls)
}
/// The most operations one run of this plan can take: its sweeps, a
/// device's cut points for them, the energy of every start's state on the
/// lowered program, the energy of the answer on the program as declared
/// (and on its weights), and its checks.
pub fn run_bound(self, p: &Program) -> u64 {
let n = p.landscape.len() as u64;
let energies = (2 * p.landscape.pairs() + n).saturating_mul(self.restarts());
let declared = if p.declared.is_some() { 2 * n * n } else { n * n };
let checks: u64 = p.notes.checks.iter().map(|c| check_items(c) as u64).sum();
self.operations_bound(p).saturating_add(energies).saturating_add(declared).saturating_add(checks)
self.operations_bound(p).saturating_add(self.tables_bound(p)).saturating_add(energies).saturating_add(declared).saturating_add(checks)
}
}
@ -198,8 +211,9 @@ pub fn job_operations(p: &Program, runs: &[(Plan, u64)]) -> u64 {
/// The most memory a job can hold, known before any of it runs: the program's
/// text, its couplings as declared and as rows, its weights where it gives
/// them, one batch of chains, `lines` lines of a replay, and a margin for
/// everything of fixed size.
/// them, one sweep's cut points as a device member holds them (at most
/// `2·nonzeros + n` of 16 bytes: `device::reach`), one batch of chains, `lines`
/// lines of a replay, and a margin for everything of fixed size.
pub fn job_bytes(p: &Program, lines: u64) -> u64 {
let n = p.landscape.len() as u64;
let nonzeros = 2 * p.landscape.pairs();
@ -209,6 +223,7 @@ pub fn job_bytes(p: &Program, lines: u64) -> u64 {
.saturating_add(n * n)
.saturating_add(weights)
.saturating_add(16 * nonzeros + 64 * n)
.saturating_add(16 * (2 * nonzeros + n))
.saturating_add(batch)
.saturating_add(lines.saturating_mul(256))
.saturating_add(1 << 20)
@ -279,13 +294,27 @@ fn replay_digest(p: &Program, seed: u64, plan: Plan, state: &[i8], energy: i64)
}
/// What the kernels leave, over every start of a run.
struct Found {
#[derive(Clone)]
pub(crate) struct Found {
state: Vec<i8>,
energy: i64,
accumulates: u64,
draws: u64,
sweeps: u64,
term: Option<Term>,
/// A meter's reading over the member's calls for this run alone; none
/// where a call covered other runs too, or a call gave no reading.
metered: Option<member::Metered>,
}
/// Sum the readings of a run's calls, figures and windows: none if any call
/// gave none, or two calls' readings are not one reading.
fn add_reading(total: &mut Option<Option<member::Metered>>, reading: Option<member::Metered>) {
*total = match (total.take(), reading) {
(None, r) => Some(r),
(Some(Some(a)), Some(b)) => Some(a.and(b)),
_ => Some(None),
};
}
fn execute(job: &Job, plan: Plan, seed: u64, c: &Certified<'_>) -> Result<Found, Refusal> {
@ -295,10 +324,12 @@ fn execute(job: &Job, plan: Plan, seed: u64, c: &Certified<'_>) -> Result<Found,
let refused = |e: String| Refusal::general(format!("The member {} refused or failed: {e}.", m.id()));
match plan {
Plan::Settle { sweeps } => {
// The host gives a member only what it takes, however it was reached.
m.takes(w, false).map_err(refused)?;
let o = m.settle(w, &vec![0i8; n], sweeps).map_err(refused)?;
member::check_settle(&o, w, sweeps).map_err(refused)?;
let energy = job.program.landscape.energy(&o.state);
Ok(Found { energy, accumulates: o.accumulates, draws: o.draws, sweeps: u64::from(o.sweeps), term: Some(o.term), state: o.state })
Ok(Found { energy, accumulates: o.accumulates, draws: o.draws, sweeps: u64::from(o.sweeps), term: Some(o.term), state: o.state, metered: m.meter() })
}
Plan::Quench { restarts, .. } | Plan::Anneal { restarts, .. } => run_chains(job, &plan.temperatures(), restarts, seed, c),
}
@ -306,19 +337,34 @@ fn execute(job: &Job, plan: Plan, seed: u64, c: &Certified<'_>) -> Result<Found,
/// Every start of a run, a batch at a time, keeping the first state of lowest energy.
fn run_chains(job: &Job, temperatures: &[f64], restarts: u32, seed: u64, c: &Certified<'_>) -> Result<Found, Refusal> {
Ok(chains_for_seeds(job, temperatures, restarts, &[seed], c, true)?.remove(0))
}
/// One seed's run as its batches come back: the best state and its energy, the
/// accumulates, the draws, and the readings of its calls.
type SeedRun = (Option<(Vec<i8>, i64)>, u64, u64, Option<Option<member::Metered>>);
/// The runs of several seeds, their starts handed to the member together a
/// batch at a time, so that a device can run them side by side. Each seed's run
/// is the one it would be alone: the same chains, reduced in the same order.
/// A reading is kept only where every call covered one run (`alone`).
fn chains_for_seeds(job: &Job, temperatures: &[f64], restarts: u32, seeds: &[u64], c: &Certified<'_>, alone: bool) -> Result<Vec<Found>, Refusal> {
let w = &job.workload;
let n = w.n();
let m = c.member();
let refused = |e: String| Refusal::general(format!("The member {} refused or failed: {e}.", m.id()));
let batch = (CHUNK_CELLS / n.max(1)).clamp(1, CHUNK_CHAINS) as u64;
let (mut accumulates, mut draws) = (0u64, 0u64);
let mut best: Option<(Vec<i8>, i64)> = None;
let mut r = 0u64;
while r < u64::from(restarts) {
let end = (r + batch).min(u64::from(restarts));
let chains: Vec<Chain> = (r..end)
m.takes(w, temperatures.iter().any(|&t| t > 0.0)).map_err(refused)?;
let batch = (CHUNK_CELLS / n.max(1)).clamp(1, CHUNK_CHAINS);
let r = u64::from(restarts);
// Every (seed, start) in order: seed after seed, start after start.
let total = seeds.len() as u64 * r;
let mut found: Vec<SeedRun> = vec![(None, 0, 0, None); seeds.len()];
let mut at = 0u64;
while at < total {
let end = (at + batch as u64).min(total);
let chains: Vec<Chain> = (at..end)
.map(|k| {
let s = sample::chain_seed(seed, k);
let s = sample::chain_seed(seeds[(k / r) as usize], k % r);
Chain { start: sample::start(s, n), seed: s }
})
.collect();
@ -326,20 +372,54 @@ fn run_chains(job: &Job, temperatures: &[f64], restarts: u32, seed: u64, c: &Cer
if outcomes.len() != chains.len() {
return Err(refused(format!("it returned {} chains for {}", outcomes.len(), chains.len())));
}
for o in outcomes {
let reading = m.meter();
let one_run = alone || (at / r == (end - 1) / r && at.is_multiple_of(r) && end.is_multiple_of(r));
for (k, o) in (at..end).zip(outcomes) {
member::check_chain(&o, w, temperatures).map_err(refused)?;
accumulates += o.accumulates;
draws += o.draws;
let f = &mut found[(k / r) as usize];
f.1 += o.accumulates;
f.2 += o.draws;
let e = w.energy(&o.state);
if best.as_ref().is_none_or(|b| e < b.1) {
best = Some((o.state, e));
if f.0.as_ref().is_none_or(|b| e < b.1) {
f.0 = Some((o.state, e));
}
}
r = end;
for s in (at / r)..=((end - 1) / r) {
add_reading(&mut found[s as usize].3, if one_run { reading.clone() } else { None });
}
at = end;
}
let (state, _) = best.expect("a run has at least one start");
let energy = job.program.landscape.energy(&state);
Ok(Found { state, energy, accumulates, draws, sweeps: u64::from(restarts) * temperatures.len() as u64, term: None })
Ok(found
.into_iter()
.map(|(best, accumulates, draws, metered)| {
let (state, _) = best.expect("a run has at least one start");
let energy = job.program.landscape.energy(&state);
Found { state, energy, accumulates, draws, sweeps: u64::from(restarts) * temperatures.len() as u64, term: None, metered: metered.flatten() }
})
.collect())
}
/// Every seed's run of a plan, the starts of several seeds handed to the member
/// together. A settle has no randomness: it runs once, and every seed has its answer.
pub(crate) fn execute_seeds(job: &Job, plan: Plan, seeds: &[u64], c: &Certified<'_>) -> Result<Vec<Found>, Refusal> {
match plan {
Plan::Settle { .. } => {
let mut f = execute(job, plan, seeds.first().copied().unwrap_or(DEFAULT_SEED), c)?;
// One call answered every seed, so its reading stands for none of them alone.
if seeds.len() > 1 {
f.metered = None;
}
Ok(vec![f; seeds.len()])
}
Plan::Quench { restarts, .. } | Plan::Anneal { restarts, .. } => chains_for_seeds(job, &plan.temperatures(), restarts, seeds, c, seeds.len() == 1),
}
}
/// How many seeds go to a member together: as many as one batch of chains holds.
pub(crate) fn seeds_a_batch(job: &Job, plan: Plan) -> usize {
let n = job.workload.n().max(1);
let chains = (CHUNK_CELLS / n).clamp(1, CHUNK_CHAINS) as u64;
(chains / plan.restarts().max(1)).max(1) as usize
}
fn truth_json(t: &Truth) -> Value {
@ -447,10 +527,15 @@ fn declared_json(real: &Weights, p: &Program, state: &[i8], exact: &Exact) -> (V
/// Run a program on a member: the answer of `extensions/quantum-landscape.md` §3.
pub fn run_on(c: &Certified<'_>, job: &Job, seed: u64, temperature: Option<f64>) -> Result<Value, Refusal> {
let plan = Plan::of(&job.program, temperature)?;
let found = execute(job, plan, seed, c)?;
Ok(answer_of(c, job, plan, seed, &found))
}
/// The answer of `extensions/quantum-landscape.md` §3, from what a run's kernels left.
pub(crate) fn answer_of(c: &Certified<'_>, job: &Job, plan: Plan, seed: u64, found: &Found) -> Value {
let (p, exact) = (&job.program, &job.exact);
let m = c.member();
let plan = Plan::of(p, temperature)?;
let found = execute(job, plan, seed, c)?;
let n = p.landscape.len();
let reached = exact.truth.as_ref().map(|t| found.energy as f64 == t.lowest);
let (declared, declared_reached, declared_energy) = match &p.declared {
@ -461,7 +546,7 @@ pub fn run_on(c: &Certified<'_>, job: &Job, seed: u64, temperature: Option<f64>)
None => (None, None, None),
};
let checks = grade_all(p, &found.state);
let (verdict, summary) = judge(&found, exact, reached, declared_reached, declared_energy, &checks);
let (verdict, summary) = judge(found, exact, reached, declared_reached, declared_energy, &checks);
let groups: Vec<Value> = [(&p.notes.plus, 1i8), (&p.notes.minus, -1), (&p.notes.zero, 0)]
.iter()
.filter_map(|(name, value)| name.as_ref().map(|name| json!({ "name": name, "items": (0..n).filter(|&i| found.state[i] == *value).collect::<Vec<_>>() })))
@ -478,13 +563,18 @@ pub fn run_on(c: &Certified<'_>, job: &Job, seed: u64, temperature: Option<f64>)
};
let rests = job.workload.rests(&found.state);
let mut joules_statement = price::priced(found.accumulates, plan.mode(), found.sweeps);
joules_statement["metered"] = match m.meter() {
Some(r) => json!({ "joules_micro": r.joules_micro, "class": r.class }),
joules_statement["metered"] = match &found.metered {
Some(r) => json!({ "joules_micro": r.joules_micro, "window_us": r.window_us, "class": r.class, "scope": r.scope }),
None => Value::Null,
};
Ok(json!({
// The label claims a reading only for a run that has one, and says what it covers.
let label = match &found.metered {
Some(r) => format!("{} Beside the priced figures is a reading of {}.", m.label(), r.scope.as_deref().unwrap_or("this run alone")),
None => m.label(),
};
json!({
"protocol": PROTOCOL,
"label": m.label(),
"label": label,
"member": member::describe(c, plan.kernel()),
"program_digest": blake3::hash(p.source.as_bytes()).to_hex().to_string(),
"title": p.notes.title,
@ -521,7 +611,7 @@ pub fn run_on(c: &Certified<'_>, job: &Job, seed: u64, temperature: Option<f64>)
"joules": joules_statement,
"state_hash": format!("{:016x}", state_hash(&found.state)),
"replay": replay_digest(p, seed, plan, &found.state, found.energy),
}))
})
}
// ── the temperature sweep ────────────────────────────────────────────────────────────────────
@ -558,10 +648,15 @@ pub fn sweep_on(c: &Certified<'_>, job: &Job, temperatures: &[f64], seeds: u32,
let plans = sweep_plans(p, temperatures, seeds)?;
let default = p.landscape.default_temperature();
let mut rows = Vec::with_capacity(plans.len());
let mut kernels = Vec::new();
for (t, plan, runs) in plans {
let (mut best, mut sum, mut reached, mut passed, mut accumulates, mut sweeps) = (i64::MAX, 0i128, 0u32, 0u32, 0u64, 0u64);
for k in 0..runs {
let found = execute(job, plan, first_seed.wrapping_add(u64::from(k)), c)?;
let seeds: Vec<u64> = (0..runs).map(|k| first_seed.wrapping_add(u64::from(k))).collect();
let mut founds = Vec::with_capacity(seeds.len());
for group in seeds.chunks(seeds_a_batch(job, plan)) {
founds.extend(execute_seeds(job, plan, group, c)?);
}
for found in founds {
best = best.min(found.energy);
sum += i128::from(found.energy);
accumulates = found.accumulates;
@ -571,10 +666,13 @@ pub fn sweep_on(c: &Certified<'_>, job: &Job, temperatures: &[f64], seeds: u32,
}
let figures = price::priced(accumulates, plan.mode(), sweeps);
let left_out = figures["left_out"]["pj"].as_f64().unwrap_or(0.0);
kernels.push(plan.kernel());
rows.push(json!({
"temperature": t,
"is_default": t == default && t > 0.0,
"mode": plan.name(),
"kernel": plan.kernel(),
"parity": c.parity(plan.kernel()).label(),
"runs": runs,
"reached_lowest": exact.truth.as_ref().map(|_| reached),
"passed_checks": (!p.notes.checks.is_empty()).then_some(passed),
@ -588,7 +686,7 @@ pub fn sweep_on(c: &Certified<'_>, job: &Job, temperatures: &[f64], seeds: u32,
Ok(json!({
"protocol": PROTOCOL,
"label": m.label(),
"member": { "id": m.id(), "class": m.class() },
"member": member::describe_all(c, &kernels),
"lowest": exact.truth.as_ref().map(|t| t.lowest),
"default_temperature": default,
"rows": rows,
@ -612,8 +710,12 @@ pub fn compare_on(c: &Certified<'_>, job: &Job, seeds: u32, first_seed: u64, tem
for (name, temperature) in [(sampling.as_str(), t0), ("emulator, deterministic", 0.0)] {
let plan_here = if temperature > 0.0 { plan } else { Plan::Quench { restarts, sweeps } };
let mut t = (name.to_string(), false, Vec::new(), 0u32, 0u64, 0u64);
for k in 0..seeds {
let found = execute(job, plan_here, first_seed.wrapping_add(u64::from(k)), c)?;
let all: Vec<u64> = (0..seeds).map(|k| first_seed.wrapping_add(u64::from(k))).collect();
let mut founds = Vec::with_capacity(all.len());
for group in all.chunks(seeds_a_batch(job, plan_here)) {
founds.extend(execute_seeds(job, plan_here, group, c)?);
}
for found in founds {
t.2.push(found.energy);
t.3 += u32::from(check(&found.state));
t.5 += found.accumulates;
@ -644,6 +746,7 @@ pub fn compare_on(c: &Certified<'_>, job: &Job, seeds: u32, first_seed: u64, tem
json!({
"solver": t.0,
"classical": t.1,
"parity": (!t.1).then(|| c.parity(member::KERNEL_SAMPLE).label()),
"runs": seeds,
"reached_target": t.2.iter().filter(|&&v| v == target).count(),
"passed_checks": (!p.notes.checks.is_empty()).then_some(t.3),
@ -656,7 +759,7 @@ pub fn compare_on(c: &Certified<'_>, job: &Job, seeds: u32, first_seed: u64, tem
Ok(json!({
"protocol": PROTOCOL,
"label": c.member().label(),
"member": { "id": c.member().id(), "class": c.member().class() },
"member": member::describe_all(c, &[member::KERNEL_SAMPLE]),
"target": target,
"target_is_exact": exact_target,
"arrangements": exact.truth.as_ref().map(|t| t.arrangements),
@ -740,7 +843,7 @@ mod tests {
CpuReference.chains(w, chains, temperatures)
}
fn meter(&self) -> Option<member::Metered> {
self.meter.then(|| member::Metered { joules_micro: 12, class: json!({ "class": "CalibratedInstrument", "calibration_ref": "a test" }) })
self.meter.then(|| member::Metered { joules_micro: 12, window_us: 5, class: json!({ "class": "CalibratedInstrument", "calibration_ref": "a test" }), scope: None })
}
}
@ -755,8 +858,9 @@ mod tests {
assert_eq!(board["member"]["id"], "fpga:test-board");
assert_eq!(board["member"]["class"], "fpga");
assert_eq!(board["member"]["parity"], "bit-exact", "its kernels give the reference's bits, and the host found so");
assert_eq!(board["label"], "A run on a test board.");
assert_eq!(board["label"], "A run on a test board. Beside the priced figures is a reading of this run alone.", "the label claims a reading where the run has one");
assert_eq!(board["joules"]["metered"]["joules_micro"], 12);
assert_eq!(board["joules"]["metered"]["window_us"], 5);
assert_eq!(board["joules"]["kind"], "priced", "a reading stands beside the price");
assert_eq!(board["replay"], reference["replay"]);
assert_eq!(board["state"], reference["state"]);
@ -764,6 +868,36 @@ mod tests {
let unmetered = Board { meter: false };
let cold = run_on(&named(&unmetered).unwrap(), &j, 5, Some(0.0)).unwrap();
assert_eq!(cold["member"]["parity"], "uncertified", "asked for by name, it is not certified");
assert_eq!(cold["label"], "A run on a test board.", "and with no reading, the label claims none");
}
/// A run of several calls adds up their figures and their windows; calls
/// whose readings are not one reading give the run none; and a settle that
/// answered several seeds from one call gives none of them a reading.
#[test]
fn a_run_s_reading_is_its_calls_added_up() {
use crate::quantum_landscape::member::Metered;
let class = |ppm: u64, window: u64| json!({ "energy_provenance": "OnChipCounter", "energy_uncertainty": { "relative_ppm": ppm, "window_us": window }, "energy_counter_filtering": "Unknown" });
let reading = |uj: u64, window: u64, ppm: u64| Some(Metered { joules_micro: uj, window_us: window, class: class(ppm, window), scope: Some("a rail".into()) });
let mut total = None;
for (uj, window, ppm) in [(2_000_000, 1_000_000, 50_000), (3_000_000, 1_500_000, 70_000), (1_000_000, 800_000, 60_000)] {
add_reading(&mut total, reading(uj, window, ppm));
}
let r = total.flatten().unwrap();
assert_eq!((r.joules_micro, r.window_us), (6_000_000, 3_300_000));
assert_eq!(r.class["energy_uncertainty"], json!({ "relative_ppm": 70_000, "window_us": 3_300_000 }));
let mut other = None;
add_reading(&mut other, reading(1, 1, 1));
add_reading(&mut other, Some(Metered { scope: None, ..reading(1, 1, 1).unwrap() }));
assert_eq!(other.flatten(), None, "two scopes are not one reading");
// A settle replayed over seeds: one call, so no seed's run has a reading of its own.
let j = job("vars 3\nagree 0 1\nfield 0 2\nminimize\n");
let metered = Board { meter: true };
let c = certify(&metered).unwrap();
let plan = Plan::of(&j.program, None).unwrap();
let found = execute_seeds(&j, plan, &[1, 2, 3], &c).unwrap();
assert!(found.iter().all(|f| f.metered.is_none()));
assert!(execute_seeds(&j, plan, &[1], &c).unwrap()[0].metered.is_some());
}
/// A member cannot pass off a broken answer: the host checks the state and works out its energy.

634
wai/wai-quantum/src/quantum_landscape/webgpu.rs Normal file
View file

@ -0,0 +1,634 @@
//! A GPU member of the landscape fabric, through WebGPU: the kernels of
//! `landscape.wgsl` on a platform GPU (Metal, Vulkan or DX12), natively.
//!
//! - **The settle kernel** runs one invocation a cell, one dispatch a sweep,
//! with a one-invocation dispatch after each that records whether the sweep
//! came to rest, closed a two-cycle or reached the bound. Its fields are
//! whole numbers in 32-bit lanes, so it gives the reference's bits.
//! - **The sampling kernel** runs one invocation a chain, its cells in order,
//! the reference's draws (SplitMix64 by counter, in pairs of 32-bit words)
//! and the reference's cut points, which the host works out for the fields
//! the cells can reach ([`device::reach`]) and ships as a table a window of
//! sweeps at a time, waiting for each window before it sends the next.
//! So it gives the reference's bits too, and many chains run side by side:
//! the starts of a run, and the seeds of a sweep, a comparison or a replay.
//! - **The coloured kernel** ([`WebGpuLandscape::coloured`]) moves a colour
//! class of cells at once, every chain in one dispatch: no two cells of a
//! class are coupled, but the order is not the reference's, so its parity
//! can only be statistical, and it has to pass the battery to get that.
//!
//! A workload whose field bound does not fit a 32-bit lane is declined before
//! anything runs, with the reason ([`Member::takes`]): the host then runs it on
//! the next member it was offered, and the answer says so. A cost above every
//! field is held at the lane's top, where it does what it does at its own size.
//!
//! The member says nothing of its own parity: the host certifies it
//! ([`super::member::certify`]). It is native only: a page's WebGPU is
//! asynchronous, and the member's calls block.
use super::device::{self, MeterSlot, Reach, SETTLE_BATCH};
use super::member::{Chain, Member, Metered};
use super::settle::{Outcome, Term, Workload};
use crate::quantum_fabric::RunMeter;
use crate::quantum_fabric_webgpu::blocking::resolve;
use std::borrow::Cow;
use std::sync::atomic::{AtomicU64, Ordering};
use std::sync::{Arc, Mutex};
use std::time::Duration;
/// The kernels.
pub const SHADER: &str = include_str!("landscape.wgsl");
/// How long one wait on the device may take: one window of sweeps, or one
/// batch of a settle, never a whole run.
const POLL_TIMEOUT: Duration = Duration::from_secs(300);
/// The largest storage binding this member uses, whatever an adapter offers:
/// its cell indices are 32-bit words, and at 4 GiB a binding's size no longer
/// fits the word a shader's bounds check is made from.
const MAX_BINDING: u64 = 1 << 31;
/// The binding limit of a device that offers `offered`.
fn binding_limit(offered: u64) -> u64 {
offered.min(MAX_BINDING)
}
/// How the sampling kernel orders its cells.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub enum Order {
/// Each chain's cells one at a time, in order: the reference's order.
InOrder,
/// A colour class at a time: statistical at most.
Coloured,
}
#[derive(Default)]
struct Sink {
errors: Vec<String>,
lost: Option<String>,
}
#[derive(Clone)]
struct Pipelines {
settle_sweep: wgpu::ComputePipeline,
settle_advance: wgpu::ComputePipeline,
in_order: wgpu::ComputePipeline,
coloured: wgpu::ComputePipeline,
copy: wgpu::ComputePipeline,
}
/// One WebGPU adapter as a landscape member.
pub struct WebGpuLandscape {
device: wgpu::Device,
queue: wgpu::Queue,
pipelines: Pipelines,
id: String,
class: &'static str,
adapter: String,
backend: &'static str,
order: Order,
max_binding: u64,
/// How long one wait may take, in nanoseconds.
wait_at_most: AtomicU64,
sink: Arc<Mutex<Sink>>,
meter: MeterSlot,
}
impl std::fmt::Debug for WebGpuLandscape {
fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
f.debug_struct("WebGpuLandscape").field("id", &self.id).field("adapter", &self.adapter).field("order", &self.order).finish()
}
}
fn lock<T>(m: &Mutex<T>) -> std::sync::MutexGuard<'_, T> {
m.lock().unwrap_or_else(|p| p.into_inner())
}
fn backend_label(b: wgpu::Backend) -> Option<&'static str> {
match b {
wgpu::Backend::Metal => Some("metal"),
wgpu::Backend::Vulkan => Some("vulkan"),
wgpu::Backend::Dx12 => Some("dx12"),
_ => None,
}
}
/// A member id from what the adapter reports ([`device::id_of`]): `gpu:wgpu-<backend>.…`,
/// so that it never reads as a member of the native backend of the same name.
fn member_id(class: &str, backend: &str, info: &wgpu::AdapterInfo) -> String {
device::id_of(class, &format!("wgpu-{backend}"), &info.name, &format!("{}/{}/{}/{}", info.vendor, info.device, info.driver, info.driver_info))
}
/// Every adapter of the platform's primary GPU interfaces, opened, each a
/// member or the reason it is not one.
pub fn discover() -> Vec<Result<WebGpuLandscape, String>> {
discover_with(SHADER)
}
/// [`discover`] with kernels other than [`SHADER`]: how a test shows that a
/// fault in the kernels loses the grade it would otherwise earn on this device.
#[doc(hidden)]
pub fn discover_with(source: &str) -> Vec<Result<WebGpuLandscape, String>> {
if wgpu::Instance::enabled_backend_features().is_empty() {
return Vec::new();
}
let mut d = wgpu::InstanceDescriptor::new_without_display_handle();
d.backends = wgpu::Backends::PRIMARY;
d.flags = wgpu::InstanceFlags::empty();
let instance = wgpu::Instance::new(d);
let Ok(adapters) = resolve(instance.enumerate_adapters(wgpu::Backends::PRIMARY)) else { return Vec::new() };
adapters.into_iter().map(|a| open(a, source)).collect()
}
fn open(adapter: wgpu::Adapter, source: &str) -> Result<WebGpuLandscape, String> {
let info = adapter.get_info();
let backend = backend_label(info.backend).ok_or_else(|| format!("{}: backend {:?} is not one this member runs on", info.name, info.backend))?;
if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
return Err(format!("{}: no compute shaders", info.name));
}
let offered = adapter.limits();
let mut short = None;
wgpu::Limits::downlevel_defaults().check_limits_with_fail_fn(&offered, true, |name, needed, has| {
short.get_or_insert_with(|| format!("{}: limit {name}: the kernels need {needed}, the adapter offers {has}", info.name));
});
if let Some(why) = short {
return Err(why);
}
let required_limits = wgpu::Limits { max_storage_buffer_binding_size: offered.max_storage_buffer_binding_size, max_buffer_size: offered.max_buffer_size, ..wgpu::Limits::downlevel_defaults() };
let descriptor = wgpu::DeviceDescriptor { label: Some("wai-quantum landscape"), required_features: wgpu::Features::empty(), required_limits, ..Default::default() };
let (device, queue) = resolve(adapter.request_device(&descriptor)).map_err(|e| e.to_string())?.map_err(|e| format!("{}: {e}", info.name))?;
let sink = Arc::new(Mutex::new(Sink::default()));
let errors = Arc::clone(&sink);
device.on_uncaptured_error(Arc::new(move |e: wgpu::Error| lock(&errors).errors.push(e.to_string())));
let lost = Arc::clone(&sink);
device.set_device_lost_callback(move |reason, message| {
let mut s = lock(&lost);
if s.lost.is_none() {
s.lost = Some(format!("device lost ({reason:?}): {message}"));
}
});
let scope = device.push_error_scope(wgpu::ErrorFilter::Validation);
let module = device.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("wai-quantum landscape"), source: wgpu::ShaderSource::Wgsl(Cow::Owned(source.to_string())) });
let pipeline = |entry: &str| {
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { label: Some(entry), layout: None, module: &module, entry_point: Some(entry), compilation_options: Default::default(), cache: None })
};
let pipelines = Pipelines {
settle_sweep: pipeline("settle_sweep"),
settle_advance: pipeline("settle_advance"),
in_order: pipeline("chains_in_order"),
coloured: pipeline("chains_coloured"),
copy: pipeline("chains_copy"),
};
if let Some(e) = resolve(scope.pop()).map_err(|e| e.to_string())? {
return Err(format!("{}: the kernels did not compile: {e}", info.name));
}
let class = if info.device_type == wgpu::DeviceType::Cpu { "cpu" } else { "gpu" };
let max_binding = binding_limit(device.limits().max_storage_buffer_binding_size);
Ok(WebGpuLandscape {
id: member_id(class, backend, &info),
class,
adapter: info.name.clone(),
backend,
device,
queue,
pipelines,
order: Order::InOrder,
max_binding,
wait_at_most: AtomicU64::new(POLL_TIMEOUT.as_nanos() as u64),
sink,
meter: MeterSlot::default(),
})
}
fn words(v: &[u32]) -> Vec<u8> {
v.iter().flat_map(|w| w.to_le_bytes()).collect()
}
/// The workload in one buffer of words: rows, columns, values, fields and each
/// cell's table offset; and where each part after the first starts.
fn graph(w: &Workload, reach: &Reach) -> (Vec<u32>, [u32; 4]) {
let r = device::rows(w);
let at_cols = r.rows.len() as u32;
let at_vals = at_cols + r.cols.len() as u32;
let at_fields = at_vals + r.vals.len() as u32;
let at_offsets = at_fields + r.fields.len() as u32;
let mut out = r.rows;
out.extend(r.cols);
out.extend(r.vals.iter().map(|&v| v as u32));
out.extend(r.fields.iter().map(|&f| f as i32 as u32));
out.extend(reach.offsets.iter().map(|&o| o as u32));
(out, [at_cols, at_vals, at_fields, at_offsets])
}
/// A greedy colouring of the coupling graph, in order of cell: each cell the
/// lowest colour none of its earlier neighbours has.
fn colouring(w: &Workload) -> Vec<Vec<u32>> {
let n = w.n();
let mut colour = vec![usize::MAX; n];
let mut classes: Vec<Vec<u32>> = Vec::new();
for i in 0..n {
let taken: Vec<usize> = w.row(i).map(|(j, _)| colour[j]).collect();
let c = (0..).find(|k| !taken.contains(k)).unwrap_or(0);
colour[i] = c;
if classes.len() <= c {
classes.resize(c + 1, Vec::new());
}
classes[c].push(i as u32);
}
classes
}
/// The coloured kernel's id: the member's own with `-coloured` after it, its name
/// shortened where it must be so that the id stays within 64 characters and keeps
/// the hex that tells two devices of one name apart.
fn coloured_id(id: &str) -> String {
const SUFFIX: &str = "-coloured";
let (stem, hex) = id.rsplit_once('.').unwrap_or((id, ""));
let room = 64usize.saturating_sub(SUFFIX.len() + hex.len() + 1);
let stem: String = stem.chars().take(room).collect();
format!("{}.{hex}{SUFFIX}", stem.trim_end_matches(['-', '.']))
}
impl WebGpuLandscape {
/// The same device running the coloured kernel: a member of its own.
pub fn coloured(&self) -> WebGpuLandscape {
let id = coloured_id(&self.id);
WebGpuLandscape {
device: self.device.clone(),
queue: self.queue.clone(),
pipelines: self.pipelines.clone(),
id,
class: self.class,
adapter: self.adapter.clone(),
backend: self.backend,
order: Order::Coloured,
max_binding: self.max_binding,
wait_at_most: AtomicU64::new(self.wait_at_most.load(Ordering::Relaxed)),
sink: Arc::clone(&self.sink),
meter: MeterSlot::default(),
}
}
/// Read a meter around each call: its reading stands beside the priced figures.
pub fn with_meter(mut self, meter: Box<dyn RunMeter + Send>) -> WebGpuLandscape {
self.meter = MeterSlot::new(meter);
self
}
/// Another limit on one wait: for a test of what a wait covers, and of
/// what the host does when one runs out.
#[doc(hidden)]
pub fn set_wait(&self, wait_at_most: Duration) {
self.wait_at_most.store(wait_at_most.as_nanos().max(1) as u64, Ordering::Relaxed);
}
/// The adapter's name, as the platform reports it.
pub fn adapter(&self) -> &str {
&self.adapter
}
/// `metal`, `vulkan` or `dx12`.
pub fn backend(&self) -> &'static str {
self.backend
}
pub fn order(&self) -> Order {
self.order
}
fn settled(&self) -> Result<(), String> {
let mut s = lock(&self.sink);
if let Some(why) = &s.lost {
return Err(why.clone());
}
if s.errors.is_empty() { Ok(()) } else { Err(std::mem::take(&mut s.errors).join("; ")) }
}
fn wait(&self, index: wgpu::SubmissionIndex) -> Result<(), String> {
let poll = || self.device.poll(wgpu::PollType::Wait { submission_index: Some(index), timeout: Some(Duration::from_nanos(self.wait_at_most.load(Ordering::Relaxed))) });
match std::panic::catch_unwind(std::panic::AssertUnwindSafe(poll)) {
Ok(Ok(_)) => Ok(()),
Ok(Err(e)) => Err(format!("the device did not finish: {e}")),
Err(_) => {
let mut s = lock(&self.sink);
s.lost.get_or_insert_with(|| "the device poll failed".into());
Err("the device poll failed".into())
}
}
}
fn buffer(&self, label: &str, bytes: &[u8], usage: wgpu::BufferUsages) -> wgpu::Buffer {
let b = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some(label), size: (bytes.len() as u64).max(16).next_multiple_of(4), usage: usage | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false });
if !bytes.is_empty() {
self.queue.write_buffer(&b, 0, bytes);
}
b
}
/// The words of `src`, `count` of them, read back.
fn read(&self, src: &wgpu::Buffer, count: u64) -> Result<Vec<u32>, String> {
let bytes = count * 4;
let staging = self.device.create_buffer(&wgpu::BufferDescriptor { label: Some("wai-quantum landscape readback"), size: bytes.max(4), usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false });
let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("readback") });
encoder.copy_buffer_to_buffer(src, 0, &staging, 0, bytes);
let index = self.queue.submit([encoder.finish()]);
let done = Arc::new(Mutex::new(None));
let filled = Arc::clone(&done);
staging.map_async(wgpu::MapMode::Read, 0..bytes, move |r| *lock(&filled) = Some(r));
self.wait(index)?;
match lock(&done).take() {
Some(Ok(())) => {}
Some(Err(e)) => return Err(e.to_string()),
None => return Err("the readback did not complete".into()),
}
let out = {
let view = staging.get_mapped_range(0..bytes).map_err(|e| e.to_string())?;
view.as_chunks::<4>().0.iter().map(|w| u32::from_le_bytes(*w)).collect()
};
staging.unmap();
self.settled()?;
Ok(out)
}
fn bind(&self, pipeline: &wgpu::ComputePipeline, entries: &[(u32, &wgpu::Buffer)]) -> wgpu::BindGroup {
let layout = pipeline.get_bind_group_layout(0);
let entries: Vec<wgpu::BindGroupEntry<'_>> = entries.iter().map(|&(binding, b)| wgpu::BindGroupEntry { binding, resource: b.as_entire_binding() }).collect();
self.device.create_bind_group(&wgpu::BindGroupDescriptor { label: Some("wai-quantum landscape"), layout: &layout, entries: &entries })
}
fn dispatch(encoder: &mut wgpu::CommandEncoder, pipeline: &wgpu::ComputePipeline, bind: &wgpu::BindGroup, invocations: u64) {
let groups = invocations.div_ceil(64).max(1);
let (x, y) = if groups <= 65_535 { (groups as u32, 1) } else { (65_535, groups.div_ceil(65_535) as u32) };
let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("wai-quantum landscape"), timestamp_writes: None });
pass.set_pipeline(pipeline);
pass.set_bind_group(0, bind, &[]);
pass.dispatch_workgroups(x, y, 1);
}
fn shape(w: &Workload, at: [u32; 4], chains: u32, bound: u32, extra: [u32; 8]) -> Vec<u8> {
// A cost above the lane holds every cell at 0, as the lane's top does:
// every field is at most 2³¹ − 1 here, or the workload was declined.
let cost = w.cost().min(i64::from(i32::MAX)) as i32 as u32;
let mut s = vec![w.n() as u32, cost, chains, bound];
s.extend(extra[..2].iter());
s.extend(&at[..3]);
s.extend(extra[2..].iter());
s.push(at[3]);
words(&s)
}
/// Whether the workload fits the kernels: every field in a 32-bit lane,
/// the workload within a binding, and, for sampling, a table of one sweep's
/// cut points within [`device::MAX_TABLE_ENTRIES`]. Only the first can
/// decline a workload a job gives (`extensions/quantum-landscape.md` §2.0).
fn fits(&self, w: &Workload, sampling: bool) -> Result<(), String> {
let bound = w.field_bound();
if bound > i64::from(i32::MAX) {
return Err(format!("its field bound is {bound}, and this member's 32-bit lanes hold at most {}", i32::MAX));
}
let n = w.n() as u64;
if 3 * n * 4 > self.max_binding || (w.nonzeros() * 2 + 3 * n + 1) * 4 > self.max_binding {
return Err("the workload is larger than a storage binding on this device".into());
}
if sampling { device::table_fits(&device::reach(w)) } else { Ok(()) }
}
fn settle_on_device(&self, w: &Workload, start: &[i8], bound: u32) -> Result<Outcome, String> {
self.settled()?;
self.fits(w, false)?;
let n = w.n();
let (g, at) = graph(w, &device::reach(w));
let graph = self.buffer("graph", &words(&g), wgpu::BufferUsages::STORAGE);
let shape = self.buffer("shape", &Self::shape(w, at, 0, bound, [0; 8]), wgpu::BufferUsages::UNIFORM);
let mut ring_words = vec![0u32; 3 * n];
for (i, &s) in start.iter().enumerate() {
ring_words[i] = i32::from(s) as u32;
}
let ring = self.buffer("ring", &words(&ring_words), wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC);
let status = self.buffer("status", &words(&[0; 4]), wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC);
let sweep = self.bind(&self.pipelines.settle_sweep, &[(0, &graph), (1, &shape), (2, &ring), (3, &status)]);
let advance = self.bind(&self.pipelines.settle_advance, &[(1, &shape), (3, &status)]);
let (mut done, mut ended) = (0u32, 0u32);
while ended == 0 && done < bound {
let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("settle") });
for _ in 0..SETTLE_BATCH.min(bound - done) {
Self::dispatch(&mut encoder, &self.pipelines.settle_sweep, &sweep, n as u64);
Self::dispatch(&mut encoder, &self.pipelines.settle_advance, &advance, 1);
}
self.queue.submit([encoder.finish()]);
let s = self.read(&status, 4)?;
(done, ended) = (s[0], s[1]);
}
let ring_words = self.read(&ring, 3 * n as u64)?;
let third = (done % 3) as usize * n;
let state: Vec<i8> = ring_words[third..third + n].iter().map(|&v| v as i32 as i8).collect();
let term = match ended {
1 => Term::FixedPoint,
2 => Term::TwoCycle,
_ => Term::Bound,
};
Ok(Outcome { state, sweeps: done, term, accumulates: w.nonzeros() * u64::from(done), draws: 0 })
}
fn chains_on_device(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String> {
let sampling = temperatures.iter().any(|&t| t > 0.0);
self.settled()?;
self.fits(w, sampling)?;
let n = w.n();
let count = chains.len();
if count == 0 {
return Ok(Vec::new());
}
if (2 * count * n * 4) as u64 > self.max_binding {
return Err("these chains are larger than a storage binding on this device".into());
}
let reach = device::reach(w);
let window = device::window(w, &reach, sampling) as u64;
let (g, at) = graph(w, &reach);
let graph = self.buffer("graph", &words(&g), wgpu::BufferUsages::STORAGE);
let mut cell_words = vec![0u32; 2 * count * n];
for (c, chain) in chains.iter().enumerate() {
if chain.start.len() != n {
return Err("a chain's start is not one cell a variable".into());
}
for (i, &s) in chain.start.iter().enumerate() {
cell_words[c * n + i] = i32::from(s) as u32;
}
}
let cells = self.buffer("cells", &words(&cell_words), wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC);
drop(cell_words);
let seeds: Vec<u32> = chains.iter().flat_map(|c| [c.seed as u32, (c.seed >> 32) as u32]).collect();
let classes = if self.order == Order::Coloured { colouring(w) } else { Vec::new() };
let members: Vec<u32> = classes.iter().flatten().copied().collect();
let mut k0 = 0usize;
while k0 < temperatures.len() {
let k1 = (k0 + window as usize).min(temperatures.len());
// The window's draws: seeds, the sweeps, the colour classes' cells, the table.
let (sweeps, table) = device::window_sweeps(w, &reach, &temperatures[k0..k1]);
let mut d = seeds.clone();
let at_sweeps = d.len() as u32;
d.extend(&sweeps);
let at_members = d.len() as u32;
d.extend(&members);
let at_table = d.len() as u32;
d.extend(table.iter().flat_map(|&v| [v as u32, (v >> 32) as u32]));
if (d.len() as u64) * 4 > self.max_binding {
return Err("a window's cut points are larger than a storage binding on this device".into());
}
let draws = self.buffer("draws", &words(&d), wgpu::BufferUsages::STORAGE);
// Each window is waited for before the next is sent, so that one wait
// covers one window, what a failed wait leaves queued is one window,
// and only one window's table is held at a time.
let last = match self.order {
Order::InOrder => {
let shape = self.buffer("shape", &Self::shape(w, at, count as u32, 0, [k0 as u32, k1 as u32, at_sweeps, at_table, 0, 0, 0, 0]), wgpu::BufferUsages::UNIFORM);
let bind = self.bind(&self.pipelines.in_order, &[(0, &graph), (1, &shape), (4, &cells), (5, &draws)]);
let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("chains") });
Self::dispatch(&mut encoder, &self.pipelines.in_order, &bind, count as u64);
Some(self.queue.submit([encoder.finish()]))
}
Order::Coloured => {
let mut last = None;
for k in k0..k1 {
let kind = d[at_sweeps as usize + k - k0];
let passes: Vec<(u32, u32)> = if kind == 0 {
vec![(0, 0)]
} else {
let mut first = at_members;
classes
.iter()
.map(|c| {
let p = (first, c.len() as u32);
first += c.len() as u32;
p
})
.collect()
};
for (first, len) in passes {
let extra = [k0 as u32, k1 as u32, at_sweeps, at_table, first, len, k as u32, kind];
let shape = self.buffer("shape", &Self::shape(w, at, count as u32, 0, extra), wgpu::BufferUsages::UNIFORM);
let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("coloured") });
let bind = self.bind(&self.pipelines.coloured, &[(0, &graph), (1, &shape), (4, &cells), (5, &draws)]);
let invocations = if kind == 0 { (count * n) as u64 } else { count as u64 * u64::from(len) };
Self::dispatch(&mut encoder, &self.pipelines.coloured, &bind, invocations);
if kind == 0 {
let copy = self.bind(&self.pipelines.copy, &[(1, &shape), (4, &cells)]);
Self::dispatch(&mut encoder, &self.pipelines.copy, &copy, (count * n) as u64);
}
last = Some(self.queue.submit([encoder.finish()]));
}
}
last
}
};
if let Some(index) = last {
self.wait(index)?;
}
self.settled()?;
k0 = k1;
}
let got = self.read(&cells, (count * n) as u64)?;
let draws_due = (n as u64) * temperatures.iter().filter(|&&t| t > 0.0).count() as u64;
Ok(got
.chunks_exact(n.max(1))
.take(count)
.map(|c| {
let state: Vec<i8> = c.iter().map(|&v| v as i32 as i8).collect();
let term = if w.rests(&state) { Term::FixedPoint } else { Term::Bound };
Outcome { state, sweeps: temperatures.len() as u32, term, accumulates: w.nonzeros() * temperatures.len() as u64, draws: draws_due }
})
.collect())
}
}
impl Member for WebGpuLandscape {
fn id(&self) -> String {
self.id.clone()
}
fn class(&self) -> &'static str {
self.class
}
fn label(&self) -> String {
let order = match self.order {
Order::InOrder => "its cells in the reference's order",
Order::Coloured => "a colour class of cells at a time",
};
device::label(self.class, &self.adapter, &format!("WebGPU ({})", self.backend), order)
}
fn takes(&self, w: &Workload, sampling: bool) -> Result<(), String> {
self.fits(w, sampling)
}
fn settle(&self, w: &Workload, start: &[i8], bound: u32) -> Result<Outcome, String> {
self.meter.around(|| self.settle_on_device(w, start, bound))
}
fn chains(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String> {
self.meter.around(|| self.chains_on_device(w, chains, temperatures))
}
fn meter(&self) -> Option<Metered> {
self.meter.last()
}
}
#[cfg(test)]
mod tests {
use super::*;
#[test]
fn a_greedy_colouring_never_puts_two_coupled_cells_together() {
let mut l = super::super::energy::Landscape::new(6);
for (a, b) in [(0, 1), (1, 2), (0, 2), (2, 3), (3, 4), (4, 5), (0, 5)] {
l.couple(a, b, 1);
}
let w = l.lower();
for class in colouring(&w) {
for &i in &class {
assert!(w.row(i as usize).all(|(j, _)| !class.contains(&(j as u32))));
}
}
}
/// The coloured kernel's id is a member id that keeps its device's hex.
#[test]
fn a_coloured_id_keeps_its_device_s_hex() {
let long = device::id_of("gpu", "wgpu-vulkan", "Radeon 8060S Graphics (RADV STRIX_HALO)", "4098/5510/radv/Mesa 25");
let c = coloured_id(&long);
assert!(c.len() <= 64, "{c}");
assert!(c.ends_with(&format!(".{}-coloured", long.rsplit_once('.').unwrap().1)), "{c}");
assert!(c.starts_with("gpu:wgpu-vulkan.radeon-8060s"), "{c}");
assert_eq!(coloured_id("gpu:wgpu-metal.apple-m3-ultra.3cc6"), "gpu:wgpu-metal.apple-m3-ultra.3cc6-coloured");
}
/// A binding this member uses stays under the size a 32-bit bounds check
/// can hold, whatever the adapter offers.
#[test]
fn a_binding_stays_below_four_gibibytes() {
assert_eq!(binding_limit(179_357_827_072), 1 << 31);
assert_eq!(binding_limit(128 << 20), 128 << 20);
}
/// A table offset is the same word whatever a field's size: every field at
/// the lanes' edge lands inside the table.
#[test]
fn a_cell_s_entry_fits_its_word_at_the_lanes_edge() {
let mut l = super::super::energy::Landscape::new(2);
l.couple(0, 1, 1).set_field(0, i64::from(i32::MAX) - 1).set_field(1, -i64::from(i32::MAX) + 1);
let w = l.lower();
assert_eq!(w.field_bound(), i64::from(i32::MAX));
let r = device::reach(&w);
let (g, at) = graph(&w, &r);
for i in 0..2 {
for s in [-1i8, 0, 1] {
let state = if i == 0 { [0, s] } else { [s, 0] };
let h = w.field(&state, i) as i32;
let e = (h as u32).wrapping_add(g[at[3] as usize + i]);
assert!(u64::from(e) < r.entries, "cell {i} field {h} entry {e}");
}
}
}
}

366
wai/wai-quantum/tests/landscape_conformance.rs
View file

@ -5,9 +5,10 @@ use serde_json::Value;
use wai_quantum::quantum_landscape::conformance::{chi_square, run_battery, vectors};
use wai_quantum::quantum_landscape::energy::Landscape;
use wai_quantum::quantum_landscape::job;
use wai_quantum::quantum_landscape::member::{Chain, CpuReference, Member, Parity, certify, KERNEL_SAMPLE};
use wai_quantum::quantum_landscape::member::{Chain, CpuReference, Member, Parity, certify, KERNEL_SAMPLE, KERNEL_SETTLE};
use wai_quantum::quantum_landscape::reach;
use wai_quantum::quantum_landscape::run::state_text;
use wai_quantum::quantum_landscape::sample::{anneal_with, cuts, schedule};
use wai_quantum::quantum_landscape::sample::{anneal_with, cuts, out, pick, schedule};
use wai_quantum::quantum_landscape::settle::{Outcome, Term, Workload};
fn text() -> String {
@ -88,12 +89,38 @@ fn the_exponential_schedule_and_cut_points_hold() {
}
}
/// A chain's sweeps: cooling from `t0`, or runs of sweeps held at a temperature.
fn temperatures_of(c: &Value) -> Vec<f64> {
match c["held"].as_array() {
Some(runs) => runs.iter().flat_map(|r| std::iter::repeat_n(bits(&r["temperature_bits"]), r["sweeps"].as_u64().unwrap() as usize)).collect(),
None => schedule(bits(&c["t0_bits"]), c["sweeps"].as_u64().unwrap() as u32),
}
}
/// The generator's outputs, worked out here in the text's own words.
#[test]
fn the_draw_vectors_hold() {
let f = file();
let cases = f["draws"].as_array().unwrap();
assert!(cases.len() >= 20);
for c in cases {
let (s, j): (u64, u64) = (c["seed"].as_str().unwrap().parse().unwrap(), c["j"].as_u64().unwrap());
let mut x = s.wrapping_add(j.wrapping_mul(0x9E37_79B9_7F4A_7C15));
x ^= x >> 30;
x = x.wrapping_mul(0xBF58_476D_1CE4_E5B9);
x ^= x >> 27;
x = x.wrapping_mul(0x94D0_49BB_1331_11EB);
x ^= x >> 31;
assert_eq!(x.to_string(), c["out"].as_str().unwrap());
}
}
#[test]
fn the_chain_vectors_hold() {
let f = file();
for c in f["chains"].as_array().unwrap() {
let l = landscape(&c["landscape"]);
let temperatures = schedule(bits(&c["t0_bits"]), c["sweeps"].as_u64().unwrap() as u32);
let temperatures = temperatures_of(c);
for ch in c["chains"].as_array().unwrap() {
let chain: u64 = ch["chain"].as_str().unwrap().parse().unwrap();
let o = anneal_with(&l.lower(), &cells(ch["start"].as_str().unwrap()), &temperatures, chain);
@ -127,28 +154,71 @@ fn the_job_vectors_hold() {
}
/// The battery's distributions are the formula, worked out here with the
/// platform's own exponential, from the file alone.
/// platform's own exponential, from the file alone: a block's law, in a
/// landscape of one block or of several.
#[test]
fn the_battery_s_distributions_are_the_formula() {
let f = file();
for c in f["statistical"]["workloads"].as_array().unwrap() {
let l = landscape(&c["landscape"]);
let n = l.len();
let whole = landscape(&c["landscape"]);
let n = c["block"].as_u64().unwrap() as usize;
let mut l = Landscape::new(n);
for i in 0..n {
for j in (i + 1)..n {
l.couple(i, j, whole.coupling(i, j));
}
l.set_field(i, whole.field()[i]);
}
l.set_cost(whole.cost());
for b in 1..whole.len() / n {
for i in 0..n {
assert_eq!(whole.field()[b * n + i], whole.field()[i], "{}: blocks are copies", c["name"]);
for j in 0..whole.len() {
let same = if j / n == b { whole.coupling(i, j % n) } else { 0 };
if j != b * n + i {
assert_eq!(whole.coupling(b * n + i, j), same, "{}: blocks are copies, none coupled to another", c["name"]);
}
}
}
}
let temperatures: Vec<f64> = c["temperatures"].as_array().unwrap().iter().map(|t| t.as_f64().unwrap()).collect();
let t = temperatures.iter().rev().copied().find(|&t| t > 0.0).unwrap();
let states: Vec<Vec<i8>> = (0..3usize.pow(n as u32)).map(|mut k| (0..n).map(|_| { let s = (k % 3) as i8 - 1; k /= 3; s }).collect()).collect();
let e: Vec<f64> = states.iter().map(|s| (-(l.energy(s) as f64) / t).exp()).collect();
let z: f64 = e.iter().sum();
let mut p: Vec<f64> = e.iter().map(|w| w / z).collect();
if *temperatures.last().unwrap() == 0.0 {
let w = l.lower();
let mut q = vec![0.0; p.len()];
for (k, s) in states.iter().enumerate() {
let next = w.sweep(s);
q[next.iter().rev().fold(0, |acc, &v| acc * 3 + (v + 1) as usize)] += p[k];
let state = |mut k: usize| -> Vec<i8> { (0..n).map(|_| { let s = (k % 3) as i8 - 1; k /= 3; s }).collect() };
let index = |s: &[i8]| s.iter().rev().fold(0, |acc, &v| acc * 3 + (v + 1) as usize);
let states: Vec<Vec<i8>> = (0..3usize.pow(n as u32)).map(state).collect();
let w = l.lower();
let p = if c["law"] == "walked" {
let mut p = vec![1.0 / states.len() as f64; states.len()];
for &t in &temperatures {
for i in 0..n {
let mut q = vec![0.0; states.len()];
for (k, s) in states.iter().enumerate() {
let h = w.field(s, i) as f64;
let e = [((-h - w.cost() as f64) / t).exp(), 1.0, ((h - w.cost() as f64) / t).exp()];
let z: f64 = e.iter().sum();
for (v, ev) in [-1i8, 0, 1].into_iter().zip(e) {
let mut u = s.clone();
u[i] = v;
q[index(&u)] += p[k] * ev / z;
}
}
p = q;
}
}
p = q;
}
p
} else {
let e: Vec<f64> = states.iter().map(|s| (-(l.energy(s) as f64) / t).exp()).collect();
let z: f64 = e.iter().sum();
let mut p: Vec<f64> = e.iter().map(|w| w / z).collect();
if *temperatures.last().unwrap() == 0.0 {
let mut q = vec![0.0; p.len()];
for (k, s) in states.iter().enumerate() {
q[index(&w.sweep(s))] += p[k];
}
p = q;
}
p
};
for (want, got) in p.iter().zip(c["exact"].as_array().unwrap()) {
assert!((want - got.as_f64().unwrap()).abs() < 1e-12, "{}", c["name"]);
}
@ -172,6 +242,16 @@ enum Mutant {
NoFinalZero,
/// A proper colouring of the coupling graph: correct, so it is granted `statistical`.
ProperColouring,
/// Cells from the fourth on left where they started.
FrozenFromFour,
/// Every sweep above 0 at the schedule's last temperature above 0.
OnlyTheEndOfTheSchedule,
/// Fields held in 16 bits.
Fields16,
/// Fields held in 32 bits, and nothing declined.
Fields32,
/// Cells taken last to first: correct, so it is granted `statistical`.
LastToFirst,
}
impl Member for Mutant {
@ -206,6 +286,7 @@ impl Mutant {
let colour: Vec<usize> = match self {
Mutant::AllAtOnce => vec![0; n],
Mutant::EvenThenOdd => (0..n).map(|i| i % 2).collect(),
Mutant::LastToFirst => (0..n).rev().collect(),
Mutant::ProperColouring => {
let mut colour = vec![usize::MAX; n];
for i in 0..n {
@ -222,6 +303,7 @@ impl Mutant {
let t = match self {
Mutant::IgnoresSchedule => temperatures[0],
Mutant::NoFinalZero if k == last => temperatures[last.saturating_sub(1)],
Mutant::OnlyTheEndOfTheSchedule if t > 0.0 => temperatures.iter().rev().copied().find(|&t| t > 0.0).unwrap(),
_ => t,
};
if t <= 0.0 {
@ -231,9 +313,15 @@ impl Mutant {
for g in 0..groups {
let old = s.clone();
for i in (0..n).filter(|&i| colour[i] == g) {
if matches!(self, Mutant::FrozenFromFour) && i >= 4 {
continue;
}
let mut h = w.field(&old, i);
if matches!(self, Mutant::ClampedTable) {
h = h.clamp(-5, 5);
match self {
Mutant::ClampedTable => h = h.clamp(-5, 5),
Mutant::Fields16 => h = i64::from(h as i16),
Mutant::Fields32 => h = i64::from(h as i32),
_ => {}
}
let (a, b) = cuts(h, w.cost(), 1.0 / t);
let u = (draw() * (1u64 << 53) as f64) as u64;
@ -260,17 +348,243 @@ fn the_battery_tells_a_sampler_from_a_wrong_one() {
let reference = run_battery(&CpuReference);
show("reference", &reference);
assert!(reference.iter().all(|r| r.passed), "the reference's own kernel passes");
for (name, m) in [("all at once", Mutant::AllAtOnce), ("even, then odd", Mutant::EvenThenOdd), ("ignores schedule", Mutant::IgnoresSchedule), ("clamped table", Mutant::ClampedTable), ("no final zero", Mutant::NoFinalZero)] {
for (name, m) in [("all at once", Mutant::AllAtOnce), ("even, then odd", Mutant::EvenThenOdd), ("ignores schedule", Mutant::IgnoresSchedule), ("clamped table", Mutant::ClampedTable), ("no final zero", Mutant::NoFinalZero), ("frozen from four", Mutant::FrozenFromFour), ("only the end", Mutant::OnlyTheEndOfTheSchedule), ("fields in 16 bits", Mutant::Fields16), ("fields in 32 bits", Mutant::Fields32)] {
let results = run_battery(&m);
show(name, &results);
assert!(results.iter().any(|r| !r.passed), "{name}: a wrong kernel passed every workload");
}
assert_eq!(certify(&Mutant::EvenThenOdd).unwrap().parity(KERNEL_SAMPLE), Parity::Uncertified);
let good = Mutant::ProperColouring;
let results = run_battery(&good);
show("proper colouring", &results);
assert!(results.iter().all(|r| r.passed), "a proper colouring passes");
assert_eq!(certify(&good).unwrap().parity(KERNEL_SAMPLE), Parity::Statistical);
for m in [Mutant::FrozenFromFour, Mutant::OnlyTheEndOfTheSchedule, Mutant::Fields16, Mutant::Fields32] {
assert_eq!(certify(&m).unwrap().parity(KERNEL_SAMPLE), Parity::Uncertified);
}
for (name, good) in [("proper colouring", Mutant::ProperColouring), ("last to first", Mutant::LastToFirst)] {
let results = run_battery(&good);
show(name, &results);
assert!(results.iter().all(|r| r.passed), "{name} passes");
assert_eq!(certify(&good).unwrap().parity(KERNEL_SAMPLE), Parity::Statistical);
}
}
// ── kernels that the vectors must refuse ───────────────────────────────────────────────────
/// A kernel that gives the reference's bits but for one fault of the kind a
/// port of it to a device can have: what the certificate exists to find.
#[derive(Clone, Copy, Debug, PartialEq)]
enum Fault {
/// A 32 by 32 product that loses the carry out of its middle.
DroppedCarry,
/// A draw compared with a cut point by its high word alone.
HighWordCompare,
/// Fields held in 16 bits.
Fields16,
/// Fields held in 32 bits, and nothing declined.
Fields32,
/// Fields held in 32 bits, and a check of `|b_i|` alone that forgets the couplings.
Fields32CheckedWithoutCouplings,
/// Each chain's cells last to first.
CellsBackwards,
/// The draw counter starting again at each window of this many sweeps.
CounterRestartsEachWindow(usize),
/// The sweep index stopping after this many sweeps.
StopsAfterOneWindow(usize),
/// Only this many chains run.
OnlyTheFirstChains(usize),
/// Only this many cells of a settle move.
OnlyTheFirstCells(usize),
/// A settle that stops looking for its end after this many sweeps.
SettleEndsLate(u32),
/// Declines everything wider than this.
DeclinesAbove(i64),
}
fn mul32(a: u32, b: u32) -> (u32, u32) {
let (a0, a1, b0, b1) = (a & 0xffff, a >> 16, b & 0xffff, b >> 16);
let (p00, p01, p10, p11) = (a0 * b0, a0 * b1, a1 * b0, a1 * b1);
let mid = p01.wrapping_add(p10);
let lo = p00.wrapping_add(mid << 16);
(lo, p11.wrapping_add(mid >> 16).wrapping_add(u32::from(lo < p00)))
}
/// The low 64 bits of a product in 32-bit words, with no carry out of the middle.
fn mul64_no_carry(a: u64, b: u64) -> u64 {
let (ax, ay, bx, by) = (a as u32, (a >> 32) as u32, b as u32, (b >> 32) as u32);
let (lo, hi) = mul32(ax, bx);
u64::from(lo) | (u64::from(hi.wrapping_add(ax.wrapping_mul(by)).wrapping_add(ay.wrapping_mul(bx))) << 32)
}
fn out_no_carry(s: u64, j: u64) -> u64 {
let mut x = s.wrapping_add(j.wrapping_mul(0x9E37_79B9_7F4A_7C15));
x ^= x >> 30;
x = mul64_no_carry(x, 0xBF58_476D_1CE4_E5B9);
x ^= x >> 27;
x = mul64_no_carry(x, 0x94D0_49BB_1331_11EB);
x ^ (x >> 31)
}
impl Fault {
fn field(self, w: &Workload, s: &[i8], i: usize) -> i64 {
let h = w.field(s, i);
match self {
Fault::Fields16 => i64::from(h as i16),
Fault::Fields32 | Fault::Fields32CheckedWithoutCouplings => i64::from(h as i32),
_ => h,
}
}
fn sweep(self, w: &Workload, s: &[i8]) -> Vec<i8> {
let c = w.cost();
(0..w.n()).map(|i| {
let h = self.field(w, s, i);
if h > c { 1 } else if h < -c { -1 } else { 0 }
}).collect()
}
fn chain(self, w: &Workload, c: &Chain, temperatures: &[f64]) -> Outcome {
let n = w.n();
let mut s = c.start.clone();
for (k, &t) in temperatures.iter().enumerate() {
if matches!(self, Fault::StopsAfterOneWindow(at) if k >= at) {
break;
}
if t <= 0.0 {
s = self.sweep(w, &s);
continue;
}
let cells: Vec<usize> = if self == Fault::CellsBackwards { (0..n).rev().collect() } else { (0..n).collect() };
for i in cells {
let counter = if let Fault::CounterRestartsEachWindow(at) = self { k % at } else { k };
let j = (counter * n + i + 1) as u64;
let u = if self == Fault::DroppedCarry { out_no_carry(c.seed, j) } else { out(c.seed, j) } >> 11;
let (a, b) = cuts(self.field(w, &s, i), w.cost(), 1.0 / t);
s[i] = if self == Fault::HighWordCompare {
if u >> 32 < a >> 32 { -1 } else if u >> 32 < b >> 32 { 0 } else { 1 }
} else {
pick(u, a, b)
};
}
}
let draws = (n as u64) * temperatures.iter().filter(|&&t| t > 0.0).count() as u64;
let term = if w.rests(&s) { Term::FixedPoint } else { Term::Bound };
Outcome { state: s, sweeps: temperatures.len() as u32, term, accumulates: w.nonzeros() * temperatures.len() as u64, draws }
}
}
impl Member for Fault {
fn id(&self) -> String {
"gpu:fault".into()
}
fn class(&self) -> &'static str {
"gpu"
}
fn label(&self) -> String {
"A kernel under test.".into()
}
fn settle(&self, w: &Workload, start: &[i8], bound: u32) -> Result<Outcome, String> {
if let Fault::SettleEndsLate(at) = *self {
let o = w.settle(start, bound);
return Ok(if o.sweeps > at { Outcome { sweeps: o.sweeps - 1, accumulates: o.accumulates - w.nonzeros(), ..o } } else { o });
}
let mut s = start.to_vec();
let mut back: Option<Vec<i8>> = None;
for k in 1..=bound {
let mut next = self.sweep(w, &s);
if let Fault::OnlyTheFirstCells(at) = *self {
next[at.min(s.len())..].copy_from_slice(&s[at.min(s.len())..]);
}
let term = if next == s { Some(Term::FixedPoint) } else if back.as_ref() == Some(&next) { Some(Term::TwoCycle) } else if k == bound { Some(Term::Bound) } else { None };
back = Some(std::mem::replace(&mut s, next));
if let Some(term) = term {
return Ok(Outcome { state: s, sweeps: k, term, accumulates: w.nonzeros() * u64::from(k), draws: 0 });
}
}
Err("no sweeps".into())
}
fn chains(&self, w: &Workload, chains: &[Chain], temperatures: &[f64]) -> Result<Vec<Outcome>, String> {
Ok(chains.iter().enumerate().map(|(k, c)| match *self {
Fault::OnlyTheFirstChains(at) if k >= at => {
let term = if w.rests(&c.start) { Term::FixedPoint } else { Term::Bound };
Outcome { state: c.start.clone(), sweeps: temperatures.len() as u32, term, accumulates: w.nonzeros() * temperatures.len() as u64, draws: (w.n() as u64) * temperatures.iter().filter(|&&t| t > 0.0).count() as u64 }
}
f => f.chain(w, c, temperatures),
}).collect())
}
fn takes(&self, w: &Workload, _sampling: bool) -> Result<(), String> {
let zero = vec![0i8; w.n()];
match *self {
Fault::DeclinesAbove(most) if w.field_bound() > most => Err(format!("its field bound is {}, above {most}", w.field_bound())),
Fault::Fields32CheckedWithoutCouplings if (0..w.n()).any(|i| w.field(&zero, i).abs() > i64::from(i32::MAX)) => Err("its fields are beyond 32 bits".into()),
_ => Ok(()),
}
}
}
/// With no fault the stand-in is the reference, so each fault below is the only difference.
#[test]
fn every_fault_loses_the_grade_of_the_kernel_it_is_in() {
let clean = Fault::DeclinesAbove(i64::MAX);
let c = certify(&clean).unwrap();
assert_eq!((c.parity(KERNEL_SETTLE), c.parity(KERNEL_SAMPLE)), (Parity::BitExact, Parity::BitExact));
let settle = [Fault::Fields16, Fault::Fields32, Fault::Fields32CheckedWithoutCouplings, Fault::OnlyTheFirstCells(64), Fault::OnlyTheFirstCells(256), Fault::SettleEndsLate(64), Fault::SettleEndsLate(256)];
let sample = [
Fault::DroppedCarry,
Fault::HighWordCompare,
Fault::Fields16,
Fault::Fields32,
Fault::Fields32CheckedWithoutCouplings,
Fault::CellsBackwards,
Fault::OnlyTheFirstChains(64),
];
for f in settle {
assert_ne!(certify(&f).unwrap().parity(KERNEL_SETTLE), Parity::BitExact, "{f:?}");
}
// A window's faults, at every window a vector crosses on a device that takes its table from the host.
let mut windows = Vec::new();
for c in wai_quantum::quantum_landscape::conformance::chain_cases() {
let w = c.landscape.lower();
let at = reach::window(&w, &reach::reach(&w), true);
if at < c.temperatures.each().len() {
windows.push(at);
}
}
windows.sort_unstable();
windows.dedup();
assert!(windows.len() >= 2, "windows by table and by work: {windows:?}");
let sample = sample.into_iter().chain(windows.iter().flat_map(|&at| [Fault::CounterRestartsEachWindow(at), Fault::StopsAfterOneWindow(at)]));
for f in sample {
assert!(!wai_quantum::quantum_landscape::conformance::chain_vectors_hold(&f), "{f:?}");
}
}
/// A member takes every vector at or under the line, and only there may it decline.
#[test]
fn a_member_that_declines_below_the_line_earns_nothing_and_one_at_it_keeps_its_grade() {
for most in [41, 70_000, i64::from(i32::MAX) - 1] {
let narrow = Fault::DeclinesAbove(most);
let c = certify(&narrow).unwrap();
assert_eq!(c.parity(KERNEL_SETTLE), Parity::Uncertified, "declines above {most}");
assert_ne!(c.parity(KERNEL_SAMPLE), Parity::BitExact, "declines above {most}");
}
let lanes32 = Fault::DeclinesAbove(i64::from(i32::MAX));
let c = certify(&lanes32).unwrap();
assert_eq!((c.parity(KERNEL_SETTLE), c.parity(KERNEL_SAMPLE)), (Parity::BitExact, Parity::BitExact));
}
/// The set reaches what large jobs reach: a settle past a workgroup and past a
/// batch, cells and chains past a workgroup, and hot chains at every width.
#[test]
fn the_set_reaches_what_large_jobs_reach() {
use wai_quantum::quantum_landscape::conformance::{chain_cases, settle_cases};
let settles = settle_cases();
let chains = chain_cases();
assert!(settles.iter().any(|c| c.landscape.len() > 256 && c.landscape.lower().settle(&c.start, c.bound).sweeps > 256));
let lane = i64::from(i32::MAX);
for bound in [lane, lane + 2] {
assert!(settles.iter().any(|c| c.landscape.lower().field_bound() == bound), "a settle vector at {bound}");
assert!(chains.iter().any(|c| c.landscape.lower().field_bound() == bound && c.temperatures.each().iter().all(|&t| t > 0.0)), "a hot chain at {bound}");
assert!(chains.iter().any(|c| c.landscape.lower().field_bound() == bound && c.temperatures.each().iter().all(|&t| t == 0.0)), "a chain at 0 at {bound}");
}
assert!(chains.iter().any(|c| c.restarts > 64));
assert!(chains.iter().filter(|c| c.temperatures.each().last().is_some_and(|&t| t > 0.0)).count() >= 8, "chains that stop hot");
}
#[test]

454
wai/wai-quantum/tests/landscape_devices.rs Normal file
View file

@ -0,0 +1,454 @@
//! The landscape kernels on every GPU this machine has, through each backend
//! this build holds (`quantum_landscape::{webgpu, cuda, metal}`), held to the
//! reference as a host holds a member: certified on the conformance set, then
//! every worked example over many seeds and jobs past every window and
//! workgroup, bit for bit; and each fault a port of the kernels can have, put
//! into the kernels on this device, shown to cost them their grade.
//!
//! - `WAI_REQUIRE_GPU`: a machine where no backend finds a device fails.
//! - `WAI_LANDSCAPE_EXPECT`: the members this machine must have, as id
//! prefixes separated by commas (`gpu:metal.,gpu:wgpu-metal.`). A prefix no
//! member has fails, and so does a backend that names one and found none,
//! with why: a card off its bus, or kernels that no longer compile.
//! - `WAI_LANDSCAPE_REPORT`: a file each test that certifies a member appends
//! its certificates to, one line each, for a workflow to read.
#![cfg(not(target_arch = "wasm32"))]
use serde_json::{Value, json};
use std::io::Write;
use std::time::{Duration, Instant};
use wai_quantum::quantum_landscape::MAX_JOB_BYTES;
use wai_quantum::quantum_landscape::energy::Landscape;
use wai_quantum::quantum_landscape::job;
use wai_quantum::quantum_landscape::member::{self, Certified, Chain, CpuReference, KERNEL_SAMPLE, KERNEL_SETTLE, Member, Parity};
use wai_quantum::quantum_landscape::sample::{chain_seed, out, schedule, start};
type Found = Vec<Result<Box<dyn Member>, String>>;
fn boxed<M: Member + 'static>(found: Vec<Result<M, String>>) -> Found {
found.into_iter().map(|r| r.map(|m| Box::new(m) as Box<dyn Member>)).collect()
}
/// Every backend this build holds, its name, and what it finds with these kernels.
#[allow(unused_mut, unused_variables)]
fn backends(faults: &[(&str, &str, &str)]) -> Vec<(&'static str, Found)> {
let apply = |source: &str, backend: &str| {
let mut s = source.to_string();
for (b, from, to) in faults {
if *b == backend {
assert!(s.contains(from), "{backend}: the kernel text no longer holds {from:?}");
s = s.replace(from, to);
}
}
s
};
let mut out: Vec<(&'static str, Found)> = Vec::new();
#[cfg(feature = "quantum_landscape_webgpu")]
out.push(("webgpu", boxed(wai_quantum::quantum_landscape::webgpu::discover_with(&apply(wai_quantum::quantum_landscape::webgpu::SHADER, "webgpu")))));
#[cfg(feature = "quantum_landscape_cuda")]
out.push(("cuda", boxed(wai_quantum::quantum_landscape::cuda::discover_with(&apply(wai_quantum::quantum_landscape::cuda::SOURCE, "cuda")))));
#[cfg(feature = "quantum_landscape_metal")]
out.push(("metal", boxed(wai_quantum::quantum_landscape::metal::discover_with(&apply(wai_quantum::quantum_landscape::metal::SOURCE, "metal")))));
out
}
/// The prefixes `WAI_LANDSCAPE_EXPECT` names.
fn expected() -> Vec<String> {
std::env::var("WAI_LANDSCAPE_EXPECT").map(|v| v.split(',').map(|p| p.trim().to_string()).filter(|p| !p.is_empty()).collect()).unwrap_or_default()
}
/// The backends this build holds.
fn built() -> Vec<&'static str> {
[cfg!(feature = "quantum_landscape_webgpu").then_some("webgpu"), cfg!(feature = "quantum_landscape_cuda").then_some("cuda"), cfg!(feature = "quantum_landscape_metal").then_some("metal")].into_iter().flatten().collect()
}
/// The backend a member's id names.
fn backend_of(id: &str) -> &'static str {
if id.contains(":wgpu-") {
"webgpu"
} else if id.contains(":cuda.") {
"cuda"
} else {
"metal"
}
}
/// Every member this build's backends find, held to what the machine expects.
fn members() -> Vec<Box<dyn Member>> {
let expect = expected();
let mut found: Vec<Box<dyn Member>> = Vec::new();
let mut not = Vec::new();
for (backend, results) in backends(&[]) {
for r in results {
match r {
Ok(m) => found.push(m),
Err(why) => {
eprintln!("{backend}: not a member: {why}");
not.push(format!("{backend}: {why}"));
}
}
}
}
let missing: Vec<&String> = expect.iter().filter(|p| !found.iter().any(|m| m.id().starts_with(p.as_str()))).collect();
assert!(missing.is_empty(), "WAI_LANDSCAPE_EXPECT is not met: no member's id starts with {missing:?}; found {:?}; not members: {not:?}", found.iter().map(|m| m.id()).collect::<Vec<_>>());
if found.is_empty() {
assert!(std::env::var_os("WAI_REQUIRE_GPU").is_none(), "no device here, and WAI_REQUIRE_GPU is set");
eprintln!("no device here: nothing to hold");
}
found
}
/// A line of the report a workflow reads.
fn report(line: &str) {
eprintln!("{line}");
if let Some(path) = std::env::var_os("WAI_LANDSCAPE_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 file takes a line");
}
}
#[test]
fn every_device_is_certified_bit_exact_for_both_kernels() {
for m in members() {
let c = member::certify(m.as_ref()).unwrap();
report(&format!("certified {}: settle {}, sample {}", m.id(), c.parity(KERNEL_SETTLE).label(), c.parity(KERNEL_SAMPLE).label()));
assert_eq!(c.parity(KERNEL_SETTLE), Parity::BitExact, "{}", m.id());
assert_eq!(c.parity(KERNEL_SAMPLE), Parity::BitExact, "{}", m.id());
}
}
/// Every worked example at three temperatures over 50 seeds, the starts of
/// many seeds on the device together: the reference's bits.
#[test]
fn every_worked_example_gives_the_reference_s_bits() {
let want = job::corpus_bits_on(&Certified::reference(), 50);
for m in members() {
let c = member::certify(m.as_ref()).unwrap();
assert_eq!(job::corpus_bits_on(&c, 50), want, "{}", m.id());
let talks = wai_quantum::quantum_landscape::EXAMPLES[0].1;
let req = json!({ "program": talks, "seed": "7" });
let cpu = job::call("run", &req).unwrap();
let gpu = job::call_certified(&c, "run", &req, MAX_JOB_BYTES).unwrap();
for key in ["state", "energy", "verdict", "work", "replay", "truth", "joules"] {
assert_eq!(cpu[key], gpu[key], "{}: {key}", m.id());
}
assert_eq!(gpu["member"]["id"], m.id());
assert_eq!(gpu["member"]["parity"], "bit-exact");
}
}
#[cfg(feature = "quantum_landscape_webgpu")]
#[test]
fn the_coloured_kernel_earns_statistical_and_no_more() {
for r in wai_quantum::quantum_landscape::webgpu::discover() {
let Ok(m) = r else { continue };
let coloured = m.coloured();
let c = member::certify(&coloured).unwrap();
for r in wai_quantum::quantum_landscape::conformance::run_battery(&coloured) {
eprintln!("{} {:<28} chi-square {:>8.1} against {:>6.1}: {}", coloured.id(), r.name, r.chi_square, r.critical, if r.passed { "pass" } else { "FAIL" });
}
report(&format!("certified {}: settle {}, sample {}", coloured.id(), c.parity(KERNEL_SETTLE).label(), c.parity(KERNEL_SAMPLE).label()));
assert_eq!(c.parity(KERNEL_SETTLE), Parity::BitExact, "{}", coloured.id());
assert_eq!(c.parity(KERNEL_SAMPLE), Parity::Statistical, "{}", coloured.id());
}
}
/// A workload whose field bound is past 32 bits: the WebGPU member declines it
/// and the job runs on the reference, which the answer names with the refusal
/// beside it; the CUDA and Metal members take it and give the reference's bits.
/// A narrow one runs on every device.
#[test]
fn a_workload_a_device_declines_goes_to_the_reference_and_the_answer_says_so() {
for m in members() {
let placed = [member::certify(m.as_ref()).unwrap(), Certified::reference()];
let wide = json!({ "program": "vars 2\nfield 0 1099511627776\nagree 0 1\nsolve t0=0 restarts=4\n" });
let a = job::call_placed(&placed, "run", &wide, MAX_JOB_BYTES).unwrap();
let reference = job::call("run", &wide).unwrap();
assert_eq!((&a["state"], &a["replay"]), (&reference["state"], &reference["replay"]), "{}", m.id());
if m.id().contains(":wgpu-") {
assert_eq!(a["member"]["id"], "cpu:reference");
assert_eq!(a["refusals"][0]["member"], m.id());
assert!(a["refusals"][0]["why"].as_str().unwrap().contains("field bound is 1099511627777"), "{a}");
} else {
assert_eq!(a["member"]["id"], m.id(), "a 64-bit member takes it");
assert!(a.get("refusals").is_none());
}
let narrow = json!({ "program": "vars 2\nfield 0 3\nagree 0 1\nsolve\n" });
let a = job::call_placed(&placed, "run", &narrow, MAX_JOB_BYTES).unwrap();
assert_eq!(a["member"]["id"], m.id());
assert!(a.get("refusals").is_none());
}
}
/// `n` cells with about `degree` couplings each, of both signs.
fn sparse(n: usize, degree: usize, seed: u64) -> Landscape {
let mut l = Landscape::new(n);
let mut k = 0u64;
for i in 0..n {
for _ in 0..degree / 2 {
k += 1;
let j = (out(seed, k) % n as u64) as usize;
if j != i {
l.couple(i, j, if k % 3 == 0 { -1 } else { 1 });
}
}
}
l.set_cost(1);
l
}
/// A program's text for a landscape and a run line.
fn program(l: &Landscape, run: &str) -> String {
let n = l.len();
let mut text = format!("vars {n}\ncost {}\n", l.cost());
for i in 0..n {
for j in (i + 1)..n {
match l.coupling(i, j) {
1 => text.push_str(&format!("agree {i} {j}\n")),
-1 => text.push_str(&format!("oppose {i} {j}\n")),
_ => {}
}
}
if l.field()[i] != 0 {
text.push_str(&format!("field {i} {}\n", l.field()[i]));
}
}
text.push_str(run);
text.push('\n');
text
}
/// Jobs past every window and workgroup a device splits them into, through the
/// host as a placed job: the reference's bits, from the device.
#[test]
fn jobs_past_every_window_and_workgroup_give_the_reference_s_bits() {
let mut wide = sparse(300, 6, 0x51DE);
for i in 0..300 {
wide.set_field(i, (out(0xF1E1D, i as u64 + 1) % 200_001) as i64 - 100_000);
}
let jobs: Vec<(&str, String)> = vec![
("a settle of 4096 cells", program(&sparse(4096, 8, 0x5E77), "minimize steps=200")),
("a sampling run of many sweeps", program(&sparse(1024, 6, 0x7177), "solve t0=3 restarts=40 steps=600")),
("wide fields, every window hot", program(&wide, "solve t0=150000 restarts=40 steps=120")),
("more chains than a batch", program(&sparse(64, 6, 0xC4A1), "solve restarts=5000 steps=20")),
];
for m in members() {
let placed = [member::certify(m.as_ref()).unwrap(), Certified::reference()];
for (name, text) in &jobs {
let req = json!({ "program": text, "seed": "11" });
let cpu = job::call("run", &req).unwrap();
let gpu = job::call_placed(&placed, "run", &req, MAX_JOB_BYTES).unwrap();
assert_eq!(gpu["member"]["id"], m.id(), "{}: {name}", m.id());
for key in ["state", "energy", "verdict", "work", "replay"] {
assert_eq!(cpu[key], gpu[key], "{}: {name}: {key}", m.id());
}
}
}
}
/// The faults a port of the kernels can have, as text put into each backend's
/// kernels: `(name, kernel it is in, [(backend, from, to)])`.
fn faults() -> Vec<(&'static str, &'static str, Vec<(&'static str, &'static str, &'static str)>)> {
vec![
("a dropped carry in a 32 by 32 product", KERNEL_SAMPLE, vec![("webgpu", "let mid_carry = select(0u, 0x10000u, mid < p01);", "let mid_carry = 0u;")]),
("a draw compared with a cut point by its high word", KERNEL_SAMPLE, vec![("webgpu", "return u.y < a.y || (u.y == a.y && u.x < a.x);", "return u.y < a.y;")]),
(
"fields held in 16 bits",
KERNEL_SETTLE,
vec![("webgpu", " h = h + bitcast<i32>(graph[shape.vals + k]) * ring[base + graph[shape.cols + k]];\n }\n return h;", " h = h + bitcast<i32>(graph[shape.vals + k]) * ring[base + graph[shape.cols + k]];\n }\n return (h << 16u) >> 16u;")],
),
(
"sampled fields held in 16 bits",
KERNEL_SAMPLE,
vec![("webgpu", " h = h + bitcast<i32>(graph[shape.vals + k]) * cells[base + graph[shape.cols + k]];\n }\n return h;", " h = h + bitcast<i32>(graph[shape.vals + k]) * cells[base + graph[shape.cols + k]];\n }\n return (h << 16u) >> 16u;")],
),
(
"fields held in 32 bits",
KERNEL_SETTLE,
vec![
("metal", " h += long(vals[k]) * long(s[cols[k]]);\n }\n return h;", " h += long(vals[k]) * long(s[cols[k]]);\n }\n return long(int(h));"),
("cuda", " h += (i64)vals[k] * (i64)s[cols[k]];\n }\n return h;", " h += (i64)vals[k] * (i64)s[cols[k]];\n }\n return (i64)(int)h;"),
],
),
(
"a draw counter that starts again at each window",
KERNEL_SAMPLE,
vec![
("webgpu", "pick(draw(seed, k, i), entry_of(kind, i, h))", "pick(draw(seed, k - shape.k0, i), entry_of(kind, i, h))"),
("metal", "splitmix(seed, ulong(k) * n + i + 1)", "splitmix(seed, ulong(k - sh.k0) * n + i + 1)"),
("cuda", "splitmix(seed, (u64)k * n + i + 1)", "splitmix(seed, (u64)(k - k0) * n + i + 1)"),
],
),
(
"cells taken last to first",
KERNEL_SAMPLE,
vec![
("webgpu", " for (var i = 0u; i < n; i = i + 1u) {\n let h = cell_field(base, i);", " for (var r = 0u; r < n; r = r + 1u) {\n let i = n - 1u - r;\n let h = cell_field(base, i);"),
("metal", " for (uint i = 0; i < n; ++i) {\n long h = field_of(", " for (uint r = 0; r < n; ++r) {\n uint i = n - 1u - r;\n long h = field_of("),
("cuda", " for (unsigned i = 0; i < n; ++i) {\n i64 h = field_of(", " for (unsigned r = 0; r < n; ++r) {\n unsigned i = n - 1u - r;\n i64 h = field_of("),
],
),
(
"a settle whose cells past one workgroup never move",
KERNEL_SETTLE,
vec![
("webgpu", "if (i >= shape.n || atomicLoad(&status[1]) != 0u) { return; }", "if (i >= min(shape.n, 64u) || atomicLoad(&status[1]) != 0u) { return; }"),
("metal", "if (i >= sh.n || atomic_load_explicit(&status[1], memory_order_relaxed) != 0u) return;", "if (i >= min(sh.n, 64u) || atomic_load_explicit(&status[1], memory_order_relaxed) != 0u) return;"),
("cuda", "if (i >= n || st[1] != 0u) return;", "if (i >= (n < 64u ? n : 64u) || st[1] != 0u) return;"),
],
),
]
}
/// Each fault, put into the kernels on this device, costs the kernel it is in
/// its grade: the certificate the host finds for it is not bit-exact. Where the
/// fault leaves a correct sampler in another order, the battery may still
/// grant it `statistical`, which is what it is.
#[test]
fn every_fault_in_a_device_s_kernels_loses_its_grade() {
let present: Vec<String> = members().iter().map(|m| m.id()).collect();
let have = built();
for (name, kernel, edits) in faults() {
let edits: Vec<(&str, &str, &str)> = edits.into_iter().filter(|(b, _, _)| have.contains(b)).collect();
if edits.is_empty() {
continue;
}
for (backend, found) in backends(&edits) {
if !edits.iter().any(|(b, _, _)| *b == backend) {
continue;
}
let faulted: Vec<Box<dyn Member>> = found.into_iter().filter_map(Result::ok).collect();
let clean: Vec<&String> = present.iter().filter(|id| backend_of(id) == backend).collect();
assert_eq!(faulted.len(), clean.len(), "{backend}, {name}: the faulted kernels opened on {} devices, the clean ones on {}", faulted.len(), clean.len());
for m in faulted {
let c = member::certify(m.as_ref()).unwrap();
report(&format!("fault {name} in {}: settle {}, sample {}", m.id(), c.parity(KERNEL_SETTLE).label(), c.parity(KERNEL_SAMPLE).label()));
assert!(present.contains(&m.id()), "{} is not a member without the fault", m.id());
assert_ne!(c.parity(kernel), Parity::BitExact, "{}: {name} kept {kernel} bit-exact", m.id());
}
}
}
}
/// A WebGPU member waits once a window, so a run longer than one wait but made
/// of windows shorter than it finishes; and a job a member cannot finish in its
/// wait goes to the next member, the answer naming the one that failed.
/// What one wait may take while the host certifies a member here: the member's own limit.
#[cfg(feature = "quantum_landscape_webgpu")]
const POLL_FOR_CERTIFYING: Duration = Duration::from_secs(300);
#[cfg(feature = "quantum_landscape_webgpu")]
#[test]
fn a_wait_covers_one_window_and_a_job_a_member_cannot_finish_goes_on() {
use wai_quantum::quantum_landscape::reach;
let l = sparse(2048, 64, 0xA17);
let w = l.lower();
let per_window = reach::window(&w, &reach::reach(&w), true);
let chains: Vec<Chain> = (0..4u64).map(|r| Chain { start: start(chain_seed(5, r), w.n()), seed: chain_seed(5, r) }).collect();
for r in wai_quantum::quantum_landscape::webgpu::discover() {
let Ok(m) = r else { continue };
let one = vec![1.5; per_window];
let _ = m.chains(&w, &chains, &one);
let mut took: Vec<f64> = (0..3)
.map(|_| {
let t = Instant::now();
m.chains(&w, &chains, &one).unwrap();
t.elapsed().as_secs_f64()
})
.collect();
took.sort_by(f64::total_cmp);
let wait = Duration::from_secs_f64((10.0 * took[1]).max(0.05));
m.set_wait(wait);
let long = vec![1.5; 40 * per_window];
let t = Instant::now();
let got = m.chains(&w, &chains, &long).unwrap_or_else(|e| panic!("{}: a run of 40 windows with a wait of {wait:?} each: {e}", m.id()));
let ran = t.elapsed();
report(&format!("waits {}: 40 windows in {:.2} s, one wait at most {:.3} s", m.id(), ran.as_secs_f64(), wait.as_secs_f64()));
assert!(ran > wait, "{}: the run should be longer than one wait to show anything ({ran:?} against {wait:?})", m.id());
assert_eq!(got, CpuReference.chains(&w, &chains, &long).unwrap(), "{}", m.id());
m.set_wait(POLL_FOR_CERTIFYING);
let placed = [member::certify(&m).unwrap(), Certified::reference()];
assert_eq!(placed[0].parity(KERNEL_SAMPLE), Parity::BitExact, "{}", m.id());
m.set_wait(Duration::from_nanos(1));
let text = program(&sparse(2048, 64, 0xA17), "solve t0=1.5 restarts=4 steps=400");
let a = job::call_placed(&placed, "run", &json!({ "program": text, "seed": "5" }), MAX_JOB_BYTES).unwrap();
assert_eq!(a["member"]["id"], "cpu:reference");
assert_eq!(a["replay"], job::call("run", &json!({ "program": text, "seed": "5" })).unwrap()["replay"]);
let refusals = a["refusals"].as_array().unwrap();
assert_eq!(refusals.len(), 1, "{a}");
assert_eq!(refusals[0]["member"], m.id());
let why = refusals[0]["why"].as_str().unwrap();
assert!(why.contains("did not finish"), "{why}");
}
}
fn time<T>(f: impl FnOnce() -> T) -> (T, f64) {
let t = Instant::now();
let v = f();
(v, t.elapsed().as_secs_f64())
}
/// Wall-clock seconds on this machine, the CPU reference on one thread against
/// each device, for the same work, with whether the bits are the same, each the
/// middle of three runs. The bits are asserted; the times are printed, with the
/// host's load beside them where the platform gives it. Run with
/// `--ignored --nocapture`.
#[test]
#[ignore]
fn timings() {
let load = std::fs::read_to_string("/proc/loadavg").ok().or_else(|| std::process::Command::new("sysctl").args(["-n", "vm.loadavg"]).output().ok().map(|o| String::from_utf8_lossy(&o.stdout).into_owned()));
println!("host load: {}", load.as_deref().map(str::trim).unwrap_or("not given here"));
let middle = |f: &mut dyn FnMut() -> f64| {
let mut t = [f(), f(), f()];
t.sort_by(f64::total_cmp);
t[1]
};
let ms = members();
for (n, degree, chains, sweeps) in [(256usize, 6usize, 4096usize, 40u32), (1024, 6, 4096, 40), (4096, 8, 1024, 40)] {
let l = sparse(n, degree, 0x7177);
let w = l.lower();
let temperatures = schedule(l.default_temperature(), sweeps);
let cs: Vec<Chain> = (0..chains as u64)
.map(|r| {
let c = chain_seed(9, r);
Chain { start: start(c, n), seed: c }
})
.collect();
let cpu = CpuReference.chains(&w, &cs, &temperatures).unwrap();
let t_cpu = middle(&mut || time(|| CpuReference.chains(&w, &cs, &temperatures).unwrap()).1);
println!("chains n={n} degree≈{degree} chains={chains} sweeps={sweeps}: cpu:reference, one thread, {t_cpu:.3} s");
for m in &ms {
let _ = m.chains(&w, &cs[..64.min(cs.len())], &temperatures);
let gpu = m.chains(&w, &cs, &temperatures).unwrap();
assert_eq!(gpu, cpu, "{}", m.id());
let t = middle(&mut || time(|| m.chains(&w, &cs, &temperatures).unwrap()).1);
println!(" {} {t:.3} s, the same bits", m.id());
}
}
for n in [1024usize, 4096] {
let l = sparse(n, 32, 0x5e77);
let w = l.lower();
let from: Vec<i8> = start(3, n);
let cpu = CpuReference.settle(&w, &from, 200).unwrap();
let t_cpu = middle(&mut || time(|| CpuReference.settle(&w, &from, 200).unwrap()).1);
println!("settle n={n} degree≈32: cpu:reference, one thread, {t_cpu:.3} s ({} sweeps, {:?})", cpu.sweeps, cpu.term);
for m in &ms {
assert_eq!(m.settle(&w, &from, 200).unwrap(), cpu, "{}", m.id());
let t = middle(&mut || time(|| m.settle(&w, &from, 200).unwrap()).1);
println!(" {} {t:.3} s, the same bits", m.id());
}
}
let talks = wai_quantum::quantum_landscape::EXAMPLES[0].1;
let req = json!({ "program": talks, "seed": "1", "seeds": 2000 });
let cpu = job::call("replay", &req).unwrap();
let t_cpu = middle(&mut || time(|| job::call("replay", &req).unwrap()).1);
println!("replay of the talks, 2,000 seeds: cpu:reference, one thread, {t_cpu:.3} s");
let columns = |a: &Value| a["lines"].as_str().unwrap().lines().map(|l| l.split(' ').skip(2).collect::<Vec<_>>().join(" ")).collect::<Vec<_>>();
for m in &ms {
let c = member::certify(m.as_ref()).unwrap();
let gpu = job::call_certified(&c, "replay", &req, MAX_JOB_BYTES).unwrap();
assert_eq!(columns(&gpu), columns(&cpu), "{}", m.id());
let t = middle(&mut || time(|| job::call_certified(&c, "replay", &req, MAX_JOB_BYTES).unwrap()).1);
println!(" {} {t:.3} s, the same states, energies, verdicts, work and replay digests", m.id());
}
}

27
wai/wai-rs/Cargo.lock generated
View file

@ -163,7 +163,7 @@ version = "0.38.0+1.3.281"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f"
dependencies = [
"libloading",
"libloading 0.8.9",
]
[[package]]
@ -1047,6 +1047,15 @@ dependencies = [
"subtle",
]
[[package]]
name = "cudarc"
version = "0.19.10"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7359bf1de037ddada72729c3acec9b040f3468400c6e9e740b01ba0cd2e89562"
dependencies = [
"libloading 0.9.0",
]
[[package]]
name = "curve25519-dalek"
version = "4.1.3"
@ -2340,6 +2349,16 @@ dependencies = [
"windows-link",
]
[[package]]
name = "libloading"
version = "0.9.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "754ca22de805bb5744484a5b151a9e1a8e837d5dc232c2d7d8c2e3492edc8b60"
dependencies = [
"cfg-if",
"windows-link",
]
[[package]]
name = "libm"
version = "0.2.16"
@ -4779,7 +4798,11 @@ name = "wai-quantum"
version = "0.8.1"
dependencies = [
"blake3",
"cudarc",
"ed25519-dalek 2.2.0",
"objc2",
"objc2-foundation",
"objc2-metal",
"serde",
"serde_json",
"wgpu",
@ -5001,7 +5024,7 @@ dependencies = [
"gpu-allocator",
"hashbrown 0.17.1",
"libc",
"libloading",
"libloading 0.8.9",
"log",
"naga",
"naga-types",

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

@ -1032,6 +1032,12 @@ quantum_fabric = ["quantum", "quantum_ops", "wai-quantum/quantum_fabric"]
# `wai_quantum_vectors verify --webgpu`. Opt-in: it is the one feature that
# reaches a platform interface.
quantum_fabric_webgpu = ["quantum_fabric", "wai-quantum/quantum_fabric_webgpu"]
# The engine's energy-landscape kernels on a GPU (wai_quantum::quantum_landscape::{webgpu, cuda,
# metal}), which sink_meter::RunCounter meters with an on-chip counter (with sink_energy).
quantum_landscape = ["quantum_fabric", "wai-quantum/quantum_landscape"]
quantum_landscape_webgpu = ["quantum_landscape", "quantum_fabric_webgpu", "wai-quantum/quantum_landscape_webgpu"]
quantum_landscape_cuda = ["quantum_landscape", "wai-quantum/quantum_landscape_cuda"]
quantum_landscape_metal = ["quantum_landscape", "wai-quantum/quantum_landscape_metal"]
# Variational quantum machine learning (extensions/quantum-sim § VML,
# `wai.quantum.vml`): the interference-based ML the thermodynamic-computing pivot
# left behind — a Quantum Circuit Born Machine (P_θ(x)=|ψ_θ(x)|²) trained by the
@ -2369,6 +2375,12 @@ name = "sink_energy"
path = "tests/sink_energy.rs"
required-features = ["sink_energy"]
# The energy-landscape GPU members metered by an on-chip counter: hardware only, run by hand.
[[test]]
name = "landscape_meter"
path = "tests/landscape_meter.rs"
required-features = ["sink_energy", "quantum_landscape"]
[[test]]
name = "browser_sink"
path = "tests/browser_sink.rs"

268
wai/wai-rs/src/sink_meter.rs
View file

@ -59,6 +59,10 @@ pub trait Counter: Send {
fn wrap_nj(&self) -> Option<u64> {
None
}
/// The step its register counts in, in nanojoules: a read is a whole number of these.
fn step_nj(&self) -> u64 {
1
}
}
/// One interval of a bracket (sink-energy §4.3): the counter's advance in nanojoules across
@ -540,6 +544,9 @@ impl Counter for Nvml {
let mut mj = 0u64;
(unsafe { (self.energy)(self.dev, &mut mj) } == 0).then_some(mj.saturating_mul(1_000_000))
}
fn step_nj(&self) -> u64 {
1_000_000
}
}
// ---------------------------------------------------------------------------------------
@ -617,6 +624,199 @@ pub fn survey(children: &Children, every: Duration, intervals: u32) -> (Vec<Surv
(out, live)
}
// ---------------------------------------------------------------------------------------
// A counter as the meter of a fabric run
// ---------------------------------------------------------------------------------------
/// A counter as the meter of a fabric run (`wai_quantum::quantum_fabric::RunMeter`), which a
/// device member of the engine reads around each of its calls.
///
/// - **Live through the call.** It reads the counter at the call's start and end and every
/// sampling interval between, on a thread of its own, and gives a figure only where the
/// counter advanced over the window as a live rail does (energy-measurement §3.1): two
/// reads alone never pass. A window shorter than the profile's minimum gives none either.
/// - **Its uncertainty.** The relative uncertainty the operator declares for the counter at
/// these windows (§2.3, as [`crate::energy_meter::CounterMeter`] takes it), combined with
/// two terms the meter evaluates itself: the register's step, read at each end, and the
/// counter's refresh, measured when the meter is made, at each edge of the window.
/// - **What it covers.** What its profile covers, over the call's window (§9.1, attribution
/// none): the whole graphics rail or board, its idle draw, and any other work on the device
/// in that window. The meter says so ([`RunMeter::scope`]), and an answer carries it beside
/// the figure. Nothing here subtracts a baseline or claims a share for the run.
/// - **NVML** is metered only through [`RunCounter::nvml`], which holds the counter to §9.3's
/// plausibility before it gives a meter, and every window to its ceiling after.
///
/// [`RunMeter::scope`]: wai_quantum::quantum_fabric::RunMeter::scope
#[cfg(feature = "quantum_fabric")]
pub struct RunCounter {
counter: Arc<Mutex<Box<dyn Counter>>>,
profile: &'static str,
step_nj: u64,
wrap_nj: Option<u64>,
relative_ppm: u32,
minimum: Duration,
every: Duration,
refresh: Duration,
ceiling_mw: Option<u64>,
span: Option<RunSpan>,
}
#[cfg(feature = "quantum_fabric")]
struct RunSpan {
reads: Arc<Mutex<Vec<(u64, Instant)>>>,
done: Arc<std::sync::atomic::AtomicBool>,
sampler: std::thread::JoinHandle<()>,
}
/// How long a counter takes to refresh: the longest gap between its changes while it is
/// polled every millisecond, up to `within`. `None` when it changed fewer than three times.
#[cfg(feature = "quantum_fabric")]
fn refresh_of(c: &mut dyn Counter, within: Duration) -> Option<Duration> {
let start = Instant::now();
let mut last = c.read_nj()?;
let mut changes: Vec<Instant> = Vec::new();
while start.elapsed() < within && changes.len() < 6 {
std::thread::sleep(Duration::from_millis(1));
let now = c.read_nj()?;
if now != last {
changes.push(Instant::now());
last = now;
}
}
(changes.len() >= 3).then(|| changes.windows(2).map(|w| w[1] - w[0]).max().unwrap_or_default())
}
#[cfg(feature = "quantum_fabric")]
impl RunCounter {
/// A meter over `counter`, whose figures carry the relative uncertainty `relative_ppm`, in
/// parts per million, that the operator declares for the counter at the windows it will be
/// read over, with the meter's own terms beside it. Refused for an uncertainty of 0, for a
/// counter whose profile energy-measurement §9 does not give, for an NVML counter (see
/// [`RunCounter::nvml`]), and for a counter that does not refresh often enough to be shown
/// live in the profile's minimum window.
pub fn new(counter: Box<dyn Counter>, relative_ppm: u32) -> Result<RunCounter, String> {
if counter.profile() == "linux-nvml" {
return Err("an NVML counter is metered through RunCounter::nvml, which holds it to §9.3's plausibility first".into());
}
RunCounter::open(counter, relative_ppm, None)
}
/// A meter over NVML, once the counter has shown itself plausible (energy-measurement
/// §9.3): over five one-second intervals its rate stayed within 1.5 times the enforced
/// power limit and twice the mean power reading plus 5 W. Every window after is held to
/// the first of those, and gives no figure above it.
pub fn nvml(mut counter: Nvml, relative_ppm: u32) -> Result<RunCounter, String> {
let (ok, seen) = counter.plausible(5, Duration::from_secs(1));
if !ok {
let shown: Vec<String> = seen.iter().map(|(rate, power)| format!("{rate} mW against {power} mW")).collect();
return Err(format!("NVML is implausible here (energy-measurement §9.3), so it gives no figure: its counter's rate against its power reading, each second: {}", shown.join(", ")));
}
let ceiling = counter.limit_mw().map(|l| u64::from(l) * 3 / 2);
RunCounter::open(Box::new(counter), relative_ppm, ceiling)
}
fn open(mut counter: Box<dyn Counter>, relative_ppm: u32, ceiling_mw: Option<u64>) -> Result<RunCounter, String> {
if relative_ppm == 0 {
return Err("an on-chip counter's figure declares its uncertainty".into());
}
let profile = counter.profile();
let (_, minimum_us) = crate::sink_energy::profile(profile).ok_or_else(|| format!("{profile}: no profile gives this counter's minimum window"))?;
let minimum = Duration::from_micros(minimum_us);
let refresh = refresh_of(counter.as_mut(), Duration::from_secs(2)).ok_or_else(|| format!("{profile}: the counter did not change three times in 2 s, so it cannot be shown live"))?;
let every = (2 * refresh).max(minimum / 8);
if every > minimum / 5 {
return Err(format!("{profile}: the counter refreshes every {} ms, too seldom to be shown live in a {} ms window", refresh.as_millis(), minimum.as_millis()));
}
let (step_nj, wrap_nj) = (counter.step_nj().max(1), counter.wrap_nj());
Ok(RunCounter { counter: Arc::new(Mutex::new(counter)), profile, step_nj, wrap_nj, relative_ppm, minimum, every, refresh, ceiling_mw, span: None })
}
/// The same meter with another minimum window and sampling interval, for tests.
#[doc(hidden)]
pub fn with_timing(mut self, minimum: Duration, every: Duration) -> RunCounter {
self.minimum = minimum;
self.every = every;
self
}
/// The same meter with a ceiling on a window's rate, in milliwatts, for tests.
#[doc(hidden)]
pub fn with_ceiling_mw(mut self, ceiling_mw: u64) -> RunCounter {
self.ceiling_mw = Some(ceiling_mw);
self
}
/// How long the counter took to refresh when the meter was made.
pub fn refresh(&self) -> Duration {
self.refresh
}
}
#[cfg(feature = "quantum_fabric")]
impl wai_quantum::quantum_fabric::RunMeter for RunCounter {
fn start(&mut self) -> Result<(), String> {
let c = self.counter.lock().unwrap_or_else(|p| p.into_inner()).read_nj().ok_or_else(|| format!("{}: the counter did not read", self.profile))?;
let reads = Arc::new(Mutex::new(vec![(c, Instant::now())]));
let done = Arc::new(std::sync::atomic::AtomicBool::new(false));
let (r, d, counter, every) = (Arc::clone(&reads), Arc::clone(&done), Arc::clone(&self.counter), self.every);
let sampler = std::thread::spawn(move || {
while !d.load(std::sync::atomic::Ordering::Acquire) {
std::thread::park_timeout(every);
if d.load(std::sync::atomic::Ordering::Acquire) {
break;
}
let read = counter.lock().unwrap_or_else(|p| p.into_inner()).read_nj();
match read {
Some(v) => r.lock().unwrap_or_else(|p| p.into_inner()).push((v, Instant::now())),
// A read that fails ends the samples, and the window cannot be shown live.
None => break,
}
}
});
self.span = Some(RunSpan { reads, done, sampler });
Ok(())
}
fn stop(&mut self) -> Result<Option<(u64, wai_quantum::quantum_energy::EnergyClass)>, String> {
use wai_quantum::quantum_energy::{CounterFiltering, EnergyClass, Uncertainty};
let Some(span) = self.span.take() else { return Err("the meter was stopped without being started".into()) };
span.done.store(true, std::sync::atomic::Ordering::Release);
span.sampler.thread().unpark();
let _ = span.sampler.join();
let last = self.counter.lock().unwrap_or_else(|p| p.into_inner()).read_nj();
let Some(last) = last else { return Ok(None) };
let mut reads = std::mem::take(&mut *span.reads.lock().unwrap_or_else(|p| p.into_inner()));
reads.push((last, Instant::now()));
let window = reads[reads.len() - 1].1 - reads[0].1;
let window_us = window.as_micros() as u64;
let deltas: Vec<u64> = reads.windows(2).map(|w| advance(w[0].0, w[1].0, self.wrap_nj)).collect();
if window < self.minimum || !crate::energy_report::Liveness::from_advances(deltas.iter().map(|&d| d > 0)).is_live_over(window_us) {
return Ok(None);
}
let nj: u64 = deltas.iter().sum();
if nj < 1000 || self.ceiling_mw.is_some_and(|c| nj / window_us.max(1) > c) {
return Ok(None);
}
// The operator's term, the register's step at each end, and the counter's refresh at
// each edge of the window: each an error the reads cannot see, combined as independent.
let e_uj = nj as f64 / 1000.0;
let step = self.step_nj as f64 / 1000.0 / 6f64.sqrt();
let edge = e_uj * self.refresh.as_secs_f64() / window.as_secs_f64() / 6f64.sqrt();
let rel = ((f64::from(self.relative_ppm) / 1e6).powi(2) + (step / e_uj).powi(2) + (edge / e_uj).powi(2)).sqrt();
let relative_ppm = u32::try_from((rel * 1e6).ceil() as u64).unwrap_or(u32::MAX);
let uncertainty = Uncertainty { relative_ppm, window_us };
Ok(Some((nj / 1000, EnergyClass::OnChipCounter { uncertainty, filtering: CounterFiltering::Unknown })))
}
fn scope(&self) -> Option<String> {
Some(match self.profile {
"macos-gpu" => "the graphics processor's whole rail as IOReport's GPU Energy channel reports it, over the windows of this run's calls: this run's work, the rail's idle draw and any other work on the graphics processor in those windows, with no share attributed to this run".into(),
"linux-nvml" => "the whole board as NVML reports it, over the windows of this run's calls: this run's work, the board's idle draw and any other work on the board in those windows, with no share attributed to this run".into(),
p => format!("what the {p} profile covers (energy-measurement §9), over the windows of this run's calls, with no share attributed to this run"),
})
}
}
#[cfg(test)]
mod tests {
use super::*;
@ -654,6 +854,74 @@ mod tests {
assert_eq!(liveness(&mut Scripted(failing, 0), 10, d), CounterStatus::Frozen, "a read that fails is no reading");
}
/// A counter as a run's meter: a figure only from a window in which the counter was live,
/// at least the profile's minimum and under any ceiling; its uncertainty never below what
/// it declares, with the meter's own terms; and what it covers, said. Two reads alone, a
/// counter that moves in one lump, one that never changes, and an NVML counter offered as
/// any other, give nothing.
#[cfg(feature = "quantum_fabric")]
#[test]
fn a_counter_meters_a_run_only_where_it_was_live_and_says_what_it_covers() {
use wai_quantum::quantum_energy::{CounterFiltering, EnergyClass};
use wai_quantum::quantum_fabric::RunMeter;
/// A counter that rises `.1` nJ a read once it has been read `.3` times; `.4` counts reads.
struct Steps(u64, u64, &'static str, u64, u64);
impl Counter for Steps {
fn profile(&self) -> &'static str {
self.2
}
fn detail(&self) -> String {
"steps".into()
}
fn read_nj(&mut self) -> Option<u64> {
self.4 += 1;
if self.4 >= self.3 {
self.0 += self.1;
}
Some(self.0)
}
fn step_nj(&self) -> u64 {
1000
}
}
let rising = |profile| Box::new(Steps(0, 2_500_000, profile, 0, 0)) as Box<dyn Counter>;
assert!(RunCounter::new(rising("macos-gpu"), 0).is_err());
assert!(RunCounter::new(rising("no-such-profile"), 10_000).is_err());
assert!(RunCounter::new(rising("linux-nvml"), 10_000).err().unwrap().contains("RunCounter::nvml"));
assert!(RunCounter::new(Box::new(Steps(7, 0, "macos-gpu", 0, 0)), 50_000).err().unwrap().contains("cannot be shown live"));
let fast = |m: RunCounter| m.with_timing(Duration::ZERO, Duration::from_millis(20));
let mut m = fast(RunCounter::new(rising("macos-gpu"), 50_000).unwrap());
assert!(m.stop().is_err(), "stopped before it started");
// Two reads, start and stop: one interval, never live.
m.start().unwrap();
assert_eq!(m.stop().unwrap(), None);
// A live window of about 400 ms.
m.start().unwrap();
std::thread::sleep(Duration::from_millis(400));
let (uj, class) = m.stop().unwrap().expect("a reading");
assert!(uj >= 2_500 * 10, "{uj}");
let EnergyClass::OnChipCounter { uncertainty, filtering } = class else { panic!("{class:?}") };
assert_eq!(filtering, CounterFiltering::Unknown);
assert!(uncertainty.relative_ppm > 50_000, "the meter's own terms are added: {}", uncertainty.relative_ppm);
assert!(uncertainty.window_us >= 400_000, "{}", uncertainty.window_us);
assert!(m.scope().unwrap().contains("idle draw and any other work on the graphics processor"));
// Above the ceiling: no figure.
let mut hot = fast(RunCounter::new(rising("macos-gpu"), 50_000).unwrap()).with_ceiling_mw(1);
hot.start().unwrap();
std::thread::sleep(Duration::from_millis(400));
assert_eq!(hot.stop().unwrap(), None, "a window above the ceiling gives nothing");
// Still until late in the window, then one lump: not live.
let mut lump = fast(RunCounter::new(rising("macos-gpu"), 50_000).unwrap());
*lump.counter.lock().unwrap() = Box::new(Steps(0, 2_500_000_000, "macos-gpu", 18, 0));
lump.start().unwrap();
std::thread::sleep(Duration::from_millis(400));
assert_eq!(lump.stop().unwrap(), None, "one lump is one advance");
// A window under the profile's minimum gives none.
let mut short = RunCounter::new(rising("macos-gpu"), 50_000).unwrap();
short.start().unwrap();
assert_eq!(short.stop().unwrap(), None, "a window under the profile's 1 s gives no figure");
}
/// An advance across a counter's wrap, and none where a counter with no wrap goes back.
#[test]
fn an_advance_crosses_the_wrap() {

1
wai/wai-rs/tests/NOT-IN-CI.txt
View file

@ -36,3 +36,4 @@ jpegai_scope_ort reads the JPEG AI assets root (wai/jpegai-vm-assets, not commi
jpegai_sigma_index_calibration reads a fixture of the reference decoder's intermediate values, named by an environment variable and held outside the tree; without it the test returns without checking anything
jpegai_synthesis_calibration reads a fixture of the reference decoder's intermediate values, named by an environment variable and held outside the tree; without it the test returns without checking anything
jpegai_uv_calibration reads a fixture of the reference decoder's intermediate values, named by an environment variable and held outside the tree; without it the test returns without checking anything
landscape_meter hardware only: a GPU and its on-chip counter, each test ignored and run by hand with --ignored; its bridge, sink_meter::RunCounter, is tested in the library

88
wai/wai-rs/tests/landscape_meter.rs Normal file
View file

@ -0,0 +1,88 @@
//! The engine's energy-landscape GPU members metered by this machine's own on-chip
//! counter (`sink_meter::RunCounter`): the GPU rail through IOReport on macOS, the board
//! through NVML on Linux. Hardware only, run by hand with `--ignored --nocapture`.
//!
//! The operator's term these tests declare for the counter, 100,000 ppm, is a coarse choice
//! for a hand run, not an evaluation of the counter's own error; the meter adds its own terms
//! beside it, and nothing here seals a figure. The figure is the whole rail or board over the
//! call's window, idle draw and any other work on the device included, as the answer says.
#![allow(dead_code, unused_imports)]
use std::time::Instant;
use wai::sink_meter::RunCounter;
use wai_quantum::quantum_landscape::energy::Landscape;
use wai_quantum::quantum_landscape::member::{Chain, CpuReference, Member};
use wai_quantum::quantum_landscape::sample::{chain_seed, out, schedule, start};
use wai_quantum::quantum_landscape::settle::Workload;
const DECLARED_PPM: u32 = 100_000;
fn work(n: usize, chains: u64, sweeps: u32) -> (Workload, Vec<Chain>, Vec<f64>) {
let mut l = Landscape::new(n);
let mut k = 0u64;
for i in 0..n {
for _ in 0..4 {
k += 1;
let j = (out(0x7177, k) % n as u64) as usize;
if j != i {
l.couple(i, j, if k % 3 == 0 { -1 } else { 1 });
}
}
}
l.set_cost(1);
let cs = (0..chains)
.map(|r| {
let c = chain_seed(9, r);
Chain { start: start(c, n), seed: c }
})
.collect();
(l.lower(), cs, schedule(l.default_temperature(), sweeps))
}
fn metered(m: &dyn Member, w: &Workload, cs: &[Chain], temperatures: &[f64]) {
let t = Instant::now();
let got = m.chains(w, cs, temperatures).unwrap();
let wall = t.elapsed().as_secs_f64();
let reading = m.meter();
println!("{}: {} chains, {} cells, {} sweeps in {wall:.2} s wall-clock", m.id(), cs.len(), w.n(), temperatures.len());
match &reading {
Some(r) => println!(" read: {} µJ over {} µs, class {}\n covering {}", r.joules_micro, r.window_us, r.class, r.scope.as_deref().unwrap_or("this run alone")),
None => println!(" no reading"),
}
let r = reading.expect("a call longer than the profile's minimum window, on a live counter, gives a reading");
assert_eq!(r.class["energy_provenance"], "OnChipCounter", "{}", r.class);
assert!(r.scope.as_deref().is_some_and(|s| s.contains("idle draw")), "{:?}", r.scope);
// The same bits as the reference, on a slice of the chains.
assert_eq!(got[..4], CpuReference.chains(w, &cs[..4], temperatures).unwrap()[..]);
}
#[cfg(all(target_os = "macos", feature = "quantum_landscape_metal"))]
#[test]
#[ignore]
fn a_metal_member_reads_the_gpu_rail() {
let counter = wai::sink_meter::IoReportGpu::open().expect("IOReport's GPU Energy channel");
let m = wai_quantum::quantum_landscape::metal::discover().into_iter().find_map(Result::ok).expect("a Metal device");
let m = m.with_meter(Box::new(RunCounter::new(Box::new(counter), DECLARED_PPM).unwrap()));
let (w, cs, t) = work(4096, 1024, 400);
metered(&m, &w, &cs, &t);
}
#[cfg(all(target_os = "linux", feature = "quantum_landscape_cuda"))]
#[test]
#[ignore]
fn a_cuda_member_reads_the_board() {
let counter = wai::sink_meter::Nvml::open().expect("NVML");
// energy-measurement §9.3: the bridge itself refuses an implausible counter, so no figure
// can come from one. This hand run then has nothing to read, and says so.
let meter = match RunCounter::nvml(counter, DECLARED_PPM) {
Ok(m) => m,
Err(why) => {
println!("no meter: {why}");
return;
}
};
let m = wai_quantum::quantum_landscape::cuda::discover().into_iter().find_map(Result::ok).expect("a CUDA device");
let m = m.with_meter(Box::new(meter));
let (w, cs, t) = work(4096, 1024, 1200);
metered(&m, &w, &cs, &t);
}