.github/workflows/wai-ops-gpu.yml

@@ -5,19 +5,25 @@ name: wai-ops-gpu integer operators on GPUs
 # (64- and 128-bit arithmetic built from 32-bit words), and the integer decoders with
 # their graphs on the device (wai-rs `ops_gpu`).
 #
-#   - metal (the `macos` runner, whose host has a GPU): the arithmetic probe on a
+#   - metal (the `macos` runner, whose host has a GPU): WGSL's integer semantics
+#     (wai/wgsl-conformance, determinism-tiers §5, Claiming the tier on a GPU), the
+#     committed vectors and SplitMix64 cases of every operation, with an edited kernel
+#     refused at discovery; the arithmetic probe on a
 #     million cases; every graph vector of wai/ops-conformance that the subset covers,
 #     at both widths, against the committed results and the interpreter; the shader
 #     mutants, every one of which the corpus must catch; every committed vector and
 #     refusal fixture of every capability that runs a wai-ops/1 graph, the jpegai/v1
 #     inputs included, against the decode on the CPU; and the corpus again in a
 #     headless browser's WebGPU, the kernels compiled by the browser's own WGSL
-#     compiler. WAI_OPS_GPU_EXPECT declares the one adapter the host must offer, and
+#     compiler, with the semantics suite there and its vectors held to the reference
+#     compiled to wasm32, in the page and under Node. WAI_OPS_GPU_EXPECT declares the one adapter the host must offer, and
 #     WAI_OPS_GPU_REQUIRE makes an absent adapter a failure, never a skip; each step
 #     then requires the report lines with their counts, so a filtered or empty run
 #     fails too.
-#   - checks (`linux`, a container with no GPU): the crate and its tests build, the
+#   - checks (`linux`, a container with no GPU): the crate and its tests build; the
-#     device tests print SKIPPED and pass, and the same tests under
+#     WGSL rules on naga's IR (tests/wgsl_rules.rs) and the committed semantics vectors
+#     against the reference, which need no adapter, pass; the device tests print SKIPPED
+#     and pass, and the same tests under
 #     WAI_OPS_GPU_REQUIRE=1 fail: an absent adapter is not a pass. The browser module
 #     builds for wasm32-unknown-unknown; the bench tool builds.
 #
@@ -36,6 +42,7 @@ on:
       - 'wai/wai-ops-gpu/**'
       - 'wai/wai-ops/**'
       - 'wai/ops-conformance/**'
+      - 'wai/wgsl-conformance/**'
       # The decoders whose graphs run on the device, and what they decode.
       - 'wai/wai-rs/Cargo.toml'
       - 'wai/wai-rs/Cargo.lock'
@@ -115,7 +122,7 @@ jobs:
             fi
           else
             grep -q '^ops-gpu corpus: SKIPPED' "$rep" || { echo "::error::the corpus test printed neither a result nor SKIPPED"; exit 1; }
-            if WAI_OPS_GPU_REQUIRE=1 cargo test --locked --release --test corpus --test arith --test mutants > "$RUNNER_TEMP/require.log" 2>&1; then
+            if WAI_OPS_GPU_REQUIRE=1 cargo test --locked --release --test corpus --test arith --test mutants --test semantics > "$RUNNER_TEMP/require.log" 2>&1; then
               echo "::error::with no adapter, the device tests passed under WAI_OPS_GPU_REQUIRE=1"; exit 1
             fi
             grep -q 'WAI_OPS_GPU_REQUIRE is set' "$RUNNER_TEMP/require.log" \
@@ -157,16 +164,21 @@ jobs:
           : > "$WAI_OPS_GPU_REPORT"
           nice -n 10 cargo test --locked --release -- --nocapture --test-threads=1 2>&1 | tee "$log"
           k=gpu-metal-0
-          # Six tests in four binaries: the library's one, arith's two, corpus's two and
+          # Sixteen tests in six binaries: the library's four, arith's two, corpus's two,
-          # mutants' one, none filtered, ignored or failed.
+          # mutants' one, semantics' three and wgsl_rules' four, none filtered, ignored or
+          # failed.
           ran=$(awk '/^test result: ok\. [0-9]+ passed; 0 failed; 0 ignored; 0 measured; 0 filtered out/ { n += $4 } END { print n + 0 }' "$log")
-          [ "$ran" = 6 ] || { echo "::error::$ran tests ran and passed, not 6"; exit 1; }
+          [ "$ran" = 16 ] || { echo "::error::$ran tests ran and passed, not 16"; exit 1; }
           ! grep -q '^test result: FAILED' "$log" || { echo "::error::a test failed"; exit 1; }
           grep -qxF "ops-gpu adapters: as declared ($k)" "$WAI_OPS_GPU_REPORT" || { echo "::error::the adapters found are not $k"; exit 1; }
           grep -qE "^ops-gpu arith: key=$k .* cases=[0-9]+ words=[0-9]+ equal$" "$WAI_OPS_GPU_REPORT" || { echo "::error::no probe result for $k"; exit 1; }
           grep -qE "^ops-gpu corpus: key=$k .* acc32 vectors=[1-9][0-9]* decodes=[1-9][0-9]*; acc64 vectors=[1-9][0-9]* decodes=[1-9][0-9]*; .* mismatches=0 errors=0 " "$WAI_OPS_GPU_REPORT" \
             || { echo "::error::no equal corpus run for $k"; exit 1; }
           grep -qxF "ops-gpu mutants: key=$k killed 18 of 18" "$WAI_OPS_GPU_REPORT" || { echo "::error::a mutant survived on $k, or the mutant test did not run"; exit 1; }
+          grep -qE "^ops-gpu semantics: key=$k .* committed=[1-9][0-9]* cases=[1-9][0-9]* guarded=[1-9][0-9]* differ=0 " "$WAI_OPS_GPU_REPORT" \
+            || { echo "::error::no run of WGSL's integer semantics on $k with every guarded case equal"; exit 1; }
+          grep -qE "^ops-gpu semantics mutant: key=$k refused: the WGSL semantics suite: rem_i32 " "$WAI_OPS_GPU_REPORT" \
+            || { echo "::error::the edited semantics kernel was not refused on $k"; exit 1; }
       - name: The jpegai/v1 inputs
         shell: bash
         run: bash wai/tools/wai_jpegai_fetch_e2e_inputs.sh "$RUNNER_TEMP/jpegai"
@@ -214,7 +226,14 @@ jobs:
           nice -n 10 cargo build --locked --release --target wasm32-unknown-unknown
           t="${CARGO_TARGET_DIR:-target}"
           wasm-bindgen --target web --out-dir pkg --out-name wai_ops_gpu_browser "$t/wasm32-unknown-unknown/release/wai_ops_gpu_browser.wasm"
+          # The committed semantics vectors against the reference compiled to wasm32,
+          # under Node, with no GPU.
+          node semantics_reference.mjs
           log="$RUNNER_TEMP/browser.log"
           nice -n 10 node run.mjs 2>&1 | tee "$log"
           grep -qE '^browser webgpu: key=gpu-webgpu-[0-9a-f]+ .* acc32 vectors=[1-9][0-9]* decodes=[1-9][0-9]*; acc64 vectors=[1-9][0-9]* decodes=[1-9][0-9]*; .* mismatches=0 errors=0$' "$log" \
             || { echo "::error::no equal corpus run on a GPU adapter in the browser"; exit 1; }
+          grep -qE '^wasm reference: target=wasm32 .* cases=[1-9][0-9]* equal$' "$log" \
+            || { echo "::error::the semantics vectors are not the wasm32 reference's in the page"; exit 1; }
+          grep -qE '^browser webgpu semantics: key=gpu-webgpu-[0-9a-f]+ .* guarded=[1-9][0-9]* differ=0 ' "$log" \
+            || { echo "::error::no run of WGSL's integer semantics on a GPU adapter in the browser with every guarded case equal"; exit 1; }

wai/.gitignore

@@ -53,6 +53,8 @@ wai-web/demo/pkg/
 wai-web/demo/pkg-node/
 # The WebGPU member's bundle (demo/quantum_webgpu.html; built with --features quantum_webgpu).
 wai-web/demo/pkg-webgpu/
+# The platform decode bundle (demo/platform_decode.html; built with --features platform_decode).
+wai-web/demo/pkg-platform/
 # int_hyper conditional-codec demo artifacts (PTQ model + bitstream, large —
 # regenerated by tools/wai_sota_int_synth_export.py into demo/int_hyper/).
 wai-web/demo/int_hyper/

wai/CHANGELOG.md

@@ -6,6 +6,61 @@ wire format may change and such changes are recorded here.
 
 ## v0.4.0 — current
 
+### GPUs: when `decode-equivalence` may be claimed, and WGSL's integer semantics on every adapter
+
+- **Claiming the tier on a GPU** ([determinism-tiers](extensions/determinism-tiers.md)
+  §5). A decoder claims `decode-equivalence` for output whose graphs ran on a GPU only
+  when every shader module that ran them is integer only (no floating-point type, so no
+  reassociation or fusion can change a result), uses no subgroup or quad operation and
+  enables no extension, has no result that depends on scheduling (one atomic kind, its
+  result unused; no out-of-bounds access), and is translated with the translator's integer
+  guards on: signed arithmetic through unsigned, guarded division and remainder, `abs` and
+  negation of the most negative value. Where the translator does not produce WGSL's
+  result, the module does not reach the operation: naga 30.0.1 passes shift amounts to
+  MSL, HLSL and SPIR-V unreduced and writes integer `clamp` as MSL's `clamp`, so every
+  shift amount is masked below 32 and every `clamp` has literal bounds in order. Before an
+  adapter runs a graph it passes WGSL's integer semantics vectors and the implementation's
+  probe; the claim names the adapter and the translator.
+- **WGSL's integer semantics vectors** (`wgsl-conformance/`). Every integer operation WGSL
+  defines on `i32`, `u32` and their two-component vectors, at the cases WGSL defines at run
+  time and MSL, HLSL and SPIR-V leave open, each result written from the WGSL text
+  (Candidate Recommendation Draft, 21 October 2026). Cases the kernels never reach are
+  marked `unguarded`: reported, not deciding.
+- **`wai-ops-gpu`.** Discovery runs the semantics suite on each adapter after the
+  arithmetic probe, and refuses an adapter that differs on a guarded case. Every run-time
+  shift amount in the kernels is now masked (`& 31u`) and the two run-time clamps are
+  `min(max(…))`; WGSL gives the same result either way, and the translated code is then
+  defined on every backend. `tests/wgsl_rules.rs` holds every shipped module to the
+  conditions on naga's IR and in its MSL and HLSL translations, with planted violations
+  refused; the browser page runs the suite through the browser's WGSL compiler and checks
+  the vectors against the reference compiled to wasm32. Measured on Metal, a browser's
+  WebGPU, Vulkan on two vendors' GPUs on Linux and Windows, and DX12 on both on Windows:
+  every guarded case equal, every unguarded one too, and the capability corpora decoding
+  with their graphs on the device to the CPU decode's bytes.
+
+### Dispatch: what a web platform says it decodes, before a learned reconstruction
+
+- **Platform decode capabilities in dispatch** ([platform-decode](extensions/platform-decode.md)).
+  A sink in a web platform queries Media Capabilities `decodingInfo()` and WebCodecs
+  `isConfigSupported()` with a configuration read from the entry's payload (for AV1, the
+  full codec string of the AV1 ISOBMFF binding from the sequence header; for HEVC, the
+  string of ISO/IEC 14496-15 Annex E.3 from the SPS), never from what an encoder declared.
+  `supported: false` makes the capability unsupported for that object; `smooth` and
+  `powerEfficient` only order the entries, and are never written as energy. The default
+  order, a deployer policy SPEC §7.1 allows, tries classical decode the platform marks
+  power-efficient, and learned decode that carries a qualifying energy budget (an
+  `OnChipCounter` figure measured on a reference sink, as energy-measurement defines
+  it), first; then smooth classical decode; then any other classical decode; then learned
+  decode without a budget. For a `WAI1` envelope, a learned capability without a budget is
+  not advertised when its fallback is a power-efficient classical decode (SPEC §4).
+- **Reference and vectors.** `wai::platform_decode`, with `platform-decode-conformance/`
+  (the order, the `WAI1` rule, and queries from packaging-corpus floors and an SPS a
+  hardware encoder wrote), and a `platform_decode` bundle of `wai-web` with a demonstration
+  page. In a browser on an arm64 desktop, every AV1 floor of the packaging corpus decoded
+  through WebCodecs with the derived configuration, and an HEVC stream from the platform's
+  own encoder, whose SPS states a level other than the one its declared codec string
+  repeats, decoded back with the codec string read from the stream.
+
 ### Packaging: each step on the fabric that does it with the least energy
 
 - **Steps and fabrics** ([packaging](extensions/packaging.md) §11.6). The computations of

wai/SPEC.md

@@ -539,7 +539,10 @@ listing names, or one held file); a file the decode needs outside them is
 the algorithm above runs unchanged, so a sink MUST NOT report a file inert because
 of module authorisation while a fallback it can decode without that module remains,
 and an encoder emitting a capability whose decode needs a module MUST declare such a
-fallback (§3). A sink configured with no key log holds no module under §7.9, so it
+fallback (§3). Advertising is per envelope in one other place:
+[platform-decode](extensions/platform-decode.md) §4.3, under which a sink does not
+advertise a learned capability without an energy budget for an envelope whose fallback a
+web platform decodes power-efficiently. A sink configured with no key log holds no module under §7.9, so it
 runs no module obtained apart from an envelope, and it does not run check 5. The
 reference dispatcher in `wai-rs` runs no module obtained apart from an envelope; it
 runs check 5 when the sink is configured with a key log (`key_log::SinkKeyLog`,
@@ -2478,6 +2481,9 @@ A v1.1-aware sink MUST:
    skipped as if the sink did not advertise its capability, and the sink keeps the
    reason. A deployer policy MAY reorder or restrict the entries the sink
    considers; it MUST NOT pick an entry skipped for its digest or its pin.
+   [platform-decode](extensions/platform-decode.md) §4 defines one: an order from
+   what a web platform says it can decode, smoothly and power-efficiently, and from
+   the energy budgets of learned entries.
 3. Read the corresponding rendition's bytes from the payload block.
 4. Dispatch as in v1.0 §4, with the entry's `capability` and `prior` in place of
    `model_requirement.capability` and `conditioning.prior`, and no fallback. A

wai/extensions/determinism-tiers.md

@@ -523,31 +523,42 @@ shows, by fabric:
   (integer-operators §7.3). The host makes the interpreter's run-time checks before the
   device runs anything, and a `check` tensor that leaves its range refuses the decode at
   its operator. The driver's stage operators stay on the CPU, and so does a graph with an
-  operator outside `wai-ops/1-i32`. The tier is met for those graphs on each adapter and
+  operator outside `wai-ops/1-i32`. The tier is claimed for those graphs only under the
-  backend below, where a run reproduced, at both widths, every graph vector of
+  conditions of *Claiming the tier on a GPU* (below), which the crate meets, and only on
-  integer-operators §14.1 and §14.2 that `wai-ops/1-i32` covers (§6):
+  an adapter where a run, on the kernel text of this revision, gave WGSL's result for
+  every guarded case of the integer semantics suite, the host's result for every word of
+  the arithmetic probe, and the interpreter's outputs, at both widths, for every graph
+  vector of integer-operators §14.1 and §14.2 that `wai-ops/1-i32` covers (§6). That is
+  each adapter and backend below:
   - Metal, on the integrated GPU of an arm64 desktop (macOS 27.0.1);
   - a browser's WebGPU on that same GPU, the kernels compiled by the browser's own WGSL
     compiler (Chromium 148 and a Chromium-based browser at 155, headless);
   - Vulkan, on a discrete GPU of PCI vendor 0x10de, device 0x28e1 (driver 595.84), and
     on an integrated GPU of PCI vendor 0x8086, device 0xa7a0 (Mesa 25.2.8), in one Linux
     x86_64 host;
-  - Vulkan, on an integrated GPU of PCI vendor 0x1002, device 0x1586 (Mesa 26.0.8),
+  - Vulkan and DX12, on a discrete GPU of PCI vendor 0x10de, device 0x1f15 (Vulkan driver
-    Linux x86_64.
+    591.74, DX12 driver 32.0.15.9174), and on an integrated GPU of PCI vendor 0x8086,
+    device 0x9a49 (Vulkan, and DX12 driver 31.0.101.3222), in one Windows x86_64 host. The
+    DX12 backend translates the kernels to HLSL, compiled for shader model 5.1.
 
   On each of these but the browser, the capability corpora (integer-operators §14.4)
-  also decoded with their graphs on the device: every committed vector and refusal
+  also decoded with their graphs on the device, as §6 records: every committed vector and
-  fixture of `wai.neural.int_hyper`, `wai.video.int_hyper`, `wai.audio.int_codec`,
+  refusal fixture of `wai.neural.int_hyper`, `wai.video.int_hyper`, `wai.audio.int_codec`,
   `wai.splat.int_codec`, `wai.neural.int_synth` and `wai.neural.int_mlicpp` gave the
   bytes, or the refusal code, of the decode on the CPU, and every graph the device ran
-  gave the interpreter's outputs. A software Vulkan driver (llvmpipe, on both Linux hosts)
+  gave the interpreter's outputs. Software adapters, which run on the CPU, gave the same
-  gave the same results; it runs on the CPU. Nothing is claimed for an adapter, a backend
+  results: a software Vulkan driver (llvmpipe) on the Linux host, and the software DX12
-  or a driver version that is not listed; DX12 was not run. `wai.image.jpegai`'s scale
+  adapter (PCI vendor 0x1414) on the Windows host. An integrated GPU of PCI vendor 0x1002,
-  graphs (`hs_y`, `hs_uv`, at `acc32`) and hyper decoders and context stages (`rec_y`,
+  device 0x1586 (Mesa 26.0.8, Linux x86_64) reproduced the graph vectors and the
-  `rec_uv`, at `acc64`) ran on each of these devices but the browser with the
+  capability corpora on the kernel text before this revision, which differs from it only
-  interpreter's outputs, so its `entropy-consistency` claim (§5) holds with its scale
+  in the shift masks and the two clamps that *Claiming the tier on a GPU* requires; the
-  path there. Its synthesis graphs use operators outside `wai-ops/1-i32` and ran on the
+  semantics suite has not run there, and the tier is not claimed on it until a run of
-  CPU.
+  this revision's kernels passes there. Nothing is claimed for an adapter, a backend or
+  a driver version that is not listed. `wai.image.jpegai`'s scale graphs (`hs_y`,
+  `hs_uv`, at `acc32`) and hyper decoders and context stages (`rec_y`, `rec_uv`, at
+  `acc64`) ran on each of these devices but the browser with the interpreter's outputs,
+  so its `entropy-consistency` claim (§5) holds with its scale path there. Its synthesis
+  graphs use operators outside `wai-ops/1-i32` and ran on the CPU.
 - **NPU.** No decoder of these capabilities runs on one in this revision, and the tier is
   not claimed there. A network evaluated in floating point on a GPU or NPU does not meet
   it (§1).
@@ -560,6 +571,105 @@ digest, a reconstruction or an energy figure to it. Content that the application
 decrypts itself, as in Lane B of SPEC §6.5, is decoded like any plaintext, and the
 tier holds for it as it does for plaintext.
 
+### Claiming the tier on a GPU
+
+`decode-equivalence` holds on a GPU only where the shader that runs a decoder's graphs,
+the translation of that shader into the backend's language, and the device all keep the
+integer arithmetic the capability's registration fixes. The texts read here are WGSL, W3C
+Candidate Recommendation Draft, 21 October 2026, and naga 30.0.1, the translator wgpu
+30.0.1 compiles WGSL with.
+
+WGSL defines each integer operation at run time, the cases that are errors when the
+operands are known before the shader runs included (WGSL §8.8, §8.10, §17.5): addition,
+subtraction, multiplication and negation of a concrete integer type are modulo 2^32;
+`e1 / 0` is `e1`, and `e1 % 0` is 0; the most negative `i32` divided by −1 is itself, and
+its remainder 0; a shift amount is taken modulo 32; an integer `clamp(e, low, high)` is
+`min(max(e, low), high)`; a bit field's offset and count are clamped to the word. The
+languages a WGSL module is translated into leave some of these open: MSL and HLSL leave
+signed overflow undefined, and MSL, HLSL and SPIR-V leave a zero divisor, the most
+negative value divided by −1, a shift by the width or more, and a `clamp` whose bounds are
+out of order undefined or to the implementation. Floating point is open even in WGSL: an
+implementation may reassociate operations, and may fuse them where the result is at least
+as accurate (WGSL §15.7.5). A subgroup operation's result depends on its subgroup, whose
+size is a power of two from 4 to 128 that the device compiler chooses, whose invocations
+have no defined relationship to `local_invocation_index`, and of whose invocations only
+the active ones take part (WGSL §15.5, §15.6.3). An out-of-bounds access has several
+permitted outcomes (WGSL §6.5.7).
+
+So a decoder **MAY** claim `decode-equivalence` for output whose graphs ran on a GPU only
+when every condition below holds for every shader module that ran them, on that adapter,
+in that process. Otherwise the output is the GPU's, and no tier is claimed for it.
+
+1. **Integer only.** Every scalar type in the module, of every declaration and of every
+   expression, is `bool`, `i32` or `u32`; every literal is of one of them; every
+   conversion is between them. No floating-point type appears for any purpose, an index
+   or a loop bound included. With no floating-point operation, no reassociation or fusion
+   an implementation may make changes a result: integer addition and multiplication
+   modulo 2^32 give the same value in any order.
+2. **No subgroup or quad operation, and no extension.** The module enables nothing
+   (`enable`) and calls no subgroup or quad built-in.
+3. **No result depends on scheduling.** Each output word is written by one invocation and
+   read by no other invocation of the same dispatch. An atomic operation is one whose
+   final value does not depend on the order in which invocations reach it (`atomicOr`,
+   `atomicAnd`, `atomicMin`, `atomicMax`, `atomicAdd`, `atomicSub`), and the value it
+   returns is not used. No access is out of bounds.
+4. **The translation gives WGSL's results.** A translation into a language that leaves an
+   integer result open produces WGSL's result for each operation the module uses: signed
+   addition, subtraction, multiplication and negation evaluated in unsigned arithmetic
+   and reinterpreted; division and remainder through a guard that substitutes the divisor
+   1 where WGSL's result is the dividend, and remainder 0; the absolute value of the most
+   negative value, itself. naga 30.0.1 does all of these in MSL and HLSL (its `naga_div`,
+   `naga_mod`, `naga_neg` and `naga_abs`, and signed arithmetic through unsigned); in
+   SPIR-V, whose integer addition, subtraction, multiplication and negation wrap, it
+   guards division and remainder. In MSL and SPIR-V the division guards are written only
+   while its integer-division checks are on; in HLSL always. A host
+   **MUST NOT** turn a translator's guards off: in wgpu, a module created unchecked, or
+   `ShaderRuntimeChecks::int_div_checks` false. Where a translator does not produce
+   WGSL's result for an operation, the module **MUST NOT** reach that operation on the
+   operands concerned. naga 30.0.1 passes a shift amount to MSL, HLSL and SPIR-V without
+   reducing it modulo 32, and writes an integer `clamp` as MSL's `clamp`; so every shift
+   amount in the module is a literal below 32 or masked below 32 (`e & 31u`), and every
+   `clamp` has literal bounds in order, another bound being written
+   `min(max(e, low), high)`.
+5. **The adapter passes the conformance suite first.** Before an adapter runs any graph of
+   the decode, in the same process, it gives WGSL's result for every case of the integer
+   semantics vectors (`wai/wgsl-conformance/integer.jsonl`) that the file marks guarded,
+   and for further cases of each operation that the implementation draws, and the host's
+   result for every word of the implementation's probe of the wider arithmetic it builds
+   from 32-bit words. And the implementation has reproduced, on that adapter, backend and
+   driver, every graph vector of integer-operators §14.1 and §14.2 its subset covers, at
+   each width it runs (integer-operators §12, *Evidence for a fabric*). An adapter that
+   differs on any guarded case or probe word runs no graph.
+6. **The claim names what ran.** A receipt or report of a GPU decode at this tier names
+   the adapter (backend, PCI vendor and device, driver) and the translator and its
+   release.
+
+The semantics vectors cover every integer operation WGSL defines on `i32`, `u32` and
+two-component vectors of them, with the edge cases above: every pair of a set of values
+at and near the edges of both types for the arithmetic, division and remainder, every
+shift amount from 0 to 40 and the extremes past it, and bit fields at and past the word.
+They mark as unguarded the cases item 4 keeps a module from, a shift amount of 32 or more
+and `clamp` bounds out of order, whose result is the hardware's through naga 30.0.1. A run
+reports what each adapter gives on them; a difference there does not refuse the adapter,
+since no module at this tier reaches one. The file is written from WGSL's definitions by
+`wai_ops_gpu::semantics::reference` and checked against it natively, and against the same
+reference compiled to wasm32, under Node and in a browser page.
+
+`wai/wai-ops-gpu` meets these conditions. `tests/wgsl_rules.rs` holds every shipped module
+to items 1–4 on naga's IR: the types and literals; a closed list of expressions and
+built-in functions; the one atomic operation, `atomicOr`, and its unused result; the
+bindings, with no workgroup or private memory and no barrier; the shift masks and `clamp`
+bounds; that every integer operation a kernel uses is one the semantics suite runs; and
+naga's guard helpers in each module's MSL and HLSL translations, which disappear from the
+MSL with its division checks turned off. It also holds the crate to the checked
+module-creation call, and shows each planted violation (a floating-point value, `f16`, a
+subgroup reduction, an `atomicAdd`, an atomic's result read, workgroup memory and a
+barrier, an unmasked shift, a `clamp` with run-time bounds) refused by the rule it
+breaks. Discovery runs item 5's suite and probe on each adapter before it runs a graph;
+`tests/semantics.rs` shows every adapter refused when the suite's remainder is written by
+hand, `a - (a / b) * b`, which gives `a` where WGSL gives 0. In a browser,
+`wai-ops-gpu/browser` runs the same suite through the browser's own WGSL compiler.
+
 ## 6. Evidence status
 
 What the repository demonstrates, claim by claim, as of 2026-09-27 (main at
@@ -595,7 +705,7 @@ task list has no Windows task after that day. The Windows results below are hist
 | `wai.neural.int_hyper` and `wai.video.int_hyper` at full model scale (the pretrained `mbt2018_mean`, exported to integers): the Rust decode reproduces the export tool's integer reference byte for byte (synthesis, hyper-decoder, the whole conditional bitstream, an eight-frame clip), and costs at most 1 dB against the same model's float decode on one pinned photograph | `wai-sota-vectors.yml` regenerates the vectors on the runner from pinned inputs (the photograph's pixels and the checkpoint's bytes by SHA-256, the Python stack by version) and runs both tests with `WAI_REQUIRE_SOTA_VECTORS=1`, under which a missing vector fails instead of skipping. The job also fails on any skip banner in the test log, and runs `require_vectors_rule.rs`, which checks that without the switch a variable naming absent or incomplete inputs fails as well. Rust is checked against the export of the same run on the same machine: the vectors come from float computation, and some of them differed between machines and between runs, so none of their hashes is pinned and none is compared across machines. The workflow has run on the forge's Linux runner (x86_64 hardware) and passed on every run since 2026-09-28 | `wai-rs/tests/int_hyper_synth.rs`, `int_video_sota_keyframe.rs`, `require_vectors_rule.rs`; `.github/workflows/wai-sota-vectors.yml` |
 | `wai.audio.int_codec`, `wai.splat.int_codec` | committed vectors that never skip. Their frozen hashes are the output of an independent Python integer reference, computed when the vectors were exported (`tools/wai_int_audio_export.py`, `tools/wai_int_splat_export.py`); no Python runs in CI. The conformance workflow's fast-kernel steps enable `audio_int` and `splat_int` and, in debug and release builds, decode both vectors against those hashes under the reference kernels and every fast-kernel policy (`tests/int_kernels_vectors.rs`); that passed on the forge's Linux x86_64, Linux arm64 and macOS arm64 runners for the landing chain's head `7cfcecfa` (run 745, 2026-09-30). The same steps also name `tests/int_audio.rs` and `tests/int_splat.rs`, which check each vector against the same hash and against its own metadata (sample count; attribute count, tensor size and grid side); Cargo declares each with its feature as a required feature, so a step that names it without that feature fails instead of running nothing. This row records no forge run of those two | `wai-rs/tests/int_audio.rs`, `int_splat.rs`, `int_kernels_vectors.rs`; `.github/workflows/wai-conformance.yml` |
 | `wai.neural.int_mlicpp` on its committed vector (synthetic seeded weights; a 64×64 image, 8 latent channels in 4 slices): the decode reproduces the export tool's independent integer reference byte for byte — the RGB, and every scale bucket and symbol of the latent stream in coding order (512 symbols over all 16 buckets of the scale table, 47 of the symbols coded by the rANS escape) — under the reference kernels and under the fast kernels on 1, 2, 3, 5 and all available threads; one integer unit changed in each of four chosen parameters, one on each path that feeds a value (a hyper-decoder weight, an entropy-network weight, a mean-head bias, a synthesis weight), changes the image; a header whose shapes the model cannot produce, or on which a network would form a buffer over the cap, is refused before anything is allocated (`tests/int_mlicpp_preflight.rs`, which counts allocations); damaged payloads are refused or decode, and never panic in an overflow-checked build | measured 2026-09-29 with rustc 1.98.0 on the source of the change that added this row: macOS arm64 native (the wai-rs tests, with overflow checks on and off; the `byte-exact-conformance` golden, release) and wasm32-wasip2 under wasmtime 49.0.1 (the same tests and golden, release; on wasm the fast kernels run single-threaded as portable scalar code, and `usize` is 32-bit). All pinned hashes matched on both. Not measured: x86_64 — the development machine has no x86_64 translation layer installed, and an x86_64-apple-darwin build of the golden compiles but cannot execute there — nor Linux or Windows. The golden runs on every leg of the byte-exact matrix (`wai-conformance.yml`, `wai-byte-equality.yml`) and `tests/int_mlicpp.rs` on every leg of `wai-conformance.yml` in both profiles. Both passed on the forge's Linux x86_64, Linux arm64 and macOS arm64 runners for the landing chain's head `7cfcecfa` (run 745, 2026-09-30); there is no Windows leg. The dispatch through `decode_envelope_strict` (the prior pinning the one-file model bundle, the strict gate) runs in `wai-neural-gates.yml`. The weights are synthetic, so nothing here measures compression quality | `wai-rs/tests/int_mlicpp.rs`, `wai-rs/tests/int_mlicpp_preflight.rs`, `byte-exact-conformance/src/int_mlicpp_golden.rs`, `wai-rs/tests/neural_native_decode.rs`; `tools/wai_int_mlicpp_export.py` |
-| `wai-ops/1-i32` graphs on GPUs (`wai/wai-ops-gpu`, §5): each device's arithmetic agrees with 128-bit host arithmetic on every probe case; every graph vector of `ops-conformance/` that the subset covers gives the interpreter's outputs or refusal code, at `acc32` and at `acc64`; the integer decoders, with their graphs on the device, give the CPU decode's bytes or refusal codes on every committed vector and refusal fixture, and every graph the device runs gives the interpreter's outputs | measured 2026-10-20 and 2026-10-21 with rustc 1.98.0 and wgpu 30.0.1, release builds, on the source of the change that added this row, on the adapters §5 lists: Metal on macOS 27.0.1 (arm64); headless Chromium 148, and a Chromium-based browser at 155, on that GPU; Vulkan on Linux x86_64 with the 0x10de and 0x8086 adapters (driver 595.84, Mesa 25.2.8) and, on another host, with the 0x1002 adapter (Mesa 26.0.8); and llvmpipe on both Linux hosts. On each, the probe (`tests/arith.rs`) agreed on every word of 1,051,757 cases, and the corpus (`tests/corpus.rs`) ran 813 vectors (1,104 decodes) on the 32-bit kernels and 1,426 (1,921 decodes) on the 64-bit ones, 212 of the decodes refusals, with no difference; 209 graphs use an operator outside the subset and ran on the CPU, and 95 vectors are refused at load. Run on the 32-bit kernels, 330 of the 613 graphs whose certificate does not meet `acc32` gave other results, so the width a graph runs at is the certificate's. Eighteen edits of one kernel expression each (a rounding shift or its addend off by one, a tie rule, a clamp bound moved by one, a dropped carry or cross term, a sign rule, a post bias, an index rule; `tests/mutants.rs`) were each caught by the corpus on every adapter, and an edited product by the probe. `wai-rs/tests/ops_gpu.rs` decoded the committed vectors and refusal fixtures of the six capabilities §5 names, the world renderer's two vectors and `wai.image.jpegai`'s `jpegai/v1` codestream, comparing every graph the device ran with the interpreter as well as each decode's bytes; none differed. The browser leg ran the corpus, not the decoders. `wai-ops-gpu.yml` runs the Metal and browser legs on the `macos` runner, where an absent adapter fails; this row records no forge run of it yet. No other workflow runs a GPU leg of these tests | `wai-ops-gpu/tests/`, `wai-ops-gpu/browser/`, `wai-rs/tests/ops_gpu.rs`; `.github/workflows/wai-ops-gpu.yml` |
+| `wai-ops/1-i32` graphs on GPUs (`wai/wai-ops-gpu`, §5, *Claiming the tier on a GPU*): every shipped module meets items 1–4 on naga's IR; each device gives WGSL's result for every guarded case of the integer semantics suite and agrees with 128-bit host arithmetic on every probe case; every graph vector of `ops-conformance/` that the subset covers gives the interpreter's outputs or refusal code, at `acc32` and at `acc64`; the integer decoders, with their graphs on the device, give the CPU decode's bytes or refusal codes on every committed vector and refusal fixture, and every graph the device runs gives the interpreter's outputs | measured 2026-10-22 with rustc 1.98.0 and wgpu 30.0.1 (naga 30.0.1), release builds, on the source of the change that added *Claiming the tier on a GPU*, on the adapters §5 lists: Metal on macOS 27.0.1 (arm64); headless Chromium 148, and a Chromium-based browser at 155, on that GPU; Vulkan on Linux x86_64 with the 0x10de and 0x8086 adapters (driver 595.84, Mesa 25.2.8) and llvmpipe; Vulkan and DX12 on Windows x86_64 with the 0x10de and 0x8086 adapters, and DX12 with the software adapter (PCI vendor 0x1414). `tests/wgsl_rules.rs` passed, its nine planted violations each refused by the rule named. On each native adapter, the semantics suite (`tests/semantics.rs`) gave WGSL's result for every one of 254,464 cases, the 25,088 committed and 4,096 SplitMix64 cases of each operation: 231,690 guarded, and the 22,774 unguarded too (a shift amount of 32 or more, `clamp` bounds out of order), so on these adapters the hardware's result there was WGSL's; the browser legs ran 82,432 cases with none differing, and the committed file reproduced under the reference compiled to wasm32, in the page and under Node 26. The probe (`tests/arith.rs`) agreed on every word of 1,051,757 cases; the corpus (`tests/corpus.rs`) ran 813 vectors (1,104 decodes) on the 32-bit kernels and 1,426 (1,921 decodes) on the 64-bit ones, 212 of the decodes refusals, with no difference; 209 graphs use an operator outside the subset and ran on the CPU, and 95 vectors are refused at load. Run on the 32-bit kernels, the graphs whose certificate does not meet `acc32` gave other results on some of them, so the width a graph runs at is the certificate's. Eighteen edits of one kernel expression each (`tests/mutants.rs`) were each caught by the corpus on every adapter, an edited product by the probe, and a remainder written by hand by the semantics suite at discovery. `wai-rs/tests/ops_gpu.rs` decoded the committed vectors and refusal fixtures of the six capabilities §5 names, the world renderer's two vectors and `wai.image.jpegai`'s `jpegai/v1` codestream on Metal and on the Linux host's Vulkan adapters, comparing every graph the device ran with the interpreter as well as each decode's bytes; none differed; the same on the Windows host's Vulkan and DX12 adapters and on its software DX12 adapter. The browser legs ran the corpus and the semantics suite, not the decoders. Before this change the kernels differed only in the shift masks and the two clamps; on that text, the same tests but the semantics suite passed on 2026-10-20 and 2026-10-21 on Metal (macOS 27.0.1, arm64), the browser at 148 and 155, the Linux host's adapters, and an integrated GPU of PCI vendor 0x1002, device 0x1586 (Mesa 26.0.8), on which nothing has run since. `wai-ops-gpu.yml` runs the Metal and browser legs on the `macos` runner, where an absent adapter fails; this row records no forge run of it yet. No other workflow runs a GPU leg of these tests | `wai-ops-gpu/tests/`, `wai-ops-gpu/browser/`, `wai-rs/tests/ops_gpu.rs`, `wgsl-conformance/`; `.github/workflows/wai-ops-gpu.yml` |
 
 So at `entropy-consistency`, WAI's JPEG AI symbol decode is demonstrated across
 architectures (x86_64 against arm64) once the distribution indices are given; the
@@ -615,9 +725,10 @@ arm64 and on wasm32-wasip2 (a 32-bit target), measured on the development machin
 it also passed on the byte-exact matrix's Linux x86_64, Linux arm64 and macOS arm64
 legs in run 745. On GPUs, the committed vectors of `int_hyper` (and its video keyframe),
 `int_codec` audio and splat, `int_synth` and `int_mlicpp` decode with their graphs on
-the device to the CPU decode's bytes on each adapter §5 lists, measured by hand on small
+the device to the CPU decode's bytes on each adapter §5 lists but the browser, measured by
-vectors; the `wai-ops-gpu` workflow repeats it on one of them, with no forge run
+hand on small vectors, after each adapter gave WGSL's integer semantics on every guarded
-recorded yet.
+case; the `wai-ops-gpu` workflow repeats it on one of them, with no forge run recorded
+yet.
 
 For the JPEG AI exact decoder's pinned reconstruction (§5), identity across
 machines rests on construction, and on one comparison: a single codestream on the

wai/extensions/platform-decode.md

@@ -0,0 +1,239 @@
+# WAI Extension: Platform Decode Capabilities in Dispatch
+
+> Status: Draft. Maps two web platform interfaces, Media Capabilities `decodingInfo()`
+> and WebCodecs `isConfigSupported()` with its hardware-acceleration preference, onto WAI
+> capability dispatch ([SPEC.md §4](../SPEC.md)) and rendition selection (SPEC §7.1):
+> what a sink asks the platform about an entry, what each answer means for selection,
+> and a default order, a deployer policy SPEC §7.1 allows, that tries classical decode
+> the platform marks power-efficient before any learned reconstruction that carries no
+> energy budget. No manifest member and no wire format: the order is a function of the
+> manifest, the payloads, the platform's answers and the budgets.
+> Keywords **MUST**, **MUST NOT**, **SHOULD**, **MAY** are RFC 2119/8174.
+
+## 1. Scope
+
+A sink running in a web platform can decode a standard video codec through the
+platform's own decoders, which may be fixed-function hardware, or through a decoder of its
+own: WebAssembly, or a learned reconstruction such as WAI's integer codecs. The platform
+says, for a configuration, whether it can decode it, whether it can do so without
+dropping frames, and whether it judges the decode power-efficient (Media Capabilities
+§2.2). A learned reconstruction runs on the sink's CPU, GPU or neural engine, and what it
+draws there is known only where it has been measured. This extension lets a WAI sink use
+what the platform says when it chooses among the entries an envelope offers, and lets a
+learned entry rank with the platform's power-efficient decode only on a measured figure.
+
+It reads these documents:
+
+- Media Capabilities, W3C Working Draft, 18 October 2026: `MediaDecodingConfiguration`,
+  `VideoConfiguration`, `MediaCapabilitiesInfo` (§2.1, §2.2), the decoding-info algorithm
+  (§2.3.2), and the privacy considerations (§3);
+- WebCodecs, W3C Working Draft, 19 October 2026: `VideoDecoderConfig` (§7.6),
+  `isConfigSupported()` and Check Configuration Support (§4, §7.1), and
+  `HardwareAcceleration` (§7.9);
+- the WebCodecs AV1 codec registration (Group Note Draft, 13 October 2026) and HEVC codec
+  registration (Group Note Draft, 8 June 2026), which define the codec strings by the AV1
+  ISOBMFF binding (AV1 Codec ISO Media File Format Binding v1.3.0, §5) and by ISO/IEC
+  14496-15 Annex E.3.
+
+ISO/IEC 14496-15 was not read. The HEVC codec string §2 writes follows the form the HEVC
+registration describes (four dot-separated fields of Annex E.3 and the constraint bytes
+after them), checked against an open-source parser of that form (Chromium's
+`ParseHEVCCodecId`) and against a platform that decodes with it (§6).
+
+## 2. The query
+
+For an entry whose capability the table maps, the sink forms the query from the entry's
+payload. It **MUST NOT** form it from anything the envelope declares apart from the
+payload: the codec string a payload's encoder declared can differ from the stream
+(§6 records one that did).
+
+| capability | codec string | size | frame rate | bit rate |
+|---|---|---|---|---|
+| `wai.video.av1`, `wai.video.av1.lossless` | from the first temporal unit's sequence header OBU: `av01.`, the profile, the level (two digits) and the tier (`M` or `H`), the bit depth (two digits); then the monochrome flag, the chroma subsampling (three digits), the colour primaries, transfer characteristics and matrix coefficients (two digits each, 01 each where the sequence header carries no colour description), and the full-range flag; the part after the bit depth omitted where it is exactly `.0.110.01.01.01.0` (AV1-ISOBMFF §5) | `w` × `h` of SPEC §5's AV1 framing | `fps_num` / `fps_den` of the framing | the payload's bits × the frame rate ÷ `n`, rounded up |
+| `wai.video.hevc` | from the first SPS NAL unit (H.265 §7.3.2.2, §7.3.3): `hev1.`; the profile space as nothing, `A`, `B` or `C`, then `general_profile_idc` in decimal; the 32 `general_profile_compatibility_flag`s as one hexadecimal value with flag *j* at bit *j*; `L` or `H` for `general_tier_flag`, then `general_level_idc` in decimal; then the six bytes from `general_progressive_source_flag` on, each in hexadecimal, trailing zero bytes omitted. `hev1` because an Annex B stream carries its parameter sets in band | `pic_width_in_luma_samples` × `pic_height_in_luma_samples`, less the conformance window; the coded size before it | none in the payload: the rate the sink will present at, where it has one | the payload's bits × the frame rate ÷ the pictures the stream holds, rounded up |
+
+A capability the table does not map is not queried: the sink decodes it, or does not, as
+it otherwise would, and the order of §4 places it by its class alone.
+
+The sink asks:
+
+- `navigator.mediaCapabilities.decodingInfo()` with a `MediaDecodingConfiguration` of
+  `type` `"media-source"` and a `video` member: `contentType`
+  `video/mp4; codecs="<codec string>"`, the size, the bit rate and the frame rate. Where
+  the query has no frame rate it does not ask: `framerate` and `bitrate` are required
+  members, and a sink **MUST NOT** invent them.
+- `VideoDecoder.isConfigSupported()` with a `VideoDecoderConfig` of the codec string, the
+  coded size as `codedWidth` and `codedHeight`, and the `hardwareAcceleration` the sink
+  will configure, when it decodes the entry through WebCodecs.
+
+WebCodecs gives no smoothness or power answer, and Media Capabilities has no WebCodecs
+decoding type. A sink that decodes through WebCodecs reads `smooth` and `powerEfficient`
+from the `media-source` query of the same configuration: the platform's statement about
+that configuration, which the order uses as a ranking signal and nothing more (§3).
+
+## 3. The answers
+
+**`supported`.** A `decodingInfo()` answer of `supported: false`, or an
+`isConfigSupported()` answer of `supported: false` for the preference the sink will
+configure, makes the capability unsupported for that object through the platform. It is
+not a decode error and not a reason to try the decode anyway
+([capability-constraints](capability-constraints.md) §2 reads an unmet constraint set the
+same way). A sink that holds a decoder of its own for the capability decodes the entry
+with it, as an entry with no platform answer; otherwise the entry is skipped with the
+reason `missing_capability` (SPEC §7.1 step 5).
+
+**`smooth` and `powerEfficient`.** Both order the entries a sink can decode (§4) and do
+nothing else. `powerEfficient` is the user agent's judgement that the decode's power draw
+is optimal, its definition left to the user agent, which **SHOULD NOT** mark hardware
+decoding power-efficient by default (Media Capabilities §2.2). It is not a measurement. A
+sink **MUST NOT** write it, or anything derived from it, into a receipt, a measurement
+claim or an energy figure, and **MUST NOT** present it as an acquisition class
+([energy-measurement](energy-measurement.md) §2).
+
+**The hardware-acceleration preference.** `prefer-hardware` and `prefer-software` are
+hints a user agent may ignore for any reason (WebCodecs §7.9). An `isConfigSupported()`
+answer of `supported: true` with `prefer-hardware` says the platform accepts that
+configuration with that preference; it is not evidence that the decode runs on hardware,
+and a sink **MUST NOT** report it as such.
+
+**Privacy.** The answers can be used to fingerprint the sink (Media Capabilities §3.2). A
+sink **SHOULD NOT** send them, or the order they give, to any party; the order is computed
+where the answers are. A platform that fakes its capabilities to resist fingerprinting
+(§3.2 recommends a fixed set rather than always yes or always no) gives answers the order
+uses as given: the result is a choice that platform made, and the decode still conforms.
+
+## 4. The default order
+
+### 4.1 Classical and learned entries
+
+An entry is **learned** when its capability's registry class (SPEC §7) is one of the
+neural, generative or behavioural classes: `NeuralIntegerExact`, `NeuralEntropyExact`,
+`NeuralFloat`, `GenerativeDeterministic`, `Behavioral`. It is **classical** when the class
+is one of the standard classes: `StandardDefined`, `StandardTolerance`,
+`StandardUnclassified`, the WAI-native integer codecs that are not learned among them. A
+capability this revision does not register is ranked as learned: nothing says it is
+classical.
+
+### 4.2 The ranks
+
+Each entry SPEC §7.1 step 2 considers has a rank, tried in this order:
+
+1. **first**: a classical entry the platform marks `supported`, `smooth` and
+   `powerEfficient`; and a learned entry whose energy budget qualifies (§5);
+2. **smooth**: a classical entry the platform marks `supported` and `smooth`, and not
+   `powerEfficient`;
+3. **classical**: any other classical entry the sink can decode: one the platform marks
+   not `smooth`, one it has no answer for, and one the sink decodes with a decoder of its
+   own after the platform answered `supported: false`;
+4. **learned without a budget**: a learned entry whose budget is absent or does not
+   qualify.
+
+Within a rank, entries keep the manifest's order, the encoder's preference. An entry the
+sink cannot decode (§3) is skipped as `missing_capability`. The order is a deployer policy
+in SPEC §7.1's sense: it reorders the entries the sink considers and leaves out the ones it
+cannot decode, and every rule of §7.1 still holds, so an entry skipped for its digest or
+its pin is skipped whatever its rank.
+
+The order therefore decodes through a classical entry before a learned one without a
+budget whenever the sink can decode both, and through a power-efficient classical entry
+before a learned one with a budget only where the manifest lists it first. A learned
+reconstruction is dispatched above a classical one by default only when it carries a
+measured energy figure (§5).
+
+### 4.3 A `WAI1` envelope
+
+SPEC §4 tries `model_requirement.capability`, then `model_requirement.fallback`. Under the
+default order a sink does not advertise, for an envelope, a learned capability without a
+qualifying budget when the envelope's fallback is a classical capability the platform
+marks `supported`, `smooth` and `powerEfficient`: step 3 then does not decode, and step 5
+decodes the fallback. In every other case the capability is advertised as the sink
+otherwise would. Advertising for an envelope is per envelope, as module authorisation is
+(SPEC §4, *Module authorisation*).
+
+### 4.4 What the order is not
+
+The order is a choice among decodes the sink can make, not a statement about energy. It
+claims no figure for the entry it picks, and a receipt over that entry's decode states its
+energy by [energy-measurement](energy-measurement.md)'s rules, or states none. A deployer
+**MAY** use another order (SPEC §7.1 step 2); it **MUST NOT** describe another order as
+this extension's.
+
+## 5. Energy budgets
+
+A learned entry's **energy budget** is a figure the energy-measurement extension defines
+for a learned capability and parameter set: an `OnChipCounter` figure (energy-measurement
+§2) measured on a reference sink. How a budget is carried and checked is that extension's;
+this extension reads only these members:
+
+| member | meaning |
+|---|---|
+| `class` | the acquisition class; the budget qualifies only when it is `OnChipCounter` |
+| `joules_micro` | the figure, in microjoules, for one unit of decode |
+| `per` | the unit: `frame`, `picture` or `second` of output |
+| `relative_ppm`, `window_us` | the declared standard uncertainty and the integration window (energy-measurement §2.3) |
+| `reference_sink` | the reference sink it was measured on |
+
+A budget qualifies when its class is `OnChipCounter`, `joules_micro`, `relative_ppm` and
+`window_us` are each non-zero, `reference_sink` names a sink, `per` is one of the units
+above, and, where the deployer sets a ceiling for that unit, `joules_micro` is at or below
+it. A budget of another class does not qualify: a `ModelBased` or `Estimator` figure is not
+a measurement of the reconstruction. Until a learned capability has a budget that
+qualifies, it ranks last (§4.2).
+
+## 6. Conformance
+
+The reference implementation is `wai::platform_decode` (`wai/wai-rs/src/platform_decode.rs`):
+`query` forms §2's query, `rank` and `order` give §4.2's order, `select` runs SPEC §7.1
+selection in it, and `advertise_for_envelope` applies §4.3. It calls no platform: the
+answers are inputs. The vectors are in `wai/platform-decode-conformance/`:
+
+- `order.json`: manifests of classical and learned entries with the platform's answers and
+  budgets, and each entry's rank, the order, the entries skipped with their reason, and
+  the entry SPEC §7.1 picks; and `WAI1` envelopes with whether their capability is
+  advertised (§4.3);
+- `queries.json`: two floor objects of the packaging corpus, whose codec string is the one
+  the packager wrote into its HLS master, and an HEVC SPS a platform's hardware encoder
+  wrote through WebCodecs.
+
+An implementation conforms when it gives every vector's result. The reference's tests
+(`platform_decode::tests`) also hold the AV1 codec string to the AV1-ISOBMFF binding's own
+examples, derive every floor of the packaging corpus to the codec string and size the
+packager wrote into the pack's HLS master, read constructed SPSs with sub-layers and with
+emulation-prevention bytes, and turn each rule of §4 off in turn to show that a vector
+catches it.
+
+`wai-web`'s `platform_decode` bundle binds the same code for a page
+(`platformQuery`, `platformOrder`), and `wai-web/demo/platform_decode.html`
+(`verify_platform_decode.mjs`) runs it against a browser's platform: every AV1 floor object
+of the packaging corpus, queried and decoded through WebCodecs with the configuration the
+query gives; the same codec family at 1920×1080; an HEVC stream the platform's own encoder
+writes, queried and decoded back; and the order over a learned entry and an AV1 entry from
+the answers the platform gave.
+
+Measured 2026-10-22 with that page on an arm64 desktop (macOS 27.0.1), in headless
+Chromium 148 and a Chromium-based browser at 155, with the same results in both:
+
+- every AV1 floor object of the packaging corpus (64×36, `av01.0.00M.08`) decoded every
+  temporal unit through WebCodecs. The platform answered `supported` and `smooth`, not
+  `powerEfficient`, and did not accept `prefer-hardware` at that size; at 1920×1080 it
+  answered `supported`, `smooth` and `powerEfficient`, and accepted `prefer-hardware`;
+- the platform's HEVC encoder, configured as `hvc1.1.6.L93.B0` for 640×360, wrote an SPS
+  with `general_level_idc` 63, and the codec string it declared for its output repeated the
+  configured one. The query read `hev1.1.6.L63.B0` from the stream; the platform accepted
+  that configuration with both preferences, answered `supported`, `smooth` and
+  `powerEfficient`, and decoded every picture back with it;
+- with those answers, the order put the AV1 entry first and the learned entry without a
+  budget last, and picked the AV1 entry; with a qualifying budget on the learned entry,
+  both ranked first and the manifest's order picked the learned entry.
+
+Not measured: a platform without hardware AV1 or HEVC decode, a platform that fakes its
+answers, and any engine but Chromium's.
+
+## 7. Security and privacy considerations
+
+The answers are fingerprinting surface (§3, *Privacy*). The order reads them where they
+are produced and sends nothing. The query reads only the payload's headers; a malformed
+payload gives no query, and the entry is decoded, or not, as it would be without this
+extension. An encoder cannot raise an entry's rank by what it declares: the codec string
+is read from the payload, the class from the registry, and a budget qualifies only as §5
+says.

wai/platform-decode-conformance/README.md

@@ -0,0 +1,32 @@
+# platform-decode-conformance — platform decode capabilities in dispatch
+
+Vectors for [`extensions/platform-decode.md`](../extensions/platform-decode.md): the query
+a sink reads from an entry's payload for a web platform's Media Capabilities and WebCodecs
+(§2), and the default order the platform's answers and the entries' energy budgets give
+SPEC §7.1 selection and SPEC §4 advertising (§4, §5).
+
+**Provenance.** `order.json` is written by hand, one case per rule of §4 and §5, each with
+the reason it exists. `queries.json` names two floor objects of the packaging corpus
+(`../pack-conformance/packs/`), whose codec strings are the ones the packager wrote into
+the packs' HLS masters, and carries one HEVC SPS that a platform's hardware encoder wrote
+through WebCodecs (§6 of the extension records the run).
+
+## Layout
+
+- `order.json` — `cases`: a `WAI2` manifest's `entries`, each a `capability`, optional
+  `answers` (`decodingInfo`: `supported`, `smooth`, `powerEfficient`; `webcodecs`:
+  `hardwareAcceleration`, `supported`; `ownDecoder`) and an optional `budget` (§5's
+  members), and optional `rules` (`ceiling`, microjoules); `expect`: each entry's rank
+  (`first`, `smooth`, `classical`, `learned-without-budget`, or `null` for an entry the
+  sink cannot decode), the `order`, the entries `refused` with their reason, and the entry
+  `picked` at a sink that advertises every capability, checks no digest and holds no
+  parameter set. `wai1`: a `capability` and a `fallback` entry, and whether the capability
+  is advertised for the envelope (§4.3).
+- `queries.json` — `cases`: a `capability`, a `payload_file` (a `WAI1` envelope, relative
+  to `wai/`) or an `sps_hex`, and the query's members.
+
+## Equivalence
+
+An implementation conforms when it gives every vector's ranks, order, refusals, pick and
+advertising decision, and every query's members. The reference is `wai::platform_decode`
+(`wai-rs/src/platform_decode.rs`), whose tests read these files.

wai/platform-decode-conformance/order.json

@@ -0,0 +1,178 @@
+{
+  "about": "platform-decode §4: the default order over a WAI2 manifest's entries, from each entry's capability, the platform's answers (Media Capabilities decodingInfo, WebCodecs isConfigSupported) and any energy budget; and §4.3, whether a WAI1 envelope's capability is advertised. Every entry's capability is advertised, no entry states a sha256 or a prior, so the pick is the first entry in the order. A budget is the energy-measurement extension's: here only the members this extension reads.",
+  "cases": [
+    {
+      "id": "power-efficient-classical-before-learned",
+      "why": "A learned reconstruction without a budget ranks below the AV1 rendition the platform marks supported, smooth and power-efficient, though the manifest lists it first.",
+      "entries": [
+        { "capability": "wai.neural.int_synth" },
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "first"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "learned-with-a-budget-keeps-manifest-order",
+      "why": "With an OnChipCounter budget measured on a reference sink, the learned entry ranks with power-efficient classical decode, and the manifest order decides between them.",
+      "entries": [
+        { "capability": "wai.neural.int_hyper", "budget": { "class": "OnChipCounter", "joules_micro": 41000, "per": "frame", "relative_ppm": 20000, "window_us": 5000000, "reference_sink": "reference sink: arm64 desktop, SoC energy counter" } },
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } }
+      ],
+      "expect": { "ranks": ["first", "first"], "order": [0, 1], "refused": [], "picked": 0 }
+    },
+    {
+      "id": "a-modelled-budget-does-not-qualify",
+      "why": "A budget of the Estimator class is not a measurement: the learned entry ranks last.",
+      "entries": [
+        { "capability": "wai.neural.int_hyper", "budget": { "class": "Estimator", "joules_micro": 41000, "per": "frame", "relative_ppm": 20000, "window_us": 5000000, "reference_sink": "reference sink: arm64 desktop" } },
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "first"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "a-budget-without-its-uncertainty-does-not-qualify",
+      "why": "An OnChipCounter figure with no declared uncertainty is not one energy-measurement §2.3 admits.",
+      "entries": [
+        { "capability": "wai.neural.int_hyper", "budget": { "class": "OnChipCounter", "joules_micro": 41000, "per": "frame", "relative_ppm": 0, "window_us": 5000000, "reference_sink": "reference sink: arm64 desktop" } },
+        { "capability": "wai.video.hevc", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "first"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "a-budget-above-the-deployers-ceiling",
+      "why": "The deployer sets a ceiling of 10000 microjoules a frame; a budget above it does not qualify.",
+      "rules": { "ceiling": 10000 },
+      "entries": [
+        { "capability": "wai.neural.int_hyper", "budget": { "class": "OnChipCounter", "joules_micro": 41000, "per": "frame", "relative_ppm": 20000, "window_us": 5000000, "reference_sink": "reference sink: arm64 desktop" } },
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "first"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "learned-below-a-software-classical-decode",
+      "why": "A classical entry the sink decodes with a decoder of its own, with no platform answer, still ranks above a learned entry without a budget.",
+      "entries": [
+        { "capability": "wai.neural.int_synth" },
+        { "capability": "wai.video.av1" }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "classical"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "unsupported-by-the-platform-and-no-decoder-of-its-own",
+      "why": "decodingInfo answers supported: false and the sink holds no AV1 decoder of its own: the entry is skipped as missing_capability, and the learned entry is picked.",
+      "entries": [
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": false, "smooth": false, "powerEfficient": false }, "ownDecoder": false } },
+        { "capability": "wai.neural.int_synth" }
+      ],
+      "expect": { "ranks": [null, "learned-without-budget"], "order": [1], "refused": [[0, "missing_capability"]], "picked": 1 }
+    },
+    {
+      "id": "unsupported-by-the-platform-with-a-decoder-of-its-own",
+      "why": "The platform cannot decode it, the sink's own decoder can: a classical entry with no platform answer, still above learned decode without a budget.",
+      "entries": [
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": false, "smooth": false, "powerEfficient": false }, "ownDecoder": true } },
+        { "capability": "wai.neural.int_synth" }
+      ],
+      "expect": { "ranks": ["classical", "learned-without-budget"], "order": [0, 1], "refused": [], "picked": 0 }
+    },
+    {
+      "id": "webcodecs-refuses-the-configuration",
+      "why": "Media Capabilities marks the configuration power-efficient, but the sink decodes through WebCodecs and isConfigSupported refuses it with the preference the sink configures: unsupported through the platform.",
+      "entries": [
+        { "capability": "wai.video.hevc", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true }, "webcodecs": { "hardwareAcceleration": "prefer-hardware", "supported": false } } },
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": false }, "webcodecs": { "hardwareAcceleration": "no-preference", "supported": true } } }
+      ],
+      "expect": { "ranks": [null, "smooth"], "order": [1], "refused": [[0, "missing_capability"]], "picked": 1 }
+    },
+    {
+      "id": "smooth-and-power-efficient-before-smooth",
+      "why": "Two classical entries: the power-efficient one ranks first though listed second.",
+      "entries": [
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": false } } },
+        { "capability": "wai.video.hevc", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } }
+      ],
+      "expect": { "ranks": ["smooth", "first"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "smooth-before-power-efficient-but-not-smooth",
+      "why": "A decode the platform calls power-efficient but not smooth drops frames: it ranks below a smooth one.",
+      "entries": [
+        { "capability": "wai.video.hevc", "answers": { "decodingInfo": { "supported": true, "smooth": false, "powerEfficient": true } } },
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": false } } }
+      ],
+      "expect": { "ranks": ["classical", "smooth"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "manifest-order-within-a-rank",
+      "why": "Two power-efficient classical entries keep the manifest's order.",
+      "entries": [
+        { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } },
+        { "capability": "wai.video.hevc", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } }
+      ],
+      "expect": { "ranks": ["first", "first"], "order": [0, 1], "refused": [], "picked": 0 }
+    },
+    {
+      "id": "an-unregistered-capability-ranks-as-learned",
+      "why": "Nothing says a capability this revision does not register is classical.",
+      "entries": [
+        { "capability": "x-example.video.codec" },
+        { "capability": "wai.video.av1" }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "classical"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "a-float-learned-decode-and-an-integer-one",
+      "why": "Two learned entries, one with a budget: it ranks first.",
+      "entries": [
+        { "capability": "wai.neural.mbt2018_mean" },
+        { "capability": "wai.neural.int_hyper", "budget": { "class": "OnChipCounter", "joules_micro": 52000, "per": "picture", "relative_ppm": 15000, "window_us": 2000000, "reference_sink": "reference sink: x86_64 desktop, package energy counter" } }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "first"], "order": [1, 0], "refused": [], "picked": 1 }
+    },
+    {
+      "id": "lossless-av1-is-classical",
+      "why": "wai.video.av1.lossless is a standard codec: with no answer it ranks as classical decode.",
+      "entries": [
+        { "capability": "wai.neural.int_synth" },
+        { "capability": "wai.video.av1.lossless" }
+      ],
+      "expect": { "ranks": ["learned-without-budget", "classical"], "order": [1, 0], "refused": [], "picked": 1 }
+    }
+  ],
+  "wai1": [
+    {
+      "id": "learned-capability-power-efficient-fallback",
+      "why": "SPEC §4 with a learned capability and no budget, over an AV1 fallback the platform marks power-efficient: the capability is not advertised for this envelope, and step 5 decodes the fallback.",
+      "capability": { "capability": "wai.neural.int_hyper" },
+      "fallback": { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } },
+      "expect": { "advertise": false, "decodes": "fallback" }
+    },
+    {
+      "id": "learned-capability-with-a-budget",
+      "why": "With a qualifying budget the capability is advertised as usual, and step 3 decodes it.",
+      "capability": { "capability": "wai.neural.int_hyper", "budget": { "class": "OnChipCounter", "joules_micro": 41000, "per": "frame", "relative_ppm": 20000, "window_us": 5000000, "reference_sink": "reference sink: arm64 desktop" } },
+      "fallback": { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": true } } },
+      "expect": { "advertise": true, "decodes": "capability" }
+    },
+    {
+      "id": "learned-capability-fallback-not-power-efficient",
+      "why": "The rule removes a learned capability only in favour of a fallback the platform marks power-efficient.",
+      "capability": { "capability": "wai.neural.int_hyper" },
+      "fallback": { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": true, "powerEfficient": false } } },
+      "expect": { "advertise": true, "decodes": "capability" }
+    },
+    {
+      "id": "classical-capability-learned-fallback",
+      "why": "A classical capability is advertised as usual.",
+      "capability": { "capability": "wai.video.av1", "answers": { "decodingInfo": { "supported": true, "smooth": false, "powerEfficient": false } } },
+      "fallback": { "capability": "wai.neural.int_synth" },
+      "expect": { "advertise": true, "decodes": "capability" }
+    },
+    {
+      "id": "learned-capability-no-fallback",
+      "why": "With no fallback there is nothing to prefer: the capability is advertised as usual.",
+      "capability": { "capability": "wai.neural.int_synth" },
+      "fallback": null,
+      "expect": { "advertise": true, "decodes": "capability" }
+    }
+  ]
+}

wai/platform-decode-conformance/queries.json

@@ -0,0 +1,23 @@
+{
+  "about": "platform-decode §2: the query a sink reads from an entry's payload. payload_file names a WAI1 envelope under wai/ whose payload is the entry's; sps_hex is an HEVC SPS NAL unit, header included, as an Annex B stream carries it. The two AV1 objects are floor objects of the packaging corpus (SPEC §5's AV1 framing); their codec string is the one the packager wrote into the pack's HLS master. The SPS is the one a platform's hardware HEVC encoder wrote through WebCodecs, configured as hvc1.1.6.L93.B0 for 640x360: it states general_level_idc 63 (level 2.1), and the codec string the encoder declared for its output repeated the configured one.",
+  "cases": [
+    {
+      "id": "corpus-floor-object-60-units",
+      "capability": "wai.video.av1",
+      "payload_file": "pack-conformance/packs/video-base-01/package/objects/35c7a45323b7afe979c123fa885f35cc55fabd2a1954759aebe4b95dbc76361f.wai",
+      "expect": { "codec": "av01.0.00M.08", "width": 64, "height": 36, "codedWidth": 64, "codedHeight": 36, "framerate": [25, 1], "bitrate": 7230, "bitDepth": 8 }
+    },
+    {
+      "id": "corpus-floor-object-40-units",
+      "capability": "wai.video.av1",
+      "payload_file": "pack-conformance/packs/video-base-01/package/objects/761c4649427c51796a2b699568775765b209ee80d2f4fc1cf040838c6724aa2d.wai",
+      "expect": { "codec": "av01.0.00M.08", "width": 64, "height": 36, "codedWidth": 64, "codedHeight": 36, "framerate": [25, 1], "bitrate": 8715, "bitDepth": 8 }
+    },
+    {
+      "id": "hardware-encoder-sps",
+      "capability": "wai.video.hevc",
+      "sps_hex": "420101016000000300b0000003000003003fa005020171f2e206ee45914bff2e7f13fa9b8086808040",
+      "expect": { "codec": "hev1.1.6.L63.B0", "width": 640, "height": 360, "codedWidth": 640, "codedHeight": 368, "bitDepth": 8 }
+    }
+  ]
+}

wai/wai-ops-gpu/Cargo.lock

@@ -705,6 +705,7 @@ version = "0.1.0"
 name = "wai-ops-gpu"
 version = "0.1.0"
 dependencies = [
+ "naga",
  "wai-ops",
  "wgpu",
 ]

wai/wai-ops-gpu/Cargo.toml

@@ -29,5 +29,11 @@ wgpu = { version = "30.0.1", default-features = false, features = ["std", "wgsl"
 [target.'cfg(all(target_arch = "wasm32", target_os = "unknown"))'.dependencies]
 wgpu = { version = "30.0.1", default-features = false, features = ["std", "wgsl", "webgpu"] }
 
+# The WGSL rules test (tests/wgsl_rules.rs) reads the shipped modules on the IR of the
+# translator wgpu compiles them with: naga, the same release wgpu 30.0.1 depends on, so
+# the lockfile gains no package.
+[target.'cfg(not(target_arch = "wasm32"))'.dev-dependencies]
+naga = { version = "30.0.1", default-features = false, features = ["wgsl-in", "msl-out", "hlsl-out", "spv-out"] }
+
 [profile.release]
 debug = false

wai/wai-ops-gpu/browser/index.html

@@ -5,9 +5,10 @@
 <pre id="report"></pre>
 <script type="module">
 // Loads every graph-vector file of ops-conformance/ (the list run.mjs serves as
-// files.json), runs them through wai-ops-gpu on this browser's WebGPU, and leaves the
+// files.json), runs them through wai-ops-gpu on this browser's WebGPU, then WGSL's
-// report in window.opsGpu for run.mjs to read.
+// integer semantics vectors of wgsl-conformance/, and leaves the report in
-import init, { run_corpus } from "./pkg/wai_ops_gpu_browser.js";
+// window.opsGpu for run.mjs to read.
+import init, { run_corpus, run_semantics } from "./pkg/wai_ops_gpu_browser.js";
 window.opsGpu = { done: false, report: "" };
 try {
   await init();
@@ -19,6 +20,11 @@ try {
     bundle += `### ${f}\n` + await r.text();
   }
   window.opsGpu.report = await run_corpus(bundle);
+  // WGSL's integer semantics: the committed vectors, held to the reference in this
+  // page's wasm32 and to the browser's WebGPU.
+  const sv = await fetch("/wgsl-conformance/integer.jsonl");
+  if (!sv.ok) throw new Error(`wgsl-conformance/integer.jsonl: ${sv.status}`);
+  window.opsGpu.report += "\n" + await run_semantics(await sv.text());
 } catch (e) {
   window.opsGpu.report = `browser webgpu: page error: ${e}`;
 }

wai/wai-ops-gpu/browser/run.mjs

@@ -12,10 +12,12 @@
 //   [WAI_OPS_GPU_REQUIRE=1] [WAI_OPS_GPU_BROWSER_CHANNEL=chrome] \
 //     PLAYWRIGHT_HOST=<dir with node_modules/playwright>/ node wai/wai-ops-gpu/browser/run.mjs
 //
-// It serves this directory and ../../ops-conformance on 127.0.0.1, opens index.html in
+// It serves this directory, ../../ops-conformance and ../../wgsl-conformance on
-// a headless browser and prints the page's report. It fails on a page error, a
+// 127.0.0.1, opens index.html in a headless browser and prints the page's report. It
-// mismatch or a device error, and, under WAI_OPS_GPU_REQUIRE=1, when the page found no
+// fails on a page error, a mismatch or a device error, and, under WAI_OPS_GPU_REQUIRE=1,
-// adapter or the report has no `browser webgpu: key=` line with mismatches=0 errors=0.
+// when the page found no adapter, or the report lacks a `browser webgpu: key=` line
+// with mismatches=0 errors=0, the `wasm reference:` line saying the committed WGSL
+// semantics vectors are equal, or a `browser webgpu semantics: key=` line with differ=0.
 // Headless, macOS offers a page a hardware adapter only when the browser draws through
 // Metal (`--use-angle=metal`); WAI_OPS_GPU_BROWSER_ARGS (space-separated) replaces the
 // launch arguments.
@@ -27,6 +29,7 @@ import { createRequire } from "node:module";
 
 const HERE = path.dirname(fileURLToPath(import.meta.url));
 const CORPUS = path.resolve(HERE, "../../ops-conformance");
+const WGSL_VECTORS = path.resolve(HERE, "../../wgsl-conformance");
 const REQUIRE = process.env.WAI_OPS_GPU_REQUIRE === "1";
 const PORT = Number(process.env.PORT || 8797);
 const TIMEOUT_MS = Number(process.env.WAI_OPS_GPU_TIMEOUT_MS || 20 * 60 * 1000);
@@ -37,7 +40,7 @@ const { chromium } = createRequire(process.env.PLAYWRIGHT_HOST || HERE + "/")("p
 
 const files = ["vectors", "graphs"].flatMap((d) =>
   fs.readdirSync(path.join(CORPUS, d)).filter((f) => f.endsWith(".jsonl")).sort().map((f) => `${d}/${f}`));
-const ROOTS = { "/ops-conformance/": CORPUS, "/browser/": HERE };
+const ROOTS = { "/ops-conformance/": CORPUS, "/wgsl-conformance/": WGSL_VECTORS, "/browser/": HERE };
 const MIME = { ".js": "text/javascript", ".wasm": "application/wasm", ".html": "text/html", ".json": "application/json", ".jsonl": "text/plain" };
 const server = http.createServer((req, res) => {
   const url = decodeURIComponent(req.url.split("?")[0]);
@@ -82,9 +85,13 @@ try {
   console.log(`adapter info (navigator.gpu): ${info}`);
   console.log(report);
   for (const e of errs) console.log(`page error: ${e}`);
-  const ok = report.split("\n").some((l) => /^browser webgpu: key=\S+ .* mismatches=0 errors=0$/.test(l));
+  const lines = report.split("\n");
+  const ok = lines.some((l) => /^browser webgpu: key=\S+ .* mismatches=0 errors=0$/.test(l)) &&
+    lines.some((l) => /^wasm reference: target=wasm32 .* cases=[1-9][0-9]* equal$/.test(l)) &&
+    lines.some((l) => /^browser webgpu semantics: key=\S+ .* guarded=[1-9][0-9]* differ=0 /.test(l));
   const bad = errs.length > 0 || /MISMATCH|DEVICE ERROR|page error|did not parse/.test(report) ||
-    report.split("\n").some((l) => /^browser webgpu: key=/.test(l) && !/mismatches=0 errors=0$/.test(l));
+    lines.some((l) => /^browser webgpu: key=/.test(l) && !/mismatches=0 errors=0$/.test(l)) ||
+    lines.some((l) => /^browser webgpu semantics: key=/.test(l) && !/ differ=0 /.test(l));
   if (bad || (REQUIRE && !ok)) process.exitCode = 1;
 } finally {
   if (browser) await browser.close();

wai/wai-ops-gpu/browser/semantics_reference.mjs

@@ -0,0 +1,26 @@
+// WGSL's integer semantics vectors (wai/wgsl-conformance/integer.jsonl) against the
+// reference compiled to wasm32, under Node, with no GPU: every committed result must be
+// the one `wai_ops_gpu::semantics::reference` gives in WebAssembly.
+//
+// Build the page's module first (run.mjs's header says how), then:
+//
+//   node wai/wai-ops-gpu/browser/semantics_reference.mjs
+//
+// Prints `wasm reference: target=wasm32 ... cases=N equal` and exits 0, or prints the
+// first difference and exits 1.
+import fs from "node:fs";
+import path from "node:path";
+import { fileURLToPath } from "node:url";
+
+const HERE = path.dirname(fileURLToPath(import.meta.url));
+const pkg = path.join(HERE, "pkg");
+if (!fs.existsSync(path.join(pkg, "wai_ops_gpu_browser.js"))) {
+  console.error("no pkg/wai_ops_gpu_browser.js; build it as run.mjs's header says");
+  process.exit(2);
+}
+const mod = await import(path.join(pkg, "wai_ops_gpu_browser.js"));
+mod.initSync({ module: fs.readFileSync(path.join(pkg, "wai_ops_gpu_browser_bg.wasm")) });
+const vectors = fs.readFileSync(path.resolve(HERE, "../../wgsl-conformance/integer.jsonl"), "utf8");
+const out = mod.semantics_reference(vectors);
+console.log(`node ${process.version}: ${out}`);
+if (!/^wasm reference: target=wasm32 .* cases=[1-9][0-9]* equal$/.test(out)) process.exit(1);

wai/wai-ops-gpu/browser/src/lib.rs

@@ -6,6 +6,12 @@
 //! `### <name>`, and returns the report: one line per adapter, `browser webgpu: …`,
 //! then a line per mismatch or device error (at most 20), or `browser webgpu: no adapter
 //! (<why>)`.
+//!
+//! [`run_semantics`] takes WGSL's integer semantics vectors
+//! (`wai/wgsl-conformance/integer.jsonl`) and holds them twice: to the reference
+//! (`wai_ops_gpu::semantics::reference`) compiled to wasm32 and run in the page, and to
+//! every adapter's WebGPU, the kernel compiled by the browser's own WGSL compiler.
+//! [`semantics_reference`] is the first half alone, for an engine with no WebGPU.
 
 use wai_ops::json::{self, Value, unhex};
 use wai_ops::vectors::read_tensor;
@@ -132,3 +138,77 @@ pub async fn run_corpus(bundle: String) -> String {
     }
     lines.join("\n")
 }
+
+/// The committed semantics vectors against the reference compiled to wasm32: one line,
+/// `wasm reference: cases=N equal`, or the first difference.
+#[wasm_bindgen]
+pub fn semantics_reference(vectors: String) -> String {
+    use wai_ops_gpu::semantics;
+    let cases = match semantics::parse_vector_file(&vectors) {
+        Ok(c) => c,
+        Err(e) => return format!("wasm reference: the vectors did not parse: {e}"),
+    };
+    let ptr_bits = usize::BITS;
+    for (c, want) in &cases {
+        let got = semantics::reference(c);
+        if got != *want {
+            return format!("wasm reference: MISMATCH {} a={:#010x} b={:#010x} c={:#010x} d={:#010x}: the file says {want:#010x}, the reference {got:#010x}", c.op.name(), c.a, c.b, c.c, c.d);
+        }
+    }
+    format!("wasm reference: target=wasm32 usize_bits={ptr_bits} cases={} equal", cases.len())
+}
+
+/// [`semantics_reference`], then the committed cases and SplitMix64 cases of every
+/// operation on every adapter the page is offered, against the reference. One line per
+/// adapter, `browser webgpu semantics: key=…`.
+#[wasm_bindgen]
+pub async fn run_semantics(vectors: String) -> String {
+    use wai_ops_gpu::semantics;
+    let mut lines = vec![semantics_reference(vectors.clone())];
+    let Ok(committed) = semantics::parse_vector_file(&vectors) else {
+        return lines.join("\n");
+    };
+    let mut cases: Vec<semantics::Case> = committed.iter().map(|(c, _)| *c).collect();
+    cases.extend(semantics::random_cases(1024, 0xA417_5E01));
+    let expect: Vec<u32> = cases.iter().map(semantics::reference).collect();
+    let d = wai_ops_gpu::discover().await;
+    for (a, why) in &d.refused {
+        lines.push(format!("browser webgpu semantics: adapter refused: key={} {} reason={why:?}", a.key(), a.describe()));
+    }
+    if d.found.is_empty() {
+        lines.push(format!("browser webgpu semantics: no adapter ({} refused)", d.refused.len()));
+        return lines.join("\n");
+    }
+    for g in &d.found {
+        let got = match semantics::run(g, &cases).await {
+            Ok(v) => v,
+            Err(e) => {
+                lines.push(format!("browser webgpu semantics: DEVICE ERROR key={}: {e}", g.adapter().key()));
+                continue;
+            }
+        };
+        let o = semantics::compare(&cases, &expect, &got);
+        for x in o.guarded_differ.iter().take(20) {
+            lines.push(format!("browser webgpu semantics: MISMATCH {x}"));
+        }
+        let mut shown = std::collections::BTreeSet::new();
+        for x in &o.unguarded_differ {
+            if shown.insert(x.case.op) {
+                lines.push(format!("browser webgpu semantics: unguarded {x}"));
+            }
+        }
+        lines.push(format!(
+            "browser webgpu semantics: key={} {} committed={} cases={} guarded={} differ={} unguarded={} unguarded_differ={} unguarded_ops={}",
+            g.adapter().key(),
+            g.adapter().describe(),
+            committed.len(),
+            cases.len(),
+            o.guarded,
+            o.guarded_differ.len(),
+            o.unguarded,
+            o.unguarded_differ.len(),
+            if o.unguarded_ops().is_empty() { "none".to_string() } else { o.unguarded_ops().join(",") },
+        ));
+    }
+    lines.join("\n")
+}

wai/wai-ops-gpu/src/bin/wai-wgsl-vectors.rs

@@ -0,0 +1,28 @@
+//! Writes WGSL's integer semantics vectors (`wai/wgsl-conformance/integer.jsonl`):
+//! every edge case of `wai_ops_gpu::semantics`, with the result the WGSL text gives.
+//!
+//! ```text
+//! wai-wgsl-vectors <out-file>
+//! ```
+//!
+//! `tests/semantics.rs` holds the committed file to this output, byte for byte.
+
+use std::process::ExitCode;
+
+use wai_ops_gpu::semantics;
+
+fn main() -> ExitCode {
+    let Some(out) = std::env::args().nth(1) else {
+        eprintln!("usage: wai-wgsl-vectors <out-file>");
+        return ExitCode::from(2);
+    };
+    let cases = semantics::edge_cases();
+    let text = semantics::vector_file(&cases);
+    if let Err(e) = std::fs::write(&out, &text) {
+        eprintln!("wai-wgsl-vectors: {out}: {e}");
+        return ExitCode::FAILURE;
+    }
+    let unguarded = cases.iter().filter(|c| !c.guarded()).count();
+    eprintln!("wai-wgsl-vectors: wrote {out}: {} operations, {} cases ({unguarded} unguarded)", semantics::ALL.len(), cases.len());
+    ExitCode::SUCCESS
+}

wai/wai-ops-gpu/src/device.rs

@@ -2,9 +2,16 @@
 //!
 //! [`discover`] opens every adapter of the platform's primary GPU interfaces (Metal,
 //! Vulkan, DX12; a browser's WebGPU on `wasm32-unknown-unknown`), compiles the kernels
-//! on it and runs the arithmetic probe ([`crate::probe`]) against the host. An adapter
+//! on it and runs the arithmetic probe ([`crate::probe`]) and WGSL's integer semantics
-//! that cannot host the kernels, or whose arithmetic disagrees with the host's on any
+//! ([`crate::semantics`]) against the host. An adapter that cannot host the kernels,
-//! word, is listed with the reason and never runs a graph.
+//! whose arithmetic disagrees with the host's on any word, or that gives another result
+//! than WGSL's for an operation the kernels may use, is listed with the reason and never
+//! runs a graph.
+//!
+//! Every module is created with `create_shader_module`, which keeps the translator's
+//! run-time checks on: bounds checks, and the guarded integer division and remainder
+//! (`wgpu::ShaderRuntimeChecks::int_div_checks`). The unchecked form is never used
+//! (`tests/wgsl_rules.rs` holds the crate to that).
 
 use std::borrow::Cow;
 use std::collections::HashMap;
@@ -126,6 +133,7 @@ pub(crate) struct Pipelines {
     pub ops64: Vec<wgpu::ComputePipeline>,
     pub mv: wgpu::ComputePipeline,
     pub probe: wgpu::ComputePipeline,
+    pub semantics: wgpu::ComputePipeline,
 }
 
 /// One adapter opened as a device, its kernels compiled and its arithmetic probed.
@@ -169,7 +177,7 @@ pub async fn discover() -> Discovery {
 }
 
 /// [`discover`], with the kernels compiled from `src` and labelled `label` (a test's
-/// mutant). The probe runs `src.probe`.
+/// mutant). The probe runs `src.probe`, and the semantics suite `src.semantics`.
 pub async fn discover_with(src: &Sources, label: &str) -> Discovery {
     let mut d = Discovery::default();
     if wgpu::Instance::enabled_backend_features().is_empty() {
@@ -264,7 +272,10 @@ async fn open(adapter: wgpu::Adapter, src: &Sources, label: &str) -> Result<Gpu,
         pool: Mutex::new(HashMap::new()),
         label: label.to_string(),
     };
-    match crate::probe::check(&gpu).await {
+    if let Err(why) = crate::probe::check(&gpu).await {
+        return Err(Box::new((gpu.adapter.clone(), why)));
+    }
+    match crate::semantics::check(&gpu).await {
         Ok(()) => Ok(gpu),
         Err(why) => Err(Box::new((gpu.adapter.clone(), why))),
     }
@@ -328,6 +339,7 @@ async fn pipelines(device: &wgpu::Device, src: &Sources) -> Result<Pipelines, Gp
     let m64 = module("wai-ops-gpu ops64", &src.ops64);
     let mmv = module("wai-ops-gpu move", &src.mv);
     let mpr = module("wai-ops-gpu probe", &src.probe);
+    let msem = module("wai-ops-gpu semantics", &src.semantics);
     let pipe = |m: &wgpu::ShaderModule, entry: &str| {
         device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
             label: Some(entry),
@@ -342,8 +354,9 @@ async fn pipelines(device: &wgpu::Device, src: &Sources) -> Result<Pipelines, Gp
     let ops64 = OP_ENTRIES.iter().map(|e| pipe(&m64, e)).collect();
     let mv = pipe(&mmv, "mv");
     let probe = pipe(&mpr, "probe");
+    let semantics = pipe(&msem, "semantics");
     scopes.close().await?;
-    Ok(Pipelines { layout, ops32, ops64, mv, probe })
+    Ok(Pipelines { layout, ops32, ops64, mv, probe, semantics })
 }
 
 impl Gpu {
@@ -531,4 +544,9 @@ pub mod blocking {
     pub fn probe(g: &crate::Gpu, recs: &[crate::probe::ProbeIn]) -> Result<Vec<[u32; crate::probe::OUT_WORDS]>, GpuError> {
         resolve(crate::probe::run(g, recs))?
     }
+
+    /// [`crate::semantics::run`].
+    pub fn semantics(g: &crate::Gpu, cases: &[crate::semantics::Case]) -> Result<Vec<u32>, GpuError> {
+        resolve(crate::semantics::run(g, cases))?
+    }
 }

wai/wai-ops-gpu/src/kernels.rs

@@ -11,7 +11,9 @@
 //!   (`wgsl/move.wgsl`).
 //!
 //! `probe` (`wgsl/probe.wgsl`) evaluates every arithmetic function on records the host
-//! gives, so a device's arithmetic can be held to the host's before it runs a graph.
+//! gives, so a device's arithmetic can be held to the host's before it runs a graph, and
+//! `semantics` (`wgsl/semantics.wgsl`) evaluates WGSL's integer operations themselves,
+//! the cases a translation has to guard included ([`crate::semantics`]).
 
 pub const ARITH: &str = include_str!("wgsl/arith.wgsl");
 pub const COMMON: &str = include_str!("wgsl/common.wgsl");
@@ -19,6 +21,7 @@ pub const OPS32: &str = include_str!("wgsl/ops32.wgsl");
 pub const OPS64: &str = include_str!("wgsl/ops64.wgsl");
 pub const MOVE: &str = include_str!("wgsl/move.wgsl");
 pub const PROBE: &str = include_str!("wgsl/probe.wgsl");
+pub const SEMANTICS: &str = include_str!("wgsl/semantics.wgsl");
 
 /// The text of one module: the arithmetic, the bindings, then `body`.
 pub fn module(body: &str) -> String {
@@ -31,19 +34,25 @@ pub fn module(body: &str) -> String {
     s
 }
 
-/// The sources of the four modules, in the order [`Sources::ALL`] names them, so a test
+/// The sources of the five modules, so a test can run a device on an edited text (a
-/// can run a device on an edited text (a mutant) and see the edit caught.
+/// mutant) and see the edit caught.
 #[derive(Clone, Debug, PartialEq, Eq)]
 pub struct Sources {
     pub ops32: String,
     pub ops64: String,
     pub mv: String,
     pub probe: String,
+    pub semantics: String,
 }
 
 impl Sources {
     pub fn shipped() -> Sources {
-        Sources { ops32: module(OPS32), ops64: module(OPS64), mv: module(MOVE), probe: module(PROBE) }
+        Sources { ops32: module(OPS32), ops64: module(OPS64), mv: module(MOVE), probe: module(PROBE), semantics: module(SEMANTICS) }
+    }
+
+    /// Every module's name and text, in the order above.
+    pub fn named(&self) -> [(&'static str, &str); 5] {
+        [("ops32", &self.ops32), ("ops64", &self.ops64), ("move", &self.mv), ("probe", &self.probe), ("semantics", &self.semantics)]
     }
 }
 

wai/wai-ops-gpu/src/lib.rs

@@ -25,6 +25,7 @@
 //!
 //! - [`device`]: discovery, devices, the pipelines and waiting ([`blocking`] natively);
 //! - [`probe`]: the arithmetic probe each device passes before it runs a graph;
+//! - [`semantics`]: WGSL's integer semantics, which each device also passes first;
 //! - [`prepare`]: where a graph runs, and its parameters on the device;
 //! - [`exec`]: one decode;
 //! - [`kernels`]: the shader sources and the record each dispatch reads.
@@ -34,6 +35,7 @@ pub mod exec;
 pub mod kernels;
 pub mod prepare;
 pub mod probe;
+pub mod semantics;
 
 pub use device::{Adapter, Discovery, Gpu, GpuError, discover, discover_with};
 #[cfg(not(target_arch = "wasm32"))]

wai/wai-ops-gpu/src/semantics.rs

@@ -0,0 +1,925 @@
+//! WGSL's integer semantics, held on the device: the conformance suite every adapter
+//! passes before it runs a graph (determinism-tiers §5, *Claiming the tier on a GPU*).
+//!
+//! WGSL defines every integer operation at run time, the cases that are errors when the
+//! operands are known before the shader runs included: addition, subtraction,
+//! multiplication and negation wrap modulo 2^32; `e1 / 0` is `e1` and `e1 % 0` is 0; the
+//! most negative i32 over −1 is itself, with remainder 0; a shift amount is taken modulo
+//! 32; `clamp` on an integer type is `min(max(e, low), high)`; a bit field's offset and
+//! count are clamped to the word (WGSL §8.7, §8.9, §17.5). MSL and HLSL leave signed
+//! overflow undefined, and MSL, HLSL and SPIR-V leave a zero divisor, an out-of-range
+//! shift and an out-of-range bit field undefined or implementation-defined, so a
+//! translator that targets them has to produce WGSL's result itself. This module states
+//! each operation's result from the WGSL text ([`reference`]), and [`run`] evaluates the
+//! same operations in a kernel (`wgsl/semantics.wgsl`) whose operands the device reads
+//! from memory, so nothing is folded before the shader runs.
+//!
+//! [`Op::guarded`] says which operations the shipped kernels may use: those that the
+//! translators this crate runs through give WGSL's result on every adapter the suite
+//! has been run on. [`Op::Shl`] and the other operations listed in [`UNGUARDED`] are
+//! defined by WGSL but reach the backend unchanged (naga 30.0.1 masks no shift amount),
+//! so their results on an amount of 32 or more, or bounds out of order, are the
+//! hardware's. The suite runs them on every adapter and reports what it gives; the
+//! kernels never depend on them (integer-operators §6: every shift that reaches the
+//! width is handled explicitly), and `tests/wgsl_rules.rs` holds the kernels to the
+//! guarded set.
+//!
+//! The committed vectors (`wai/wgsl-conformance/integer.jsonl`) are [`edge_cases`] with
+//! their results, written by `src/bin/wai-wgsl-vectors.rs`; a test holds the file to
+//! this module, the browser page holds it to the same module compiled to wasm32, and the
+//! device runs hold every adapter to it.
+
+use crate::device::{Gpu, GpuError, Scopes};
+use crate::kernels::{Rec, WG_PROBE, f, grid};
+
+/// Words of one case's input: the operation, then `a`, `b`, `c`, `d`.
+pub const IN_WORDS: usize = 5;
+
+/// One integer operation of WGSL, in the order of `sem` in `wgsl/semantics.wgsl`, whose
+/// code is its index.
+#[derive(Clone, Copy, Debug, PartialEq, Eq, PartialOrd, Ord, Hash)]
+pub enum Op {
+    AddU32,
+    SubU32,
+    MulU32,
+    AddI32,
+    SubI32,
+    MulI32,
+    NegI32,
+    DivU32,
+    RemU32,
+    DivI32,
+    RemI32,
+    ShlU32,
+    ShrU32,
+    ShlI32,
+    ShrI32,
+    And,
+    Or,
+    Xor,
+    Not,
+    LtI32,
+    LtU32,
+    LeI32,
+    LeU32,
+    Eq,
+    Ne,
+    SelectU32,
+    SelectI32,
+    MinI32,
+    MaxI32,
+    MinU32,
+    MaxU32,
+    ClampI32,
+    ClampU32,
+    AbsI32,
+    SignI32,
+    CountLeadingZerosU32,
+    CountTrailingZerosU32,
+    CountOneBitsU32,
+    ReverseBitsU32,
+    FirstLeadingBitU32,
+    FirstLeadingBitI32,
+    FirstTrailingBitU32,
+    CountLeadingZerosI32,
+    ExtractBitsU32,
+    ExtractBitsI32,
+    InsertBitsU32,
+    InsertBitsI32,
+    U32OfI32,
+    I32OfU32,
+    U32OfBool,
+    AddVec2U32,
+    MulVec2U32,
+    MulVec2I32,
+    ShrVec2U32,
+    ShrVec2I32,
+    SelectVec2U32,
+}
+
+/// Every operation, in code order.
+pub const ALL: [Op; 56] = [
+    Op::AddU32,
+    Op::SubU32,
+    Op::MulU32,
+    Op::AddI32,
+    Op::SubI32,
+    Op::MulI32,
+    Op::NegI32,
+    Op::DivU32,
+    Op::RemU32,
+    Op::DivI32,
+    Op::RemI32,
+    Op::ShlU32,
+    Op::ShrU32,
+    Op::ShlI32,
+    Op::ShrI32,
+    Op::And,
+    Op::Or,
+    Op::Xor,
+    Op::Not,
+    Op::LtI32,
+    Op::LtU32,
+    Op::LeI32,
+    Op::LeU32,
+    Op::Eq,
+    Op::Ne,
+    Op::SelectU32,
+    Op::SelectI32,
+    Op::MinI32,
+    Op::MaxI32,
+    Op::MinU32,
+    Op::MaxU32,
+    Op::ClampI32,
+    Op::ClampU32,
+    Op::AbsI32,
+    Op::SignI32,
+    Op::CountLeadingZerosU32,
+    Op::CountTrailingZerosU32,
+    Op::CountOneBitsU32,
+    Op::ReverseBitsU32,
+    Op::FirstLeadingBitU32,
+    Op::FirstLeadingBitI32,
+    Op::FirstTrailingBitU32,
+    Op::CountLeadingZerosI32,
+    Op::ExtractBitsU32,
+    Op::ExtractBitsI32,
+    Op::InsertBitsU32,
+    Op::InsertBitsI32,
+    Op::U32OfI32,
+    Op::I32OfU32,
+    Op::U32OfBool,
+    Op::AddVec2U32,
+    Op::MulVec2U32,
+    Op::MulVec2I32,
+    Op::ShrVec2U32,
+    Op::ShrVec2I32,
+    Op::SelectVec2U32,
+];
+
+/// What WGSL defines that the translators this crate runs through pass to the backend
+/// unchanged, and where the backend's language leaves the result undefined or to the
+/// implementation: a shift by 32 or more (MSL, HLSL and SPIR-V), and an integer `clamp`
+/// whose low bound is above its high one (MSL's `clamp`; naga writes SPIR-V's as `min`
+/// and `max`). On these operands the result is the hardware's; [`Op::guarded_case`]
+/// is false for them.
+pub const UNGUARDED: &str = "a shift amount of 32 or more; an integer clamp whose low bound is above its high bound";
+
+impl Op {
+    /// The operation's code: its index in [`ALL`] and its `case` in `sem`.
+    pub fn code(self) -> u32 {
+        ALL.iter().position(|&o| o == self).expect("in ALL") as u32
+    }
+
+    pub fn from_code(c: u32) -> Option<Op> {
+        ALL.get(c as usize).copied()
+    }
+
+    /// The name the vector file gives the operation.
+    pub fn name(self) -> &'static str {
+        use Op::*;
+        match self {
+            AddU32 => "add_u32",
+            SubU32 => "sub_u32",
+            MulU32 => "mul_u32",
+            AddI32 => "add_i32",
+            SubI32 => "sub_i32",
+            MulI32 => "mul_i32",
+            NegI32 => "neg_i32",
+            DivU32 => "div_u32",
+            RemU32 => "rem_u32",
+            DivI32 => "div_i32",
+            RemI32 => "rem_i32",
+            ShlU32 => "shl_u32",
+            ShrU32 => "shr_u32",
+            ShlI32 => "shl_i32",
+            ShrI32 => "shr_i32",
+            And => "and",
+            Or => "or",
+            Xor => "xor",
+            Not => "not",
+            LtI32 => "lt_i32",
+            LtU32 => "lt_u32",
+            LeI32 => "le_i32",
+            LeU32 => "le_u32",
+            Eq => "eq",
+            Ne => "ne",
+            SelectU32 => "select_u32",
+            SelectI32 => "select_i32",
+            MinI32 => "min_i32",
+            MaxI32 => "max_i32",
+            MinU32 => "min_u32",
+            MaxU32 => "max_u32",
+            ClampI32 => "clamp_i32",
+            ClampU32 => "clamp_u32",
+            AbsI32 => "abs_i32",
+            SignI32 => "sign_i32",
+            CountLeadingZerosU32 => "count_leading_zeros_u32",
+            CountTrailingZerosU32 => "count_trailing_zeros_u32",
+            CountOneBitsU32 => "count_one_bits_u32",
+            ReverseBitsU32 => "reverse_bits_u32",
+            FirstLeadingBitU32 => "first_leading_bit_u32",
+            FirstLeadingBitI32 => "first_leading_bit_i32",
+            FirstTrailingBitU32 => "first_trailing_bit_u32",
+            CountLeadingZerosI32 => "count_leading_zeros_i32",
+            ExtractBitsU32 => "extract_bits_u32",
+            ExtractBitsI32 => "extract_bits_i32",
+            InsertBitsU32 => "insert_bits_u32",
+            InsertBitsI32 => "insert_bits_i32",
+            U32OfI32 => "u32_of_i32",
+            I32OfU32 => "i32_of_u32",
+            U32OfBool => "u32_of_bool",
+            AddVec2U32 => "add_vec2_u32",
+            MulVec2U32 => "mul_vec2_u32",
+            MulVec2I32 => "mul_vec2_i32",
+            ShrVec2U32 => "shr_vec2_u32",
+            ShrVec2I32 => "shr_vec2_i32",
+            SelectVec2U32 => "select_vec2_u32",
+        }
+    }
+
+    pub fn from_name(n: &str) -> Option<Op> {
+        ALL.iter().copied().find(|o| o.name() == n)
+    }
+
+    /// The WGSL the kernel evaluates, with `a`, `b`, `c`, `d` the operands as u32 and
+    /// `ia`, `ib`, `ic`, `id` their bits as i32.
+    pub fn wgsl(self) -> &'static str {
+        use Op::*;
+        match self {
+            AddU32 => "a + b",
+            SubU32 => "a - b",
+            MulU32 => "a * b",
+            AddI32 => "ia + ib",
+            SubI32 => "ia - ib",
+            MulI32 => "ia * ib",
+            NegI32 => "-ia",
+            DivU32 => "a / b",
+            RemU32 => "a % b",
+            DivI32 => "ia / ib",
+            RemI32 => "ia % ib",
+            ShlU32 => "a << b",
+            ShrU32 => "a >> b",
+            ShlI32 => "ia << b",
+            ShrI32 => "ia >> b",
+            And => "a & b",
+            Or => "a | b",
+            Xor => "a ^ b",
+            Not => "~a",
+            LtI32 => "u32(ia < ib)",
+            LtU32 => "u32(a < b)",
+            LeI32 => "u32(ia <= ib)",
+            LeU32 => "u32(a <= b)",
+            Eq => "u32(a == b)",
+            Ne => "u32(a != b)",
+            SelectU32 => "select(b, a, c != 0u)",
+            SelectI32 => "select(ib, ia, c != 0u)",
+            MinI32 => "min(ia, ib)",
+            MaxI32 => "max(ia, ib)",
+            MinU32 => "min(a, b)",
+            MaxU32 => "max(a, b)",
+            ClampI32 => "clamp(ia, ib, ic)",
+            ClampU32 => "clamp(a, b, c)",
+            AbsI32 => "abs(ia)",
+            SignI32 => "sign(ia)",
+            CountLeadingZerosU32 => "countLeadingZeros(a)",
+            CountTrailingZerosU32 => "countTrailingZeros(a)",
+            CountOneBitsU32 => "countOneBits(a)",
+            ReverseBitsU32 => "reverseBits(a)",
+            FirstLeadingBitU32 => "firstLeadingBit(a)",
+            FirstLeadingBitI32 => "firstLeadingBit(ia)",
+            FirstTrailingBitU32 => "firstTrailingBit(a)",
+            CountLeadingZerosI32 => "countLeadingZeros(ia)",
+            ExtractBitsU32 => "extractBits(a, b, c)",
+            ExtractBitsI32 => "extractBits(ia, b, c)",
+            InsertBitsU32 => "insertBits(a, b, c, d)",
+            InsertBitsI32 => "insertBits(ia, ib, c, d)",
+            U32OfI32 => "u32(ia)",
+            I32OfU32 => "i32(a)",
+            U32OfBool => "u32(a < b)",
+            AddVec2U32 => "vec2(a, b) + vec2(c, d), folded",
+            MulVec2U32 => "vec2(a, b) * vec2(c, d), folded",
+            MulVec2I32 => "vec2(ia, ib) * vec2(ic, id), folded",
+            ShrVec2U32 => "vec2(a, b) >> vec2(c, d), folded",
+            ShrVec2I32 => "vec2(ia, ib) >> vec2(c, d), folded",
+            SelectVec2U32 => "select(vec2(a, b), vec2(c, d), vec2(a < c, b < d)), folded",
+        }
+    }
+
+    /// Whether the translators this crate runs through give WGSL's result for this
+    /// operation on these operands ([`UNGUARDED`] lists what they do not).
+    pub fn guarded_case(self, c: &Case) -> bool {
+        use Op::*;
+        match self {
+            ShlU32 | ShrU32 | ShlI32 | ShrI32 => c.b < 32,
+            ShrVec2U32 | ShrVec2I32 => c.c < 32 && c.d < 32,
+            ClampI32 => (c.b as i32) <= (c.c as i32),
+            ClampU32 => c.b <= c.c,
+            _ => true,
+        }
+    }
+}
+
+/// One case: an operation and its four operands (an operation reads the ones it names).
+#[derive(Clone, Copy, Debug, PartialEq, Eq, PartialOrd, Ord, Hash)]
+pub struct Case {
+    pub op: Op,
+    pub a: u32,
+    pub b: u32,
+    pub c: u32,
+    pub d: u32,
+}
+
+impl Case {
+    pub fn new(op: Op, a: u32, b: u32, c: u32, d: u32) -> Case {
+        Case { op, a, b, c, d }
+    }
+
+    /// Whether the translators give WGSL's result here ([`Op::guarded_case`]).
+    pub fn guarded(&self) -> bool {
+        self.op.guarded_case(self)
+    }
+}
+
+/// The two components of a vector result folded into one word, as `sem` folds them.
+fn fold(x: u32, y: u32) -> u32 {
+    x ^ y.wrapping_mul(0x9E37_79B9)
+}
+
+/// `(o, c)` of a bit field: `o = min(offset, 32)`, `c = min(count, 32 − o)`.
+fn field(offset: u32, count: u32) -> (u32, u32) {
+    let o = offset.min(32);
+    (o, count.min(32 - o))
+}
+
+/// The low `c` bits set, for `c` in 0..=32.
+fn low_mask(c: u32) -> u32 {
+    if c >= 32 { u32::MAX } else { (1u32 << c) - 1 }
+}
+
+/// The position of the most significant 1 bit, or `None` for 0.
+fn msb(x: u32) -> Option<u32> {
+    (x != 0).then(|| 31 - x.leading_zeros())
+}
+
+/// The result WGSL defines for `c`, as the bits of a u32. Written from the WGSL text
+/// (the sections cited in the module documentation), not from the shader: concrete
+/// integer arithmetic is modulo 2^32; `e1 / e2` truncates, is `e1` when `e2` is zero and
+/// when `e1` is the most negative i32 and `e2` is −1; `e1 % e2` is
+/// `e1 − truncate(e1 ÷ e2) × e2`, and 0 in those two cases; a shift's amount is `e2`
+/// modulo 32, and a right shift of an i32 copies the sign bit.
+pub fn reference(c: &Case) -> u32 {
+    use Op::*;
+    let (a, b, cc, d) = (c.a, c.b, c.c, c.d);
+    let (ia, ib, ic, id) = (a as i32, b as i32, cc as i32, d as i32);
+    let bit = |x: bool| x as u32;
+    match c.op {
+        AddU32 => a.wrapping_add(b),
+        SubU32 => a.wrapping_sub(b),
+        MulU32 => a.wrapping_mul(b),
+        AddI32 => ia.wrapping_add(ib) as u32,
+        SubI32 => ia.wrapping_sub(ib) as u32,
+        MulI32 => ia.wrapping_mul(ib) as u32,
+        NegI32 => ia.wrapping_neg() as u32,
+        // e1 / 0 is e1; e1 % 0 is 0.
+        DivU32 => a.checked_div(b).unwrap_or(a),
+        RemU32 => a.checked_rem(b).unwrap_or(0),
+        DivI32 => {
+            if ib == 0 || (ia == i32::MIN && ib == -1) {
+                a
+            } else {
+                (ia / ib) as u32
+            }
+        }
+        RemI32 => {
+            if ib == 0 || (ia == i32::MIN && ib == -1) {
+                0
+            } else {
+                (ia - (ia / ib) * ib) as u32
+            }
+        }
+        ShlU32 => a << (b % 32),
+        ShrU32 => a >> (b % 32),
+        ShlI32 => (ia << (b % 32)) as u32,
+        ShrI32 => (ia >> (b % 32)) as u32,
+        And => a & b,
+        Or => a | b,
+        Xor => a ^ b,
+        Not => !a,
+        LtI32 => bit(ia < ib),
+        LtU32 => bit(a < b),
+        LeI32 => bit(ia <= ib),
+        LeU32 => bit(a <= b),
+        Eq => bit(a == b),
+        Ne => bit(a != b),
+        SelectU32 | SelectI32 => {
+            if cc != 0 {
+                a
+            } else {
+                b
+            }
+        }
+        MinI32 => ia.min(ib) as u32,
+        MaxI32 => ia.max(ib) as u32,
+        MinU32 => a.min(b),
+        MaxU32 => a.max(b),
+        // min(max(e, low), high), whatever the order of the bounds.
+        ClampI32 => ia.max(ib).min(ic) as u32,
+        ClampU32 => a.max(b).min(cc),
+        // The absolute value of the most negative i32 is itself.
+        AbsI32 => {
+            if ia == i32::MIN {
+                a
+            } else {
+                ia.unsigned_abs()
+            }
+        }
+        SignI32 => ia.signum() as u32,
+        CountLeadingZerosU32 | CountLeadingZerosI32 => a.leading_zeros(),
+        CountTrailingZerosU32 => a.trailing_zeros(),
+        CountOneBitsU32 => a.count_ones(),
+        ReverseBitsU32 => a.reverse_bits(),
+        // T(−1) when e is zero; otherwise the position of the most significant 1.
+        FirstLeadingBitU32 => msb(a).unwrap_or(u32::MAX),
+        // −1 when e is 0 or −1; otherwise the position of the most significant bit
+        // that differs from the sign bit.
+        FirstLeadingBitI32 => {
+            if ia == 0 || ia == -1 {
+                u32::MAX
+            } else if ia < 0 {
+                msb(!a).unwrap_or(u32::MAX)
+            } else {
+                msb(a).unwrap_or(u32::MAX)
+            }
+        }
+        FirstTrailingBitU32 => {
+            if a == 0 {
+                u32::MAX
+            } else {
+                a.trailing_zeros()
+            }
+        }
+        // Bits o..o+c−1 of e; 0 when c is 0.
+        ExtractBitsU32 => {
+            let (o, n) = field(b, cc);
+            if n == 0 { 0 } else { (a >> o) & low_mask(n) }
+        }
+        // The same, sign-extended from bit c−1.
+        ExtractBitsI32 => {
+            let (o, n) = field(b, cc);
+            if n == 0 {
+                0
+            } else {
+                let v = (a >> o) & low_mask(n);
+                if (v >> (n - 1)) & 1 == 1 { v | !low_mask(n) } else { v }
+            }
+        }
+        // Bits o..o+c−1 from bits 0..c−1 of newbits, the others from e; e when c is 0.
+        InsertBitsU32 | InsertBitsI32 => {
+            let (o, n) = field(cc, d);
+            if n == 0 {
+                a
+            } else {
+                let m = low_mask(n) << o;
+                (a & !m) | ((b << o) & m)
+            }
+        }
+        // A conversion between i32 and u32 keeps the bit pattern.
+        U32OfI32 | I32OfU32 => a,
+        U32OfBool => bit(a < b),
+        AddVec2U32 => fold(a.wrapping_add(cc), b.wrapping_add(d)),
+        MulVec2U32 => fold(a.wrapping_mul(cc), b.wrapping_mul(d)),
+        MulVec2I32 => fold(ia.wrapping_mul(ic) as u32, ib.wrapping_mul(id) as u32),
+        ShrVec2U32 => fold(a >> (cc % 32), b >> (d % 32)),
+        ShrVec2I32 => fold((ia >> (cc % 32)) as u32, (ib >> (d % 32)) as u32),
+        SelectVec2U32 => fold(if a < cc { cc } else { a }, if b < d { d } else { b }),
+    }
+}
+
+/// The 32-bit values at the edges of i32 and u32, and near them, and at the boundaries
+/// of shift amounts and bit fields.
+pub fn edge_values() -> Vec<u32> {
+    let mut v: Vec<u32> = vec![
+        0,
+        1,
+        2,
+        3,
+        5,
+        7,
+        31,
+        32,
+        33,
+        63,
+        64,
+        255,
+        0x7FFF,
+        0x8000,
+        0xFFFF,
+        0x1_0000,
+        0x5555_5555,
+        0x7FFF_FFFE,
+        0x7FFF_FFFF,
+        0x8000_0000,
+        0x8000_0001,
+        0xAAAA_AAAA,
+        0xDEAD_BEEF,
+    ];
+    for x in [1i32, 2, 3, 7, 31, 32, 0x8000, 0x1_0000] {
+        v.push((-x) as u32);
+    }
+    v.push(u32::MAX);
+    v.sort_unstable();
+    v.dedup();
+    v
+}
+
+/// Shift amounts and bit-field offsets and counts: every value from 0 to 40, then
+/// amounts past 32 that are 0 or 1 modulo 32, and the extremes of u32.
+pub fn amounts() -> Vec<u32> {
+    let mut v: Vec<u32> = (0..=40).collect();
+    v.extend([63, 64, 65, 96, 127, 128, 255, 256, 0x7FFF_FFFF, 0x8000_0000, u32::MAX - 31, u32::MAX]);
+    v
+}
+
+/// Whether the operation is a shift whose amount is its operand `b`.
+fn shift_by_b(op: Op) -> bool {
+    matches!(op, Op::ShlU32 | Op::ShrU32 | Op::ShlI32 | Op::ShrI32)
+}
+
+/// The edge cases of every operation, in code order and then by operands: the cases the
+/// committed vector file holds. The arithmetic operations, division and remainder take
+/// every pair of [`edge_values`]; shifts every edge value by every one of [`amounts`];
+/// the bitwise operations, comparisons, `min` and `max` every pair of a short list;
+/// `clamp` and `select` every triple of it; bit fields a few values at every (offset,
+/// count) pair around 0, 16 and 32 and the extremes.
+pub fn edge_cases() -> Vec<Case> {
+    use Op::*;
+    let e = edge_values();
+    let s = amounts();
+    let short: Vec<u32> = vec![0, 1, 7, 0x7FFF_FFFF, 0x8000_0000, 0xFFFF_FFF9, u32::MAX];
+    let fields: Vec<u32> = vec![0, 1, 2, 15, 16, 17, 30, 31, 32, 33, 64, u32::MAX];
+    let mut out = Vec::new();
+    for op in ALL {
+        match op {
+            _ if shift_by_b(op) => {
+                for &a in &e {
+                    for &b in &s {
+                        out.push(Case::new(op, a, b, 0, 0));
+                    }
+                }
+            }
+            NegI32 | Not | AbsI32 | SignI32 | CountLeadingZerosU32 | CountTrailingZerosU32 | CountOneBitsU32 | ReverseBitsU32
+            | FirstLeadingBitU32 | FirstLeadingBitI32 | FirstTrailingBitU32 | CountLeadingZerosI32 | U32OfI32 | I32OfU32 => {
+                for &a in &e {
+                    out.push(Case::new(op, a, 0, 0, 0));
+                }
+                // Every single bit, and every run of low ones.
+                for k in 0..32u32 {
+                    out.push(Case::new(op, 1u32 << k, 0, 0, 0));
+                    out.push(Case::new(op, low_mask(k + 1), 0, 0, 0));
+                    out.push(Case::new(op, !low_mask(k), 0, 0, 0));
+                }
+            }
+            ClampI32 | ClampU32 | SelectU32 | SelectI32 => {
+                for &a in &short {
+                    for &b in &short {
+                        for &c in &short {
+                            out.push(Case::new(op, a, b, c, 0));
+                        }
+                    }
+                }
+            }
+            ExtractBitsU32 | ExtractBitsI32 => {
+                for &a in &[0u32, 1, 0x8000_0000, 0xDEAD_BEEF, 0x7FFF_FFFF, u32::MAX, 0x5555_5555] {
+                    for &o in &fields {
+                        for &n in &fields {
+                            out.push(Case::new(op, a, o, n, 0));
+                        }
+                    }
+                }
+            }
+            InsertBitsU32 | InsertBitsI32 => {
+                for &(a, b) in &[(0u32, u32::MAX), (u32::MAX, 0), (0xDEAD_BEEF, 0x1234_5678), (0x8000_0000, 1)] {
+                    for &o in &fields {
+                        for &n in &fields {
+                            out.push(Case::new(op, a, b, o, n));
+                        }
+                    }
+                }
+            }
+            AddVec2U32 | MulVec2U32 | MulVec2I32 | SelectVec2U32 => {
+                for &a in &short {
+                    for &b in &short {
+                        for &c in &short {
+                            out.push(Case::new(op, a, b, c, a ^ b ^ c));
+                        }
+                    }
+                }
+            }
+            ShrVec2U32 | ShrVec2I32 => {
+                for &a in &[0u32, 1, 0x8000_0000, 0xDEAD_BEEF, u32::MAX] {
+                    for &c in &s {
+                        out.push(Case::new(op, a, !a, c, c.wrapping_add(1)));
+                    }
+                }
+            }
+            AddU32 | SubU32 | MulU32 | AddI32 | SubI32 | MulI32 | DivU32 | RemU32 | DivI32 | RemI32 => {
+                for &a in &e {
+                    for &b in &e {
+                        out.push(Case::new(op, a, b, 0, 0));
+                    }
+                }
+            }
+            _ => {
+                for &a in &short {
+                    for &b in &short {
+                        out.push(Case::new(op, a, b, 0, 0));
+                    }
+                }
+            }
+        }
+    }
+    out
+}
+
+struct SplitMix(u64);
+
+impl SplitMix {
+    fn next(&mut self) -> u64 {
+        self.0 = self.0.wrapping_add(0x9E37_79B9_7F4A_7C15);
+        let mut z = self.0;
+        z = (z ^ (z >> 30)).wrapping_mul(0xBF58_476D_1CE4_E5B9);
+        z = (z ^ (z >> 27)).wrapping_mul(0x94D0_49BB_1331_11EB);
+        z ^ (z >> 31)
+    }
+}
+
+/// `per_op` SplitMix64 cases of every operation, from `seed`. Shift amounts and
+/// bit-field bounds are drawn below 40 half the time, so the boundary is reached.
+pub fn random_cases(per_op: usize, seed: u64) -> Vec<Case> {
+    let mut r = SplitMix(seed);
+    let mut out = Vec::with_capacity(per_op * ALL.len());
+    for op in ALL {
+        for k in 0..per_op {
+            let w = r.next();
+            let x = r.next();
+            let (mut a, mut b, mut c, mut d) = (w as u32, (w >> 32) as u32, x as u32, (x >> 32) as u32);
+            if k % 2 == 0 {
+                b %= 40;
+                c %= 40;
+                d %= 40;
+            }
+            if k % 8 == 1 {
+                a = 0x8000_0000;
+                b = u32::MAX;
+            }
+            out.push(Case::new(op, a, b, c, d));
+        }
+    }
+    out
+}
+
+/// The device's word for each case.
+pub async fn run(g: &Gpu, cases: &[Case]) -> Result<Vec<u32>, GpuError> {
+    let mut out = Vec::with_capacity(cases.len());
+    // At most 2^20 cases per dispatch, to keep each buffer small.
+    for chunk in cases.chunks(1 << 20) {
+        out.extend(run_chunk(g, chunk).await?);
+    }
+    Ok(out)
+}
+
+async fn run_chunk(g: &Gpu, cases: &[Case]) -> Result<Vec<u32>, GpuError> {
+    let n = cases.len() as u32;
+    if n == 0 {
+        return Ok(Vec::new());
+    }
+    let mut inw: Vec<u8> = Vec::with_capacity(cases.len() * IN_WORDS * 4);
+    for c in cases {
+        for x in [c.op.code(), c.a, c.b, c.c, c.d] {
+            inw.extend_from_slice(&x.to_le_bytes());
+        }
+    }
+    let out_bytes = (cases.len() * 4) as u64;
+    let (gx, gy, row) = grid(n, WG_PROBE).ok_or_else(|| GpuError::Device("too many semantics cases".into()))?;
+    let mut rec = Rec::new();
+    rec.set(f::N, n).set(f::ROW, row);
+    let scopes = Scopes::push(&g.device);
+    let uni = g.buffer("wai-ops-gpu semantics record", g.stride, wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST);
+    g.queue.write_buffer(&uni, 0, &rec.bytes());
+    let ins = g.buffer("wai-ops-gpu semantics in", inw.len() as u64, wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST);
+    g.queue.write_buffer(&ins, 0, &inw);
+    let out = g.buffer("wai-ops-gpu semantics out", out_bytes, wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC);
+    let flags = g.buffer("wai-ops-gpu semantics flags", 16, wgpu::BufferUsages::STORAGE);
+    let staging = g.buffer("wai-ops-gpu semantics staging", out_bytes, wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST);
+    let bind = g.device.create_bind_group(&wgpu::BindGroupDescriptor {
+        label: Some("wai-ops-gpu semantics"),
+        layout: &g.pipes.layout,
+        entries: &[
+            wgpu::BindGroupEntry {
+                binding: 0,
+                resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding { buffer: &uni, offset: 0, size: core::num::NonZeroU64::new(crate::kernels::REC_WORDS as u64 * 4) }),
+            },
+            wgpu::BindGroupEntry { binding: 1, resource: ins.as_entire_binding() },
+            wgpu::BindGroupEntry { binding: 2, resource: g.dummy.as_entire_binding() },
+            wgpu::BindGroupEntry { binding: 3, resource: g.dummy.as_entire_binding() },
+            wgpu::BindGroupEntry { binding: 4, resource: out.as_entire_binding() },
+            wgpu::BindGroupEntry { binding: 5, resource: flags.as_entire_binding() },
+        ],
+    });
+    let mut enc = g.device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("wai-ops-gpu semantics") });
+    {
+        let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("wai-ops-gpu semantics"), timestamp_writes: None });
+        pass.set_pipeline(&g.pipes.semantics);
+        pass.set_bind_group(0, &bind, &[]);
+        pass.dispatch_workgroups(gx, gy, 1);
+    }
+    enc.copy_buffer_to_buffer(&out, 0, &staging, 0, out_bytes);
+    let index = g.queue.submit([enc.finish()]);
+    scopes.close().await?;
+    let bytes = g.read(&staging, out_bytes, index).await?;
+    g.settled()?;
+    Ok(bytes.as_chunks::<4>().0.iter().map(|b| u32::from_le_bytes(*b)).collect())
+}
+
+/// One case on which a result differs from WGSL's.
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+pub struct Difference {
+    pub case: Case,
+    pub want: u32,
+    pub got: u32,
+}
+
+impl core::fmt::Display for Difference {
+    fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result {
+        let c = &self.case;
+        write!(f, "{} ({}) a={:#010x} b={:#010x} c={:#010x} d={:#010x}: WGSL gives {:#010x}, got {:#010x}", c.op.name(), c.op.wgsl(), c.a, c.b, c.c, c.d, self.want, self.got)
+    }
+}
+
+/// What a run of the suite found: the cases compared, and every difference, the
+/// guarded ones apart from the rest.
+#[derive(Clone, Debug, Default, PartialEq, Eq)]
+pub struct Outcome {
+    pub guarded: usize,
+    pub unguarded: usize,
+    pub guarded_differ: Vec<Difference>,
+    pub unguarded_differ: Vec<Difference>,
+}
+
+impl Outcome {
+    /// The operations with an unguarded case that differs, by name, each once.
+    pub fn unguarded_ops(&self) -> Vec<&'static str> {
+        let mut v: Vec<&'static str> = self.unguarded_differ.iter().map(|d| d.case.op.name()).collect();
+        v.dedup();
+        v
+    }
+}
+
+/// `got[i]` against `expect[i]` for every case.
+pub fn compare(cases: &[Case], expect: &[u32], got: &[u32]) -> Outcome {
+    let mut o = Outcome::default();
+    for (i, c) in cases.iter().enumerate() {
+        let (want, have) = (expect[i], got.get(i).copied().unwrap_or(!expect[i]));
+        let g = c.guarded();
+        if g {
+            o.guarded += 1;
+        } else {
+            o.unguarded += 1;
+        }
+        if want != have {
+            let d = Difference { case: *c, want, got: have };
+            if g { o.guarded_differ.push(d) } else { o.unguarded_differ.push(d) }
+        }
+    }
+    o
+}
+
+/// The vector file: one JSON object per operation, in code order, with its WGSL and its
+/// cases, each `[a, b, c, d, result]` as eight lowercase hexadecimal digits, and
+/// `guarded` false where [`Op::guarded_case`] is.
+pub fn vector_file(cases: &[Case]) -> String {
+    let mut s = String::new();
+    for op in ALL {
+        let mine: Vec<&Case> = cases.iter().filter(|c| c.op == op).collect();
+        if mine.is_empty() {
+            continue;
+        }
+        s.push_str(&format!("{{\"op\":\"{}\",\"code\":{},\"wgsl\":\"{}\",\"cases\":[", op.name(), op.code(), op.wgsl()));
+        for (k, c) in mine.iter().enumerate() {
+            if k > 0 {
+                s.push(',');
+            }
+            s.push_str(&format!("[\"{:08x}\",\"{:08x}\",\"{:08x}\",\"{:08x}\",\"{:08x}\"", c.a, c.b, c.c, c.d, reference(c)));
+            if !c.guarded() {
+                s.push_str(",\"unguarded\"");
+            }
+            s.push(']');
+        }
+        s.push_str("]}\n");
+    }
+    s
+}
+
+/// The cases and committed results of a vector file, in file order.
+pub fn parse_vector_file(text: &str) -> Result<Vec<(Case, u32)>, String> {
+    use wai_ops::json::{self, Value};
+    let mut out = Vec::new();
+    for (ln, line) in text.split_terminator('\n').enumerate() {
+        let v = json::parse(line.as_bytes()).map_err(|e| format!("line {}: JSON at byte {}", ln + 1, e.at))?;
+        let name = v.get("op").and_then(Value::as_str).ok_or(format!("line {}: no op", ln + 1))?;
+        let op = Op::from_name(name).ok_or(format!("line {}: unknown op {name}", ln + 1))?;
+        for c in v.get("cases").and_then(Value::as_arr).ok_or(format!("line {}: no cases", ln + 1))? {
+            let w = c.as_arr().ok_or(format!("line {}: a case is not an array", ln + 1))?;
+            let word = |k: usize| -> Result<u32, String> {
+                let s = w.get(k).and_then(Value::as_str).ok_or(format!("line {}: a case has no word {k}", ln + 1))?;
+                if s.len() != 8 {
+                    return Err(format!("line {}: {s:?} is not eight hexadecimal digits", ln + 1));
+                }
+                u32::from_str_radix(s, 16).map_err(|_| format!("line {}: {s:?} is not hexadecimal", ln + 1))
+            };
+            let case = Case::new(op, word(0)?, word(1)?, word(2)?, word(3)?);
+            let unguarded = w.get(5).and_then(Value::as_str) == Some("unguarded");
+            if unguarded == case.guarded() {
+                return Err(format!("line {}: {name} {:?}: the file's guard mark is not the module's", ln + 1, case));
+            }
+            out.push((case, word(4)?));
+        }
+    }
+    Ok(out)
+}
+
+/// The discovery check: [`edge_cases`] and SplitMix64 cases of every operation, on the
+/// device, against [`reference`]. `Err` names the first guarded case that differs; an
+/// unguarded case that differs is reported by [`run`]'s callers, and does not refuse the
+/// adapter, since no kernel depends on it.
+pub(crate) async fn check(g: &Gpu) -> Result<(), String> {
+    let mut cases = edge_cases();
+    cases.extend(random_cases(256, 0x5EED_0006));
+    let expect: Vec<u32> = cases.iter().map(reference).collect();
+    let got = run(g, &cases).await.map_err(|e| format!("the WGSL semantics suite did not run: {e}"))?;
+    let o = compare(&cases, &expect, &got);
+    match o.guarded_differ.first() {
+        None => Ok(()),
+        Some(d) => Err(format!("the WGSL semantics suite: {d}")),
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    /// The reference on the cases the WGSL text works through, and on the values of
+    /// each special rule.
+    #[test]
+    fn the_reference_is_the_wgsl_text() {
+        let r = |op, a: u32, b: u32| reference(&Case::new(op, a, b, 0, 0));
+        let min = i32::MIN as u32;
+        assert_eq!(r(Op::DivI32, 7, 0), 7, "e1 / 0 is e1");
+        assert_eq!(r(Op::DivU32, 7, 0), 7);
+        assert_eq!(r(Op::RemI32, 7, 0), 0, "e1 % 0 is 0");
+        assert_eq!(r(Op::RemU32, 7, 0), 0);
+        assert_eq!(r(Op::DivI32, min, u32::MAX), min, "the most negative i32 over -1 is itself");
+        assert_eq!(r(Op::RemI32, min, u32::MAX), 0);
+        assert_eq!(r(Op::DivI32, (-7i32) as u32, 2), (-3i32) as u32, "division truncates");
+        assert_eq!(r(Op::RemI32, (-7i32) as u32, 2), (-1i32) as u32, "the remainder has the dividend's sign");
+        assert_eq!(r(Op::ShlU32, 1, 32), 1, "the amount is taken modulo 32");
+        assert_eq!(r(Op::ShlU32, 1, 33), 2);
+        assert_eq!(r(Op::ShrI32, min, 63), u32::MAX);
+        assert_eq!(r(Op::NegI32, min, 0), min, "negation wraps");
+        assert_eq!(r(Op::AbsI32, min, 0), min, "abs of the most negative i32 is itself");
+        assert_eq!(r(Op::MulI32, 0x7FFF_FFFF, 2), (-2i32) as u32, "multiplication wraps");
+        assert_eq!(reference(&Case::new(Op::ClampI32, 5, 10, 0, 0)), 0, "clamp(5, 10, 0) = min(max(5, 10), 0)");
+        assert_eq!(r(Op::FirstLeadingBitI32, u32::MAX, 0), u32::MAX, "firstLeadingBit(-1) is -1");
+        assert_eq!(r(Op::FirstLeadingBitI32, (-2i32) as u32, 0), 0);
+        assert_eq!(r(Op::FirstLeadingBitU32, 0, 0), u32::MAX);
+        assert_eq!(reference(&Case::new(Op::ExtractBitsI32, 0x80, 4, 4, 0)), u32::MAX - 7, "bit c-1 extends");
+        assert_eq!(reference(&Case::new(Op::ExtractBitsU32, u32::MAX, 30, 10, 0)), 3, "the count is clamped to 32 - offset");
+        assert_eq!(reference(&Case::new(Op::ExtractBitsU32, u32::MAX, 40, 10, 0)), 0, "an offset past 32 extracts nothing");
+        assert_eq!(reference(&Case::new(Op::InsertBitsU32, 0, u32::MAX, 28, 8)), 0xF000_0000);
+        assert_eq!(reference(&Case::new(Op::InsertBitsU32, 0x1234, u32::MAX, 32, 8)), 0x1234, "e when c is 0");
+    }
+
+    #[test]
+    fn every_operation_has_a_code_name_and_edge_cases() {
+        for (k, op) in ALL.iter().enumerate() {
+            assert_eq!(op.code() as usize, k);
+            assert_eq!(Op::from_code(k as u32), Some(*op));
+            assert_eq!(Op::from_name(op.name()), Some(*op));
+            assert!(edge_cases().iter().any(|c| c.op == *op), "{op:?} has no edge case");
+        }
+        assert_eq!(Op::from_code(ALL.len() as u32), None);
+    }
+
+    #[test]
+    fn the_vector_file_round_trips() {
+        let cases = edge_cases();
+        let text = vector_file(&cases);
+        let back = parse_vector_file(&text).unwrap();
+        assert_eq!(back.len(), cases.len());
+        for ((c, r), want) in back.iter().zip(&cases) {
+            assert_eq!(c, want);
+            assert_eq!(*r, reference(want));
+        }
+    }
+}

wai/wai-ops-gpu/src/wgsl/arith.wgsl

@@ -5,9 +5,13 @@
 // operation here uses only what every WGSL implementation defines identically:
 // `u32` addition, subtraction and multiplication, which wrap; bitwise operations;
 // comparisons; `bitcast`; and shifts, whose amount WGSL takes modulo 32, so every
-// shift of 32 or more is handled explicitly (integer-operators §6). No floating
+// shift of 32 or more is handled explicitly (integer-operators §6). Every shift by an
-// point, and no signed operation that could overflow: signed values are read from
+// amount known only at run time is also written with its amount masked, `& 31u`:
-// a word's bits only to compare or to shift right arithmetically.
+// WGSL gives the same result either way, and a translator that passes the amount to
+// a language where a shift of 32 or more is undefined (MSL, HLSL, SPIR-V; naga 30.0.1
+// masks none) then never passes one (determinism-tiers §5, *Claiming the tier on a
+// GPU*). No floating point, and no signed operation that could overflow: signed values
+// are read from a word's bits only to compare or to shift right arithmetically.
 //
 // Each function's result is exact modulo 2^64 (W) or 2^128 (Q). A kernel uses one
 // only where the graph's certificate places the exact value inside that width
@@ -118,10 +122,10 @@ fn w_sar(a: W, s: u32) -> W {
         return a;
     }
     if (s < 32u) {
-        return W((a.x >> s) | (a.y << (32u - s)), bitcast<u32>(bitcast<i32>(a.y) >> s));
+        return W((a.x >> (s & 31u)) | (a.y << ((32u - s) & 31u)), bitcast<u32>(bitcast<i32>(a.y) >> (s & 31u)));
     }
     if (s < 64u) {
-        return W(bitcast<u32>(bitcast<i32>(a.y) >> (s - 32u)), sg);
+        return W(bitcast<u32>(bitcast<i32>(a.y) >> ((s - 32u) & 31u)), sg);
     }
     return W(sg, sg);
 }
@@ -129,9 +133,9 @@ fn w_sar(a: W, s: u32) -> W {
 // 2^k for k in 0..=63 (2^63 is the bit pattern of −2^63).
 fn w_pow2(k: u32) -> W {
     if (k < 32u) {
-        return W(1u << k, 0u);
+        return W(1u << (k & 31u), 0u);
     }
-    return W(0u, 1u << (k - 32u));
+    return W(0u, 1u << ((k - 32u) & 31u));
 }
 
 // R_round(p, s) (§2) on 64 bits, for s in 0..=63, mode 0 floor, 1 half_up, 2
@@ -159,7 +163,7 @@ fn sar32(a: u32, s: u32) -> u32 {
     if (s >= 32u) {
         return bitcast<u32>(bitcast<i32>(a) >> 31u);
     }
-    return bitcast<u32>(bitcast<i32>(a) >> s);
+    return bitcast<u32>(bitcast<i32>(a) >> (s & 31u));
 }
 
 // isqrt of a 32-bit value, digit by digit.
@@ -324,7 +328,7 @@ fn q_shr_fill(a: Q, s: u32, f: u32) -> Q {
     for (var i = 0u; i < 4u; i++) {
         let lo = q_word(a, i + k, f);
         let hi = q_word(a, i + k + 1u, f);
-        r[i] = select((lo >> b) | (hi << (32u - b)), lo, b == 0u);
+        r[i] = select((lo >> b) | (hi << ((32u - b) & 31u)), lo, b == 0u);
     }
     return Q(r[0], r[1], r[2], r[3]);
 }

wai/wai-ops-gpu/src/wgsl/ops32.wgsl

@@ -199,7 +199,10 @@ fn lut16(@builtin(global_invocation_id) g: vec3<u32>) {
 }
 
 // lo and hi are the host's clamp bounds held to i32, which gives the same result
-// on an i32 value (a bound past i32 is not reached by any i32).
+// on an i32 value (a bound past i32 is not reached by any i32). Load refuses a clamp
+// whose lo is above its hi; the bound is applied as min(max(v, lo), hi), WGSL's
+// definition of an integer clamp, which every translation keeps whatever the bounds,
+// where a backend's own clamp can leave bounds out of order undefined.
 @compute @workgroup_size(256)
 fn ew(@builtin(global_invocation_id) g: vec3<u32>) {
     let t = tid(g);
@@ -209,7 +212,7 @@ fn ew(@builtin(global_invocation_id) g: vec3<u32>) {
     let v = A[t];
     switch P.mode {
         case 0u: {
-            O[t] = bitcast<u32>(clamp(bitcast<i32>(v), bitcast<i32>(P.lo_x), bitcast<i32>(P.hi_x)));
+            O[t] = bitcast<u32>(min(max(bitcast<i32>(v), bitcast<i32>(P.lo_x)), bitcast<i32>(P.hi_x)));
         }
         case 1u: {
             O[t] = select(v, 0u, lt32(v, 0u));
@@ -239,7 +242,7 @@ fn policy(@builtin(global_invocation_id) g: vec3<u32>) {
     let lo = bitcast<i32>(P.lo_x);
     let hi = bitcast<i32>(P.hi_x);
     if (P.mode == 0u) {
-        O[t] = bitcast<u32>(clamp(v, lo, hi));
+        O[t] = bitcast<u32>(min(max(v, lo), hi));
     } else if (v < lo || v > hi) {
         atomicOr(&F[P.flag], 1u);
     }

wai/wai-ops-gpu/src/wgsl/semantics.wgsl

@@ -0,0 +1,125 @@
+// WGSL's integer semantics on the device: each record of A (five words: an
+// operation, then the operands a, b, c and d) gives one word of O, the operation's
+// result, its bits as a u32. The host computes the same operations from the text of
+// the WGSL specification (`src/semantics.rs`, `reference`) and compares every word.
+//
+// Every operand is read from a storage buffer, so no result is known before the
+// shader runs: WGSL makes a division by zero, a shift of 32 or more and an
+// out-of-range bit field errors only where the operands are const- or
+// override-expressions, and defines a result for each at run time. Those results
+// are what this kernel checks, since a translation into a language that leaves them
+// undefined (MSL, HLSL, SPIR-V) has to produce them itself. Integer only: no
+// floating-point type appears in this module.
+
+const SEM_IN: u32 = 5u;
+
+fn sem(op: u32, a: u32, b: u32, c: u32, d: u32) -> u32 {
+    let ia = bitcast<i32>(a);
+    let ib = bitcast<i32>(b);
+    let ic = bitcast<i32>(c);
+    var r = 0xDEADBEEFu;
+    switch op {
+        // Wrapping arithmetic, modulo 2^32.
+        case 0u: { r = a + b; }
+        case 1u: { r = a - b; }
+        case 2u: { r = a * b; }
+        case 3u: { r = bitcast<u32>(ia + ib); }
+        case 4u: { r = bitcast<u32>(ia - ib); }
+        case 5u: { r = bitcast<u32>(ia * ib); }
+        case 6u: { r = bitcast<u32>(-ia); }
+        // Division and remainder, with WGSL's run-time results for a zero divisor
+        // and for the most negative i32 over -1.
+        case 7u: { r = a / b; }
+        case 8u: { r = a % b; }
+        case 9u: { r = bitcast<u32>(ia / ib); }
+        case 10u: { r = bitcast<u32>(ia % ib); }
+        // Shifts: the amount is taken modulo 32.
+        case 11u: { r = a << b; }
+        case 12u: { r = a >> b; }
+        case 13u: { r = bitcast<u32>(ia << b); }
+        case 14u: { r = bitcast<u32>(ia >> b); }
+        // Bitwise.
+        case 15u: { r = a & b; }
+        case 16u: { r = a | b; }
+        case 17u: { r = a ^ b; }
+        case 18u: { r = ~a; }
+        // Comparisons, signed and unsigned.
+        case 19u: { r = select(0u, 1u, ia < ib); }
+        case 20u: { r = select(0u, 1u, a < b); }
+        case 21u: { r = select(0u, 1u, ia <= ib); }
+        case 22u: { r = select(0u, 1u, a <= b); }
+        case 23u: { r = select(0u, 1u, a == b); }
+        case 24u: { r = select(0u, 1u, a != b); }
+        // select(f, t, cond).
+        case 25u: { r = select(b, a, c != 0u); }
+        case 26u: { r = bitcast<u32>(select(ib, ia, c != 0u)); }
+        // min, max and clamp: clamp(e, low, high) is min(max(e, low), high) for an
+        // integer type, low above high included.
+        case 27u: { r = bitcast<u32>(min(ia, ib)); }
+        case 28u: { r = bitcast<u32>(max(ia, ib)); }
+        case 29u: { r = min(a, b); }
+        case 30u: { r = max(a, b); }
+        case 31u: { r = bitcast<u32>(clamp(ia, ib, ic)); }
+        case 32u: { r = clamp(a, b, c); }
+        // abs and sign of an i32; the absolute value of the most negative i32 is
+        // itself.
+        case 33u: { r = bitcast<u32>(abs(ia)); }
+        case 34u: { r = bitcast<u32>(sign(ia)); }
+        // Bit counts and positions.
+        case 35u: { r = countLeadingZeros(a); }
+        case 36u: { r = countTrailingZeros(a); }
+        case 37u: { r = countOneBits(a); }
+        case 38u: { r = reverseBits(a); }
+        case 39u: { r = firstLeadingBit(a); }
+        case 40u: { r = bitcast<u32>(firstLeadingBit(ia)); }
+        case 41u: { r = firstTrailingBit(a); }
+        case 42u: { r = bitcast<u32>(countLeadingZeros(ia)); }
+        // Bit fields: o = min(offset, 32), c = min(count, 32 - o).
+        case 43u: { r = extractBits(a, b, c); }
+        case 44u: { r = bitcast<u32>(extractBits(ia, b, c)); }
+        case 45u: { r = insertBits(a, b, c, d); }
+        case 46u: { r = bitcast<u32>(insertBits(ia, ib, c, d)); }
+        // Value conversions between i32 and u32 keep the bits; bool converts to 0 or 1.
+        case 47u: { r = u32(ia); }
+        case 48u: { r = bitcast<u32>(i32(a)); }
+        case 49u: { r = u32(a < b); }
+        // The same arithmetic on two-component vectors, component by component: the
+        // kernels' 64-bit and 128-bit words are vectors.
+        case 50u: {
+            let v = vec2<u32>(a, b) + vec2<u32>(c, d);
+            r = v.x ^ (v.y * 0x9E3779B9u);
+        }
+        case 51u: {
+            let v = vec2<u32>(a, b) * vec2<u32>(c, d);
+            r = v.x ^ (v.y * 0x9E3779B9u);
+        }
+        case 52u: {
+            let v = bitcast<vec2<u32>>(vec2<i32>(ia, ib) * vec2<i32>(ic, bitcast<i32>(d)));
+            r = v.x ^ (v.y * 0x9E3779B9u);
+        }
+        case 53u: {
+            let v = vec2<u32>(a, b) >> vec2<u32>(c, d);
+            r = v.x ^ (v.y * 0x9E3779B9u);
+        }
+        case 54u: {
+            let v = bitcast<vec2<u32>>(vec2<i32>(ia, ib) >> vec2<u32>(c, d));
+            r = v.x ^ (v.y * 0x9E3779B9u);
+        }
+        case 55u: {
+            let v = select(vec2<u32>(a, b), vec2<u32>(c, d), vec2<bool>(a < c, b < d));
+            r = v.x ^ (v.y * 0x9E3779B9u);
+        }
+        default: {}
+    }
+    return r;
+}
+
+@compute @workgroup_size(64)
+fn semantics(@builtin(global_invocation_id) g: vec3<u32>) {
+    let t = tid(g);
+    if (t >= P.n) {
+        return;
+    }
+    let i = t * SEM_IN;
+    O[t] = sem(A[i], A[i + 1u], A[i + 2u], A[i + 3u], A[i + 4u]);
+}

wai/wai-ops-gpu/tests/mutants.rs

@@ -30,7 +30,7 @@ const MUTANTS: &[(&str, Module, &str, &str)] = &[
     ("half_even ties to odd (acc64)", Module::Ops64, "(q_eq(rem, h) && (q.x & 1u) == 1u)", "(q_eq(rem, h) && (q.x & 1u) == 0u)"),
     ("half_even ties to odd (acc32)", Module::Ops32, "(w_eq(rem, h) && (q.x & 1u) == 1u)", "(w_eq(rem, h) && (q.x & 1u) == 0u)"),
     ("clamp operator saturates one below its top (acc64)", Module::Ops64, "put64(t, w_clamp(a64(t), lo, hi));", "put64(t, w_clamp(a64(t), lo, w_sub(hi, W(1u, 0u))));"),
-    ("clamp policy saturates one above its bottom (acc32)", Module::Ops32, "O[t] = bitcast<u32>(clamp(v, lo, hi));", "O[t] = bitcast<u32>(clamp(v, lo + 1, hi));"),
+    ("clamp policy saturates one above its bottom (acc32)", Module::Ops32, "O[t] = bitcast<u32>(min(max(v, lo), hi));", "O[t] = bitcast<u32>(min(max(v, lo + 1), hi));"),
     ("check policy misses values above its range (acc64)", Module::Ops64, "} else if (w_lt(v, lo) || w_lt(hi, v)) {", "} else if (w_lt(v, lo)) {"),
     // In the 32-bit kernels the 32 × 32-bit product only ever sees magnitudes of at most
     // 2^31, for which neither carry of its middle sum can be set (a0·b1 + a1·b0 stays

wai/wai-ops-gpu/tests/semantics.rs

@@ -0,0 +1,113 @@
+//! WGSL's integer semantics on every adapter (`src/semantics.rs`): the committed vectors
+//! of `wai/wgsl-conformance/` are the WGSL text's results, and every adapter found gives
+//! them, with SplitMix64 cases of every operation besides. A case the translators do not
+//! guard (`semantics::UNGUARDED`) is run and reported, and does not fail the adapter: no
+//! kernel depends on one (`tests/wgsl_rules.rs`). Then an edited kernel is shown to be
+//! caught: discovery refuses every adapter.
+//!
+//! Each adapter's result is one line, `ops-gpu semantics: key=…`, appended to
+//! `WAI_OPS_GPU_REPORT` when it is set. No adapter is not a pass: the test prints
+//! SKIPPED, and fails under `WAI_OPS_GPU_REQUIRE`.
+
+use std::io::Write;
+use std::path::PathBuf;
+
+use wai_ops_gpu::semantics::{self, Case};
+use wai_ops_gpu::{blocking, kernels::Sources};
+
+fn report(line: &str) {
+    println!("{line}");
+    if let Ok(p) = std::env::var("WAI_OPS_GPU_REPORT") {
+        let mut f = std::fs::OpenOptions::new().create(true).append(true).open(&p).unwrap_or_else(|e| panic!("{p}: {e}"));
+        writeln!(f, "{line}").unwrap();
+    }
+}
+
+fn committed() -> (PathBuf, String) {
+    let p = PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("../wgsl-conformance/integer.jsonl");
+    let text = std::fs::read_to_string(&p).unwrap_or_else(|e| panic!("{}: {e}", p.display()));
+    (p, text)
+}
+
+/// The committed file is what `wai-wgsl-vectors` writes, byte for byte, and each of its
+/// results is the reference's.
+#[test]
+fn the_committed_vectors_are_the_wgsl_texts_results() {
+    let (p, text) = committed();
+    let want = semantics::vector_file(&semantics::edge_cases());
+    assert!(text == want, "{} is not what `cargo run --bin wai-wgsl-vectors -- {}` writes", p.display(), p.display());
+    let cases = semantics::parse_vector_file(&text).unwrap();
+    assert_eq!(cases.len(), semantics::edge_cases().len());
+    for (c, r) in &cases {
+        assert_eq!(*r, semantics::reference(c), "{c:?}");
+    }
+    let unguarded = cases.iter().filter(|(c, _)| !c.guarded()).count();
+    assert!(unguarded > 0 && unguarded < cases.len() / 4, "{unguarded} unguarded cases");
+}
+
+#[test]
+fn every_adapter_gives_wgsls_integer_semantics() {
+    let d = blocking::discover();
+    for (a, why) in &d.refused {
+        report(&format!("ops-gpu semantics adapter refused: key={} {} reason={why:?}", a.key(), a.describe()));
+    }
+    if d.found.is_empty() {
+        report("ops-gpu semantics: SKIPPED: no adapter");
+        assert!(std::env::var_os("WAI_OPS_GPU_REQUIRE").is_none(), "WAI_OPS_GPU_REQUIRE is set and no adapter was found");
+        return;
+    }
+    let (_, text) = committed();
+    let mut cases: Vec<Case> = semantics::parse_vector_file(&text).unwrap().into_iter().map(|(c, _)| c).collect();
+    let n_committed = cases.len();
+    cases.extend(semantics::random_cases(4096, 0xA417_5E00));
+    let expect: Vec<u32> = cases.iter().map(semantics::reference).collect();
+    for g in &d.found {
+        let got = blocking::semantics(g, &cases).unwrap_or_else(|e| panic!("{}: {e}", g.adapter().key()));
+        let o = semantics::compare(&cases, &expect, &got);
+        for d in o.guarded_differ.iter().take(20) {
+            println!("MISMATCH {}: {d}", g.adapter().key());
+        }
+        let mut shown = std::collections::BTreeSet::new();
+        for d in &o.unguarded_differ {
+            if shown.insert(d.case.op) {
+                println!("ops-gpu semantics unguarded: key={} {d}", g.adapter().key());
+            }
+        }
+        report(&format!(
+            "ops-gpu semantics: key={} {} committed={n_committed} cases={} guarded={} differ={} unguarded={} unguarded_differ={} unguarded_ops={}",
+            g.adapter().key(),
+            g.adapter().describe(),
+            cases.len(),
+            o.guarded,
+            o.guarded_differ.len(),
+            o.unguarded,
+            o.unguarded_differ.len(),
+            if o.unguarded_ops().is_empty() { "none".to_string() } else { o.unguarded_ops().join(",") },
+        ));
+        assert!(o.guarded_differ.is_empty(), "{}: {} guarded cases differ from WGSL, the first: {}", g.adapter().key(), o.guarded_differ.len(), o.guarded_differ[0]);
+    }
+}
+
+/// A remainder written out by hand, `a - (a / b) * b`, gives `a` where WGSL gives 0 (a
+/// zero divisor): an adapter running that kernel must be refused at discovery, for the
+/// semantics suite, on every adapter.
+#[test]
+fn an_edited_semantics_kernel_is_caught() {
+    let d = blocking::discover();
+    if d.found.is_empty() {
+        println!("ops-gpu semantics mutant: SKIPPED: no adapter");
+        assert!(std::env::var_os("WAI_OPS_GPU_REQUIRE").is_none(), "WAI_OPS_GPU_REQUIRE is set and no adapter was found");
+        return;
+    }
+    let mut src = Sources::shipped();
+    let from = "case 10u: { r = bitcast<u32>(ia % ib); }";
+    assert!(src.semantics.contains(from));
+    src.semantics = src.semantics.replace(from, "case 10u: { r = bitcast<u32>(ia - (ia / ib) * ib); }");
+    let m = blocking::discover_with(&src, "mutant: a remainder by hand");
+    assert!(m.found.is_empty(), "an adapter passed the semantics suite with a remainder written by hand");
+    for (a, why) in &m.refused {
+        assert!(why.contains("WGSL semantics suite") && why.contains("rem_i32"), "{}: refused for another reason: {why}", a.key());
+        report(&format!("ops-gpu semantics mutant: key={} refused: {why}", a.key()));
+    }
+    assert_eq!(m.refused.len(), d.found.len() + d.refused.len());
+}

wai/wai-ops-gpu/tests/wgsl_rules.rs

@@ -0,0 +1,486 @@
+//! The shipped WGSL modules held, on naga 30.0.1's IR (the translator wgpu runs them
+//! through), to the rules the `decode-equivalence` tier on a GPU rests on
+//! (determinism-tiers §5, *Claiming the tier on a GPU*). No adapter is needed: these run
+//! on every leg.
+//!
+//! 1. Each module parses, and validates with every validation flag and no capability:
+//!    so no `f16`, no 64-bit type, no subgroup or quad operation, no extension.
+//! 2. Every entry point is a compute entry point; the globals are exactly the six
+//!    bindings of `wgsl/common.wgsl`, and there is no workgroup or private memory.
+//! 3. Every scalar, in the type arena and in the type of every expression, is `bool`,
+//!    `i32` or `u32`; every literal is one of those. No floating-point value exists.
+//! 4. Every expression is one of a closed list: no derivative, image, ray-query,
+//!    subgroup or cooperative-matrix expression; every built-in function is one of the
+//!    integer functions the semantics suite tests; every conversion is between `bool`,
+//!    `i32` and `u32`.
+//! 5. In the kernels (every module but `semantics`, which exists to evaluate them), every
+//!    shift amount is a literal in 0..=31 or masked to 0..=31 (`e & 31u`), and every
+//!    `clamp` has literal bounds in order: the two integer operations whose WGSL result
+//!    naga 30.0.1 does not produce itself in MSL, HLSL and SPIR-V
+//!    (`semantics::UNGUARDED`) are reached only where every backend defines them.
+//! 6. The only atomic operation is `atomicOr`, its result unused, so no result depends
+//!    on the order in which invocations reach it; there is no barrier.
+//! 7. Every integer operation the kernels use (by operator and by signedness) is one the
+//!    semantics suite runs on every adapter (`src/semantics.rs`).
+//! 8. Each module translates to MSL 2.4, HLSL shader model 5.1 and SPIR-V; the MSL and
+//!    HLSL carry naga's guards for every guarded operation the module uses (signed
+//!    arithmetic through unsigned, `naga_div`, `naga_mod`, `naga_neg`, `naga_abs`), and
+//!    the MSL guards for division vanish when `emit_int_div_checks` is turned off,
+//!    which is what the crate must never do.
+//!
+//! Then each planted violation (a floating-point value, `f16`, a subgroup operation, an
+//! `atomicAdd`, an atomic's result read, a workgroup variable and barrier, an unmasked
+//! shift, a run-time `clamp`) is shown to be refused, by the rule it breaks.
+
+use std::collections::BTreeSet;
+
+use naga::back::{hlsl, msl, spv};
+use naga::valid::{Capabilities, ModuleInfo, ValidationFlags, Validator};
+use naga::{AddressSpace, AtomicFunction, BinaryOperator, Block, Expression, Literal, MathFunction, Module, Scalar, ScalarKind, ShaderStage, Statement, StorageAccess, TypeInner, UnaryOperator};
+use wai_ops_gpu::kernels::Sources;
+use wai_ops_gpu::semantics::Op;
+
+type Check<T = ()> = Result<T, String>;
+
+/// Step 1.
+fn compile(text: &str) -> Check<(Module, ModuleInfo)> {
+    let module = naga::front::wgsl::parse_str(text).map_err(|e| format!("rule 1: parse: {}", e.emit_to_string(text)))?;
+    let info = Validator::new(ValidationFlags::all(), Capabilities::empty()).validate(&module).map_err(|e| format!("rule 1: validation: {}", e.emit_to_string(text)))?;
+    Ok((module, info))
+}
+
+/// The functions and entry points, with their expression-info tables.
+fn functions<'a>(m: &'a Module, info: &'a ModuleInfo) -> Vec<(String, &'a naga::Function, &'a naga::valid::FunctionInfo)> {
+    let mut v: Vec<(String, &naga::Function, &naga::valid::FunctionInfo)> = m.functions.iter().map(|(h, f)| (f.name.clone().unwrap_or_default(), f, &info[h])).collect();
+    for (i, e) in m.entry_points.iter().enumerate() {
+        v.push((e.name.clone(), &e.function, info.get_entry_point(i)));
+    }
+    v
+}
+
+/// Step 2.
+fn interface(m: &Module) -> Check {
+    for e in &m.entry_points {
+        if e.stage != ShaderStage::Compute {
+            return Err(format!("rule 2: entry point `{}` is a {:?} stage", e.name, e.stage));
+        }
+        if e.workgroup_size_overrides.is_some() || e.workgroup_size[1] != 1 || e.workgroup_size[2] != 1 {
+            return Err(format!("rule 2: entry point `{}` has workgroup size {:?}", e.name, e.workgroup_size));
+        }
+    }
+    let mut found = Vec::new();
+    for (_, v) in m.global_variables.iter() {
+        let name = v.name.clone().unwrap_or_default();
+        let space = match v.space {
+            AddressSpace::Uniform => "uniform",
+            AddressSpace::Storage { access } if access == StorageAccess::LOAD => "read",
+            AddressSpace::Storage { access } if access.contains(StorageAccess::LOAD | StorageAccess::STORE) => "read_write",
+            other => return Err(format!("rule 2: global `{name}` is in address space {other:?}")),
+        };
+        let b = v.binding.ok_or(format!("rule 2: global `{name}` has no binding"))?;
+        found.push((b.group, b.binding, space));
+    }
+    found.sort();
+    let want = vec![(0, 0, "uniform"), (0, 1, "read"), (0, 2, "read"), (0, 3, "read"), (0, 4, "read_write"), (0, 5, "read_write")];
+    if found != want {
+        return Err(format!("rule 2: the globals are {found:?}, not the six bindings of common.wgsl"));
+    }
+    Ok(())
+}
+
+fn scalar_ok(s: Scalar) -> bool {
+    s.kind == ScalarKind::Bool || (matches!(s.kind, ScalarKind::Sint | ScalarKind::Uint) && s.width == 4)
+}
+
+fn type_ok(t: &TypeInner) -> Check {
+    match *t {
+        TypeInner::Scalar(s) | TypeInner::Vector { scalar: s, .. } | TypeInner::Atomic(s) | TypeInner::ValuePointer { scalar: s, .. } => {
+            if scalar_ok(s) { Ok(()) } else { Err(format!("{:?} of width {}", s.kind, s.width)) }
+        }
+        TypeInner::Pointer { .. } | TypeInner::Array { .. } | TypeInner::Struct { .. } => Ok(()),
+        ref other => Err(format!("{other:?}")),
+    }
+}
+
+/// Step 3.
+fn types(m: &Module, info: &ModuleInfo) -> Check {
+    for (_, t) in m.types.iter() {
+        type_ok(&t.inner).map_err(|e| format!("rule 3: the type {} is not made of bool, i32 and u32: {e}", t.name.clone().unwrap_or_default()))?;
+    }
+    for (name, f, fi) in functions(m, info) {
+        for (h, e) in f.expressions.iter() {
+            type_ok(fi[h].ty.inner_with(&m.types)).map_err(|err| format!("rule 3: in `{name}`, {e:?} has a type that is not bool, i32 or u32: {err}"))?;
+            if let Expression::Literal(l) = e {
+                literal_ok(l).map_err(|err| format!("rule 3: in `{name}`: {err}"))?;
+            }
+        }
+    }
+    for (_, e) in m.global_expressions.iter() {
+        if let Expression::Literal(l) = e {
+            literal_ok(l).map_err(|err| format!("rule 3: at module scope: {err}"))?;
+        }
+    }
+    Ok(())
+}
+
+fn literal_ok(l: &Literal) -> Check {
+    match l {
+        Literal::U32(_) | Literal::I32(_) | Literal::Bool(_) => Ok(()),
+        other => Err(format!("the literal {other:?}")),
+    }
+}
+
+/// The integer functions the semantics suite tests.
+fn math_ok(f: MathFunction) -> bool {
+    use MathFunction::*;
+    matches!(f, Min | Max | Clamp | Abs | Sign | CountLeadingZeros | CountTrailingZeros | CountOneBits | ReverseBits | FirstLeadingBit | FirstTrailingBit | ExtractBits | InsertBits)
+}
+
+/// The integer value of a literal, or of a constant initialised with one.
+fn literal_value(m: &Module, arena: &naga::Arena<Expression>, h: naga::Handle<Expression>) -> Option<i64> {
+    let lit = |l: &Literal| match *l {
+        Literal::U32(v) => Some(i64::from(v)),
+        Literal::I32(v) => Some(i64::from(v)),
+        _ => None,
+    };
+    match &arena[h] {
+        Expression::Literal(l) => lit(l),
+        Expression::Constant(c) => match &m.global_expressions[m.constants[*c].init] {
+            Expression::Literal(l) => lit(l),
+            _ => None,
+        },
+        _ => None,
+    }
+}
+
+/// A shift amount the kernels may use: a literal in 0..=31, or `e & k` with `k` a literal
+/// in 0..=31 (either side).
+fn amount_bounded(m: &Module, arena: &naga::Arena<Expression>, h: naga::Handle<Expression>) -> bool {
+    if let Some(v) = literal_value(m, arena, h) {
+        return (0..=31).contains(&v);
+    }
+    match arena[h] {
+        Expression::Binary { op: BinaryOperator::And, left, right } => [left, right].iter().any(|&x| literal_value(m, arena, x).is_some_and(|v| (0..=31).contains(&v))),
+        _ => false,
+    }
+}
+
+/// The scalar kind of an expression and whether it is a vector.
+fn kind_of(m: &Module, fi: &naga::valid::FunctionInfo, h: naga::Handle<Expression>) -> (Option<ScalarKind>, bool) {
+    let t = fi[h].ty.inner_with(&m.types);
+    (t.scalar_kind(), matches!(t, TypeInner::Vector { .. }))
+}
+
+/// Steps 4, 5 and 7: the closed list of expressions, the kernel rules on shifts and
+/// `clamp`, and the semantics operation each integer operation is. Returns the
+/// operations used.
+fn expressions(m: &Module, info: &ModuleInfo, kernel: bool) -> Check<BTreeSet<Op>> {
+    let mut used = BTreeSet::new();
+    for (name, f, fi) in functions(m, info) {
+        let arena = &f.expressions;
+        for (h, e) in arena.iter() {
+            let at = |what: String| format!("in `{name}`, {what}");
+            match *e {
+                Expression::Literal(_) | Expression::Constant(_) | Expression::ZeroValue(_) | Expression::Compose { .. } | Expression::Access { .. } | Expression::AccessIndex { .. } | Expression::Splat { .. } | Expression::Swizzle { .. } | Expression::FunctionArgument(_) | Expression::GlobalVariable(_) | Expression::LocalVariable(_) | Expression::Load { .. } | Expression::CallResult(_) | Expression::ArrayLength(_) => {}
+                Expression::Relational { fun, .. } => {
+                    if !matches!(fun, naga::RelationalFunction::All | naga::RelationalFunction::Any) {
+                        return Err(format!("rule 4: {}", at(format!("the relational function {fun:?}"))));
+                    }
+                }
+                // naga gives every atomic call a result; the call statement does not count
+                // as a reference to it, so a count above zero is a read.
+                Expression::AtomicResult { .. } => {
+                    if fi[h].ref_count > 0 {
+                        return Err(format!("rule 6: {}", at("an atomic operation's result is read".into())));
+                    }
+                }
+                Expression::Unary { op, expr } => {
+                    let (k, _) = kind_of(m, fi, expr);
+                    match (op, k) {
+                        (UnaryOperator::Negate, Some(ScalarKind::Sint)) => {
+                            used.insert(Op::NegI32);
+                        }
+                        (UnaryOperator::BitwiseNot, Some(ScalarKind::Sint | ScalarKind::Uint)) => {
+                            used.insert(Op::Not);
+                        }
+                        (UnaryOperator::LogicalNot, Some(ScalarKind::Bool)) | (UnaryOperator::BitwiseNot, Some(ScalarKind::Bool)) => {}
+                        (op, k) => return Err(format!("rule 4: {}", at(format!("the unary {op:?} on {k:?}")))),
+                    }
+                }
+                Expression::Binary { op, left, right } => {
+                    let (k, vector) = kind_of(m, fi, left);
+                    use BinaryOperator as B;
+                    let s = k == Some(ScalarKind::Sint);
+                    let int = matches!(k, Some(ScalarKind::Sint | ScalarKind::Uint));
+                    let v = vector;
+                    let op_used = match op {
+                        B::Add if int && v && !s => Some(Op::AddVec2U32),
+                        B::Multiply if int && v => Some(if s { Op::MulVec2I32 } else { Op::MulVec2U32 }),
+                        B::ShiftRight if v => Some(if s { Op::ShrVec2I32 } else { Op::ShrVec2U32 }),
+                        B::Add if int && !v => Some(if s { Op::AddI32 } else { Op::AddU32 }),
+                        B::Subtract if int && !v => Some(if s { Op::SubI32 } else { Op::SubU32 }),
+                        B::Multiply if int && !v => Some(if s { Op::MulI32 } else { Op::MulU32 }),
+                        B::Divide if int && !v => Some(if s { Op::DivI32 } else { Op::DivU32 }),
+                        B::Modulo if int && !v => Some(if s { Op::RemI32 } else { Op::RemU32 }),
+                        B::ShiftLeft if int && !v => Some(if s { Op::ShlI32 } else { Op::ShlU32 }),
+                        B::ShiftRight if int && !v => Some(if s { Op::ShrI32 } else { Op::ShrU32 }),
+                        B::And if int && !v => Some(Op::And),
+                        B::InclusiveOr if int && !v => Some(Op::Or),
+                        B::ExclusiveOr if int && !v => Some(Op::Xor),
+                        B::Less | B::Greater if int && !v => Some(if s { Op::LtI32 } else { Op::LtU32 }),
+                        B::LessEqual | B::GreaterEqual if int && !v => Some(if s { Op::LeI32 } else { Op::LeU32 }),
+                        B::Equal if int && !v => Some(Op::Eq),
+                        B::NotEqual if int && !v => Some(Op::Ne),
+                        B::LogicalAnd | B::LogicalOr | B::And | B::InclusiveOr | B::Equal | B::NotEqual if k == Some(ScalarKind::Bool) => None,
+                        _ => return Err(format!("rule 7: {}", at(format!("{op:?} on {k:?} (vector {vector}), which the semantics suite does not test")))),
+                    };
+                    if kernel && matches!(op, B::ShiftLeft | B::ShiftRight) && !amount_bounded(m, arena, right) {
+                        return Err(format!("rule 5: {}", at(format!("{op:?} by an amount that is neither a literal in 0..=31 nor masked to it: {:?}", arena[right]))));
+                    }
+                    used.extend(op_used);
+                }
+                Expression::Select { accept, .. } => {
+                    let (k, vector) = kind_of(m, fi, accept);
+                    used.insert(match (k, vector) {
+                        (_, true) => Op::SelectVec2U32,
+                        (Some(ScalarKind::Sint), false) => Op::SelectI32,
+                        _ => Op::SelectU32,
+                    });
+                }
+                Expression::Math { fun, arg, arg1, arg2, .. } => {
+                    if !math_ok(fun) {
+                        return Err(format!("rule 4: {}", at(format!("the built-in function {fun:?}, not an integer function the suite tests"))));
+                    }
+                    let (k, _) = kind_of(m, fi, arg);
+                    let s = k == Some(ScalarKind::Sint);
+                    use MathFunction as F;
+                    if kernel && fun == F::Clamp {
+                        let lo = arg1.and_then(|x| literal_value(m, arena, x));
+                        let hi = arg2.and_then(|x| literal_value(m, arena, x));
+                        match (lo, hi) {
+                            (Some(lo), Some(hi)) if lo <= hi => {}
+                            _ => return Err(format!("rule 5: {}", at("a clamp whose bounds are not literals in order: write min(max(e, low), high)".into()))),
+                        }
+                    }
+                    used.insert(match fun {
+                        F::Min => if s { Op::MinI32 } else { Op::MinU32 },
+                        F::Max => if s { Op::MaxI32 } else { Op::MaxU32 },
+                        F::Clamp => if s { Op::ClampI32 } else { Op::ClampU32 },
+                        F::Abs => Op::AbsI32,
+                        F::Sign => Op::SignI32,
+                        F::CountLeadingZeros => if s { Op::CountLeadingZerosI32 } else { Op::CountLeadingZerosU32 },
+                        F::CountTrailingZeros => Op::CountTrailingZerosU32,
+                        F::CountOneBits => Op::CountOneBitsU32,
+                        F::ReverseBits => Op::ReverseBitsU32,
+                        F::FirstLeadingBit => if s { Op::FirstLeadingBitI32 } else { Op::FirstLeadingBitU32 },
+                        F::FirstTrailingBit => Op::FirstTrailingBitU32,
+                        F::ExtractBits => if s { Op::ExtractBitsI32 } else { Op::ExtractBitsU32 },
+                        F::InsertBits => if s { Op::InsertBitsI32 } else { Op::InsertBitsU32 },
+                        _ => unreachable!(),
+                    });
+                }
+                Expression::As { expr, kind, convert } => {
+                    let (from, _) = kind_of(m, fi, expr);
+                    let ok = |k: Option<ScalarKind>| matches!(k, Some(ScalarKind::Sint | ScalarKind::Uint | ScalarKind::Bool));
+                    if !ok(Some(kind)) || !ok(from) || convert.is_some_and(|w| w != 4) {
+                        return Err(format!("rule 4: {}", at(format!("a conversion from {from:?} to {kind:?} ({convert:?})"))));
+                    }
+                    if convert.is_some() {
+                        used.insert(match (from, kind) {
+                            (Some(ScalarKind::Bool), _) => Op::U32OfBool,
+                            (Some(ScalarKind::Sint), ScalarKind::Uint) => Op::U32OfI32,
+                            (Some(ScalarKind::Uint), ScalarKind::Sint) => Op::I32OfU32,
+                            _ => continue,
+                        });
+                    }
+                }
+                ref other => return Err(format!("rule 4: {}", at(format!("the expression {other:?}")))),
+            }
+        }
+    }
+    Ok(used)
+}
+
+/// Step 6, recursively.
+fn statements(block: &Block, name: &str) -> Check {
+    for s in block.iter() {
+        match s {
+            // Whether its result is read is the expression rule's (rule 6 there).
+            Statement::Atomic { fun, .. } => {
+                if !matches!(fun, AtomicFunction::InclusiveOr) {
+                    return Err(format!("rule 6: in `{name}`, the atomic operation {fun:?}"));
+                }
+            }
+            Statement::ControlBarrier(_) | Statement::MemoryBarrier(_) | Statement::WorkGroupUniformLoad { .. } => return Err(format!("rule 6: in `{name}`, a barrier")),
+            Statement::Block(b) => statements(b, name)?,
+            Statement::If { accept, reject, .. } => {
+                statements(accept, name)?;
+                statements(reject, name)?;
+            }
+            Statement::Switch { cases, .. } => {
+                for c in cases {
+                    statements(&c.body, name)?;
+                }
+            }
+            Statement::Loop { body, continuing, .. } => {
+                statements(body, name)?;
+                statements(continuing, name)?;
+            }
+            Statement::Emit(_) | Statement::Break | Statement::Continue | Statement::Return { .. } | Statement::Store { .. } | Statement::Call { .. } => {}
+            other => return Err(format!("rule 4: in `{name}`, the statement {other:?}")),
+        }
+    }
+    Ok(())
+}
+
+/// What one module was found to use.
+struct Found {
+    module: Module,
+    info: ModuleInfo,
+    used: BTreeSet<Op>,
+}
+
+/// Rules 1–7 on one module's text; `kernel` is false only for the semantics module.
+fn check(text: &str, kernel: bool) -> Check<Found> {
+    let (module, info) = compile(text)?;
+    interface(&module)?;
+    types(&module, &info)?;
+    let used = expressions(&module, &info, kernel)?;
+    for (name, f, _) in functions(&module, &info) {
+        statements(&f.body, &name)?;
+    }
+    Ok(Found { module, info, used })
+}
+
+fn msl_text(f: &Found, int_div_checks: bool) -> Check<String> {
+    let options = msl::Options { lang_version: (2, 4), emit_int_div_checks: int_div_checks, ..Default::default() };
+    msl::write_string(&f.module, &f.info, &options, &msl::PipelineOptions::default()).map(|(s, _)| s).map_err(|e| format!("rule 8: MSL: {e}"))
+}
+
+fn hlsl_text(f: &Found) -> Check<String> {
+    let mut s = String::new();
+    hlsl::Writer::new(&mut s, &hlsl::Options::default(), &hlsl::PipelineOptions::default()).write(&f.module, &f.info, None).map_err(|e| format!("rule 8: HLSL: {e}"))?;
+    Ok(s)
+}
+
+/// Rule 8 on one module: the translations, and the guards each must carry.
+fn translations(name: &str, f: &Found) -> Check {
+    let m = msl_text(f, true)?;
+    let h = hlsl_text(f)?;
+    spv::write_vec(&f.module, &f.info, &spv::Options::default(), None).map_err(|e| format!("rule 8: {name}: SPIR-V: {e}"))?;
+    let u = &f.used;
+    let need = |cond: bool, text: &str, pat: &str, lang: &str| -> Check {
+        if cond && !text.contains(pat) {
+            return Err(format!("rule 8: {name}: the {lang} lacks `{pat}`, the guard for an operation the module uses"));
+        }
+        Ok(())
+    };
+    let signed = u.contains(&Op::AddI32) || u.contains(&Op::SubI32) || u.contains(&Op::MulI32) || u.contains(&Op::MulVec2I32);
+    need(signed, &m, "as_type<uint>(", "MSL")?;
+    need(signed, &h, "asuint(", "HLSL")?;
+    let div = [Op::DivU32, Op::DivI32].iter().any(|o| u.contains(o));
+    let rem = [Op::RemU32, Op::RemI32].iter().any(|o| u.contains(o));
+    for (text, lang) in [(&m, "MSL"), (&h, "HLSL")] {
+        need(div, text, "naga_div(", lang)?;
+        need(rem, text, "naga_mod(", lang)?;
+        need(u.contains(&Op::NegI32), text, "naga_neg(", lang)?;
+        need(u.contains(&Op::AbsI32), text, "naga_abs(", lang)?;
+    }
+    if div || rem {
+        let raw = msl_text(f, false)?;
+        if raw.contains("naga_div(") || raw.contains("naga_mod(") {
+            return Err(format!("rule 8: {name}: the MSL keeps its division guards with emit_int_div_checks off, so this check does not see the switch"));
+        }
+    }
+    Ok(())
+}
+
+#[test]
+fn every_shipped_module_meets_the_rules() {
+    let src = Sources::shipped();
+    let mut all = BTreeSet::new();
+    for (name, text) in src.named() {
+        let kernel = name != "semantics";
+        let f = check(text, kernel).unwrap_or_else(|e| panic!("{name}: {e}"));
+        translations(name, &f).unwrap_or_else(|e| panic!("{e}"));
+        let ops: Vec<&str> = f.used.iter().map(|o| o.name()).collect();
+        println!("ops-gpu wgsl rules: {name}: rules 1-8 hold; integer operations used: {}", ops.join(" "));
+        if kernel {
+            all.extend(f.used);
+        }
+    }
+    let ops: Vec<&str> = all.iter().map(|o| o.name()).collect();
+    println!("ops-gpu wgsl rules: the kernels use {} of the suite's operations: {}", all.len(), ops.join(" "));
+}
+
+/// Rule 7 the other way: the semantics module evaluates every operation of the suite,
+/// so a kernel operation the suite names is one the device is held to.
+#[test]
+fn the_semantics_module_evaluates_every_operation_of_the_suite() {
+    let src = Sources::shipped();
+    let f = check(&src.semantics, false).unwrap();
+    for op in wai_ops_gpu::semantics::ALL {
+        // The conversions u32(i32) and i32(u32) are bit-preserving, as bitcast is; the
+        // comparisons with > and >= are those with < and <= on swapped operands.
+        assert!(f.used.contains(&op), "the semantics module does not evaluate {op:?}");
+    }
+}
+
+/// The crate creates every module through `create_shader_module`, which keeps wgpu's
+/// run-time checks, the division guards among them, on; never through the unchecked
+/// form or with checks of its own choosing.
+#[test]
+fn the_crate_never_turns_the_translators_checks_off() {
+    let dir = std::path::PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("src");
+    let mut seen = 0;
+    for entry in std::fs::read_dir(&dir).unwrap() {
+        let p = entry.unwrap().path();
+        if p.extension().is_some_and(|e| e == "rs") {
+            let text = std::fs::read_to_string(&p).unwrap();
+            for (n, line) in text.lines().enumerate() {
+                let code = line.split("//").next().unwrap_or("");
+                for bad in ["create_shader_module_trusted", "ShaderRuntimeChecks", "create_shader_module_passthrough", "create_shader_module_spirv"] {
+                    assert!(!code.contains(bad), "{}:{}: {bad}", p.display(), n + 1);
+                }
+                seen += code.matches("create_shader_module(").count();
+            }
+        }
+    }
+    assert!(seen >= 1, "no module is created through create_shader_module");
+}
+
+/// Each violation, planted in a shipped module, is refused by the rule it breaks.
+#[test]
+fn each_planted_violation_is_refused() {
+    let src = Sources::shipped();
+    let ops32 = &src.ops32;
+    let anchor = "    let v = A[t];\n    switch P.mode {";
+    assert_eq!(ops32.matches(anchor).count(), 1);
+    let plant = |code: &str| ops32.replacen(anchor, &format!("{code}\n{anchor}"), 1);
+    let cases: Vec<(&str, String, &str)> = vec![
+        ("a floating-point value", plant("    let fl = f32(P.n) * 0.5;\n    O[t] = u32(fl);"), "rule 3"),
+        ("an f16 value", format!("enable f16;\n{}", plant("    let h = f16(A[t]);\n    O[t] = u32(h);")), "rule 1"),
+        ("a subgroup reduction", format!("enable subgroups;\n{}", plant("    O[t] = subgroupAdd(A[t]);")), "rule 1"),
+        ("an atomicAdd", plant("    atomicAdd(&F[P.flag], 1u);"), "rule 6"),
+        ("an atomic's result read", plant("    let old = atomicOr(&F[P.flag], 1u);\n    O[t] = old;"), "rule 6"),
+        ("a workgroup variable and a barrier", format!("var<workgroup> wg: array<u32, 64>;\n{}", plant("    wg[t & 63u] = A[t];\n    workgroupBarrier();")), "rule 2"),
+        ("an unmasked shift", plant("    O[t] = A[t] << P.f;"), "rule 5"),
+        ("a clamp with run-time bounds", plant("    O[t] = clamp(A[t], P.lo_x, P.hi_x);"), "rule 5"),
+        ("a floating-point built-in", plant("    O[t] = u32(sqrt(f32(A[t])));"), "rule 3"),
+    ];
+    for (what, text, rule) in &cases {
+        assert_ne!(text, ops32, "{what}: the plant did not change the text");
+        match check(text, true) {
+            Ok(_) => panic!("{what}: the rules accepted it"),
+            Err(e) => {
+                assert!(e.starts_with(rule), "{what}: refused, but by `{e}`, not {rule}");
+                println!("ops-gpu wgsl rules: planted {what}: refused: {}", e.lines().next().unwrap_or(""));
+            }
+        }
+    }
+    // And the shift rule reads the masks the kernels carry: unmask one of arith.wgsl's.
+    let unmasked = ops32.replacen("(a.y << ((32u - s) & 31u))", "(a.y << (32u - s))", 1);
+    assert_ne!(&unmasked, ops32);
+    let e = check(&unmasked, true).err().expect("an unmasked shift in arith.wgsl was accepted");
+    assert!(e.starts_with("rule 5"), "{e}");
+}

wai/wai-rs/src/lib.rs

@@ -22,6 +22,10 @@
 
 pub mod codecs;
 pub mod container;
+// Platform decode capabilities in dispatch (extensions/platform-decode.md): the query a
+// sink makes of Media Capabilities and WebCodecs for an entry, and the default order those
+// answers and energy budgets give SPEC §4 and §7.1 selection. Pure logic, no dependency.
+pub mod platform_decode;
 // Deduplication totals (extensions/staged-measurement.md §9): objects, distinct
 // contents, referenced and stored bytes, recomputed from the objects. Pure integer,
 // no dependency, so every build (wasm32 included) has it.

wai/wai-rs/src/platform_decode.rs

@@ -0,0 +1,921 @@
+//! Platform decode capabilities in dispatch ([platform-decode](../../extensions/platform-decode.md)):
+//! what a sink asks a web platform about an entry it could decode through the platform's
+//! own decoders — Media Capabilities `decodingInfo()` and WebCodecs `isConfigSupported()` —
+//! and how the answers order and restrict WAI's selection (SPEC §4, §7.1).
+//!
+//! - [`query`]: the configuration a sink asks about for an entry (§2): the codec string,
+//!   the size, the frame rate and the bit rate, read from the payload.
+//! - [`Answers`]: what the platform answered (§3). `supported: false` makes the capability
+//!   unsupported for that object through the platform; `smooth` and `powerEfficient` only
+//!   order the entries. Neither is an energy figure.
+//! - [`EnergyBudget`]: what lets a learned reconstruction rank with power-efficient
+//!   classical decode (§5): a figure of the `OnChipCounter` class measured on a reference
+//!   sink, as [energy-measurement](../../extensions/energy-measurement.md) defines one.
+//! - [`rank`] and [`order`]: the default order (§4), a deployer policy SPEC §7.1 step 2
+//!   allows: power-efficient classical decode and learned decode with a budget first, then
+//!   smooth classical decode, then any other classical decode, then learned decode without
+//!   a budget; within a rank, manifest order.
+//! - [`select`]: SPEC §7.1 selection over a `WAI2` manifest in that order, and
+//!   [`advertise_for_envelope`]: SPEC §4 for a `WAI1` envelope, where the rule decides
+//!   whether a learned capability is advertised for an envelope whose fallback a platform
+//!   decodes power-efficiently.
+//!
+//! Nothing here calls a platform: the answers are inputs, so the order is a function of
+//! the manifest, the answers and the budgets, and its vectors
+//! (`wai/platform-decode-conformance/`) run on every target.
+
+use crate::codecs::{self, ReplicateClass};
+use crate::container::{self, ManifestV2, PinRefusal, Prior, RenditionMeta, Skip};
+
+// ---------------------------------------------------------------- the query (§2)
+
+/// A frame rate, `num / den` frames a second.
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub struct Rate {
+    pub num: u32,
+    pub den: u32,
+}
+
+/// The configuration a sink asks the platform about for one entry (§2).
+#[derive(Debug, Clone, PartialEq, Eq)]
+pub struct Query {
+    /// The codecs parameter (RFC 6381) and WebCodecs codec string.
+    pub codec: String,
+    /// The picture size: the largest frame's, after cropping where the bitstream crops.
+    pub width: u32,
+    pub height: u32,
+    /// The coded size WebCodecs' `codedWidth`/`codedHeight` take.
+    pub coded_width: u32,
+    pub coded_height: u32,
+    /// The frame rate the payload states, or the presentation rate the caller gave where
+    /// it states none; `None` when neither is known.
+    pub framerate: Option<Rate>,
+    /// The mean bit rate over the payload, bits a second, rounded up; `None` without a rate.
+    pub bitrate: Option<u64>,
+    /// Luma bit depth.
+    pub bit_depth: u8,
+}
+
+impl Query {
+    /// The `MediaDecodingConfiguration` for `decodingInfo()` (Media Capabilities §2.1),
+    /// with `type` set to `decoding_type` (`"file"`, `"media-source"` or `"webrtc"`).
+    /// `None` without a frame rate: `framerate` and `bitrate` are required members, and a
+    /// sink does not invent them (§2).
+    pub fn decoding_configuration(&self, decoding_type: &str) -> Option<serde_json::Value> {
+        let rate = self.framerate?;
+        let framerate = f64::from(rate.num) / f64::from(rate.den);
+        Some(serde_json::json!({
+            "type": decoding_type,
+            "video": {
+                "contentType": format!("video/mp4; codecs=\"{}\"", self.codec),
+                "width": self.width,
+                "height": self.height,
+                "bitrate": self.bitrate?,
+                "framerate": framerate,
+            }
+        }))
+    }
+
+    /// The `VideoDecoderConfig` for `VideoDecoder.isConfigSupported()` (WebCodecs §7.6).
+    pub fn decoder_config(&self, hardware_acceleration: HardwareAcceleration) -> serde_json::Value {
+        serde_json::json!({
+            "codec": self.codec,
+            "codedWidth": self.coded_width,
+            "codedHeight": self.coded_height,
+            "hardwareAcceleration": hardware_acceleration.as_str(),
+        })
+    }
+}
+
+/// Why no query is formed for an entry.
+#[derive(Debug, Clone, PartialEq, Eq)]
+pub enum QueryError {
+    /// §2's table does not map this capability to a platform codec string.
+    NotMapped(String),
+    /// The payload does not hold what the query is read from.
+    Malformed(String),
+}
+
+impl std::fmt::Display for QueryError {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        match self {
+            Self::NotMapped(c) => write!(f, "{c}: no platform codec string is mapped for it"),
+            Self::Malformed(w) => write!(f, "the payload: {w}"),
+        }
+    }
+}
+
+impl std::error::Error for QueryError {}
+
+fn malformed<T>(why: impl Into<String>) -> Result<T, QueryError> {
+    Err(QueryError::Malformed(why.into()))
+}
+
+/// The capabilities §2 maps to a platform codec string.
+pub const MAPPED: [&str; 3] = [codecs::CAP_VIDEO_AV1, codecs::CAP_VIDEO_AV1_LOSSLESS, codecs::CAP_VIDEO_HEVC];
+
+/// The query for an entry of `capability` whose payload is `payload` (§2). For a payload
+/// that states no frame rate (an HEVC byte stream), `presentation_rate` is the rate the
+/// sink will present at, and without one the query has no rate.
+pub fn query(capability: &str, payload: &[u8], presentation_rate: Option<Rate>) -> Result<Query, QueryError> {
+    match capability {
+        codecs::CAP_VIDEO_AV1 | codecs::CAP_VIDEO_AV1_LOSSLESS => av1_query(payload),
+        codecs::CAP_VIDEO_HEVC => hevc_query(payload, presentation_rate),
+        other => Err(QueryError::NotMapped(other.to_string())),
+    }
+}
+
+fn mean_bitrate(bytes: usize, frames: u64, rate: Rate) -> Option<u64> {
+    if frames == 0 || rate.den == 0 || rate.num == 0 {
+        return None;
+    }
+    let num = (bytes as u128) * 8 * u128::from(rate.num);
+    let den = u128::from(frames) * u128::from(rate.den);
+    u64::try_from(num.div_ceil(den)).ok()
+}
+
+// AV1 (SPEC §5's framing, AV1 §5, AV1-ISOBMFF §5)
+
+/// A bit reader, most significant bit first.
+struct Bits<'a> {
+    b: &'a [u8],
+    at: usize,
+}
+
+impl Bits<'_> {
+    fn f(&mut self, n: u32) -> Result<u32, QueryError> {
+        let mut v = 0u32;
+        for _ in 0..n {
+            let Some(byte) = self.b.get(self.at >> 3) else {
+                return malformed("a header runs past its unit");
+            };
+            v = (v << 1) | u32::from((byte >> (7 - (self.at & 7))) & 1);
+            self.at += 1;
+        }
+        Ok(v)
+    }
+
+    fn flag(&mut self) -> Result<bool, QueryError> {
+        Ok(self.f(1)? == 1)
+    }
+
+    /// `uvlc()` (AV1 §4.10.3).
+    fn uvlc(&mut self) -> Result<u32, QueryError> {
+        let mut lz = 0u32;
+        while !self.flag()? {
+            lz += 1;
+            if lz > 32 {
+                return malformed("a uvlc() past 32 leading zeros");
+            }
+        }
+        if lz >= 32 {
+            return Ok(u32::MAX);
+        }
+        Ok(self.f(lz)? + ((1u32 << lz) - 1))
+    }
+
+    /// `ue(v)` (H.265 §9.2).
+    fn ue(&mut self) -> Result<u32, QueryError> {
+        let mut lz = 0u32;
+        while !self.flag()? {
+            lz += 1;
+            if lz > 31 {
+                return malformed("an exp-Golomb code past 31 leading zeros");
+            }
+        }
+        Ok(self.f(lz)? + ((1u32 << lz) - 1))
+    }
+}
+
+/// `leb128()` (AV1 §4.10.5).
+fn leb128(b: &[u8]) -> Result<(u64, usize), QueryError> {
+    let mut v = 0u64;
+    for i in 0..8 {
+        let Some(&byte) = b.get(i) else { return malformed("an OBU size runs past the end") };
+        v |= u64::from(byte & 0x7f) << (7 * i);
+        if byte & 0x80 == 0 {
+            return Ok((v, i + 1));
+        }
+    }
+    malformed("an OBU size longer than 8 bytes")
+}
+
+/// The fields of an AV1 sequence header (AV1 §5.5) that the codecs parameter carries.
+#[derive(Debug, Clone, PartialEq, Eq)]
+pub struct Av1Sequence {
+    pub seq_profile: u8,
+    pub seq_level_idx_0: u8,
+    pub seq_tier_0: u8,
+    pub bit_depth: u8,
+    pub mono_chrome: bool,
+    pub subsampling_x: bool,
+    pub subsampling_y: bool,
+    pub chroma_sample_position: u8,
+    /// `color_primaries`, `transfer_characteristics`, `matrix_coefficients` when
+    /// `color_description_present_flag` is 1.
+    pub color_description: Option<(u8, u8, u8)>,
+    pub color_range: bool,
+}
+
+impl Av1Sequence {
+    /// Read a sequence header OBU's payload through `color_config()` (AV1 §5.5.1, §5.5.2).
+    pub fn parse(payload: &[u8]) -> Result<Av1Sequence, QueryError> {
+        let mut r = Bits { b: payload, at: 0 };
+        let seq_profile = r.f(3)? as u8;
+        if seq_profile > 2 {
+            return malformed(format!("seq_profile {seq_profile}"));
+        }
+        let _still_picture = r.flag()?;
+        let reduced = r.flag()?;
+        let (seq_level_idx_0, seq_tier_0);
+        if reduced {
+            seq_level_idx_0 = r.f(5)? as u8;
+            seq_tier_0 = 0;
+        } else {
+            let timing_info_present = r.flag()?;
+            let mut decoder_model_info_present = false;
+            let mut buffer_delay_length = 0;
+            if timing_info_present {
+                r.f(32)?;
+                r.f(32)?;
+                if r.flag()? {
+                    r.uvlc()?;
+                }
+                decoder_model_info_present = r.flag()?;
+                if decoder_model_info_present {
+                    buffer_delay_length = r.f(5)? + 1;
+                    r.f(32)?;
+                    r.f(5)?;
+                    r.f(5)?;
+                }
+            }
+            let initial_display_delay_present = r.flag()?;
+            let count = r.f(5)? + 1;
+            let (mut lvl0, mut tier0) = (0u8, 0u8);
+            for i in 0..count {
+                r.f(12)?;
+                let lvl = r.f(5)? as u8;
+                let tier = if lvl > 7 { r.f(1)? as u8 } else { 0 };
+                if decoder_model_info_present && r.flag()? {
+                    r.f(buffer_delay_length)?;
+                    r.f(buffer_delay_length)?;
+                    r.f(1)?;
+                }
+                if initial_display_delay_present && r.flag()? {
+                    r.f(4)?;
+                }
+                if i == 0 {
+                    (lvl0, tier0) = (lvl, tier);
+                }
+            }
+            seq_level_idx_0 = lvl0;
+            seq_tier_0 = tier0;
+        }
+        let wbits = r.f(4)? + 1;
+        let hbits = r.f(4)? + 1;
+        r.f(wbits)?;
+        r.f(hbits)?;
+        if !reduced && r.flag()? {
+            r.f(4)?;
+            r.f(3)?;
+        }
+        r.f(3)?; // use_128x128_superblock, enable_filter_intra, enable_intra_edge_filter
+        if !reduced {
+            r.f(4)?; // interintra compound, masked compound, warped motion, dual filter
+            let order_hint = r.flag()?;
+            if order_hint {
+                r.f(2)?; // jnt_comp, ref_frame_mvs
+            }
+            let force_sct = if r.flag()? { 2 } else { r.f(1)? };
+            if force_sct > 0 && !r.flag()? {
+                r.f(1)?;
+            }
+            if order_hint {
+                r.f(3)?;
+            }
+        }
+        r.f(3)?; // superres, cdef, restoration
+        // color_config()
+        let high_bitdepth = r.flag()?;
+        let twelve_bit = seq_profile == 2 && high_bitdepth && r.flag()?;
+        let bit_depth = match (high_bitdepth, twelve_bit) {
+            (false, _) => 8,
+            (true, false) => 10,
+            (true, true) => 12,
+        };
+        let mono_chrome = seq_profile != 1 && r.flag()?;
+        let color_description = if r.flag()? { Some((r.f(8)? as u8, r.f(8)? as u8, r.f(8)? as u8)) } else { None };
+        // CP_BT_709, TC_SRGB, MC_IDENTITY.
+        let srgb = color_description == Some((1, 13, 0));
+        let (sx, sy, csp, color_range);
+        if mono_chrome {
+            color_range = r.flag()?;
+            (sx, sy, csp) = (true, true, 0);
+        } else if srgb {
+            (sx, sy, csp, color_range) = (false, false, 0, true);
+        } else {
+            color_range = r.flag()?;
+            let (x, y) = if seq_profile == 0 {
+                (true, true)
+            } else if seq_profile == 1 {
+                (false, false)
+            } else if bit_depth == 12 {
+                let x = r.flag()?;
+                (x, if x { r.flag()? } else { false })
+            } else {
+                (true, false)
+            };
+            (sx, sy) = (x, y);
+            csp = if x && y { r.f(2)? as u8 } else { 0 };
+        }
+        Ok(Av1Sequence {
+            seq_profile,
+            seq_level_idx_0,
+            seq_tier_0,
+            bit_depth,
+            mono_chrome,
+            subsampling_x: sx,
+            subsampling_y: sy,
+            chroma_sample_position: csp,
+            color_description,
+            color_range,
+        })
+    }
+
+    /// The codecs parameter string (AV1-ISOBMFF §5): every field, then the trailing part
+    /// omitted where it is exactly the default `.0.110.01.01.01.0`. Without a colour
+    /// description the primaries, transfer and matrix are written as the defaults, 01.
+    pub fn codec_string(&self) -> String {
+        let head = format!("av01.{}.{:02}{}.{:02}", self.seq_profile, self.seq_level_idx_0, if self.seq_tier_0 == 1 { 'H' } else { 'M' }, self.bit_depth);
+        let (cp, tc, mc) = self.color_description.unwrap_or((1, 1, 1));
+        let third = if self.subsampling_x && self.subsampling_y { self.chroma_sample_position } else { 0 };
+        let tail = format!(
+            ".{}.{}{}{}.{cp:02}.{tc:02}.{mc:02}.{}",
+            u8::from(self.mono_chrome),
+            u8::from(self.subsampling_x),
+            u8::from(self.subsampling_y),
+            third,
+            u8::from(self.color_range)
+        );
+        if tail == ".0.110.01.01.01.0" { head } else { head + &tail }
+    }
+}
+
+/// SPEC §5's framing of `wai.video.av1`: `n`, `h`, `w`, `fps_num`, `fps_den`, then each
+/// temporal unit with its length. Returns the five values and the first temporal unit.
+fn av1_framing(p: &[u8]) -> Result<([u32; 5], &[u8]), QueryError> {
+    if p.len() < 24 {
+        return malformed("shorter than the AV1 framing's header and one unit length");
+    }
+    let w = |i: usize| u32::from_le_bytes([p[4 * i], p[4 * i + 1], p[4 * i + 2], p[4 * i + 3]]);
+    let head = [w(0), w(1), w(2), w(3), w(4)];
+    let len = w(5) as usize;
+    let Some(tu) = p.get(24..24 + len) else { return malformed("the first temporal unit runs past the payload") };
+    Ok((head, tu))
+}
+
+fn av1_query(p: &[u8]) -> Result<Query, QueryError> {
+    let ([n, h, w, fps_num, fps_den], tu) = av1_framing(p)?;
+    if n == 0 || h == 0 || w == 0 {
+        return malformed("an AV1 payload of no frames, or of a zero size");
+    }
+    // The first temporal unit's sequence header OBU (AV1 §5.3).
+    let mut at = 0usize;
+    let seq = loop {
+        let Some(&header) = tu.get(at) else { return malformed("the first temporal unit holds no sequence header") };
+        let obu_type = (header >> 3) & 0x0f;
+        let ext = header & 0x04 != 0;
+        if header & 0x02 == 0 {
+            return malformed("an OBU without its size field");
+        }
+        let hl = if ext { 2 } else { 1 };
+        let (size, sl) = leb128(tu.get(at + hl..).unwrap_or(&[]))?;
+        let start = at + hl + sl;
+        let Some(body) = usize::try_from(size).ok().and_then(|s| tu.get(start..start.checked_add(s)?)) else {
+            return malformed("an OBU runs past its temporal unit");
+        };
+        if obu_type == 1 {
+            break Av1Sequence::parse(body)?;
+        }
+        at = start + body.len();
+    };
+    let rate = (fps_num > 0 && fps_den > 0).then_some(Rate { num: fps_num, den: fps_den });
+    Ok(Query {
+        codec: seq.codec_string(),
+        width: w,
+        height: h,
+        coded_width: w,
+        coded_height: h,
+        framerate: rate,
+        bitrate: rate.and_then(|r| mean_bitrate(p.len(), u64::from(n), r)),
+        bit_depth: seq.bit_depth,
+    })
+}
+
+// HEVC (H.265 Annex B, §7.3.2.2, §7.3.3; ISO/IEC 14496-15 Annex E.3)
+
+/// The NAL units of an Annex B byte stream (H.265 Annex B), each without its start code.
+fn annex_b_units(s: &[u8]) -> Vec<&[u8]> {
+    let mut starts = Vec::new();
+    let mut i = 0;
+    while i + 3 <= s.len() {
+        if s[i] == 0 && s[i + 1] == 0 && s[i + 2] == 1 {
+            starts.push(i + 3);
+            i += 3;
+        } else {
+            i += 1;
+        }
+    }
+    let mut out = Vec::with_capacity(starts.len());
+    for (k, &st) in starts.iter().enumerate() {
+        let mut end = starts.get(k + 1).map_or(s.len(), |&n| n - 3);
+        // The zero_byte before a four-byte start code, and trailing_zero_8bits.
+        while end > st && s[end - 1] == 0 {
+            end -= 1;
+        }
+        out.push(&s[st..end]);
+    }
+    out
+}
+
+/// A NAL unit's RBSP: the two-byte header dropped, and each emulation-prevention byte
+/// (`0x000003`) removed (H.265 §7.3.1.1).
+fn rbsp(nal: &[u8]) -> Vec<u8> {
+    let mut out = Vec::with_capacity(nal.len());
+    let mut zeros = 0;
+    for &b in nal.iter().skip(2) {
+        if zeros >= 2 && b == 3 {
+            zeros = 0;
+            continue;
+        }
+        zeros = if b == 0 { zeros + 1 } else { 0 };
+        out.push(b);
+    }
+    out
+}
+
+/// The fields of an HEVC sequence parameter set that the codecs parameter and the query
+/// carry.
+#[derive(Debug, Clone, PartialEq, Eq)]
+pub struct HevcSps {
+    pub profile_space: u8,
+    pub tier: u8,
+    pub profile_idc: u8,
+    /// `general_profile_compatibility_flag[j]` at bit `j`.
+    pub compatibility: u32,
+    /// The 48 bits from `general_progressive_source_flag` on, as six bytes.
+    pub constraints: [u8; 6],
+    pub level_idc: u8,
+    pub chroma_format_idc: u32,
+    pub width: u32,
+    pub height: u32,
+    pub cropped_width: u32,
+    pub cropped_height: u32,
+    pub bit_depth: u8,
+}
+
+impl HevcSps {
+    /// Read an SPS NAL unit (type 33), header included, through `bit_depth_luma_minus8`.
+    pub fn parse(nal: &[u8]) -> Result<HevcSps, QueryError> {
+        if nal.len() < 3 || (nal[0] >> 1) & 0x3f != 33 {
+            return malformed("not an SPS NAL unit");
+        }
+        let b = rbsp(nal);
+        let mut r = Bits { b: &b, at: 0 };
+        r.f(4)?; // sps_video_parameter_set_id
+        let max_sub_layers_minus1 = r.f(3)?;
+        r.f(1)?; // sps_temporal_id_nesting_flag
+        // profile_tier_level(1, sps_max_sub_layers_minus1), H.265 §7.3.3
+        let profile_space = r.f(2)? as u8;
+        let tier = r.f(1)? as u8;
+        let profile_idc = r.f(5)? as u8;
+        let mut compatibility = 0u32;
+        for j in 0..32 {
+            compatibility |= r.f(1)? << j;
+        }
+        let mut constraints = [0u8; 6];
+        for c in &mut constraints {
+            *c = r.f(8)? as u8;
+        }
+        let level_idc = r.f(8)? as u8;
+        let mut present = Vec::new();
+        for _ in 0..max_sub_layers_minus1 {
+            present.push((r.flag()?, r.flag()?));
+        }
+        if max_sub_layers_minus1 > 0 {
+            for _ in max_sub_layers_minus1..8 {
+                r.f(2)?;
+            }
+        }
+        for (profile, level) in present {
+            if profile {
+                for _ in 0..11 {
+                    r.f(8)?; // 88 bits
+                }
+            }
+            if level {
+                r.f(8)?;
+            }
+        }
+        r.ue()?; // sps_seq_parameter_set_id
+        let chroma_format_idc = r.ue()?;
+        if chroma_format_idc > 3 {
+            return malformed(format!("chroma_format_idc {chroma_format_idc}"));
+        }
+        let separate = chroma_format_idc == 3 && r.flag()?;
+        let width = r.ue()?;
+        let height = r.ue()?;
+        let (mut cw, mut ch) = (width, height);
+        if r.flag()? {
+            let (l, rr, t, bt) = (r.ue()?, r.ue()?, r.ue()?, r.ue()?);
+            // SubWidthC and SubHeightC (H.265 Table 6-1); 1 for 4:4:4 and for separate planes.
+            let (sw, sh) = match (chroma_format_idc, separate) {
+                (1, _) => (2, 2),
+                (2, _) => (2, 1),
+                _ => (1, 1),
+            };
+            cw = width.checked_sub(sw * (l + rr)).ok_or(QueryError::Malformed("a conformance window wider than the picture".into()))?;
+            ch = height.checked_sub(sh * (t + bt)).ok_or(QueryError::Malformed("a conformance window taller than the picture".into()))?;
+        }
+        let bit_depth = r.ue()? + 8;
+        if bit_depth > 16 {
+            return malformed(format!("bit depth {bit_depth}"));
+        }
+        Ok(HevcSps { profile_space, tier, profile_idc, compatibility, constraints, level_idc, chroma_format_idc, width, height, cropped_width: cw, cropped_height: ch, bit_depth: bit_depth as u8 })
+    }
+
+    /// The codecs parameter (ISO/IEC 14496-15 Annex E.3, as WebCodecs' HEVC registration
+    /// cites it): `hev1.`, the profile space as nothing, `A`, `B` or `C` followed by the
+    /// profile in decimal; the compatibility flags as one hexadecimal value with flag `j`
+    /// at bit `j`; `L` or `H` and the level in decimal; then each constraint byte in
+    /// hexadecimal, trailing zero bytes omitted. `hev1` because an Annex B stream carries
+    /// its parameter sets in band.
+    pub fn codec_string(&self) -> String {
+        let space = ["", "A", "B", "C"][usize::from(self.profile_space & 3)];
+        let mut s = format!("hev1.{space}{}.{:X}.{}{}", self.profile_idc, self.compatibility, if self.tier == 1 { 'H' } else { 'L' }, self.level_idc);
+        let last = self.constraints.iter().rposition(|&c| c != 0);
+        if let Some(k) = last {
+            for c in &self.constraints[..=k] {
+                s.push_str(&format!(".{c:X}"));
+            }
+        }
+        s
+    }
+}
+
+fn hevc_query(p: &[u8], presentation_rate: Option<Rate>) -> Result<Query, QueryError> {
+    let units = annex_b_units(p);
+    let sps = units.iter().find(|u| u.len() >= 2 && (u[0] >> 1) & 0x3f == 33).ok_or(QueryError::Malformed("no SPS NAL unit".into()))?;
+    let sps = HevcSps::parse(sps)?;
+    // A picture begins at each VCL NAL unit (types 0..=31) whose
+    // first_slice_segment_in_pic_flag, the slice header's first bit, is 1.
+    let pictures = units.iter().filter(|u| u.len() >= 3 && (u[0] >> 1) & 0x3f <= 31 && u[2] & 0x80 != 0).count() as u64;
+    Ok(Query {
+        codec: sps.codec_string(),
+        width: sps.cropped_width,
+        height: sps.cropped_height,
+        coded_width: sps.width,
+        coded_height: sps.height,
+        framerate: presentation_rate,
+        bitrate: presentation_rate.and_then(|r| mean_bitrate(p.len(), pictures, r)),
+        bit_depth: sps.bit_depth,
+    })
+}
+
+// ---------------------------------------------------------------- the answers (§3)
+
+/// WebCodecs `HardwareAcceleration` (WebCodecs §7.9): a hint the user agent may ignore.
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub enum HardwareAcceleration {
+    NoPreference,
+    PreferHardware,
+    PreferSoftware,
+}
+
+impl HardwareAcceleration {
+    pub fn as_str(self) -> &'static str {
+        match self {
+            Self::NoPreference => "no-preference",
+            Self::PreferHardware => "prefer-hardware",
+            Self::PreferSoftware => "prefer-software",
+        }
+    }
+
+    pub fn parse(s: &str) -> Option<HardwareAcceleration> {
+        [Self::NoPreference, Self::PreferHardware, Self::PreferSoftware].into_iter().find(|h| h.as_str() == s)
+    }
+}
+
+/// `MediaCapabilitiesInfo` (Media Capabilities §2.2) for the entry's query.
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub struct DecodingInfo {
+    pub supported: bool,
+    pub smooth: bool,
+    /// The user agent's judgement that the decode's power draw is optimal, its definition
+    /// left to the user agent (§2.2): an ordering signal, never an energy figure.
+    pub power_efficient: bool,
+}
+
+/// What the platform answered about one entry, through the path the sink decodes it by.
+#[derive(Debug, Clone, PartialEq, Eq, Default)]
+pub struct Answers {
+    /// `decodingInfo()` for the entry's query.
+    pub decoding_info: Option<DecodingInfo>,
+    /// `VideoDecoder.isConfigSupported()` for the entry's query with the hardware
+    /// acceleration the sink will configure: `(that preference, supported)`. Present when
+    /// the sink decodes the entry through WebCodecs.
+    pub webcodecs: Option<(HardwareAcceleration, bool)>,
+    /// Whether the sink holds a decoder of its own for the capability, apart from the
+    /// platform's (a WebAssembly AV1 decoder, say), which it uses when the platform's
+    /// path is unsupported.
+    pub own_decoder: bool,
+}
+
+// ---------------------------------------------------------------- the budget (§5)
+
+/// An energy budget for a learned reconstruction (§5): the energy-measurement extension
+/// defines it and how it is carried and checked; the order reads only these members.
+#[derive(Debug, Clone, PartialEq, Eq)]
+pub struct EnergyBudget {
+    /// The acquisition class (energy-measurement §2), by its name.
+    pub class: String,
+    /// The figure: microjoules for one unit of decode.
+    pub joules_micro: u64,
+    /// What one unit is: `frame`, `picture` or `second` of output.
+    pub per: String,
+    /// The declared standard uncertainty, parts per million (energy-measurement §2.3).
+    pub relative_ppm: u64,
+    /// The integration window, microseconds (energy-measurement §2.3).
+    pub window_us: u64,
+    /// The reference sink it was measured on.
+    pub reference_sink: String,
+}
+
+/// Why a budget does not let a learned entry rank with power-efficient classical decode.
+#[derive(Debug, Clone, PartialEq, Eq)]
+pub enum BudgetRefusal {
+    /// Its class is not `OnChipCounter`.
+    Class(String),
+    /// A zero figure, uncertainty or window, or no reference sink: not a measurement.
+    NotMeasured,
+    /// A unit other than `frame`, `picture` or `second`.
+    Unit(String),
+    /// Above the ceiling the deployer set for that unit.
+    Ceiling { joules_micro: u64, ceiling: u64 },
+}
+
+impl EnergyBudget {
+    /// Whether the budget lets its entry rank first (§5): of the `OnChipCounter` class,
+    /// measured (a figure, an uncertainty and a window, each non-zero, on a named reference
+    /// sink), per a unit this extension reads, and, where the deployer sets a ceiling for
+    /// that unit, at or below it.
+    pub fn qualifies(&self, ceiling: Option<u64>) -> Result<(), BudgetRefusal> {
+        if self.class != "OnChipCounter" {
+            return Err(BudgetRefusal::Class(self.class.clone()));
+        }
+        if self.joules_micro == 0 || self.relative_ppm == 0 || self.window_us == 0 || self.reference_sink.is_empty() {
+            return Err(BudgetRefusal::NotMeasured);
+        }
+        if !matches!(self.per.as_str(), "frame" | "picture" | "second") {
+            return Err(BudgetRefusal::Unit(self.per.clone()));
+        }
+        match ceiling {
+            Some(c) if self.joules_micro > c => Err(BudgetRefusal::Ceiling { joules_micro: self.joules_micro, ceiling: c }),
+            _ => Ok(()),
+        }
+    }
+}
+
+// ---------------------------------------------------------------- the order (§4)
+
+/// Whether an entry's capability is decoded by a learned reconstruction (§4.1): its
+/// registry class is a neural, generative or behavioural one. A capability this revision
+/// does not register is ranked as learned: nothing says it is classical.
+pub fn is_learned(capability: &str) -> bool {
+    use ReplicateClass::*;
+    !matches!(codecs::capability_replicate_class(capability), Some(StandardDefined | StandardTolerance | StandardUnclassified))
+}
+
+/// An entry's rank under the default order (§4.2); lower ranks are tried first.
+#[derive(Debug, Clone, Copy, PartialEq, Eq, PartialOrd, Ord, Hash)]
+pub enum Rank {
+    /// Classical decode the platform marks supported, smooth and power-efficient; or a
+    /// learned reconstruction whose budget qualifies (§5).
+    First = 0,
+    /// Classical decode the platform marks supported and smooth, not power-efficient.
+    Smooth = 1,
+    /// Any other classical decode the sink can make: not smooth, or with no platform
+    /// answer (its own decoder).
+    Classical = 2,
+    /// A learned reconstruction without a qualifying budget.
+    LearnedWithoutBudget = 3,
+}
+
+impl Rank {
+    pub fn as_str(self) -> &'static str {
+        match self {
+            Self::First => "first",
+            Self::Smooth => "smooth",
+            Self::Classical => "classical",
+            Self::LearnedWithoutBudget => "learned-without-budget",
+        }
+    }
+}
+
+/// One entry as the order reads it.
+#[derive(Debug, Clone, Copy)]
+pub struct Entry<'a> {
+    pub capability: &'a str,
+    /// The platform's answers, for a classical entry the sink could decode through the
+    /// platform; `None` when it asks nothing (a learned entry, or no platform).
+    pub answers: Option<&'a Answers>,
+    pub budget: Option<&'a EnergyBudget>,
+}
+
+/// The rules the order applies; [`Rules::DEFAULT`] is §4's. A test turns one off to show
+/// the vectors catch it.
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub struct Rules {
+    /// §4.2: a learned entry without a qualifying budget ranks below every classical entry.
+    pub learned_below_classical: bool,
+    /// §3: `supported: false` makes the platform path unsupported.
+    pub unsupported_is_unsupported: bool,
+    /// §5: a budget qualifies only as [`EnergyBudget::qualifies`] says.
+    pub budget_must_qualify: bool,
+    /// The deployer's ceiling on a budget's figure, if any.
+    pub budget_ceiling: Option<u64>,
+}
+
+impl Rules {
+    pub const DEFAULT: Rules = Rules { learned_below_classical: true, unsupported_is_unsupported: true, budget_must_qualify: true, budget_ceiling: None };
+}
+
+/// An entry's rank, or `Err` when the sink cannot decode it at all: the platform answered
+/// `supported: false` for the path it decodes by and it holds no decoder of its own.
+pub fn rank(e: &Entry<'_>, rules: &Rules) -> Result<Rank, Skip> {
+    if is_learned(e.capability) {
+        let qualified = e.budget.is_some_and(|b| !rules.budget_must_qualify || b.qualifies(rules.budget_ceiling).is_ok());
+        return Ok(if qualified || !rules.learned_below_classical { Rank::First } else { Rank::LearnedWithoutBudget });
+    }
+    let Some(a) = e.answers else { return Ok(Rank::Classical) };
+    let platform_says_no = a.decoding_info.is_some_and(|d| !d.supported) || a.webcodecs.is_some_and(|(_, s)| !s);
+    if platform_says_no && rules.unsupported_is_unsupported {
+        return if a.own_decoder { Ok(Rank::Classical) } else { Err(Skip::NotEligible) };
+    }
+    Ok(match a.decoding_info {
+        Some(d) if d.smooth && d.power_efficient => Rank::First,
+        Some(d) if d.smooth => Rank::Smooth,
+        _ => Rank::Classical,
+    })
+}
+
+/// The default order over `entries` (§4.2): their indices by rank, manifest order within a
+/// rank, and those the sink cannot decode with their reason (`missing_capability`).
+pub fn order(entries: &[Entry<'_>], rules: &Rules) -> (Vec<usize>, Vec<(usize, Skip)>) {
+    let mut ranked: Vec<(Rank, usize)> = Vec::new();
+    let mut out = Vec::new();
+    for (i, e) in entries.iter().enumerate() {
+        match rank(e, rules) {
+            Ok(r) => ranked.push((r, i)),
+            Err(s) => out.push((i, s)),
+        }
+    }
+    ranked.sort();
+    (ranked.into_iter().map(|(_, i)| i).collect(), out)
+}
+
+/// What [`select`] gives: the picked index, what `verify` returned for it and the entries
+/// skipped before it; or every entry's index and reason.
+pub type Selection<V> = Result<(usize, Option<V>, Vec<(usize, Skip)>), Vec<(usize, Skip)>>;
+
+/// SPEC §7.1 selection in the default order: `answers` and `budgets` give each entry's,
+/// by index; `advertises` is the sink's capability test; `digest_ok` and `verify` are
+/// [`container::select_entries`]'s. An entry the platform cannot decode is skipped as
+/// `missing_capability`, and an entry the order puts later is tried later; the digest and
+/// pin checks are SPEC §7.1's, unchanged.
+#[allow(clippy::too_many_arguments)]
+pub fn select<V>(
+    manifest: &ManifestV2,
+    answers: impl Fn(usize, &RenditionMeta) -> Option<Answers>,
+    budgets: impl Fn(usize, &RenditionMeta) -> Option<EnergyBudget>,
+    advertises: impl Fn(&str) -> bool,
+    rules: &Rules,
+    digest_ok: impl Fn(usize, &RenditionMeta) -> bool,
+    verify: impl FnMut(&str, &Prior) -> Result<V, PinRefusal>,
+) -> Selection<V> {
+    let metas = &manifest.renditions;
+    let held_answers: Vec<Option<Answers>> = metas.iter().enumerate().map(|(i, m)| answers(i, m)).collect();
+    let held_budgets: Vec<Option<EnergyBudget>> = metas.iter().enumerate().map(|(i, m)| budgets(i, m)).collect();
+    let entries: Vec<Entry<'_>> = metas
+        .iter()
+        .enumerate()
+        .map(|(i, m)| Entry { capability: m.capability.as_str(), answers: held_answers[i].as_ref(), budget: held_budgets[i].as_ref() })
+        .collect();
+    let (mut ord, refused) = order(&entries, rules);
+    // The refused entries are still walked, last, so their reason is reported in place.
+    ord.extend(refused.iter().map(|(i, _)| *i));
+    let refused_at: Vec<usize> = refused.iter().map(|(i, _)| *i).collect();
+    container::select_entries(
+        manifest,
+        &ord,
+        |i, m| {
+            if refused_at.contains(&i) || !advertises(&m.capability) {
+                Err(Skip::NotEligible)
+            } else {
+                Ok(())
+            }
+        },
+        digest_ok,
+        verify,
+    )
+}
+
+/// SPEC §4 for a `WAI1` envelope under the default order: whether the sink advertises
+/// `capability` for this envelope. A learned capability without a qualifying budget is
+/// not advertised for an envelope whose `fallback` is a classical capability the platform
+/// marks supported, smooth and power-efficient, so step 5 decodes the fallback; anything
+/// else is advertised as the sink otherwise would (`advertised`).
+pub fn advertise_for_envelope(capability: &Entry<'_>, fallback: Option<&Entry<'_>>, advertised: bool, rules: &Rules) -> bool {
+    if !advertised {
+        return false;
+    }
+    let Some(fb) = fallback else { return true };
+    match (rank(capability, rules), rank(fb, rules)) {
+        (Ok(Rank::LearnedWithoutBudget), Ok(Rank::First)) => is_learned(fb.capability),
+        _ => true,
+    }
+}
+
+// ---------------------------------------------------------------- JSON forms
+
+/// [`Answers`] from the JSON the vectors and the browser binding use:
+/// `{"decodingInfo": {"supported", "smooth", "powerEfficient"}, "webcodecs":
+/// {"hardwareAcceleration", "supported"}, "ownDecoder"}`, each member optional.
+pub fn answers_from_json(v: &serde_json::Value) -> Result<Answers, String> {
+    let flag = |o: &serde_json::Value, k: &str| o.get(k).and_then(serde_json::Value::as_bool).ok_or(format!("answers: no boolean {k}"));
+    let decoding_info = match v.get("decodingInfo") {
+        None | Some(serde_json::Value::Null) => None,
+        Some(d) => Some(DecodingInfo { supported: flag(d, "supported")?, smooth: flag(d, "smooth")?, power_efficient: flag(d, "powerEfficient")? }),
+    };
+    let webcodecs = match v.get("webcodecs") {
+        None | Some(serde_json::Value::Null) => None,
+        Some(w) => {
+            let h = w.get("hardwareAcceleration").and_then(serde_json::Value::as_str).and_then(HardwareAcceleration::parse).ok_or("answers: webcodecs.hardwareAcceleration is not a HardwareAcceleration")?;
+            Some((h, flag(w, "supported")?))
+        }
+    };
+    Ok(Answers { decoding_info, webcodecs, own_decoder: v.get("ownDecoder").and_then(serde_json::Value::as_bool).unwrap_or(false) })
+}
+
+/// [`EnergyBudget`] from its JSON form: `class`, `joules_micro`, `per`, `relative_ppm`,
+/// `window_us`, `reference_sink`.
+pub fn budget_from_json(v: &serde_json::Value) -> Result<EnergyBudget, String> {
+    let s = |k: &str| v.get(k).and_then(serde_json::Value::as_str).map(str::to_string).ok_or(format!("budget: no string {k}"));
+    let n = |k: &str| v.get(k).and_then(serde_json::Value::as_u64).ok_or(format!("budget: no integer {k}"));
+    Ok(EnergyBudget { class: s("class")?, joules_micro: n("joules_micro")?, per: s("per")?, relative_ppm: n("relative_ppm")?, window_us: n("window_us")?, reference_sink: s("reference_sink")? })
+}
+
+/// A query's JSON form for a caller that asks a platform: its members, with the
+/// `decodingInfo()` configuration (type `media-source`, or `null` without a rate) and the
+/// `isConfigSupported()` configurations with `prefer-hardware` and `no-preference`.
+pub fn query_json(q: &Query) -> serde_json::Value {
+    serde_json::json!({
+        "codec": q.codec,
+        "width": q.width,
+        "height": q.height,
+        "codedWidth": q.coded_width,
+        "codedHeight": q.coded_height,
+        "framerate": q.framerate.map(|r| serde_json::json!([r.num, r.den])),
+        "bitrate": q.bitrate,
+        "bitDepth": q.bit_depth,
+        "decodingConfiguration": q.decoding_configuration("media-source"),
+        "decoderConfigPreferHardware": q.decoder_config(HardwareAcceleration::PreferHardware),
+        "decoderConfigNoPreference": q.decoder_config(HardwareAcceleration::NoPreference),
+    })
+}
+
+/// The default order over a `WAI2` manifest's entries as JSON, from each entry's answers
+/// and budget (arrays aligned with the entries, `null` where none): `ranks` (each rank's
+/// name, `null` for an entry the sink cannot decode), `order`, `refused` (`[index,
+/// reason]`), and `picked`, the entry SPEC §7.1 picks at a sink that advertises every
+/// capability, holds no parameter set and checks no digest.
+pub fn order_json(manifest: &ManifestV2, answers: &[Option<Answers>], budgets: &[Option<EnergyBudget>], rules: &Rules) -> serde_json::Value {
+    let n = manifest.renditions.len();
+    let a = |i: usize| answers.get(i).cloned().flatten();
+    let b = |i: usize| budgets.get(i).cloned().flatten();
+    let held_a: Vec<Option<Answers>> = (0..n).map(a).collect();
+    let held_b: Vec<Option<EnergyBudget>> = (0..n).map(b).collect();
+    let entries: Vec<Entry<'_>> = manifest.renditions.iter().enumerate().map(|(i, m)| Entry { capability: m.capability.as_str(), answers: held_a[i].as_ref(), budget: held_b[i].as_ref() }).collect();
+    let ranks: Vec<serde_json::Value> = entries.iter().map(|e| rank(e, rules).map_or(serde_json::Value::Null, |r| r.as_str().into())).collect();
+    let (ord, refused) = order(&entries, rules);
+    let picked = select(manifest, |i, _| a(i), |i, _| b(i), |_| true, rules, |_, _| true, container::no_parameter_sets()).ok().map(|(i, _, _)| i);
+    serde_json::json!({
+        "ranks": ranks,
+        "order": ord,
+        "refused": refused.iter().map(|(i, s)| serde_json::json!([i, s.code()])).collect::<Vec<_>>(),
+        "picked": picked,
+    })
+}
+
+#[cfg(test)]
+mod tests;

wai/wai-rs/src/platform_decode/tests.rs

@@ -0,0 +1,559 @@
+//! The queries against real AV1 floors and constructed sequence headers and SPSs, and the
+//! default order against `wai/platform-decode-conformance/order.json`; then each rule of
+//! the order, turned off, is shown to be caught by a vector.
+
+use super::*;
+use serde_json::Value;
+
+const ORDER: &str = include_str!("../../../platform-decode-conformance/order.json");
+const QUERIES: &str = include_str!("../../../platform-decode-conformance/queries.json");
+
+// ---------------------------------------------------------------- a bit writer
+
+/// Bits, most significant first, as the AV1 and H.265 headers are written.
+#[derive(Default)]
+struct W {
+    bits: Vec<bool>,
+}
+
+impl W {
+    fn f(&mut self, n: u32, v: u64) -> &mut Self {
+        for k in (0..n).rev() {
+            self.bits.push((v >> k) & 1 == 1);
+        }
+        self
+    }
+    fn ue(&mut self, v: u32) -> &mut Self {
+        let x = u64::from(v) + 1;
+        let len = 64 - x.leading_zeros();
+        self.f(len - 1, 0).f(len, x)
+    }
+    /// The bytes, padded with `trailing_bits` (a 1, then zeros).
+    fn bytes(&self) -> Vec<u8> {
+        let mut b = self.bits.clone();
+        b.push(true);
+        while b.len() % 8 != 0 {
+            b.push(false);
+        }
+        b.chunks(8).map(|c| c.iter().fold(0u8, |a, &x| (a << 1) | u8::from(x))).collect()
+    }
+}
+
+/// An RBSP with emulation-prevention bytes inserted (H.265 §7.4.2).
+fn escape(rbsp: &[u8]) -> Vec<u8> {
+    let mut out = Vec::new();
+    let mut zeros = 0;
+    for &b in rbsp {
+        if zeros >= 2 && b <= 3 {
+            out.push(3);
+            zeros = 0;
+        }
+        zeros = if b == 0 { zeros + 1 } else { 0 };
+        out.push(b);
+    }
+    out
+}
+
+// ---------------------------------------------------------------- AV1
+
+struct Seq {
+    profile: u64,
+    level: u64,
+    tier: u64,
+    high: bool,
+    twelve: bool,
+    mono: bool,
+    color: Option<(u64, u64, u64)>,
+    range: bool,
+    ss: Option<(bool, bool)>,
+    csp: u64,
+}
+
+impl Seq {
+    fn new(profile: u64, level: u64) -> Seq {
+        Seq { profile, level, tier: 0, high: false, twelve: false, mono: false, color: None, range: false, ss: None, csp: 0 }
+    }
+
+    /// A sequence header OBU's payload (AV1 §5.5): one operating point, frames of
+    /// 640×360, order hints, screen content tools chosen per frame.
+    fn payload(&self) -> Vec<u8> {
+        let mut w = W::default();
+        w.f(3, self.profile).f(1, 0).f(1, 0); // still_picture, reduced_still_picture_header
+        w.f(1, 0).f(1, 0).f(5, 0); // timing_info, initial_display_delay, one operating point
+        w.f(12, 0).f(5, self.level);
+        if self.level > 7 {
+            w.f(1, self.tier);
+        }
+        w.f(4, 15).f(4, 15).f(16, 639).f(16, 359);
+        w.f(1, 0); // frame_id_numbers_present_flag
+        w.f(3, 0b011); // 128x128 superblock, filter intra, intra edge
+        w.f(4, 0b1111).f(1, 1).f(2, 0b11); // interintra..dual filter; order hint; jnt, ref mvs
+        w.f(1, 1).f(1, 1); // seq_choose_screen_content_tools, seq_choose_integer_mv
+        w.f(3, 6); // order_hint_bits_minus_1
+        w.f(3, 0b011); // superres, cdef, restoration
+        w.f(1, u64::from(self.high));
+        if self.profile == 2 && self.high {
+            w.f(1, u64::from(self.twelve));
+        }
+        if self.profile != 1 {
+            w.f(1, u64::from(self.mono));
+        }
+        match self.color {
+            Some((cp, tc, mc)) => {
+                w.f(1, 1).f(8, cp).f(8, tc).f(8, mc);
+            }
+            None => {
+                w.f(1, 0);
+            }
+        }
+        let srgb = self.color == Some((1, 13, 0));
+        if self.mono {
+            w.f(1, u64::from(self.range));
+        } else if !srgb {
+            w.f(1, u64::from(self.range));
+            if self.profile == 2 && self.twelve {
+                let (x, y) = self.ss.unwrap_or((true, true));
+                w.f(1, u64::from(x));
+                if x {
+                    w.f(1, u64::from(y));
+                }
+            }
+            let (x, y) = match self.profile {
+                0 => (true, true),
+                1 => (false, false),
+                _ => self.ss.unwrap_or((true, false)),
+            };
+            if x && y {
+                w.f(2, self.csp);
+            }
+        }
+        if !self.mono {
+            w.f(1, 0); // separate_uv_delta_q
+        }
+        w.f(1, 0); // film_grain_params_present
+        w.bytes()
+    }
+}
+
+/// A WAI AV1 payload (SPEC §5): `n`, `h`, `w`, `fps_num`, `fps_den`, then each temporal
+/// unit with its length.
+fn framing(units: &[Vec<u8>], h: u32, w: u32, fps: (u32, u32)) -> Vec<u8> {
+    let mut p = Vec::new();
+    for v in [units.len() as u32, h, w, fps.0, fps.1] {
+        p.extend_from_slice(&v.to_le_bytes());
+    }
+    for u in units {
+        p.extend_from_slice(&(u.len() as u32).to_le_bytes());
+        p.extend_from_slice(u);
+    }
+    p
+}
+
+/// One temporal unit: a temporal delimiter, then a sequence header OBU around `seq`.
+fn unit_with(seq: &[u8]) -> Vec<u8> {
+    let mut u = vec![0x12, 0x00, 0x0a];
+    u.push(seq.len() as u8);
+    u.extend_from_slice(seq);
+    u
+}
+
+#[test]
+fn av1_codec_strings_follow_the_isobmff_binding() {
+    let cs = |s: &Seq| Av1Sequence::parse(&s.payload()).unwrap().codec_string();
+    // The binding's own examples (AV1-ISOBMFF §5).
+    let mut hdr = Seq::new(0, 4);
+    hdr.high = true;
+    hdr.color = Some((9, 16, 9));
+    hdr.csp = 2;
+    assert_eq!(cs(&hdr), "av01.0.04M.10.0.112.09.16.09.0", "the binding's example: Main, level 3.0, 10-bit, BT.2100 PQ");
+    assert_eq!(cs(&Seq::new(0, 1)), "av01.0.01M.08", "the binding's example: Main, level 2.1, 8-bit, the defaults omitted");
+    let mut high_tier = Seq::new(0, 8);
+    high_tier.tier = 1;
+    assert_eq!(cs(&high_tier), "av01.0.08H.08");
+    let mut full_range_444 = Seq::new(1, 5);
+    full_range_444.range = true;
+    assert_eq!(cs(&full_range_444), "av01.1.05M.08.0.000.01.01.01.1", "4:4:4 and full range are not the defaults");
+    let mut srgb = Seq::new(1, 5);
+    srgb.color = Some((1, 13, 0));
+    assert_eq!(cs(&srgb), "av01.1.05M.08.0.000.01.13.00.1", "sRGB implies full range and no subsampling");
+    let mut mono = Seq::new(0, 4);
+    mono.mono = true;
+    assert_eq!(cs(&mono), "av01.0.04M.08.1.110.01.01.01.0");
+    let mut p2 = Seq::new(2, 10);
+    p2.high = true;
+    p2.twelve = true;
+    p2.ss = Some((true, false));
+    assert_eq!(cs(&p2), "av01.2.10M.12.0.100.01.01.01.0", "12-bit 4:2:2");
+    let mut csp = Seq::new(0, 4);
+    csp.csp = 1;
+    assert_eq!(cs(&csp), "av01.0.04M.08.0.111.01.01.01.0", "a chroma sample position other than unknown is written");
+}
+
+#[test]
+fn an_av1_query_reads_the_framing_and_the_first_sequence_header() {
+    let mut s = Seq::new(0, 8);
+    s.high = true;
+    let seq = s.payload();
+    let units = vec![unit_with(&seq), vec![0x12, 0x00, 0x32, 0x01, 0x00]];
+    let p = framing(&units, 360, 640, (30000, 1001));
+    let q = query(codecs::CAP_VIDEO_AV1, &p, None).unwrap();
+    assert_eq!(q.codec, "av01.0.08M.10");
+    assert_eq!((q.width, q.height, q.coded_width, q.coded_height, q.bit_depth), (640, 360, 640, 360, 10));
+    assert_eq!(q.framerate, Some(Rate { num: 30000, den: 1001 }));
+    let want = ((p.len() as u128) * 8 * 30000).div_ceil(2 * 1001) as u64;
+    assert_eq!(q.bitrate, Some(want));
+    let mc = q.decoding_configuration("media-source").unwrap();
+    assert_eq!(mc["video"]["contentType"], "video/mp4; codecs=\"av01.0.08M.10\"");
+    assert_eq!(mc["type"], "media-source");
+    let wc = q.decoder_config(HardwareAcceleration::PreferHardware);
+    assert_eq!(wc["hardwareAcceleration"], "prefer-hardware");
+    assert_eq!(wc["codedWidth"], 640);
+    // A payload with a zero frame rate has a query with no rate, and so no
+    // decodingInfo() configuration.
+    let q0 = query(codecs::CAP_VIDEO_AV1, &framing(&units, 360, 640, (0, 1)), None).unwrap();
+    assert_eq!((q0.framerate, q0.bitrate), (None, None));
+    assert!(q0.decoding_configuration("file").is_none());
+    // Truncated, and a unit with no sequence header.
+    assert!(matches!(query(codecs::CAP_VIDEO_AV1, &p[..30], None), Err(QueryError::Malformed(_))));
+    let no_seq = framing(&[vec![0x12, 0x00]], 360, 640, (30, 1));
+    assert!(matches!(query(codecs::CAP_VIDEO_AV1, &no_seq, None), Err(QueryError::Malformed(_))));
+    assert_eq!(query(codecs::CAP_NEURAL_INT_SYNTH, &p, None), Err(QueryError::NotMapped(codecs::CAP_NEURAL_INT_SYNTH.into())));
+}
+
+/// The low-overhead OBU stream split into temporal units at each temporal delimiter.
+fn temporal_units(obus: &[u8]) -> Vec<Vec<u8>> {
+    let mut units: Vec<Vec<u8>> = Vec::new();
+    let mut at = 0;
+    while at < obus.len() {
+        let header = obus[at];
+        let hl = if header & 0x04 != 0 { 2 } else { 1 };
+        let (size, sl) = leb128(&obus[at + hl..]).unwrap();
+        let end = at + hl + sl + size as usize;
+        if (header >> 3) & 0x0f == 2 || units.is_empty() {
+            units.push(Vec::new());
+        }
+        units.last_mut().unwrap().extend_from_slice(&obus[at..end]);
+        at = end;
+    }
+    units
+}
+
+/// Every AV1 floor the packaging corpus supplies, wrapped in SPEC §5's framing at the
+/// pack's frame rate: its codec string is one the packager wrote into the pack's HLS
+/// master (`CODECS`), and its size one of the master's `RESOLUTION`s. The packager reads
+/// the sequence header with its own reader (`wai-pack` `av1.rs`).
+#[test]
+fn the_packaging_corpus_floors_give_their_catalogs_codec_strings() {
+    let packs = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("../pack-conformance/packs");
+    let mut checked = 0;
+    let mut names: Vec<_> = std::fs::read_dir(&packs).unwrap().filter_map(|e| e.ok()).map(|e| e.path()).collect();
+    names.sort();
+    for pack in names {
+        let master = pack.join("package/hls/classic/master.m3u8");
+        let (Ok(floors), Ok(m)) = (std::fs::read_dir(pack.join("floor")), std::fs::read_to_string(&master)) else { continue };
+        let codecs_of: Vec<String> = m.split("CODECS=\"").skip(1).filter_map(|s| s.split('"').next()).flat_map(|s| s.split(',').map(str::to_string).collect::<Vec<_>>()).filter(|c| c.starts_with("av01")).collect();
+        let sizes: Vec<(u32, u32)> = m
+            .split("RESOLUTION=")
+            .skip(1)
+            .filter_map(|s| {
+                let r: String = s.chars().take_while(|c| c.is_ascii_alphanumeric()).collect();
+                let (w, h) = r.split_once('x')?;
+                Some((w.parse().ok()?, h.parse().ok()?))
+            })
+            .collect();
+        for f in floors.filter_map(|e| e.ok()).map(|e| e.path()).filter(|p| p.extension().is_some_and(|x| x == "obu")) {
+            let obus = std::fs::read(&f).unwrap();
+            let units = temporal_units(&obus);
+            let seq = obus_sequence(&units[0]);
+            let (w, h) = sizes.iter().copied().find(|&(w, h)| (w, h) == seq).unwrap_or(seq);
+            let p = framing(&units, h, w, (25, 1));
+            let q = query(codecs::CAP_VIDEO_AV1, &p, None).unwrap_or_else(|e| panic!("{}: {e}", f.display()));
+            assert!(codecs_of.contains(&q.codec), "{}: {} is not among the master's {codecs_of:?}", f.display(), q.codec);
+            assert!(sizes.contains(&(q.width, q.height)), "{}: {}x{} is not among the master's {sizes:?}", f.display(), q.width, q.height);
+            checked += 1;
+        }
+    }
+    assert!(checked >= 4, "only {checked} floors were checked");
+}
+
+/// The max frame size the first sequence header states.
+fn obus_sequence(unit: &[u8]) -> (u32, u32) {
+    let mut at = 0;
+    loop {
+        let header = unit[at];
+        let hl = if header & 0x04 != 0 { 2 } else { 1 };
+        let (size, sl) = leb128(&unit[at + hl..]).unwrap();
+        let body = &unit[at + hl + sl..at + hl + sl + size as usize];
+        if (header >> 3) & 0x0f == 1 {
+            let mut r = Bits { b: body, at: 0 };
+            let _profile = r.f(3).unwrap();
+            r.f(1).unwrap();
+            let reduced = r.flag().unwrap();
+            assert!(!reduced);
+            assert!(!r.flag().unwrap(), "no timing info in the corpus floors");
+            assert!(!r.flag().unwrap());
+            let cnt = r.f(5).unwrap() + 1;
+            for _ in 0..cnt {
+                r.f(12).unwrap();
+                if r.f(5).unwrap() > 7 {
+                    r.f(1).unwrap();
+                }
+            }
+            let (wb, hb) = (r.f(4).unwrap() + 1, r.f(4).unwrap() + 1);
+            return (r.f(wb).unwrap() + 1, r.f(hb).unwrap() + 1);
+        }
+        at += hl + sl + size as usize;
+    }
+}
+
+// ---------------------------------------------------------------- HEVC
+
+struct Sps {
+    space: u64,
+    tier: u64,
+    idc: u64,
+    compat: u32,
+    constraints: [u8; 6],
+    level: u64,
+    sub_layers: Vec<(bool, bool)>,
+    chroma: u32,
+    width: u32,
+    height: u32,
+    window: Option<(u32, u32, u32, u32)>,
+    depth: u32,
+}
+
+impl Sps {
+    fn main(level: u64) -> Sps {
+        Sps { space: 0, tier: 0, idc: 1, compat: 0b110, constraints: [0xB0, 0, 0, 0, 0, 0], level, sub_layers: Vec::new(), chroma: 1, width: 1920, height: 1088, window: Some((0, 0, 0, 4)), depth: 8 }
+    }
+
+    /// The NAL unit, header and emulation prevention included.
+    fn nal(&self) -> Vec<u8> {
+        let mut w = W::default();
+        w.f(4, 0).f(3, self.sub_layers.len() as u64).f(1, 1);
+        w.f(2, self.space).f(1, self.tier).f(5, self.idc);
+        for j in 0..32 {
+            w.f(1, u64::from((self.compat >> j) & 1));
+        }
+        for c in self.constraints {
+            w.f(8, u64::from(c));
+        }
+        w.f(8, self.level);
+        for &(p, l) in &self.sub_layers {
+            w.f(1, u64::from(p)).f(1, u64::from(l));
+        }
+        if !self.sub_layers.is_empty() {
+            for _ in self.sub_layers.len()..8 {
+                w.f(2, 0);
+            }
+        }
+        for &(p, l) in &self.sub_layers {
+            if p {
+                w.f(32, 0x1234_5678).f(32, 0x9abc_def0).f(24, 0x00ab_cdef);
+            }
+            if l {
+                w.f(8, 90);
+            }
+        }
+        w.ue(0).ue(self.chroma);
+        if self.chroma == 3 {
+            w.f(1, 0);
+        }
+        w.ue(self.width).ue(self.height);
+        match self.window {
+            Some((l, r, t, b)) => {
+                w.f(1, 1).ue(l).ue(r).ue(t).ue(b);
+            }
+            None => {
+                w.f(1, 0);
+            }
+        }
+        w.ue(self.depth - 8).ue(self.depth - 8);
+        let mut nal = vec![33 << 1, 1];
+        nal.extend(escape(&w.bytes()));
+        nal
+    }
+}
+
+/// An Annex B stream: an SPS, then `pictures` slices, each the first of its picture.
+fn hevc_stream(sps: &Sps, pictures: usize) -> Vec<u8> {
+    let mut s = vec![0, 0, 0, 1];
+    s.extend(sps.nal());
+    for k in 0..pictures {
+        s.extend_from_slice(&[0, 0, 1, 1 << 1, 1, 0x80 | (k as u8 & 0x7f), 0xAA, 0x55]);
+    }
+    s
+}
+
+#[test]
+fn hevc_codec_strings_and_queries() {
+    let s = Sps::main(93);
+    let parsed = HevcSps::parse(&s.nal()).unwrap();
+    assert_eq!(parsed.codec_string(), "hev1.1.6.L93.B0");
+    assert_eq!((parsed.width, parsed.height, parsed.cropped_width, parsed.cropped_height, parsed.bit_depth), (1920, 1088, 1920, 1080, 8));
+    // Main 10, high tier, level 5.1, a second constraint byte.
+    let mut m10 = Sps::main(153);
+    (m10.idc, m10.compat, m10.tier, m10.depth, m10.constraints) = (2, 0b100, 1, 10, [0x90, 0x00, 0x12, 0, 0, 0]);
+    assert_eq!(HevcSps::parse(&m10.nal()).unwrap().codec_string(), "hev1.2.4.H153.90.0.12");
+    // A profile space, and sub-layers with a profile and a level present.
+    let mut spaced = Sps::main(120);
+    spaced.space = 1;
+    spaced.sub_layers = vec![(true, false), (false, true)];
+    let p = HevcSps::parse(&spaced.nal()).unwrap();
+    assert_eq!(p.codec_string(), "hev1.A1.6.L120.B0");
+    assert_eq!((p.cropped_width, p.cropped_height), (1920, 1080), "the sub-layer fields are skipped exactly");
+    // Every compatibility flag clear: the RBSP holds a run of zero bytes, which the NAL
+    // carries with emulation-prevention bytes.
+    let mut zeros = Sps::main(93);
+    (zeros.compat, zeros.constraints, zeros.window) = (0, [0; 6], None);
+    let nal = zeros.nal();
+    assert!(nal.windows(3).any(|w| w == [0, 0, 3]), "the constructed NAL needs emulation prevention");
+    let z = HevcSps::parse(&nal).unwrap();
+    assert_eq!(z.codec_string(), "hev1.1.0.L93");
+    assert_eq!((z.cropped_width, z.cropped_height), (1920, 1088));
+    // The query: the cropped size, the coded size, the presentation rate the sink gives,
+    // and the bit rate over the pictures the stream holds.
+    let stream = hevc_stream(&s, 4);
+    let q = query(codecs::CAP_VIDEO_HEVC, &stream, Some(Rate { num: 25, den: 1 })).unwrap();
+    assert_eq!((q.codec.as_str(), q.width, q.height, q.coded_width, q.coded_height), ("hev1.1.6.L93.B0", 1920, 1080, 1920, 1088));
+    assert_eq!(q.bitrate, Some(((stream.len() as u64) * 8 * 25).div_ceil(4)));
+    let unrated = query(codecs::CAP_VIDEO_HEVC, &stream, None).unwrap();
+    assert_eq!((unrated.framerate, unrated.bitrate), (None, None), "the payload states no frame rate");
+    assert!(unrated.decoding_configuration("media-source").is_none());
+    assert!(matches!(query(codecs::CAP_VIDEO_HEVC, &[0, 0, 1, 2, 1, 0x80], None), Err(QueryError::Malformed(_))), "no SPS");
+}
+
+// ---------------------------------------------------------------- the order
+
+fn answers_of(v: &Value) -> Option<Answers> {
+    v.get("answers").map(|a| answers_from_json(a).unwrap())
+}
+
+fn budget_of(v: &Value) -> Option<EnergyBudget> {
+    v.get("budget").map(|b| budget_from_json(b).unwrap())
+}
+
+fn rules_of(case: &Value, base: Rules) -> Rules {
+    Rules { budget_ceiling: case.get("rules").and_then(|r| r.get("ceiling")).and_then(Value::as_u64), ..base }
+}
+
+/// Every vector of `order.json` under `base`; `Err` names the first that differs.
+fn run_vectors(base: Rules) -> Result<usize, String> {
+    let v: Value = serde_json::from_str(ORDER).unwrap();
+    let mut n = 0;
+    for case in v["cases"].as_array().unwrap() {
+        let id = case["id"].as_str().unwrap();
+        let rules = rules_of(case, base);
+        let es = case["entries"].as_array().unwrap();
+        let answers: Vec<Option<Answers>> = es.iter().map(answers_of).collect();
+        let budgets: Vec<Option<EnergyBudget>> = es.iter().map(budget_of).collect();
+        let entries: Vec<Entry<'_>> = es.iter().enumerate().map(|(i, e)| Entry { capability: e["capability"].as_str().unwrap(), answers: answers[i].as_ref(), budget: budgets[i].as_ref() }).collect();
+        let want = &case["expect"];
+        let renditions: Vec<Value> = es.iter().map(|e| serde_json::json!({ "capability": e["capability"], "kind": "bytes" })).collect();
+        let manifest: ManifestV2 = serde_json::from_value(serde_json::json!({ "wai": "1.1", "media": "video", "intent": "display", "renditions": renditions, "target": null })).unwrap();
+        let got = order_json(&manifest, &answers, &budgets, &rules);
+        for k in ["ranks", "order", "refused", "picked"] {
+            if got[k] != want[k] {
+                return Err(format!("{id}: {k} {}, not {}", got[k], want[k]));
+            }
+        }
+        // The same ranks, entry by entry, through rank() directly.
+        for (i, e) in entries.iter().enumerate() {
+            let r = rank(e, &rules).map_or(Value::Null, |r| Value::from(r.as_str()));
+            if r != want["ranks"][i] {
+                return Err(format!("{id}: entry {i} ranks {r}"));
+            }
+        }
+        n += 1;
+    }
+    for case in v["wai1"].as_array().unwrap() {
+        let id = case["id"].as_str().unwrap();
+        let rules = rules_of(case, base);
+        let (ca, cb) = (answers_of(&case["capability"]), budget_of(&case["capability"]));
+        let cap = Entry { capability: case["capability"]["capability"].as_str().unwrap(), answers: ca.as_ref(), budget: cb.as_ref() };
+        let (fa, fbud) = if case["fallback"].is_null() { (None, None) } else { (answers_of(&case["fallback"]), budget_of(&case["fallback"])) };
+        let fb = (!case["fallback"].is_null()).then(|| Entry { capability: case["fallback"]["capability"].as_str().unwrap(), answers: fa.as_ref(), budget: fbud.as_ref() });
+        let adv = advertise_for_envelope(&cap, fb.as_ref(), true, &rules);
+        let decodes = if adv { "capability" } else { "fallback" };
+        if Value::from(adv) != case["expect"]["advertise"] || case["expect"]["decodes"] != decodes {
+            return Err(format!("{id}: advertise {adv}"));
+        }
+        n += 1;
+    }
+    Ok(n)
+}
+
+#[test]
+fn the_default_order_gives_every_vector() {
+    let n = run_vectors(Rules::DEFAULT).unwrap_or_else(|e| panic!("{e}"));
+    assert!(n >= 20, "{n} vectors");
+}
+
+/// Each rule of the order, turned off, is caught: some vector differs.
+#[test]
+fn each_rule_turned_off_is_caught_by_a_vector() {
+    let off = [
+        ("learned below classical", Rules { learned_below_classical: false, ..Rules::DEFAULT }),
+        ("unsupported is unsupported", Rules { unsupported_is_unsupported: false, ..Rules::DEFAULT }),
+        ("a budget must qualify", Rules { budget_must_qualify: false, ..Rules::DEFAULT }),
+    ];
+    for (name, rules) in off {
+        let e = run_vectors(rules).err();
+        assert!(e.is_some(), "with \"{name}\" off every vector still passed");
+        println!("platform-decode rule off: {name}: caught by {}", e.unwrap());
+    }
+}
+
+#[test]
+fn a_budget_qualifies_only_as_a_measurement() {
+    let good = EnergyBudget { class: "OnChipCounter".into(), joules_micro: 1, per: "frame".into(), relative_ppm: 1, window_us: 1, reference_sink: "r".into() };
+    assert_eq!(good.qualifies(None), Ok(()));
+    assert_eq!(good.qualifies(Some(1)), Ok(()));
+    assert_eq!(good.qualifies(Some(0)), Err(BudgetRefusal::Ceiling { joules_micro: 1, ceiling: 0 }));
+    for class in ["ModelBased", "Estimator", "HwShunt", "CalibratedInstrument", ""] {
+        assert_eq!(EnergyBudget { class: class.into(), ..good.clone() }.qualifies(None), Err(BudgetRefusal::Class(class.into())));
+    }
+    assert_eq!(EnergyBudget { window_us: 0, ..good.clone() }.qualifies(None), Err(BudgetRefusal::NotMeasured));
+    assert_eq!(EnergyBudget { reference_sink: String::new(), ..good.clone() }.qualifies(None), Err(BudgetRefusal::NotMeasured));
+    assert_eq!(EnergyBudget { per: "byte".into(), ..good.clone() }.qualifies(None), Err(BudgetRefusal::Unit("byte".into())));
+}
+
+#[test]
+fn learned_is_read_from_the_registry_class() {
+    for c in [codecs::CAP_NEURAL_INT_SYNTH, codecs::CAP_NEURAL_INT_HYPER, codecs::CAP_IMAGE_JPEGAI, codecs::CAP_NEURAL_MBT2018_MEAN, codecs::CAP_SEMANTIC_CLIP, codecs::CAP_GENERATE_SPECTRAL] {
+        assert!(is_learned(c), "{c}");
+    }
+    for c in MAPPED.iter().copied().chain([codecs::CAP_VIDEO_INT_MOTION, codecs::CAP_IMAGE_PNG, codecs::CAP_VIDEO_VP9]) {
+        assert!(!is_learned(c), "{c}");
+    }
+}
+
+/// `queries.json`: two floor objects of the packaging corpus, and an SPS a hardware
+/// encoder wrote, whose level is not the one the encoder was configured with.
+#[test]
+fn the_query_vectors() {
+    let v: Value = serde_json::from_str(QUERIES).unwrap();
+    let wai = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("..");
+    let mut n = 0;
+    for case in v["cases"].as_array().unwrap() {
+        let id = case["id"].as_str().unwrap();
+        let cap = case["capability"].as_str().unwrap();
+        let want = &case["expect"];
+        let got = if let Some(f) = case["payload_file"].as_str() {
+            let env = crate::container::Wai::from_bytes(&std::fs::read(wai.join(f)).unwrap()).unwrap();
+            assert_eq!(env.capability(), cap, "{id}");
+            query_json(&query(cap, &env.payload, None).unwrap())
+        } else {
+            let hex = case["sps_hex"].as_str().unwrap();
+            let nal: Vec<u8> = (0..hex.len()).step_by(2).map(|i| u8::from_str_radix(&hex[i..i + 2], 16).unwrap()).collect();
+            let sps = HevcSps::parse(&nal).unwrap();
+            serde_json::json!({ "codec": sps.codec_string(), "width": sps.cropped_width, "height": sps.cropped_height, "codedWidth": sps.width, "codedHeight": sps.height, "bitDepth": sps.bit_depth })
+        };
+        for (k, w) in want.as_object().unwrap() {
+            assert_eq!(&got[k.as_str()], w, "{id}: {k}");
+        }
+        n += 1;
+    }
+    assert_eq!(n, 3);
+}

wai/wai-web/Cargo.toml

@@ -38,6 +38,11 @@ quantum_webgpu = ["wai/quantum_fabric", "wai/quantum_fabric_webgpu", "dep:wasm-b
 # into the IDE's bundle (scripts/wasm-pkgs.tsv on the site), which the console
 # also loads; off by default, so the main bundle is unchanged.
 quantum_engine = ["wai/quantum_engine"]
+# Platform decode capabilities in dispatch (src/platform_decode.rs,
+# extensions/platform-decode.md): the query a page makes of Media Capabilities and
+# WebCodecs for an entry, and the default order. Built as its own bundle with its own
+# --out-name (demo/platform_decode.html); off by default, so the main bundle is unchanged.
+platform_decode = []
 
 [dev-dependencies]
 # Pinned to the wasm-bindgen release Cargo.lock holds, so the test runner and the

wai/wai-web/demo/platform_decode.html

@@ -0,0 +1,6 @@
+<!doctype html>
+<meta charset="utf-8">
+<title>Platform decode in dispatch</title>
+<p>The query the reference reads from each payload, this platform's Media Capabilities and WebCodecs answers to it, a decode through WebCodecs, and the default order those answers give (extensions/platform-decode.md).</p>
+<pre id="report">running</pre>
+<script type="module" src="platform_decode.js"></script>

wai/wai-web/demo/platform_decode.js

@@ -0,0 +1,168 @@
+// Platform decode capabilities in dispatch (extensions/platform-decode.md), in a page.
+//
+// For each AV1 floor object of the packaging corpus (WAI1 envelopes whose payload is SPEC
+// §5's AV1 framing): the query the reference reads from the payload (platformQuery), the
+// platform's answers to it (Media Capabilities decodingInfo, WebCodecs isConfigSupported
+// with prefer-hardware and no-preference), and a decode of every temporal unit through
+// WebCodecs with the configuration the query gives. Then the same codec family at a
+// broadcast size; then an HEVC Annex B stream the platform's own encoder writes, read by
+// the reference's query and decoded back; then the default order (platformOrder) over a
+// manifest of a learned entry and an AV1 entry, from the answers this platform gave.
+//
+// The report is left in window.platformDecode for verify_platform_decode.mjs to read.
+import init, { parse, platformQuery, platformOrder } from "./pkg-platform/wai_platform_decode.js";
+
+const lines = [];
+const say = (l) => lines.push(l);
+const info = (d) => d ? `supported=${d.supported} smooth=${d.smooth} powerEfficient=${d.powerEfficient}` : "none";
+
+/// SPEC §5's AV1 framing: n, h, w, fps_num, fps_den, then each unit with its length.
+function av1Units(p) {
+  const dv = new DataView(p.buffer, p.byteOffset, p.byteLength);
+  const n = dv.getUint32(0, true);
+  const units = [];
+  let at = 20;
+  for (let k = 0; k < n; k++) {
+    const len = dv.getUint32(at, true);
+    units.push(p.subarray(at + 4, at + 4 + len));
+    at += 4 + len;
+  }
+  return units;
+}
+
+/// Whether a temporal unit carries a sequence header OBU (type 1): the corpus floors start
+/// each closed group with one, before its key frame.
+function hasSequenceHeader(tu) {
+  let at = 0;
+  while (at < tu.length) {
+    const h = tu[at];
+    const ext = (h & 4) ? 1 : 0;
+    let size = 0, shift = 0, i = at + 1 + ext;
+    for (;;) { const b = tu[i++]; size += (b & 0x7f) * 2 ** shift; shift += 7; if (!(b & 0x80)) break; }
+    if (((h >> 3) & 15) === 1) return true;
+    at = i + size;
+  }
+  return false;
+}
+
+async function answers(q) {
+  const out = { decodingInfo: null, preferHardware: null, noPreference: null };
+  if (q.decodingConfiguration) {
+    const d = await navigator.mediaCapabilities.decodingInfo(q.decodingConfiguration);
+    out.decodingInfo = { supported: d.supported, smooth: d.smooth, powerEfficient: d.powerEfficient };
+  }
+  out.preferHardware = (await VideoDecoder.isConfigSupported(q.decoderConfigPreferHardware)).supported;
+  out.noPreference = (await VideoDecoder.isConfigSupported(q.decoderConfigNoPreference)).supported;
+  return out;
+}
+
+/// Decode `chunks` ([data, key]) with the configuration `config`: the frames output.
+async function decodeAll(config, chunks, frameUs) {
+  let frames = 0, error = null;
+  const dec = new VideoDecoder({ output: (f) => { frames++; f.close(); }, error: (e) => { error = String(e); } });
+  dec.configure(config);
+  chunks.forEach(([data, key], k) => dec.decode(new EncodedVideoChunk({ type: key ? "key" : "delta", timestamp: k * frameUs, data })));
+  try { await dec.flush(); } catch (e) { error = error || String(e); }
+  if (dec.state !== "closed") dec.close();
+  return { frames, error };
+}
+
+async function av1Objects(paths) {
+  let firstAnswers = null;
+  for (const path of paths) {
+    const bytes = new Uint8Array(await (await fetch(path)).arrayBuffer());
+    const env = parse(bytes);
+    if (env.capability !== "wai.video.av1") continue;
+    const payload = env.payload;
+    const q = JSON.parse(platformQuery("wai.video.av1", payload, undefined, undefined));
+    const a = await answers(q);
+    const units = av1Units(payload);
+    const config = { ...(a.preferHardware ? q.decoderConfigPreferHardware : q.decoderConfigNoPreference) };
+    const r = await decodeAll(config, units.map((u) => [u, hasSequenceHeader(u)]), Math.round(1e6 * q.framerate[1] / q.framerate[0]));
+    say(`platform av1 object: ${path.split("/").slice(-4).join("/")} codec=${q.codec} size=${q.width}x${q.height} rate=${q.framerate.join("/")} bitrate=${q.bitrate} decodingInfo=[${info(a.decodingInfo)}] isConfigSupported prefer-hardware=${a.preferHardware} no-preference=${a.noPreference} decoded=${r.frames}/${units.length} with ${config.hardwareAcceleration}${r.error ? " error=" + r.error : ""}`);
+    firstAnswers = firstAnswers || a;
+  }
+  return firstAnswers;
+}
+
+/// The same codec family at a broadcast size, through the same query form.
+async function av1Broadcast() {
+  const q = {
+    codec: "av01.0.08M.08",
+    decodingConfiguration: { type: "media-source", video: { contentType: 'video/mp4; codecs="av01.0.08M.08"', width: 1920, height: 1080, bitrate: 5000000, framerate: 30 } },
+    decoderConfigPreferHardware: { codec: "av01.0.08M.08", codedWidth: 1920, codedHeight: 1080, hardwareAcceleration: "prefer-hardware" },
+    decoderConfigNoPreference: { codec: "av01.0.08M.08", codedWidth: 1920, codedHeight: 1080, hardwareAcceleration: "no-preference" },
+  };
+  const a = await answers(q);
+  say(`platform av1 1920x1080: codec=${q.codec} decodingInfo=[${info(a.decodingInfo)}] isConfigSupported prefer-hardware=${a.preferHardware} no-preference=${a.noPreference}`);
+  return a;
+}
+
+/// The first SPS NAL unit (type 33) of an Annex B stream, as hexadecimal.
+function firstSps(s) {
+  for (let i = 0; i + 4 < s.length; i++) {
+    if (s[i] === 0 && s[i + 1] === 0 && s[i + 2] === 1 && ((s[i + 3] >> 1) & 63) === 33) {
+      let j = i + 3;
+      while (j + 3 <= s.length && !(s[j] === 0 && s[j + 1] === 0 && (s[j + 2] === 1 || (s[j + 2] === 0 && s[j + 3] === 1)))) j++;
+      return Array.from(s.subarray(i + 3, j), (b) => b.toString(16).padStart(2, "0")).join("");
+    }
+  }
+  return "";
+}
+
+async function hevcRoundTrip() {
+  const configured = "hvc1.1.6.L93.B0";
+  const [w, h, n] = [640, 360, 12];
+  const enc = { codec: configured, width: w, height: h, bitrate: 1000000, framerate: 30, hevc: { format: "annexb" }, hardwareAcceleration: "prefer-hardware" };
+  const es = await VideoEncoder.isConfigSupported(enc);
+  if (!es.supported) { say(`platform hevc encoder: not supported (${configured} ${w}x${h} prefer-hardware)`); return; }
+  const chunks = [];
+  let declared = "";
+  const e = new VideoEncoder({ output: (c, meta) => { const b = new Uint8Array(c.byteLength); c.copyTo(b); chunks.push([b, c.type === "key"]); if (meta && meta.decoderConfig) declared = meta.decoderConfig.codec; }, error: (x) => say(`platform hevc encoder: error ${x}`) });
+  e.configure(enc);
+  const rgba = new Uint8Array(w * h * 4);
+  for (let k = 0; k < n; k++) {
+    for (let y = 0; y < h; y++) for (let x = 0; x < w; x++) { const o = 4 * (y * w + x); rgba[o] = (x + 4 * k) & 255; rgba[o + 1] = (y * 2) & 255; rgba[o + 2] = (x ^ y) & 255; rgba[o + 3] = 255; }
+    const f = new VideoFrame(rgba, { format: "RGBA", codedWidth: w, codedHeight: h, timestamp: k * 33333 });
+    e.encode(f, { keyFrame: k === 0 });
+    f.close();
+  }
+  await e.flush();
+  e.close();
+  const stream = new Uint8Array(chunks.reduce((s, [b]) => s + b.length, 0));
+  let at = 0;
+  for (const [b] of chunks) { stream.set(b, at); at += b.length; }
+  const q = JSON.parse(platformQuery("wai.video.hevc", stream, 30, 1));
+  const a = await answers(q);
+  const r = await decodeAll(a.preferHardware ? q.decoderConfigPreferHardware : q.decoderConfigNoPreference, chunks, 33333);
+  say(`platform hevc encoder: configured=${configured} encoder-declares=${declared || "nothing"} derived=${q.codec} size=${q.width}x${q.height} coded=${q.codedWidth}x${q.codedHeight} depth=${q.bitDepth} bitrate=${q.bitrate} decodingInfo=[${info(a.decodingInfo)}] isConfigSupported prefer-hardware=${a.preferHardware} no-preference=${a.noPreference} decoded=${r.frames}/${chunks.length}${r.error ? " error=" + r.error : ""} sps=${firstSps(stream)}`);
+}
+
+function order(a) {
+  const manifest = { wai: "1.1", media: "video", intent: "display", target: null, renditions: [{ capability: "wai.neural.int_synth", kind: "bytes" }, { capability: "wai.video.av1", kind: "bytes" }] };
+  // The sink configures prefer-hardware when the platform accepts it, and otherwise
+  // no-preference; the answer it passes is for the preference it configures.
+  const webcodecs = a.preferHardware ? { hardwareAcceleration: "prefer-hardware", supported: true } : { hardwareAcceleration: "no-preference", supported: a.noPreference };
+  const av1 = { decodingInfo: a.decodingInfo, webcodecs, ownDecoder: false };
+  const o = JSON.parse(platformOrder(JSON.stringify(manifest), JSON.stringify([null, av1]), JSON.stringify([null, null]), undefined));
+  say(`platform order: entries=[wai.neural.int_synth, wai.video.av1] av1=[${info(a.decodingInfo)} webcodecs=${webcodecs.hardwareAcceleration}:${webcodecs.supported}] ranks=${JSON.stringify(o.ranks)} order=${JSON.stringify(o.order)} picked=${o.picked}`);
+  const budget = { class: "OnChipCounter", joules_micro: 41000, per: "frame", relative_ppm: 20000, window_us: 5000000, reference_sink: "reference sink: arm64 desktop" };
+  const ob = JSON.parse(platformOrder(JSON.stringify(manifest), JSON.stringify([null, av1]), JSON.stringify([budget, null]), undefined));
+  say(`platform order with a budget: ranks=${JSON.stringify(ob.ranks)} order=${JSON.stringify(ob.order)} picked=${ob.picked}`);
+}
+
+window.platformDecode = { done: false, report: "" };
+try {
+  await init();
+  say(`platform: userAgent=${navigator.userAgent}`);
+  const paths = await (await fetch("objects.json")).json();
+  await av1Objects(paths);
+  const broadcast = await av1Broadcast();
+  await hevcRoundTrip();
+  order(broadcast);
+} catch (e) {
+  say(`platform: page error: ${e && e.stack || e}`);
+}
+window.platformDecode.report = lines.join("\n");
+document.getElementById("report").textContent = window.platformDecode.report;
+window.platformDecode.done = true;

wai/wai-web/demo/verify_platform_decode.mjs

@@ -0,0 +1,117 @@
+// Platform decode capabilities in dispatch, in a browser (extensions/platform-decode.md).
+//
+// Build the page's bundle first, from wai/wai-web:
+//
+//   cargo build --release --target wasm32-unknown-unknown --features platform_decode
+//   wasm-bindgen --target web --out-dir demo/pkg-platform --out-name wai_platform_decode \
+//     target/wasm32-unknown-unknown/release/wai_web.wasm
+//
+// then, with Playwright installed under PLAYWRIGHT_HOST:
+//
+//   [WAI_PLATFORM_BROWSER_CHANNEL=chrome] PLAYWRIGHT_HOST=<dir with node_modules/playwright>/ \
+//     node wai/wai-web/demo/verify_platform_decode.mjs
+//
+// It serves wai-web/demo and wai/pack-conformance on 127.0.0.1, lists the AV1 floor
+// objects of the packaging corpus (WAI1 envelopes of wai.video.av1) as objects.json,
+// opens platform_decode.html in a headless browser and prints the page's report. It fails
+// on a page error, on an AV1 object that does not decode every temporal unit through
+// WebCodecs with the configuration the reference's query gives, on an HEVC stream the
+// platform's encoder wrote that does not decode back with the derived configuration, and
+// on an order that is not the one the answers give: AV1, which the platform can decode,
+// before the learned entry without a budget; the learned entry with a budget first.
+import http from "node:http";
+import fs from "node:fs";
+import path from "node:path";
+import { fileURLToPath } from "node:url";
+import { createRequire } from "node:module";
+
+const HERE = path.dirname(fileURLToPath(import.meta.url));
+const PACKS = path.resolve(HERE, "../../pack-conformance");
+const PORT = Number(process.env.PORT || 8798);
+const { chromium } = createRequire(process.env.PLAYWRIGHT_HOST || HERE + "/")("playwright");
+
+/// The AV1 floor objects: WAI1 envelopes ("WAI1", a u32 manifest length, the manifest)
+/// whose manifest's capability is wai.video.av1.
+function av1Objects() {
+  const out = [];
+  const packs = path.join(PACKS, "packs");
+  for (const p of fs.readdirSync(packs).sort()) {
+    const dir = path.join(packs, p, "package", "objects");
+    if (!fs.existsSync(dir)) continue;
+    for (const f of fs.readdirSync(dir).sort()) {
+      const b = fs.readFileSync(path.join(dir, f));
+      if (b.length < 8 || b.toString("latin1", 0, 4) !== "WAI1") continue;
+      const m = JSON.parse(b.toString("utf8", 8, 8 + b.readUInt32LE(4)));
+      if (m.model_requirement && m.model_requirement.capability === "wai.video.av1") out.push(`/pack-conformance/packs/${p}/package/objects/${f}`);
+    }
+  }
+  return out;
+}
+
+const objects = av1Objects();
+const ROOTS = { "/pack-conformance/": PACKS, "/demo/": HERE };
+const MIME = { ".js": "text/javascript", ".wasm": "application/wasm", ".html": "text/html", ".json": "application/json" };
+const server = http.createServer((req, res) => {
+  const url = decodeURIComponent(req.url.split("?")[0]);
+  if (url === "/demo/objects.json") {
+    res.writeHead(200, { "content-type": "application/json" });
+    return res.end(JSON.stringify(objects));
+  }
+  const prefix = Object.keys(ROOTS).find((p) => url.startsWith(p));
+  const root = prefix && ROOTS[prefix];
+  const f = root && path.resolve(root, "." + url.slice(prefix.length - 1));
+  if (!f || !f.startsWith(root + path.sep) || !fs.existsSync(f) || !fs.statSync(f).isFile()) {
+    res.writeHead(404);
+    return res.end("not found");
+  }
+  res.writeHead(200, { "content-type": MIME[path.extname(f)] || "application/octet-stream" });
+  fs.createReadStream(f).pipe(res);
+});
+
+let browser;
+try {
+  if (!fs.existsSync(path.join(HERE, "pkg-platform", "wai_platform_decode.js"))) {
+    console.error("no pkg-platform/wai_platform_decode.js; build it as this file's header says");
+    process.exit(2);
+  }
+  if (objects.length === 0) {
+    console.error("no AV1 floor object found under pack-conformance/packs");
+    process.exit(2);
+  }
+  await new Promise((r) => server.listen(PORT, "127.0.0.1", r));
+  const channel = process.env.WAI_PLATFORM_BROWSER_CHANNEL || undefined;
+  browser = await chromium.launch({ channel, args: process.platform === "darwin" ? ["--use-angle=metal"] : [] });
+  console.log(`browser: ${channel || "chromium"} ${browser.version()}`);
+  const page = await browser.newPage();
+  const errs = [];
+  page.on("pageerror", (e) => errs.push(String(e)));
+  await page.goto(`http://127.0.0.1:${PORT}/demo/platform_decode.html`, { waitUntil: "domcontentloaded", timeout: 30000 });
+  await page.waitForFunction(() => window.platformDecode && window.platformDecode.done === true, null, { timeout: 5 * 60 * 1000, polling: 500 });
+  const report = await page.evaluate(() => window.platformDecode.report);
+  console.log(report);
+  for (const e of errs) console.log(`page error: ${e}`);
+  const lines = report.split("\n");
+  const fail = [];
+  if (errs.length || /page error/.test(report)) fail.push("a page error");
+  const objs = lines.filter((l) => l.startsWith("platform av1 object:"));
+  if (objs.length !== objects.length) fail.push(`${objs.length} AV1 objects reported, ${objects.length} served`);
+  for (const l of objs) {
+    const m = / decoded=(\d+)\/(\d+) /.exec(l);
+    if (!m || m[1] !== m[2] || / error=/.test(l)) fail.push(`not every unit decoded: ${l}`);
+  }
+  const hevc = lines.find((l) => l.startsWith("platform hevc encoder:"));
+  if (!hevc) fail.push("no HEVC line");
+  else if (!/not supported/.test(hevc)) {
+    const m = / decoded=(\d+)\/(\d+)/.exec(hevc);
+    if (!m || m[1] !== m[2] || / error=/.test(hevc) || !/ no-preference=true /.test(hevc)) fail.push(`the derived HEVC configuration did not decode the stream back: ${hevc}`);
+  }
+  const ord = lines.find((l) => l.startsWith("platform order: "));
+  if (!ord || !/ picked=1$/.test(ord)) fail.push(`the AV1 entry was not picked over the learned entry without a budget: ${ord}`);
+  const ob = lines.find((l) => l.startsWith("platform order with a budget:"));
+  if (!ob || !/ picked=0$/.test(ob)) fail.push(`the learned entry with a budget was not picked: ${ob}`);
+  for (const f of fail) console.log(`FAIL ${f}`);
+  if (fail.length) process.exitCode = 1;
+} finally {
+  if (browser) await browser.close();
+  server.close();
+}

wai/wai-web/src/lib.rs

@@ -159,6 +159,10 @@ pub mod quantum_qbom;
 // the separate `quantum_webgpu` bundle; the main bundle does not carry it.
 #[cfg(all(feature = "quantum_webgpu", target_arch = "wasm32"))]
 mod quantum_webgpu;
+// Platform decode capabilities in dispatch (extensions/platform-decode.md), in the
+// separate `platform_decode` bundle; the main bundle does not carry it.
+#[cfg(feature = "platform_decode")]
+mod platform_decode;
 
 // The job service's engine in the page: plan, run and sign a job in the tab, and
 // check any job's receipt there, the service's included. The IDE's bundle only.

wai/wai-web/src/platform_decode.rs

@@ -0,0 +1,56 @@
+//! Platform decode capabilities in dispatch (`wai::platform_decode`,
+//! extensions/platform-decode.md), bound for JS in the separate `platform_decode` bundle.
+//!
+//! [`platform_query`] gives the configuration a page asks the platform about for an entry
+//! (`decodingInfo()`'s and `isConfigSupported()`'s), read from the payload by the same Rust
+//! a native sink runs. [`platform_order`] takes the platform's answers and any energy
+//! budgets and returns the default order and the entry SPEC §7.1 picks. The page calls the
+//! platform itself; nothing here does, so the order is the reference's for those answers.
+
+use wai::container::ManifestV2;
+use wai::platform_decode as pd;
+use wasm_bindgen::prelude::*;
+
+fn js(e: impl std::fmt::Display) -> JsError {
+    JsError::new(&e.to_string())
+}
+
+/// The query for an entry of `capability` whose payload is `payload`, as JSON
+/// (`wai::platform_decode::query_json`); `rate_num`/`rate_den` is the presentation rate,
+/// for a payload that states none.
+#[wasm_bindgen(js_name = platformQuery)]
+pub fn platform_query(capability: &str, payload: &[u8], rate_num: Option<u32>, rate_den: Option<u32>) -> Result<String, JsError> {
+    let rate = match (rate_num, rate_den) {
+        (Some(num), Some(den)) if num > 0 && den > 0 => Some(pd::Rate { num, den }),
+        _ => None,
+    };
+    let q = pd::query(capability, payload, rate).map_err(js)?;
+    Ok(pd::query_json(&q).to_string())
+}
+
+/// The default order over a `WAI2` manifest's entries (`manifest_json`), from the
+/// platform's answers and the energy budgets, each a JSON array aligned with the entries
+/// (`null` where none), as JSON (`wai::platform_decode::order_json`). `ceiling` is the
+/// deployer's ceiling on a budget's figure, in microjoules, if any.
+#[wasm_bindgen(js_name = platformOrder)]
+pub fn platform_order(manifest_json: &str, answers_json: &str, budgets_json: &str, ceiling: Option<f64>) -> Result<String, JsError> {
+    let manifest: ManifestV2 = serde_json::from_str(manifest_json).map_err(js)?;
+    let arr = |text: &str, what: &str| -> Result<Vec<serde_json::Value>, JsError> {
+        match serde_json::from_str::<serde_json::Value>(text).map_err(js)? {
+            serde_json::Value::Array(a) => Ok(a),
+            _ => Err(JsError::new(&format!("{what}: not a JSON array"))),
+        }
+    };
+    let answers = arr(answers_json, "answers")?
+        .iter()
+        .map(|v| if v.is_null() { Ok(None) } else { pd::answers_from_json(v).map(Some) })
+        .collect::<Result<Vec<_>, String>>()
+        .map_err(js)?;
+    let budgets = arr(budgets_json, "budgets")?
+        .iter()
+        .map(|v| if v.is_null() { Ok(None) } else { pd::budget_from_json(v).map(Some) })
+        .collect::<Result<Vec<_>, String>>()
+        .map_err(js)?;
+    let rules = pd::Rules { budget_ceiling: ceiling.filter(|c| *c >= 0.0).map(|c| c as u64), ..pd::Rules::DEFAULT };
+    Ok(pd::order_json(&manifest, &answers, &budgets, &rules).to_string())
+}

wai/wgsl-conformance/.gitattributes

@@ -0,0 +1,2 @@
+# Compared as bytes: a line-ending conversion on checkout would change the file.
+* -text

wai/wgsl-conformance/README.md

@@ -0,0 +1,37 @@
+# wgsl-conformance — WGSL's integer semantics on a GPU
+
+The vectors an adapter passes before it runs a graph of a decode at `decode-equivalence`
+([`extensions/determinism-tiers.md`](../extensions/determinism-tiers.md) §5, *Claiming
+the tier on a GPU*, item 5): every integer operation WGSL defines on `i32`, `u32` and
+two-component vectors of them, at the cases WGSL defines at run time and the languages it
+is translated into leave open — wrapping addition, subtraction, multiplication and
+negation; division and remainder by zero and of the most negative `i32` by −1; shift
+amounts from 0 to 40 and past them; `clamp` with its bounds in and out of order;
+`abs`, `sign`, the bit counts and positions; bit fields at and past the word; the
+conversions between `i32`, `u32` and `bool`.
+
+**Provenance.** Written by `wai/wai-ops-gpu`'s `wai-wgsl-vectors` binary from
+`wai_ops_gpu::semantics::reference`, which states each result from the text of WGSL
+(W3C Candidate Recommendation Draft, 21 October 2026, §8.8, §8.10, §17.5), not from a
+shader. No case comes from a device.
+
+## Layout
+
+- `integer.jsonl` — one JSON object per operation, in the order of its `code`: `op`, its
+  name; `code`, the operation's number in `wai-ops-gpu/src/wgsl/semantics.wgsl`; `wgsl`,
+  the expression evaluated, with `a`, `b`, `c`, `d` the operands as `u32` and `ia` … `id`
+  their bits as `i32`; and `cases`, each `[a, b, c, d, result]` as eight lowercase
+  hexadecimal digits, followed by `"unguarded"` for a case the module of item 4 never
+  reaches (a shift amount of 32 or more, `clamp` bounds out of order). A vector
+  operation's result is its two components folded into one word,
+  `x ^ (y × 0x9E3779B9)` modulo 2^32.
+
+## Equivalence
+
+An adapter conforms when, for every case not marked `unguarded`, the word it computes
+for the operation on the operands, read from memory at run time, is `result`. Results on
+`unguarded` cases are reported, and do not decide conformance.
+
+The file is checked against the reference natively (`wai-ops-gpu/tests/semantics.rs`),
+under Node (`wai-ops-gpu/browser/semantics_reference.mjs`) and in a browser page
+(`wai-ops-gpu/browser/run.mjs`), the last two running the reference compiled to wasm32.