Measurement

Workload

Two 256 MiB buffers, touched in a pattern no cache can hold.

Every pass touches each element of one buffer exactly once; consecutive passes alternate buffers. No cache can hold the reuse distance.

Scope

Measured	Not measured
Effective throughput that WebGPU compute moves through GPU-visible memory, from this browser	Manufacturer peak bandwidth. A result shows it only as a published reference.
Scattered-access throughput and a dependent-load latency proxy	Memory capacity. Browsers do not expose VRAM size.
How much throughput holds across 18 s of continuous load	Who made the memory chips. HBM never infers a memory vendor.

Buffers

Working set: Two storage buffers of 256 MiB (power of two, 16 to 512 MiB). Smaller only if the adapter caps buffer sizes; such runs are unranked.
Element: vec4<u32>
Layout: Grid-stride, 4 elements per invocation, workgroups of 256. Neighbouring invocations touch neighbouring addresses.
Dispatch: 16,384 workgroups per pass at 256 MiB, 32,768 at 512 MiB (limit 65,535).
Initial data: Both buffers are filled with hashed values before the first phase, so no page is untouched or zero.

Alternation

Re-reading inside a dispatch lets small per-workgroup footprints hit in L2, Infinity Cache or a system-level cache and inflates the result. HBM never does it. Repetition happens across dispatches, and consecutive dispatches alternate between the two buffers: READ reads A then B, WRITE writes B then A, COPY copies A to B then B to A. Any byte is reused only after 512 MiB of other traffic, about four times the largest consumer last-level cache.

Kernels

Phase	Access pattern	Bytes counted per dispatch	Reported
Read	Four independent 16-byte loads per invocation, XOR-folded	256 MiB	Median GB/s
Write	Four 16-byte stores per invocation, data varies per element	256 MiB	Median GB/s
Copy	Load from one buffer, store to the other	512 MiB (read + write)	Median GB/s
Random	1M invocations, 16 gathers each, from both buffers at xorshift addresses	256 MiB of useful bytes	Median GB/s and G accesses/s
Latency	One invocation follows a random single cycle	Dependent loads	Median ns per load
Sustain	The copy kernel, continuously, for 18 s	512 MiB	Median of the last 60 %, retention

Loads stay observable because the folded result is compared with sentinels drawn from a uniform; a match is practically impossible, but the compiler cannot prove it, so no load is removed.

Sequential read, src/lib/bench/kernels.ts

@group(0) @binding(0) var<storage, read> src: array<vec4<u32>>;
@group(0) @binding(1) var<storage, read_write> sink: array<vec4<u32>>;
@group(0) @binding(2) var<uniform> p: Params;
@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let t = gid.x;
  let T = p.threads;
  let a = src[t];
  let b = src[t + T];
  let c = src[t + 2u * T];
  let d = src[t + 3u * T];
  let r = (a ^ b) ^ (c ^ d);
  if (r.x == p.magic0 && r.y == p.magic1) {
    sink[t & 255u] = r;
  }
}

Latency proxy

The chase buffer holds a single random cycle through 4M slots spaced 64 bytes apart, 256 MiB in total, built with Sattolo's algorithm on the CPU and scattered into place on the GPU. One invocation follows it; each address depends on the previous load, so the time per hop approximates a round trip to memory, TLB misses included. The position is carried across dispatches, so no pass ever replays a warm path. Hops per sample are calibrated to about 45 ms (90 ms on the wall clock).

Dependent-load chase

@group(0) @binding(0) var<storage, read> chain: array<u32>;
@group(0) @binding(1) var<storage, read_write> state: array<u32>;
@group(0) @binding(2) var<uniform> p: Params;
@compute @workgroup_size(1)
fn main() {
  var idx = state[0];
  for (var i = 0u; i < p.iters; i++) {
    idx = chain[idx];
  }
  state[0] = idx;
}

Sustain

Sustain repeats the copy kernel for 18 seconds. It reports the median of the last 60 % of samples, and retention: that median divided by the median of the first 20 %. A laptop that throttles after ten seconds shows it; a desktop card usually holds 100 %.