Bug 323560

Created attachment 481310 [details] Self-contained repro: two entry points, same function, one call site vs three SUMMARY On iOS, a WGSL compute shader that inlines three or more calls to the same function over a 256-byte struct silently produces wrong results. No validation error, no compilation message, no device loss. Two entry points in one module compute the same value from the same read-only buffer. The one written with three inlined call sites writes all zeros; the one written with a single call site writes the correct data. Chrome 152 on macOS (Dawn/Metal) makes the two byte-identical. Reduced from a BN254 Groth16 prover, where it made every G2 multi-scalar-multiplication result a point that is not on the curve, so the proofs it produced were rejected by an independent verifier. Everything else in that prover, including four structurally identical MSMs over 128-byte G1 points, was bit-exact on the same device. ENVIRONMENT iPhone 15 Pro (iPhone16,1), iOS 26.6, Safari. Mozilla/5.0 (iPhone; CPU iPhone OS 18_7 like Mac OS X) AppleWebKit/605.1.15 (KHTML, like Gecko) Version/26.6 Mobile/15E148 Safari/604.1 Adapter: vendor apple, architecture apple. maxComputeWorkgroupStorageSize 16384, maxComputeInvocationsPerWorkgroup 256, maxStorageBufferBindingSize 1073741824. Not reproducible in Chrome 152 on macOS (Apple silicon, Dawn/Metal). STEPS TO REPRODUCE Open the attached page on the device and tap "run" (or append ?auto=1). It is self-contained: no network, no build step, no dependencies. Expected, and what Chrome on macOS prints: repro_folded: 468 ms repro_direct: 461 ms outputs MATCH device: alive Actual on the iPhone: repro_folded: 158 ms repro_direct: 519 ms all-zero output: direct true, folded false outputs DIFFER from word 0 (window 0) device: alive THE TWO KERNELS Both read the same storage buffer of PtG2, a 256-byte struct of four Fq2, each two 8-word BN254 base field elements. pt_add_g2 is unified projective XYZZ point addition: roughly 200 lines of 256-bit modular arithmetic, no loops, no memory access of its own. @compute @workgroup_size(1) fn repro_direct(@builtin(workgroup_id) wid: vec3<u32>) { let w = wid.x; var run = pt_zero_g2(); var tot = pt_zero_g2(); for (var j = P.n_buckets; j > 0u; j = j - 1u) { run = pt_add_g2(run, BUCKETS[w * P.n_buckets + (j - 1u)]); tot = pt_add_g2(tot, run); } WSUMS[w] = pt_add_g2(tot, run); } @compute @workgroup_size(1) fn repro_folded(@builtin(workgroup_id) wid: vec3<u32>) { let w = wid.x; var pts: array<PtG2, 2>; pts[0] = pt_zero_g2(); pts[1] = pt_zero_g2(); for (var step = 0u; step < 2u * P.n_buckets + 1u; step = step + 1u) { var dst = 1u; var src = pts[0]; if (step < 2u * P.n_buckets && (step & 1u) == 0u) { dst = 0u; src = BUCKETS[w * P.n_buckets + (P.n_buckets - 1u - (step >> 1u))]; } pts[dst] = pt_add_g2(pts[dst], src); } WSUMS[w] = pts[1]; } repro_direct has three inlined pt_add_g2 call sites. repro_folded has one, reached through a two-element array and a step counter: even steps fold a bucket into pts[0], odd steps fold pts[0] into pts[1], and the final step is the closing addition. The two are mathematically identical for any input, and only repro_folded returns data. WHAT THIS IS NOT Not a race. @workgroup_size(1), no workgroupBarrier, no var<workgroup>, no atomics, one dispatch per pass and one pass per submit. Not out of bounds. Both kernels index the same buffer over the same range, and WGSL clamps or discards out-of-range accesses rather than leaving them undefined. Not input-dependent branching. The buffer holds a fixed xorshift pattern, identical for both entry points, and pt_add_g2 branches only on its own intermediates. Not a shader-creation error. getCompilationInfo() returns no messages. Not the identity fast path. The buffer is filled with nonzero words precisely so the short-circuit on a zero z-coordinate cannot skip the arithmetic under test. WHERE THE EDGE IS Bisected in the original kernel by holding the phone against a laptop one construct at a time, with the count of inlined point-operation call sites as the only variable: 2 call sites bit-exact 3 call sites every output is the all-zero identity value 4 or more call sites outputs wrong, and different on every run of the same page one 4-site spelling the device is lost, reason "unknown", empty message, and nothing in the device system log The identical algorithm over 128-byte G1 points (Fq rather than Fq2, half the struct) never failed at any call-site count on the same device in the same submit. That is why this reads as a code-size or inlining-budget cliff in the WGSL-to-Metal translation rather than a logic error in one routine. The "different on every run" behaviour above is worth flagging because it cost time to interpret: the miscompiled code is deterministic, but it was being fed projective coordinates whose encoding legitimately varies between runs, so the corruption looked like a race for a while. The reduced case here removes that and fails the same way every time. WORKAROUND Route every point operation in the affected kernels through a single inlined call site, as repro_folded does. It costs roughly twice the loop trips and recovers correct results. POSSIBLY RELATED Bug 323543, "[WebGPU][WGSL] Array value constructor named through a type alias emits invalid MSL, silently producing wrong compute results". Different root cause, same failure mode of compiling clean and computing wrong. The project this was reduced from already carries a workaround for that one. SOURCE The generator that produces the WGSL, and the same probe wired into the project's own build, are at https://github.com/zemse/gpu-snark (web/static/dbg/reduce.html and crates/g16-wgpu/src/gen/points.rs).