WGSL Shader Debugging

[brendan-duncan]

It would be great to be able to debug shaders, like you can with RenderDoc or PIX. To step through a specific invocation of a compute, vertex, or fragment shader, stepping through the code, inspecting variables.

There are a couple different ways this could be done.

RenderDoc has debugger code for SPIR-V, DXIB, and DXBC bytecodes. It is provides an interrupter for the bytecodes, allowing you to step through its execution. A sourcemap can provide mapping back to the original shader code.

I want to be able to debug WGSL shader code. Some options that could be done to implement this:

• Use a WASM build of a WGSL compiler like Tint to generate a shader bytecode such as SPIR-V, implement a SPIR-V interpreter using RenderDoc's implementation as inspiration, and use a sourcemap generated from Tint to map back into the WGSL code.

• See if we can extend the work from James Price did to implement a prototype WGSL debugger in Tint, https://dawn-review.googlesource.com/c/dawn/+/118901. (Just found out about this. Very cool.)

• Write a WGSL interpreter and debugger from scratch purely in Typescript.

For the last option, I already have a Typescript WGSL parser written, https://github.com/brendan-duncan/wgsl_reflect, which is used as the basis for the reflection library used by webgpu_inspector, and by other tools.

The parser generates an AST for the WGSL shader code, and we could use that to develop an interpreter. To verify this idea I started experimenting with doing that. https://github.com/brendan-duncan/wgsl_reflect/tree/exec, implementing an interpreter via walking the AST.

The following test is working with that approach:

const shader = `
        @group(0) @binding(0) var<storage, read_write> bins: array<u32>;
        @group(0) @binding(1) var ourTexture: texture_2d<f32>;

        // from: https://www.w3.org/WAI/GL/wiki/Relative_luminance
        const kSRGBLuminanceFactors = vec3f(0.2126, 0.7152, 0.0722);
        fn srgbLuminance(color: vec3f) -> f32 {
            return saturate(dot(color, kSRGBLuminanceFactors));
        }
        @compute @workgroup_size(1) fn main() {
            let size = textureDimensions(ourTexture, 0);
            let numBins = f32(arrayLength(&bins));
            let lastBinIndex = u32(numBins - 1);
            for (var y = 0u; y < size.y; y++) {
                for (var x = 0u; x < size.x; x++) {
                    let position = vec2u(x, y);
                    let color = textureLoad(ourTexture, position, 0);
                    let v = srgbLuminance(color.rgb);
                    let bin = min(u32(v * numBins), lastBinIndex);
                    bins[bin] += 1;
                }
            }
        }`;

const numBins = 256;
const histogramBuffer = new Uint32Array(numBins);

const size = [16, 16];
const texture = new Uint8Array(size[0] * size[1] * 4);

const bg = {0: {0: histogramBuffer, 1: {texture, size}}};

// Run a real WebGPU dispatch to get the expected results.
// The webgpuDispatch function will create GPU buffers, textures, and bindgroups from the contents of bg.
const _data = await webgpuDispatch(shader, "main", 1, bg);
const webgpuData = new Uint32Array(_data);

// Run the same dispatch using the WgslExec interpreter and compare the results,
const wgsl = new WgslExec(shader);
wgsl.dispatchWorkgroups("main", 1, bg);

// verifying the interpreter produces the same results as the GPU.
test.equals(histogramBuffer, webgpuData);

An AST walking interpreter would be difficult to build a debugger on top of, as it involves recursively executing nodes from a tree structure. We need to be able to execute a single statement at a time, passing the control back to the UI. A bytecode representation of the shader would be ideal, but it could probably be done in other ways.

Once the debugger execution is sorted out, then it can be integrated into the frame Capture tool. The Capture tool records all data and commands for rendering a frame. Select a dispatchWorkgroups command and you can see all of the buffers and textures that went into and out of the compute dispatch. The idea would then be that you can select a command like dispatchWorkgroups, then debug its shader for a particular invocation id. The debugger already has all of the buffer and texture data and bindgroup descriptors, so it could set up the variables in the shader so you can see the data as you step through the code.

Most of the building blocks for being able to do this are already written. Mostly just the debug-able interpreter just needs to be finished.

[brendan-duncan]

Making some progress on shader debugging with the custom parser / AST.

I didn't want to go through the work of creating a bytecode and VM for executing the shader in a way that supports iteratively stepping through instructions needed for debugging, but I found another way that is a hybrid of executing the AST directly, and allows support for stepping through the code.

The following test is working:

await test("dispatch function call", async function (test) {
      const shader = `
          fn scale(x: f32, y: f32) -> f32 {
            return x * y;
          }
          @group(0) @binding(0) var<storage, read_write> buffer: array<f32>;
          @compute @workgroup_size(1)
          fn main(@builtin(global_invocation_id) id: vec3<u32>) {
              let i = id.x;
              buffer[i] = scale(buffer[i], 2.0);
          }`;

      // Verify the emulated dispatch has the same results as the WebGPU dispatch.
      const buffer = new Float32Array([1, 2, 6, 0]);
      const bg = {0: {0: buffer}};

      const dbg = new WgslDebug(shader);
      dbg.debugWorkgroup("main", [1, 0, 0], 4, bg);
      dbg.stepNext(); // LET: i = id.x;
      dbg.stepNext(); // CALL: scale(buffer[i], 2.0)
      dbg.stepNext(); // RETURN: x * y
      dbg.stepNext(); // ASSIGN: buffer[i] = <value> 

      // Test that we only executed the [1, 0, 0] global_invocation_id.
      test.equals(buffer, [1, 4, 6, 0]);
});

Where we start debugging a workgroup invocation, and can then step through all of the statements being executed individually, even as it goes into expression function calls. It has all the state info of where the execution currently is at each step, so we can get stacktrace info and all the variables available in the current scope.

There's a lot more work to do to fill it out, but I'm happy this proof of concept is starting to show positive results.

[brendan-duncan]

A very rough, unpolished, incomplete first version of the debugger where it does something that looks kinda allright. It makes me excited to continue down this road.