// Tile runs: one workgroup owns one tile of run.size consecutive amplitudes of
// the state and applies a run of gates to it in place, in op-log order, with a
// storage barrier between gates. Each gate either has its target inside the
// tile (tbit < size) or is diagonal (m01 = m10 = 0) and acts per amplitude.
// Uses fx64, add64, sub64, fits, row and the bindings re, im, report of
// lanes32.wgsl, which this file is appended to. No shift amount here is
// anything but a literal; sizes and indices come multiplied, not shifted. No
// workgroup memory: the tile stays in re and im, which no other workgroup of
// the dispatch touches.
//
// Uniformity: the early returns depend only on run (a uniform buffer),
// workgroup_id, num_workgroups and the state planes' lengths. The barrier sits
// at the top level of the r loop, whose bound run.count is a uniform-buffer
// value; no break, continue or return occurs inside that loop, and no branch
// that depends on a gate or an invocation holds a barrier.
struct TileGate {
row0: vec4<u32>, // m00.re, m00.im, m01.re, m01.im
row1: vec4<u32>, // m10.re, m10.im, m11.re, m11.im
tbit: u32, // 1 << target
ctrl: u32, // the controls' bits, OR-ed
diag: u32, // 1: m01 = m10 = (0, 0)
index: u32, // the gate's position in the op-log
}
struct TileRun {
first: u32, // the run's first record in tile_gates
count: u32, // gates in the run, at least 1
size: u32, // amplitudes per tile: a power of two, at most TILE_MAX
x_groups: u32, // workgroups along x
tiles: u32, // tiles in the state: 2^n / size
pad0: u32,
pad1: u32,
pad2: u32,
}
@group(0) @binding(6) var<storage, read> tile_gates: array<TileGate>;
@group(0) @binding(7) var<uniform> run: TileRun;
const TILE_MAX: u32 = 1024u;
const TILE_LANES: u32 = 256u;
const NONE: u32 = 0xffffffffu;
const ZERO2: vec2<u32> = vec2<u32>(0u, 0u);
@compute @workgroup_size(256)
fn apply_tile(@builtin(workgroup_id) wid: vec3<u32>,
@builtin(num_workgroups) groups: vec3<u32>,
@builtin(local_invocation_index) lid: u32) {
let size = run.size;
let span = run.tiles * size; // 2^n
// A malformed run record, or one whose tiles or grid do not cover the state
// exactly once, is reported and never run (uniform: run is a uniform buffer,
// num_workgroups a uniform built-in, and a buffer's length is the same for
// every invocation).
if (size == 0u || size > TILE_MAX || (size & (size - 1u)) != 0u || run.count == 0u
|| run.tiles == 0u || span != arrayLength(&re) || span != arrayLength(&im)
|| run.x_groups == 0u || groups.x != run.x_groups || groups.z != 1u
|| groups.x * groups.y < run.tiles) {
atomicOr(&report[0], 4u);
return;
}
let tile = wid.y * run.x_groups + wid.x;
if (tile >= run.tiles) { return; } // workgroup-uniform
let base = tile * size;
// A tile outside the state is reported, and nothing of it is touched.
let inside = base + size <= arrayLength(&re) && base + size <= arrayLength(&im);
if (!inside) { atomicOr(&report[0], 2u); }
var bad = NONE;
var malformed = false;
for (var r = 0u; r < run.count; r++) {
// A record outside tile_gates reads as malformed, never out of bounds.
let at = run.first + r;
let held = at < arrayLength(&tile_gates);
let h = tile_gates[select(0u, at, held)];
let shape = held && h.tbit != 0u && (h.tbit & (h.tbit - 1u)) == 0u && h.tbit < span
&& (h.ctrl & h.tbit) == 0u && (h.ctrl & ~(span - 1u)) == 0u;
var ok = true;
if (shape && h.diag == 1u) {
for (var k = 0u; k < 4u; k++) {
let e = lid + TILE_LANES * k;
let gi = base + e;
if (inside && e < size && (gi & h.ctrl) == h.ctrl) {
let a = vec2<u32>(re[gi], im[gi]);
var o: vec4<u32>;
if ((gi & h.tbit) == 0u) { o = row(h.row0, a, ZERO2); } else { o = row(h.row1, ZERO2, a); }
re[gi] = o.y;
im[gi] = o.w;
if (!(fits(o.xy) && fits(o.zw))) { ok = false; }
}
}
} else if (shape && h.tbit < size) {
for (var k = 0u; k < 2u; k++) {
let q = lid + TILE_LANES * k;
if (inside && q < (size >> 1u)) {
let low = q & (h.tbit - 1u);
let gi = base + (((q ^ low) << 1u) | low); // q with a 0 inserted at the target bit
let gj = gi | h.tbit;
if ((gi & h.ctrl) == h.ctrl) {
let a0 = vec2<u32>(re[gi], im[gi]);
let a1 = vec2<u32>(re[gj], im[gj]);
let o0 = row(h.row0, a0, a1);
let o1 = row(h.row1, a0, a1);
re[gi] = o0.y;
im[gi] = o0.w;
re[gj] = o1.y;
im[gj] = o1.w;
if (!(fits(o0.xy) && fits(o0.zw) && fits(o1.xy) && fits(o1.zw))) { ok = false; }
}
}
}
} else {
malformed = true;
}
if (!ok) { bad = select(bad, h.index, h.index < bad); }
storageBarrier(); // this gate's writes before the next gate's reads
}
if (malformed) { atomicOr(&report[0], 4u); }
if (bad != NONE) {
atomicOr(&report[0], 1u);
atomicMin(&report[1], bad);
}
}