Here's a WGSL loop:
for (var j = 0u; j < 4u; j++) {
result[j] = j + 1u;
}
Until v0.20.3 in October, naga, our shader compiler written in Go, turned it into DXIL that never ran the loop. Microsoft's validator looked at the module and said it was valid. On an Intel and an NVIDIA GPU the loop body didn't execute once.
The validator wasn't wrong. That's the interesting part.
What the validator saw
This is the DXIL that naga v0.20.2 writes for that loop, trimmed to the lines that matter, with my comments on the right:
%2 = alloca i32, align 4 ; var j
br label %3
...
%5 = load i32, i32* %2, align 4 ; read j for "j < 4"
%6 = icmp ult i32 %5, 4
br i1 %6, label %11, label %12
The variable j gets memory, and the very first thing the loop does is read it. Nothing ever stored the 0. Reading uninitialized memory is legal in LLVM IR, so every instruction here is well formed, the types match, the control flow is structured. IDxcValidator checks exactly those things and reports VALID (S_OK).
What it doesn't check is whether the module computes what the source says. No static validator does, not DXC's and not spirv-val. They answer "can a driver run this?", and the answer was yes. The driver ran it, j started as whatever was in that memory, and the loop exited before its first iteration.
naga v0.20.3 fixed three bugs of this kind in the DXIL backend:
| The shader | What the DXIL did |
|---|---|
for (var j = 0u; j < 4u; j++) |
never stored the initializer; the loop didn't run |
for (var j = i; j < n; j++) |
an optimization pass deleted the initial store, then noticed the loop also writes j; the loop started from an uninitialized j
|
v[i] = v[i] + 1u on a local vector |
a run-time index always picked v.x
|
For this article I compiled all three shaders with v0.20.2 again. All three come back VALID (S_OK).
Our own tests missed them too. We compare naga's text output with Rust naga's and its DXIL with DXC's, and none of that asks what the shader computes. What found them was a one-off experiment: run the shaders on real GPUs through Direct3D 12 and read the buffers back.
What naga is
naga reads WGSL, the WebGPU shading language, and writes SPIR-V for Vulkan, MSL for Metal, GLSL for OpenGL, HLSL for Direct3D, and DXIL for Direct3D 12. The DXIL is emitted directly, without DXC or LLVM. It's plain Go with no CGO, so it cross-compiles like any other Go package.
package main
import (
"fmt"
"log"
"github.com/gogpu/naga/dxil"
"github.com/gogpu/naga/hlsl"
"github.com/gogpu/naga/msl"
"github.com/gogpu/naga/spirv"
"github.com/gogpu/naga/valid"
"github.com/gogpu/naga/wgsl"
)
const source = `
@group(0) @binding(0) var<storage, read_write> result: array<u32>;
@compute @workgroup_size(1)
fn main() {
for (var j = 0u; j < 4u; j++) {
result[j] = j + 1u;
}
}
`
func main() {
module, _, err := wgsl.Parse(source, wgsl.ParseOptions{})
check(err)
_, err = valid.Validate(module, valid.Options{})
check(err)
spv := must(spirv.Compile(module, spirv.Options{})) // Vulkan
metal := must(msl.Compile(module, msl.Options{})) // Metal
hlslSrc := must(hlsl.Compile(module, hlsl.Options{})) // Direct3D, through FXC or DXC
blob := must(dxil.Compile(module, dxil.Options{})) // Direct3D 12, no DXC needed
fmt.Printf("SPIR-V %d bytes, MSL %d bytes, HLSL %d bytes, DXIL %d bytes\n",
len(spv), len(metal), len(hlslSrc), len(blob))
fmt.Println(metal)
}
// must returns a backend's output and stops on its error.
func must[T, I any](out T, _ I, err error) T {
check(err)
return out
}
func check(err error) {
if err != nil {
log.Fatal(err)
}
}
With naga v0.21.2 it prints SPIR-V 1240 bytes, MSL 917 bytes, HLSL 667 bytes, DXIL 2960 bytes, then the Metal source. The zero value of every options struct is the safe choice: bounds checks are on and workgroup memory is zeroed. The SPIR-V, MSL and HLSL writers also give every loop an iteration limit, because a loop the driver's compiler can prove never ends is undefined behavior, and the compiler is then free to delete it, bounds checks included. Two of those show up in the Metal output:
while(true) {
if (metal::all(loop_bound == uint2(0u))) { break; }
...
result[metal::min(unsigned(_e6), (_buffer_sizes.size0 - 0 - 4) / 4)] = _e8 + 1u;
}
Keep an eye on that metal::min. It's a fix from this month, and it comes back later.
naga takes its IR and much of its design from Rust naga, the shader compiler of Rust wgpu, which is the core of WebGPU in Firefox, Servo and Deno. We run Rust naga's 172-shader WGSL test corpus, compare our output with Rust's wherever Rust writes the same format (SPIR-V only structurally so far), and record every difference with its reason. Rust naga has no DXIL backend, so for DXIL the references are DXC and Mesa. naga's main user is gogpu/wgpu, a WebGPU implementation in Go, which compiles its WGSL shaders with naga in its default pure-Go build.
Running the output instead
After v0.20.3 we turned that experiment into a permanent test corpus: small compute shaders, each paired with the exact values a GPU has to write when it runs them. It lives in testdata/readback and has 55 cases today, mostly loops, function calls and local variables. Each case is a .wgsl file and a .json file:
{
"format": 1,
"desc": "for-loop counter declared in the loop header writes 1..4; the sentinel after them stays",
"entry": "main",
"workgroups": [1, 1, 1],
"bindings": [
{"binding": 0, "kind": "storage-rw", "len": 5, "fill": "0xDEADBEEF",
"expect": [[1, 2, 3, 4, "fill"]]}
]
}
Three habits made the corpus useful:
Sentinels. Every output buffer starts filled with 0xDEADBEEF, and the expected values say where it has to survive. A loop that runs one time too many, or a store through the wrong index, shows up as a missing sentinel.
Inputs from buffers. Most cases read their inputs from a storage buffer, so the compiler can't fold the work away at compile time, and a few, like collatz-lit.wgsl and collatz-rt.wgsl, test the same computation both ways. The WebGPU conformance test suite uses the same trick.
Known failures with a card. A case can list a backend it currently fails on, with the issue that tracks it. The run stays readable while the bug is open, and the entry goes away with the fix.
naga can't drive a GPU itself, so the work is split. naga checks that every case is well formed and compiles on every backend, and runs spirv-val, glslangValidator, dxc and IDxcValidator on the output where they're installed. wgpu runs the corpus on its GPUs before naga releases that touch a backend, with a Go workspace (go.work) pointing at the naga candidate, so a change runs on real hardware before it's tagged. That runner still lives outside the wgpu repository; moving it in is planned, and then anyone can run the corpus on their own GPU.
For v0.21.1 the corpus ran on eight backend and adapter pairs:
| Backend | Adapters |
|---|---|
| Vulkan | Intel Iris Xe, NVIDIA CMP 30HX |
| Direct3D 12, HLSL through FXC | Intel Iris Xe, NVIDIA GTX 1660 SUPER |
| Direct3D 12, naga's DXIL | Intel Iris Xe, NVIDIA GTX 1660 SUPER |
| OpenGL ES | Intel Iris Xe |
| wgpu's software backend | CPU |
There's no Mac in that list. Metal is covered by wgpu's own tests on a macOS CI runner, which is where the third bug below turned up.
What it caught
Local arrays that forgot their initializers. wgpu found this one with a vertex shader that picks its corners from local arrays. The same shape, cut down:
@vertex
fn vs(@builtin(vertex_index) v: u32) -> @builtin(position) vec4f {
var xs = array<f32, 6>(-1.0, 1.0, -1.0, -1.0, 1.0, 1.0);
var ys = array<f32, 6>(-1.0, -1.0, 1.0, 1.0, -1.0, 1.0);
return vec4f(xs[v % 6u], ys[v % 6u], 0.0, 1.0);
}
On naga v0.21.0 the DXIL for this has two alloca instructions and not a single store, so every index reads uninitialized memory, and the validator again says VALID (S_OK). wgpu's shader drew nothing on Intel. The bug was there since at least v0.17.16. Since v0.21.1 naga stores the initializers: 12 stores for this shader, checked on v0.21.2.
A loop-carried variable read as if the loop ran once. naga's DXIL backend turns local variables into SSA values where it can. When every access to a variable sat in one block inside a loop, and that block read the variable before writing it, the pass treated the block as running only once:
var x: u32;
for (var i = 0u; i < n; i++) {
x += i + 1u;
dst[i] = x; // want 1, 3, 6, 10. DXIL wrote 1, 2, 3, 4.
}
Every iteration saw the value from before the loop. This one shipped in every release since v0.17.4. The corpus case for it also checks shapes the old code got right, such as variables read after the loop or updated only in branches, so the fix can't quietly break those.
The v0.21.1 run compared every result with the same run on v0.21.0. 54 of the 55 cases ran on each of the eight pairs (one needs shader-f16, which the runner skipped), 432 results in all. Seven changed, all on DXIL, all from fail to pass:
| Backend and adapter | v0.21.0 | v0.21.1 |
|---|---|---|
| Direct3D 12 + naga's DXIL, Intel Iris Xe | 50 of 54 pass | 54 of 54 |
| Direct3D 12 + naga's DXIL, GTX 1660 SUPER | 51 of 54 pass | 54 of 54 |
| the other six pairs | unchanged | unchanged |
Not everything in that run was green, and both versions failed the same cases: two workgroup-memory cases on the CMP 30HX under Vulkan (not diagnosed yet; Intel passes them), one known OpenGL ES failure, and 15 of the 54 in wgpu's software backend (11 failures, 4 already known). For a release gate that's the point: nothing got worse.
An out-of-bounds write on Metal. wgpu has hand-written readback tests too. One of them, TestStorageBufferOutOfBoundsContained, writes past the end of a storage binding and then checks that the binding next to it is untouched. WGSL allows an implementation to clamp such an access, and wgpu's Metal backend asks naga to clamp every index into a buffer.
On the macOS CI runner the test failed. For a plain var<storage> a: array<u32>, naga's MSL used the index as it was, and the write landed in the next binding. Only a runtime-sized array at the end of a struct got clamped. v0.21.2 clamps every index into a runtime-sized array; that's the metal::min from the output above. The unclamped MSL was perfectly valid MSL.
When Rust naga and the spec disagree
While fixing the Metal bug I noticed that our bound was computed from the array's stride and Rust naga's from the element's size. I switched to Rust's form and called ours a defect.
It wasn't. WGSL defines how many elements a runtime-sized array has, and calls it NRuntime: the buffer binding's size minus the array's offset, divided by the stride, rounded down. arrayLength must return exactly that.
The two forms differ when an element is smaller than its stride and the binding ends partway through the last stride. Take array<vec3<f32>>: each element is 12 bytes and the stride is 16. On a 28-byte binding the spec says there is truncate(28 / 16) = 1 element. Rust naga's MSL computes 1 + (28 - 12) / 16 = 2, because a second vec3 would fit in bytes 16 to 27. That's true about memory and wrong about WGSL.
The correction came from the wgpu side, and wgpu's vec3 test now pins it: its comment quotes the spec's NRuntime and requires arrayLength to follow it. v0.21.2 uses the spec's count for both the clamp and arrayLength, which is what naga's HLSL, DXIL, GLSL and SPIR-V output already did. On the macOS CI runner, a 28-byte array<vec3<u32>> binding now reads arrayLength back as 1 (run result). The difference from Rust is written down in the code, in the tests that pin the bound, and in the CHANGELOG.
So this is how we decide now. Where the spec defines a value or an error, the spec decides. Rust naga decides what the spec leaves open, like the shape of the IR or the wording of an error message. When we differ from Rust, we assume the bug is ours until we've read the spec. After that it might be theirs.
What I'd do differently
Run the output from the first day of a backend. We had validators, golden files and Rust parity long before we had a GPU in the loop. All three check that the output looks right. Only running it checks that it is right.
Read a green validator as "legal", not "correct". It's a gate worth having, and it says nothing about values.
Make the expected values hard to fake. Sentinels in every slot that shouldn't change, inputs from buffers, and several shapes per bug, including the ones the old code already handled.
What's next
DXIL is opt-in in wgpu (GOGPU_DX12_DXIL=1). The default Direct3D 12 path is HLSL through FXC, so DX12 users on that path were never affected by the DXIL bugs above.
Before we call DXIL stable, the shaders of gg, our 2D graphics library, have to pass the validator, readback has to be green on Intel and NVIDIA, and no silent bug may be known. Two are known and open right now:
- more shapes of the local-initializer bug: indexed vectors, arrays of arrays, zero-initialized arrays of vectors, and struct or array variables of a helper function called inside a loop;
-
primitive_indexreads wrong in fragment shaders on Intel, probably because the DXIL doesn't mark fragment system values as non-interpolated.
If naga compiles one of your shaders wrong, the corpus format is documented. A WGSL file plus the values it should produce is the most useful bug report we can get.
naga: github.com/gogpu/naga · pkg.go.dev
The Go gopher was designed by Renée French and is licensed under CC BY 4.0.
go get github.com/gogpu/naga@v0.21.2
go install github.com/gogpu/naga/cmd/nagac@latest
nagac -o shader.spv shader.wgsl
Top comments (0)