File size: 12,616 Bytes
9ae1216
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
54e1054
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
9ae1216
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
f295800
 
 
 
9ae1216
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
// Mutation-scoring the ORACLES, using a bug taxonomy someone else wrote.
//
// Bugs ported from dipankarsarkar/gpuemu-corpus (the 10 buggy variants). Their
// kernels are Triton/numpy and mine are WGSL/JS, so nothing runs across. What
// ports is the taxonomy: each bug is specified, so it can be re-injected into a
// different codebase. The point is that the bug list has a different author than
// the checks, which is the one thing I cannot manufacture myself.
//
// The question is NOT "do my kernels pass". It is "which of my checks catch a
// bug I did not think of". Every bug an oracle waves through is a hole with a
// name on it instead of a feeling about my own coverage.
const fs = require("fs");
const path = require("path");
const V = require("./public/verified_core.js");
const L = { mul: new Int16Array(fs.readFileSync(path.join(__dirname, "public", "mul_lut.bin")).buffer.slice(0)) };

const f32 = Math.fround;
const randf = (n) => Float32Array.from({ length: n }, () => Math.random() * 2 - 1);

// ---- MY kernel, with one corpus bug injected at a time ----------------------
// Faithful to the real bgemmJS/WGSL structure, so the injected bug is the only
// difference. `bug` names map to corpus variants; see TAXONOMY below.
function kernel(bug) {
  return (Xq, Wq, rs, cs, d) => {
    const { m, k, n } = d, batch = d.batch || 1, relu = !!d.relu, mul = L.mul, raw = !!d.acc;
    const out = raw ? new Int32Array(batch * m * n) : new Float32Array(batch * m * n);
    const acc = new Int32Array(n);
    for (let bz = 0; bz < batch; bz++) {
      const xo = bz * m * k, wo = bz * k * n, oo = bz * m * n, co = bz * n;
      for (let i = 0; i < m; i++) {
        acc.fill(0);
        const xrow = xo + i * k;
        for (let p = 0; p < k; p++) {
          const au = (Xq[xrow + p] & 0xFF) * 256, wrow = wo + p * n;
          for (let j = 0; j < n; j++) {
            const prod = mul[au + (Wq[wrow + j] & 0xFF)];
            // corpus: matmul_triton_buggy — acc= instead of acc+=, "correct only when K=1"
            if (bug === "accOverwrite") acc[j] = prod; else acc[j] += prod;
          }
        }
        const orow = oo + i * n;
        if (raw) {
          for (let j = 0; j < n; j++) out[orow + j] = acc[j];
          continue;
        }
        const rscale = rs[bz * m + i];
        for (let j = 0; j < n; j++) {
          // corpus: softmax_llm_buggy — "correct only when last dim is a multiple
          // of 128". My analogue: workgroup_size(8,8), so a missing bounds guard
          // is correct only when n is a multiple of 8.
          if (bug === "boundsMult8" && j >= (n - (n % 8))) { out[orow + j] = 0; continue; }
          let v = V.epi(acc[j], rscale, cs[co + j]);
          // corpus: gelu_triton_buggy — a dropped constant factor, uniformly ~2x off
          if (bug === "droppedFactor") v = f32(v * 2);
          // corpus: leaky_relu_triton_buggy — wrong alpha (0.1 vs 0.01)
          if (relu && bug === "leakyAlpha") v = v < 0 ? f32(v * 0.1) : v;
          else if (relu && v < 0) v = 0;
          out[orow + j] = v;
        }
      }
    }
    return out;
  };
}

// what each injected bug is, and where it lives
const TAXONOMY = {
  accOverwrite:  { from: "matmul_triton_buggy",  where: "loop", desc: "acc= instead of acc+= (only correct at K=1)" },
  boundsMult8:   { from: "softmax_llm_buggy",    where: "loop", desc: "correct only when n is a multiple of the workgroup width" },
  droppedFactor: { from: "gelu_triton_buggy",    where: "math", desc: "constant factor dropped, uniformly 2x" },
  leakyAlpha:    { from: "leaky_relu_buggy",     where: "math", desc: "wrong negative-slope alpha" },
};

// ---- oracle 1: the property suite (no reference implementation) -------------
const SHAPES = { m: 6, k: 32, n: 5, batch: 1 };
function quantize(Xf, Wf, d) {
  const batch = d.batch || 1;
  const x = V.quantizeRows(Xf, batch * d.m, d.k);
  let wq, ws;
  if (batch === 1) { const w = V.quantizeCols(Wf, d.k, d.n); wq = w.q; ws = w.s; }
  else {
    wq = new Int8Array(batch * d.k * d.n); ws = new Float32Array(batch * d.n);
    for (let bz = 0; bz < batch; bz++) {
      const w = V.quantizeCols(Wf.subarray(bz * d.k * d.n, (bz + 1) * d.k * d.n), d.k, d.n);
      wq.set(w.q, bz * d.k * d.n); ws.set(w.s, bz * d.n);
    }
  }
  return { x, wq, ws };
}
const call = (K, Xf, Wf, d) => { const q = quantize(Xf, Wf, d); return K(q.x.q, q.wq, q.x.s, q.ws, d); };

function propertySuite(K) {
  // NON-TRIVIALITY. Added after the corpus scored my suite 0/4: a kernel that
  // returns all zeros satisfies every relation below, because zero is
  // permutation-equivariant, zero-row-preserving and batch-decomposable. Every
  // algebraic property suite needs this or it can be passed by doing nothing.
  {
    const d = { ...SHAPES };
    const A = randf(d.m * d.k), B = randf(d.k * d.n);
    const o = call(K, A, B, d);
    let nz = 0;
    for (let i = 0; i < o.length; i++) if (o[i] !== 0) nz++;
    if (nz < o.length / 4) return "nonTriviality";
  }
  // SENSITIVITY. Every element of A must be able to move its output row. This is
  // what catches an accumulator that overwrites instead of accumulating: the
  // structure is perfect, every relation holds, but only the last k contributes.
  //
  // Measured against the RAW accumulator, and perturbed by a sign flip that
  // leaves |value| untouched. Both matter: block scaling coupled the row absmax
  // to every output in the row, so a naive perturbation moves the output through
  // the SCALE and the test passes while proving nothing about the accumulation.
  // That is how the first version of this property fooled me.
  {
    const d = { ...SHAPES, acc: true };
    const A = randf(d.m * d.k), B = randf(d.k * d.n);
    for (let p = 0; p < d.k; p++) A[1 * d.k + p] = 0.3;
    A[1 * d.k + 0] = 1.0;                          // pin the row absmax at p=0
    const q0 = quantize(A, B, d);
    const base = K(q0.x.q, q0.wq, q0.x.s, q0.ws, d);
    for (let p = 1; p < d.k; p++) {
      const A2 = Float32Array.from(A);
      A2[1 * d.k + p] = -0.3;                      // sign flip: absmax unchanged
      const q2 = quantize(A2, B, d);
      if (q2.x.s[1] !== q0.x.s[1]) continue;       // scale moved anyway: inconclusive
      const o2 = K(q2.x.q, q2.wq, q2.x.s, q2.ws, d);
      let moved = false;
      for (let j = 0; j < d.n; j++) if (o2[1 * d.n + j] !== base[1 * d.n + j]) { moved = true; break; }
      if (!moved) return `sensitivity(A[.,${p}] ignored)`;
    }
  }
  // RELU RANGE. relu(x) >= 0 is part of the kernel's DEFINITION when the
  // fused ReLU is on — a range constraint, not a relation between calls.
  // This is what catches a wrong negative slope: the structure survives a
  // leak, the sign does not.
  {
    const d = { ...SHAPES, relu: true };
    const A = randf(d.m * d.k), B = randf(d.k * d.n);
    const o = call(K, A, B, d);
    for (let i = 0; i < o.length; i++) if (o[i] < 0) return "reluRange";
  }
  // UNIT-SCALE ANCHOR. With rs = cs = 1 the dequant is the identity, so the
  // definition pins ABSOLUTE values: out must equal the exact integer dot
  // product (f32-exact far below 2^24). No RELATION can do this — if out
  // satisfies every relation, so does c*out — so the suite needs one point
  // where the spec fixes the scale. Expected values come from plain integer
  // arithmetic: no reference implementation, no LUT, no mirror.
  {
    const d = { m: 3, k: 5, n: 4, batch: 1 };
    const Xq = new Int8Array(d.m * d.k), Wq = new Int8Array(d.k * d.n);
    for (let i = 0; i < Xq.length; i++) Xq[i] = ((i * 37 + 11) % 25) - 12;
    for (let i = 0; i < Wq.length; i++) Wq[i] = ((i * 53 + 7) % 25) - 12;
    const rs = new Float32Array(d.m).fill(1), cs = new Float32Array(d.n).fill(1);
    const o = K(Xq, Wq, rs, cs, d);
    for (let i = 0; i < d.m; i++)
      for (let j = 0; j < d.n; j++) {
        let dot = 0;
        for (let p = 0; p < d.k; p++) dot += Xq[i * d.k + p] * Wq[p * d.n + j];
        if (o[i * d.n + j] !== dot) return "unitScaleAnchor";
      }
  }
  // zero row of A -> zero row out
  {
    const d = { ...SHAPES };
    const A = randf(d.m * d.k), B = randf(d.k * d.n);
    for (let p = 0; p < d.k; p++) A[2 * d.k + p] = 0;
    const o = call(K, A, B, d);
    for (let j = 0; j < d.n; j++) if (o[2 * d.n + j] !== 0) return "zeroRow";
  }
  // permuting rows of A permutes output rows
  {
    const d = { ...SHAPES };
    const A = randf(d.m * d.k), B = randf(d.k * d.n);
    const perm = [3, 1, 5, 0, 4, 2];
    const Ap = new Float32Array(A.length);
    perm.forEach((src, dst) => Ap.set(A.subarray(src * d.k, (src + 1) * d.k), dst * d.k));
    const o = call(K, A, B, d), op = call(K, Ap, B, d);
    for (let r = 0; r < d.m; r++) for (let j = 0; j < d.n; j++)
      if (op[r * d.n + j] !== o[perm[r] * d.n + j]) return "rowPermutation";
  }
  // permuting columns of B permutes output columns
  {
    const d = { ...SHAPES };
    const A = randf(d.m * d.k), B = randf(d.k * d.n);
    const perm = [2, 0, 4, 1, 3];
    const Bp = new Float32Array(B.length);
    for (let p = 0; p < d.k; p++) perm.forEach((src, dst) => { Bp[p * d.n + dst] = B[p * d.n + src]; });
    const o = call(K, A, B, d), op = call(K, A, Bp, d);
    for (let r = 0; r < d.m; r++) for (let j = 0; j < d.n; j++)
      if (op[r * d.n + j] !== o[r * d.n + perm[j]]) return "colPermutation";
  }
  // batched == each element on its own
  {
    const d = { m: 4, k: 32, n: 5, batch: 3 };
    const A = randf(d.batch * d.m * d.k), B = randf(d.batch * d.k * d.n);
    const together = call(K, A, B, d);
    for (let bz = 0; bz < d.batch; bz++) {
      const alone = call(K, A.subarray(bz * d.m * d.k, (bz + 1) * d.m * d.k),
                            B.subarray(bz * d.k * d.n, (bz + 1) * d.k * d.n), { ...d, batch: 1 });
      for (let i = 0; i < alone.length; i++)
        if (alone[i] !== together[bz * d.m * d.n + i]) return "batchDecomposition";
    }
  }
  return null;
}

// ---- oracle 2: the exact differential gate (kernel vs the JS mirror) --------
function differentialGate(K) {
  for (const d of [{ m: 5, k: 9, n: 6, batch: 3, relu: true },
                   { m: 32, k: 64, n: 32, batch: 1, relu: false },
                   { m: 7, k: 253, n: 5, batch: 2, relu: true },
                   { m: 17, k: 33, n: 9, batch: 1, relu: true }]) {
    const A = randf(d.batch * d.m * d.k), B = randf(d.batch * d.k * d.n);
    const q = quantize(A, B, d);
    for (const acc of [true, false]) {
      const dd = { ...d, acc };
      const hw = K(q.x.q, q.wq, q.x.s, q.ws, dd);
      const ref = V.bgemmJS(q.x.q, q.wq, q.x.s, q.ws, dd, L);
      for (let i = 0; i < ref.length; i++) if (hw[i] !== ref[i]) return (acc ? "accumulator" : "epilogue");
    }
  }
  return null;
}

// the oracles are reusable: test_selfcorpus.js scores them against MY OWN bugs
module.exports = { propertySuite, differentialGate, quantize, call, kernel };
if (require.main !== module) return;

// ---- run --------------------------------------------------------------------
console.log("\nBugs from dipankarsarkar/gpuemu-corpus, re-injected into my kernel.");
console.log("Scoring the ORACLES, not the kernel.\n");
console.log("bug              lives in   properties        differential");
console.log("---------------------------------------------------------------");
const score = { prop: 0, diff: 0, total: 0 };
for (const [bug, meta] of Object.entries(TAXONOMY)) {
  const K = kernel(bug);
  const p = propertySuite(K), dgate = differentialGate(K);
  score.total++;
  if (p) score.prop++;
  if (dgate) score.diff++;
  console.log(`${bug.padEnd(16)} ${meta.where.padEnd(10)} ${(p ? "CAUGHT (" + p + ")" : "MISSED").padEnd(17)} ${dgate ? "CAUGHT (" + dgate + ")" : "MISSED"}`);
}
// control: the real kernel must pass both
const cleanP = propertySuite(kernel(null)), cleanD = differentialGate(kernel(null));
console.log(`${"(no bug)".padEnd(16)} ${"-".padEnd(10)} ${(cleanP ? "FALSE POSITIVE" : "clean").padEnd(17)} ${cleanD ? "FALSE POSITIVE" : "clean"}`);

console.log(`\nmutation score, properties:   ${score.prop}/${score.total}`);
console.log(`mutation score, differential: ${score.diff}/${score.total}`);
console.log("\nby where the bug lives:");
for (const w of ["loop", "math"]) {
  const bugs = Object.entries(TAXONOMY).filter(([, m]) => m.where === w);
  const caught = bugs.filter(([b]) => propertySuite(kernel(b)) !== null).length;
  console.log(`  ${w.padEnd(5)} properties caught ${caught}/${bugs.length}`);
}
const ok = !cleanP && !cleanD;
console.log(ok ? "\nCONTROL OK (clean kernel passes both oracles)" : "\nCONTROL FAILED");
process.exit(ok ? 0 : 1);