The op we didn't ship

Suppose you need a windowed median. In front of you: a one-dimensional series, about a million samples long, with occasional single-sample spikes. A mean filter does exactly the wrong thing there — it smears each spike across its neighbors instead of removing it. The robust, boring, correct tool is a median over a sliding window. You open the Catalog, search, and find nothing. It is simply not a builtin.

The shortcut worth refusing#

The maintainer could fix that the maintainer way. A definition file, a reference implementation, a place in the builtin set — by evening, windowed_median would be a builtin and nobody would ever know it had been missing. It would also have proven nothing. The deepest claim Scellis makes about extensibility is that the maintainer keeps only three grammars — the shape-signature language, the buffer-binding descriptor, the backward representation — plus the checker and the reference regime, and that users author every instance. Claims like that decay unless someone exercises them from the outside. So: author the op the way any user does — through the authoring forms in the Studio, as content, on your own device, with no engine patch and no maintainer key. The authoring guide walks exactly this path.

Shape first: what the op promises#

An op here is not a function body; it is a contract with a body attached. So the signature comes first: input, a 1-D f32 tensor of length n; one parameter w, an odd window width capped at 31; output, length n − w + 1. The output shape class is static — computable from shapes and parameters alone — so the planner can budget memory before anything runs. Then the interesting part: the tolerance class. Declare exact, and you are required to justify it. An odd-window median performs no arithmetic on its values at all — it compares and selects, and the output is one of the inputs, bit for bit. The checker demands that a declared class match the op's semantic category: a matmul cannot claim a loose tolerance to dodge conformance, and by the same rule a pure selection may honestly claim exactness. An even window would average the two middle elements — one rounding step — and would have to declare differently; restricting the signature to odd windows keeps the clean claim.

The reference is the contract#

The rule for user compute is the same as for the built-ins: checked against a reference, never self-attested. So the reference implementation comes before the kernel — for each output index, gather w values, sort, take the middle — and it runs on the CPU reference path, where the engine anchors numeric truth. Your kernel is not correct because you say so; it is correct exactly insofar as it matches that reference, within the declared tolerance, on cases you do not get to pick. And the pairing is worth something beyond your own session: the reference implementation and the conformance corpus live in the permissively licensed interchange ring, so “matches its reference” is a claim anyone can re-derive — not a trust-me from a vendor.

The kernel, and what the suite catches#

The WGSL itself is unglamorous, which is a feature. One invocation per output element; gather the window into a fixed-size local array (the cap of 31 is what keeps that honest); insertion-sort; take the middle. The binding layout — two storage buffers and a uniform — is written in the buffer-binding grammar the platform owns. Everything inside the braces is yours.

// user/kemal/windowed_median@1 -- odd window w <= 31, one thread per output
@group(0) @binding(0) var<storage, read>       x   : array<f32>;
@group(0) @binding(1) var<storage, read_write> dst : array<f32>;

struct Params { n: u32, w: u32 }
@group(0) @binding(2) var<uniform> p: Params;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= p.n - p.w + 1u) { return; }  // the "+ 1u" is the fix the suite forces

  var win: array<f32, 31>;
  for (var k = 0u; k < p.w; k++) { win[k] = x[i + k]; }

  // insertion sort: pure selection -- no arithmetic ever touches the values
  for (var a = 1u; a < p.w; a++) {
    let v = win[a];
    var b = a;
    while (b > 0u && win[b - 1u] > v) { win[b] = win[b - 1u]; b--; }
    win[b] = v;
  }
  dst[i] = win[p.w / 2u];
}

Then conformance. The suite generates its cases itself — shapes, windows, edge inputs — because an author grading their own exam is self-attestation with extra steps. And the case that catches the classic first-draft bug is not exotic; it is the smallest legal input. With n = 7 and w = 7 there is exactly one output element, so a guard written i >= n − w instead of i >= n − w + 1 writes nothing at all. The + 1u above is the fix. Once the kernel matches the reference on every generated case — where exact means equal, not close — the engine mints a conformance receipt, hash-bound to the kernel and scoped to the GPU profile the proof ran on. That scoping is the honest part: trust holds on the envelope that ran the proof; another machine re-conforms locally rather than importing yours. And none of this ever stands between you and your own GPU — an unconformed kernel runs on your own device from the first compile. The receipt is what lets the op travel.

Subgradients, honestly#

Now put the despike step inside a pipeline, with trainable pieces upstream of it — which means gradients must pass through the median. There is a ladder for that, and its rungs are honest. If an op is a composition of existing differentiable ops, autodiff derives the backward for free — and to be fair, a windowed median can be composed that way: unfold the windows, reduce each with a median. But the unfold materializes an n-by-w intermediate — w times the memory — and the whole reason this op exists is a flat budget. So: a dedicated kernel, and the middle rung of the ladder — a declared backward, verified against your own forward by finite differences. The bottom rung is the honest boundary: declare the op non-differentiable, and the engine refuses to train through it rather than pretend.

The VJP is max-pooling's close cousin: the median is one of the inputs, so the incoming cotangent routes to the argmedian index and is zero everywhere else. Between order changes the function is piecewise linear, and that gradient is exact — not approximately right, exact. At a tie it is not differentiable at all, and the record must say so: what you declare is a subgradient convention — route to the selected element — stated in the op's metadata rather than buried in a comment. Gradcheck applies the same honesty to itself. A probe that lands within the step size of an order crossing straddles a kink where no derivative exists, so the harness excludes it instead of averaging across the crease. Its report has this shape:

gradcheck  user/kemal/windowed_median@1   (w = 7, reference path, f64)
  forward  : author reference (gather -> sort -> middle)
  vjp      : declared backward -- cotangent routed to the argmedian index
  probes   : sampled (input, output) pairs, central differences, h = 1e-3
  skipped  : probes within h of an order crossing (no derivative exists there)
  bound    : max |analytic - numeric| must stay under the declared tolerance
  negative control : a corrupted vjp must be REJECTED -- a check that cannot
                     fail certifies nothing

Two lines in that transcript carry the weight. Central differences on a piecewise-linear function have no truncation error, so away from kinks the analytic and numeric gradients agree down to f64 rounding — the agreement is not luck, it is the shape of the function. And the negative control is not decoration: the harness re-runs with a deliberately corrupted VJP and requires the check to fail, because a check that cannot fail certifies nothing.

Published, and indistinguishable#

Publishing is the least dramatic step, which is the point of the whole design. Share the op with a workspace and it has already crossed the only gate that matters: nothing leaves private without an SPDX license attached — because anything you share is something a colleague may fork, and unlicensed shared work is legally all-rights-reserved, which would quietly void the fork flywheel the platform is built on. Going public asks for more on top: the full metadata bar, the same bar the built-in content clears. The card wears its provenance without apology — source user, trust unverified until it is earned, the exact tolerance chip, the receipt.

Then slot it into the workflow between two builtin steps and run. Here is the punchline, and it is not a figure of speech: at dispatch time the engine cannot tell a user's op from a builtin. Not “treats it fairly” — cannot tell. One lookup, one validation, one execution path; the run's provenance record pins the authored op by content hash exactly as it pins the builtin ops on either side of it. The spikes are gone. The pipeline trains through the median, gradients and all.

The op nobody shipped ships anyway — authored, checked against the reference, receipt in hand, published by a user. That the first such user might be the maintainer is a detail the engine is structurally incapable of noticing. That is the whole point.