Der Op, den wir nie ausgeliefert haben

Angenommen, du brauchst einen gleitenden Median. Vor dir: eine eindimensionale Reihe, rund eine Million Samples, mit gelegentlichen Ein-Sample-Spitzen. Ein Mittelwertfilter tut dort genau das Falsche — er verschmiert jede Spitze über ihre Nachbarn, statt sie zu entfernen. Das robuste, langweilige, richtige Werkzeug ist ein Median über ein gleitendes Fenster. Du öffnest den Catalog, suchst — und findest nichts. Er ist schlicht kein Built-in.

Die Abkürzung, die man sich verbieten sollte#

Der Maintainer könnte das auf dem Maintainer-Weg beheben. Eine Definitionsdatei, eine Referenzimplementierung, ein Platz im Built-in-Satz — bis zum Abend wäre windowed_median ein Built-in, und niemand wüsste je, dass er gefehlt hat. Bewiesen hätte das: nichts. Die tiefste Aussage, die Scellis über Erweiterbarkeit trifft, ist diese: Der Maintainer behält genau drei Grammatiken — die Shape-Signatur-Sprache, den Buffer-Binding-Deskriptor, die Backward-Repräsentation — dazu den Checker und das Referenz-Regime; jede Instanz erstellen die Nutzer. Solche Aussagen verfallen, wenn niemand sie von außen auf die Probe stellt. Also: den Op so erstellen, wie es jeder Nutzer täte — über die Authoring-Formulare im Studio, als Content, auf dem eigenen Gerät, ohne Engine-Patch und ohne Maintainer-Schlüssel. Der Authoring-Guide geht exakt diesen Weg.

Erst die Form: was der Op verspricht#

Ein Op ist hier kein Funktionskörper, sondern ein Vertrag mit angehängtem Körper. Also zuerst die Signatur: Eingabe ein 1-D-f32-Tensor der Länge n; ein Parameter w, eine ungerade Fensterbreite, begrenzt auf 31; Ausgabe der Länge n − w + 1. Die Output-Shape-Klasse ist static — allein aus Shapes und Parametern berechenbar —, damit der Planner Speicher budgetieren kann, bevor irgendetwas läuft. Dann der interessante Teil: die Toleranzklasse. Du deklarierst exact — und musst es begründen. Ein Median mit ungeradem Fenster rechnet auf seinen Werten überhaupt nicht: Er vergleicht und wählt aus, und die Ausgabe ist eine der Eingaben, Bit für Bit. Der Checker verlangt, dass die deklarierte Klasse zur semantischen Kategorie des Ops passt: Eine Matrixmultiplikation kann sich nicht mit einer loose-Toleranz an der Konformität vorbeimogeln — und nach derselben Regel darf reine Selektion ehrlich Exaktheit beanspruchen. Ein gerades Fenster würde die beiden mittleren Elemente mitteln — ein Rundungsschritt — und müsste anders deklarieren; die Beschränkung auf ungerade Fenster erhält die saubere Aussage.

Die Referenz ist der Vertrag#

Für User-Compute gilt dieselbe Regel wie für die Built-ins: gegen eine Referenz geprüft, nie selbst-attestiert. Die Referenzimplementierung kommt also vor dem Kernel — pro Ausgabeindex: w Werte einsammeln, sortieren, die Mitte nehmen — und sie läuft auf dem CPU-Referenzpfad, an dem die Engine alle numerische Wahrheit verankert. Dein Kernel ist nicht korrekt, weil du das behauptest; er ist genau insoweit korrekt, wie er diese Referenz trifft, innerhalb der deklarierten Toleranz, auf Fällen, die du dir nicht aussuchen darfst. Und die Paarung ist über deine Session hinaus etwas wert: Referenzimplementierung und Konformitätskorpus liegen im permissiv lizenzierten Interchange-Ring — „trifft seine Referenz“ ist eine Aussage, die jeder selbst nachrechnen kann. Kein Vertrau-mir von einem Anbieter.

Der Kernel — und was die Suite fängt#

Das WGSL selbst ist unspektakulär, und genau das ist ein Feature. Eine Invocation pro Ausgabeelement; das Fenster in ein lokales Array fester Größe einsammeln (die Schranke von 31 hält das ehrlich); Insertion-Sort; die Mitte nehmen. Das Binding-Layout — zwei Storage-Buffer und ein Uniform — ist in der Buffer-Binding-Grammatik geschrieben, die der Plattform gehört. Alles innerhalb der Klammern gehört dir.

// 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];
}

Dann die Konformität. Die Suite erzeugt ihre Fälle selbst — Shapes, Fenster, Randeingaben —, denn ein Autor, der seine eigene Prüfung benotet, ist Selbst-Attestierung mit Umwegen. Und der Fall, der den klassischen Erstentwurfs-Bug fängt, ist nichts Exotisches: die kleinste zulässige Eingabe. Bei n = 7 und w = 7 gibt es genau ein Ausgabeelement — ein Guard, der i >= n − w liest statt i >= n − w + 1, schreibt also gar nichts. Das + 1u oben ist der Fix. Trifft der Kernel die Referenz auf jedem erzeugten Fall — wobei exact gleich bedeutet, nicht ungefähr —, prägt die Engine ein Konformitäts-Receipt: hash-gebunden an den Kernel, gescoped auf das GPU-Profil, auf dem der Beweis lief. Dieses Scoping ist der ehrliche Teil: Vertrauen gilt auf der Umgebung, die den Beweis ausgeführt hat; eine andere Maschine konformiert lokal nach, statt deins zu importieren. Und nichts davon steht je zwischen dir und deiner eigenen GPU — ein unkonformierter Kernel läuft auf dem eigenen Gerät ab dem ersten Compile. Das Receipt ist, was den Op weitergabefähig macht.

Subgradienten, ehrlich#

Jetzt soll der Despike-Schritt in die Pipeline, mit trainierbaren Teilen davor — also müssen Gradienten durch den Median hindurch. Dafür gibt es eine Leiter, und ihre Sprossen sind ehrlich. Ist ein Op eine Komposition bestehender differenzierbarer Ops, leitet Autodiff das Backward gratis ab — und fairerweise: Ein gleitender Median lässt sich so komponieren, Fenster entfalten, jedes per Median reduzieren. Aber das Entfalten materialisiert ein Zwischenergebnis von n mal w — das w-Fache an Speicher —, und der ganze Grund für diesen Op ist ein flaches Budget. Also: ein eigener Kernel, und die mittlere Sprosse der Leiter — ein deklariertes Backward, per finiter Differenzen gegen den eigenen Forward geprüft. Die unterste Sprosse ist die ehrliche Grenze: Du deklarierst den Op als nicht differenzierbar, und die Engine weigert sich, durch ihn zu trainieren, statt so zu tun, als ob.

Die VJP ist die nahe Verwandte des Max-Poolings: Der Median ist eine der Eingaben, also fließt der ankommende Kotangens auf den Argmedian-Index — und ist überall sonst null. Zwischen Ordnungswechseln ist die Funktion stückweise linear, und dieser Gradient ist exakt — nicht ungefähr richtig, exakt. An einem Gleichstand ist sie überhaupt nicht differenzierbar, und genau das muss im Record stehen: Deklariert wird eine Subgradienten-Konvention — route zum ausgewählten Element —, festgehalten in den Metadaten des Ops, nicht in einem Kommentar vergraben. Gradcheck wendet dieselbe Ehrlichkeit auf sich selbst an. Eine Probe, die innerhalb der Schrittweite eines Ordnungswechsels landet, sitzt rittlings auf einem Knick, an dem keine Ableitung existiert — der Harness schließt sie aus, statt über die Falte hinweg zu mitteln. Sein Report hat diese Form:

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

Zwei Zeilen in diesem Transkript tragen das Gewicht. Zentrale Differenzen haben auf einer stückweise linearen Funktion keinen Trunkierungsfehler; abseits der Knicke stimmen analytischer und numerischer Gradient bis aufs f64-Runden überein — die Übereinstimmung ist kein Glück, sie ist die Gestalt der Funktion. Und die Negativkontrolle ist keine Dekoration: Der Harness läuft erneut mit einer absichtlich verfälschten VJP und verlangt, dass die Prüfung durchfällt — denn eine Prüfung, die nicht durchfallen kann, bescheinigt nichts.

Veröffentlicht — und ununterscheidbar#

Das Veröffentlichen ist der undramatischste Schritt — und genau das ist der Punkt des ganzen Designs. Teilst du den Op mit einem Workspace, hat er die einzige Schranke schon passiert, auf die es ankommt: Nichts verlässt private ohne angehängte SPDX-Lizenz — denn was du teilst, kann eine Kollegin forken, und ungelizenzierte geteilte Arbeit ist rechtlich vollständig vorbehalten, was das Fork-Schwungrad still aushebeln würde, auf dem diese Plattform steht. Der Gang an die Öffentlichkeit verlangt obendrauf mehr: die volle Metadaten-Messlatte, dieselbe, die auch der eingebaute Content nimmt. Die Karte trägt ihre Herkunft ohne Entschuldigung — Quelle user, Vertrauen unverified, bis es verdient ist, der exact-Chip, das Receipt.

Dann setzt du den Op zwischen zwei Built-in-Schritte in den Workflow und lässt laufen. Hier ist die Pointe, und sie ist keine Redewendung: Zur Dispatch-Zeit kann die Engine den Op eines Nutzers nicht von einem Built-in unterscheiden. Nicht „behandelt ihn fair“ — kann ihn nicht unterscheiden. Ein Nachschlagen, eine Validierung, ein Ausführungspfad; der Herkunfts-Record des Laufs pinnt den erstellten Op per Content-Hash exakt so, wie er die Built-in-Ops links und rechts davon pinnt. Die Spitzen sind weg. Die Pipeline trainiert durch den Median hindurch, Gradienten inklusive.

Der Op, den niemand ausgeliefert hat, erscheint trotzdem — erstellt, gegen die Referenz geprüft, Receipt in der Hand, veröffentlicht von einem Nutzer. Dass dieser erste Nutzer der Maintainer sein mag, ist ein Detail, das die Engine strukturell gar nicht bemerken kann. Genau darum ist das Ganze gebaut.