wgslsmith

wgslsmith provides a set of tools for randomized testing of WGSL shader compilers. This includes a random program generator, as well as tools for executing arbitrary compute shaders and performing automated test case reduction on WGSL programs.

The compilers that are supported for testing currently include naga via wgpu and tint via dawn.

This was developed for @hasali19's final year project at Imperial College London, and went on to win a prize. The project report is available here.

Requirements

• Rust - Latest stable toolchain

A more complete list of requirements is available here.

Building

The full set of tools requires compiling Dawn and wgpu. Full instructions can be found in the docs.

Alternatively, some of the tools can be used without any WebGPU implementations/compilers (e.g. generator and reconditioner). To compile these, run:

$ git clone --recursive https://github.com/wgslsmith/wgslsmith
$ cd wgslsmith
$ cargo build -p wgslsmith --release
$ target/release/wgslsmith --help

Usage

All the tools can be used through the wgslsmith command:

# Do some fuzzing
$ wgslsmith fuzz
# Recondition a shader
$ wgslsmith recondition /path/to/shader.wgsl
# Reduce a crash
$ wgslsmith reduce crash path/to/shader.wgsl --config wgpu:dx12:9348 --regex '...'
# Run a shader
$ wgslsmith run path/to/shader.wgsl

Some options can be configured through a config file. Run wgslsmith config to open the default config file in a text editor. You can also specify a custom config file with the --config-file option.

[reducer]
tmpdir = "/home/hasan/dev/wgslsmith/tmp"
parallelism = 24

[reducer.creduce]
path = "/optional/path/to/creduce"

[reducer.cvise]
path = "/optional/path/to/cvise"

[reducer.perses]
# You need this if you want to reduce with perses
jar = "/path/to/perses_deploy.jar"

To use perses for reduction, grab and build it from https://github.com/wgslsmith/perses, then add it to the config as above.

Development

Insta is used for snapshot testing the parser.

Install the tool with cargo install cargo-insta and use cargo insta test -p parser to run the parser tests.

Building wgslsmith

• Dawn Prerequisites
  • CMake
  • depot_tools
  • Ninja
  • Linux dependencies
• Windows
• Building
• Installing

Dawn Prerequisites

The test harness and program reduction tools depend on dawn, which has a few prerequisites for building. If you don't want to build those tools, you can skip this bit.

CMake

If you don't already have it, run the following (for Ubuntu):

$ sudo apt install cmake

Otherwise, download it from https://cmake.org/download/.

depot_tools

Dawn uses depot_tools to fetch its dependencies. Detailed usage instructions can be found here. To install it:

# Clone the depot_tools repo somewhere on your system
$ git clone https://chromium.googlesource.com/chromium/tools/depot_tools.git
# Add it to your PATH
$ export PATH=/path/to/depot_tools:$PATH

Ninja

On Ubuntu:

$ sudo apt install ninja-build

Otherwise, grab the binary from https://github.com/ninja-build/ninja/releases and add it your PATH.

Linux dependencies

If you're building for Linux, you might need a few more dependencies. On Ubuntu, you can install the following packages:

$ sudo apt install libxrandr-dev libxinerama-dev libx11-dev \
    libxcursor-dev libxi-dev libxext-dev libxcb-shm0-dev libxtst-dev \
    libx11-xcb-dev

Windows

If cross-compiling for Windows, you'll also need to follow the instructions here to set up the compiler and SDK.

Building

Make sure to clone wgslsmith recursively to fetch the submodules:

$ git clone --recursive https://github.com/wgslsmith/wgslsmith
$ cd wgslsmith

To build everything, run the following:

$ ./build.py

This will automatically build dawn, and then build wgslsmith.

You can also disable some features if you don't want to build dawn and wgpu:

$ ./build.py --no-reducer --no-harness

If cross compiling for Windows, you need to set the target explicitly:

$ ./build.py --target x86_64-pc-windows-msvc

It's possible to build the harness as a standalone tool:

$ ./build.py harness

Build output will be in target/release (or cross-target/<target>/release when cross compiling).

Installing

To make the wgslsmith command available globally, run the following (after building):

$ ./build.py install [--install-prefix <path>]

This will create a symlink to target/release/wgslsmith, so you don't need to rerun it every time you rebuild. Use the --install-prefix option to specify the directory to install in (defaults to /usr/local/bin).

Alternatively you can just put the binary somewhere on your PATH.

Cross-compiling

• Windows

Windows

Windows is supported through cross-compilation using llvm's msvc target.

Add the msvc target to your rust toolchain with rustup:

$ rustup target add x86_64-pc-windows-msvc

Install the clang and llvm packages (on Ubuntu):

$ sudo apt install clang-14 clang-tools-14 llvm-14 lld-14

You'll also need a copy of the Windows SDK headers and libraries. xwin is a super handy tool which can be used to download them on Linux. Install it from source or grab a binary from the releases, then run the following:

$ xwin splat --include-debug-libs --output /path/to/sdk

The /path/to/sdk can be anywhere on your system where you'd like to download the SDK.

Also, make sure to set these environment variables:

# If you installed llvm on Ubuntu like above, this should be `/usr/lib/llvm-14`
export LLVM_NATIVE_TOOLCHAIN="/path/to/llvm"
# This should be the path to wherever you downloaded the Windows SDK with xwin
export XWIN_CACHE="/path/to/sdk"

Generator

The program generator is able to randomly generate WGSL programs using a range of language features.

# Generate a shader
$ wgslsmith gen
# Show help text
$ wgslsmith gen --help

Note that programs produced by the generator may not always compile (despite being syntactically valid and well-typed). This is because some WGSL compilers implement additional validation such as rejecting obvious infinite loops. wgslsmith uses a technique called reconditioning (see here) to guarantee validity. You can recondition shaders by passing --recondition to the generator, or by invoking the reconditioner separately on the generated shader which allows more control over its behaviour.

The generator has various options to control the generation process. See the help text for a full list.

Pointers are currently supported as an opt-in feature (since the reconditioner may reject some shaders with invalid pointer operations). To enable them, use the --enable-pointers flag. If reconditioning (with --recondition), you can also pass --skip-pointer-checks to stop it from erroring if the program contains possible invalid pointer operations.

Reconditioner

wgslsmith uses a simple technique called reconditioning to remove certain unwanted behaviour from programs (such as out-of-bounds array accesses and infinite loops). This involves a tool called a reconditioner, which validates and transforms shaders to add safety checks around potentially dangerous operations. The advantage of using a separate tool rather than implementing this in the generator is that it enables using off-the-shelf program reduction tools such as C-Reduce and Perses.

# Recondition a shader
$ wgslsmith recondition path/to/shader.wgsl

The reconditioner can be used to guarantee loop termination, which is important for making sure that programs can be compiled as some compilers reject obvious infinite loops. If you only want to enforce loop terminate without any other runtime checks, pass --enable loop-limiters to the reconditioner.

Harness

The test harness is used to run shader programs within a WebGPU pipeline. The harness is flexible and can be used for shaders with different inputs and outputs by providing a JSON description of the shader's I/O interface.

The harness will execute the shaders against multiple WebGPU implementations/configurations and compare the outputs to detect mismatches. It can also run in server mode to enable remote execution on a separate machine.

Currently, the harness only supports running compute shaders.

Basic usage

To execute a shader, run:

$ wgslsmith run /path/to/shader.wgsl

You can also supply input data to initialize any unfiform/storage buffers, by writing a json file of the form:

{
  "{group}:{binding}": [1, 2, 3, ...]
}

This contains a mapping from the buffer ID to a byte array containing the init data. The values of group and binding should be the corresponding attribute values in the WGSL program:

struct Buffer {
    value: u32,
};

@group(0) @binding(0) // <- 0:0
var<uniform> input: Buffer;

@group(0) @binding(1) // <- 0:1
var<storage, read_write> output: Buffer;

@compute
@workgroup_size(1)
fn main() {
    output.value = input.value;
}

By default, when executing a shader with an explicit path, the harness will look for a json file with the same name and parent directory as the shader. For example, given a shader file at /path/to/shader.wgsl, the harness will look for the inputs file at /path/to/shader.json.

You can also specify the inputs file path explicitly by passing /path/to/inputs.json as the second positional argument on the command line, or even specify the json object inline: '{"0:0": [...]}'.

Configurations

The harness will execute the input shader against one or more configurations, and compare the results of the output buffers for each configuration to detect possible miscompilation bugs.

A configuration is defined as the combination of a WebGPU implementation (such as dawn or wgpu) and a graphics adapter. The graphics adapter is identified by its backend type (D3D12, Metal, Vulkan) and a platform-specific integer identifier for the device.

Use the list subcommand to get a list of available configurations on your machine.

$ wgslsmith harness list
ID             | Adapter Name
---------------+------------------------------
wgpu:vk:9348   | NVIDIA GeForce RTX 3070
wgpu:dx12:9348 | NVIDIA GeForce RTX 3070
wgpu:dx12:140  | Microsoft Basic Render Driver
dawn:dx12:9348 | NVIDIA GeForce RTX 3070
dawn:dx12:140  | Microsoft Basic Render Driver
dawn:vk:9348   | NVIDIA GeForce RTX 3070

On my machine there are three adapters available, corresponding to hardware Vulkan and D3D12 implementations as well as a D3D12 software implementation. The configuration IDs consist of the WebGPU implementation, the backend type, and the PCI ID for the adapter.

By default, the harness will attempt to find the first available adapter for each combination of WebGPU implementation and backend type. Thus, all configurations above will be selected except for the D3D12 software adapter.

To specify configurations manually, you can pass them on the command line using the -c option.

$ wgslsmith run test.wgsl -c wgpu:dx12:140 -c dawn:dx12:140 -c dawn:vk:9348
executing wgpu:dx12:140
outputs:
  0: [2, 0, 0, 0]

executing dawn:dx12:140
outputs:
  0: [2, 0, 0, 0]

executing dawn:vk:9348
outputs:
  0: [2, 0, 0, 0]

ok

Exit codes

Test case reduction tools such as c-reduce typically take an interestingness test as input, which returns 0 for a useful test case or 1 if the test case should be discarded.

The harness can produce two types of errors:

• If the actual shader execution failed, this will manifest as a panic with exit code 101.
• If the shader was successfully executed for all configurations but the outputs differ, the program will exit with code 1.

Otherwise, the program exits normally with code 0.

Normally when using this with a reduction tool to find miscompilations, you will want to discard the shader if the harness returns 0 or 101, since execution failure means that the reduction process probably produced an invalid program. Only the exits with 1 are likely to be interesting.

Remote execution

The harness supports remote execution over a TCP connection, to enable running shaders on a separate device such as an Android smartphone or across a VM boundary (e.g. WSL to Windows host). This is done by running the harness as a server, which receives requests to execute a shader against a set of configurations.

To enable this, wgslsmith should be compiled for both the target machine as well as the client. It is also possible to compile the harness as a standalone tool to run on the execution server (without the other fuzzing tools) as described here.

Use the serve subcommand to start the server.

$ wgslsmith harness serve -a 0.0.0.0:1234

The remote subcommand can then be used to interact with the server. The command syntax is similar to the normal harness command.

$ wgslsmith remote 192.168.1.23:1234 list
$ wgslsmith remote 192.168.1.23:1234 run path/to/shader.wgsl

Note that the first argument to remote is the address of the server to connect to. For convenience, wgslsmith allows you to create friendly names for addresses and to set a default address. This is done through a configuration file (open it in your editor by running wgslsmith config).

# wgslsmith.toml

# Define a remote server called 'android-phone'
[remote.android-phone]
address = "192.168.1.23:1234"

[harness]
remote = "android-phone" # Set 'android-phone' as the default remote

This will allow you to connect to the remote using:

$ wgslsmith remote android-phone run shader.wgsl
# or
$ wgslsmith remote run shader.wgsl

Reducer

wgslsmith is able to integrate with a number of off-the-shelf program reduction tools to reduce WGSL programs that exhibit compilation crashes or have mismatching output buffers during differential testing. Currently, the following tools are supported:

• C-Reduce
• C-Vise
• Perses
• Picire

Perses tends to work the best thanks to its syntax-aware approach, though C-Reduce also works very well. If you want to use Perses, you must use this fork which includes the WGSL grammar, and add the following to your wgslsmith config file (open it in your editor by running wgslsmith config):

[reducer.perses]
jar = "/path/to/perses_deploy.jar"

The other tools can be installed by following their respective documentation.

To reduce WGSL programs use the reduce command:

$ wgslsmith reduce --help

For reducing mismatches, there are no extra required arguments other than the path to the shader and input data. For crashes, wgslsmith supports two approaches described below. In both cases you must provide the --regex <REGEX> option to specify a regex to match against the crash output (e.g. an error code that you're interested in).

Using the harness

The obvious way is to use the harness to attempt to execute the shader and check if it crashes. In this case, you will need to provide the --config <CONFIG> option with a config string that produces the crash (see here).

Using standalone validation tools

Executing the shaders can be relatively slow and requires access to the platform that has the bug. Instead, there are existing tools that can be used to validate HLSL, MSL and SPIR-V shaders generated by WGSL compilers. To use this, you will need to set up the validation server as described here. Then pass the --compiler <COMPILER> and --backend <BACKEND> options to the reduce command.

Validation Server

• Overview
• Obtaining the compilers
  • FXC
  • Metal
• Toolchain setup
• Usage

Overview

The FXC and Metal compilers can be used to validate HLSL and MSL shaders. Unfortunately neither of these tools are actually available for Linux, but it's possible to run them using wine.

Running with wine has a startup cost of a few seconds, which can make program reduction significantly slower. We get around this by using a server for validating shaders. This is packaged along with a Wine installation in a docker container.

Obtaining the compilers

FXC

The FXC dll is already included in the tools subdirectory of the wgslsmith repository.

Metal

You can download the Metal Developer Tools for Windows from here. You'll need an Apple ID to get access. Copy the installer to tools/Metal_Developer_Tools.exe.

Toolchain setup

The validation server needs to be compiled for Windows. Following the instructions here to setup a cross-compilation environment on Linux.

Usage

Start the validation server using the following command. This will compile the server and then build and start a docker container containing wine and the compiler tools.

$ ./start-validation-server.sh

The default server port is 9123. Add this to your wgslsmith config file (using the wgslsmith config command).

[validator]
server = "localhost:9123"