wgpu Documentation

File: README.md Copy Chapter

wgpu

wgpu is a cross-platform, safe, pure-Rust graphics API. It runs natively on Vulkan, Metal, D3D12, and OpenGL; and on top of WebGL2 and WebGPU on wasm.

The API is based on the WebGPU standard, but is a fully native Rust library. It serves as the core of the WebGPU integration in Firefox, Servo, and Deno.

Getting Started Copy Chapter

See our examples online at https://wgpu.rs/examples/. You can see the Rust sources at examples and run them directly with cargo run --bin wgpu-examples <example>.

Learning `wgpu` Copy Section

If you are new to wgpu and graphics programming, we recommend starting with Learn Wgpu.

Additionally, WebGPU Fundamentals is a tutorial for WebGPU which is very similar to our API, minus differences between Rust and Javascript.

Wiki Copy Section

We have a wiki which has information on useful architecture patterns, debugging tips, and more getting started information.

Need Help? Want to Contribute? Copy Section

The wgpu community uses Matrix and Discord to discuss.

• - discussion of wgpu's development.
• - discussion of using the library and the surrounding ecosystem.
• - Dedicated support channel on the Rust Gamedev Discord.

Other Languages Copy Section

To use wgpu in C or dozens of other languages, look at wgpu-native. These are C bindings to wgpu and has an up-to-date list of libraries bringing support to other languages.

Learn WebGPU (for C++) is a good resource for learning how to use wgpu-native from C++.

Quick Links Copy Chapter

Contributors are welcome! See CONTRIBUTING.md for more information.

Supported Platforms Copy Chapter

✅ = First Class Support
🆗 = Downlevel/Best Effort Support
📐 = Requires the ANGLE translation layer (GL ES 3.0 only). On macOS/iOS, use the angle feature. On Windows, gles uses WGL by default; build with cfg(windows_angle) to use ANGLE instead.
🌋 = Requires the MoltenVK translation layer
🛠️ = Unsupported, though open to contributions

Environment Variables Copy Chapter

Testing, examples, and ::from_env() methods use a standardized set of environment variables to control wgpu's behavior.

• WGPU_BACKEND with a comma-separated list of the backends you want to use (vulkan, metal, dx12, or gl).
• WGPU_ADAPTER_NAME with a case-insensitive substring of the name of the adapter you want to use (ex. 1080 will match NVIDIA GeForce 1080ti).
• WGPU_DX12_COMPILER with the DX12 shader compiler you wish to use (dxc, static-dxc, or fxc). Note that dxc requires dxcompiler.dll (min v1.8.2502) to be in the working directory, and static-dxc requires the static-dxc crate feature to be enabled. Otherwise, it will fall back to fxc.

See the documentation for more environment variables.

When running the CTS, use the variables DENO_WEBGPU_ADAPTER_NAME, DENO_WEBGPU_BACKEND, DENO_WEBGPU_POWER_PREFERENCE, DENO_WEBGPU_DX12_COMPILER, and DENO_WEBGPU_STRICT_COMPLIANCE.

Repo Overview Copy Chapter

For an overview of all the components in the gfx-rs ecosystem, see the big picture.

MSRV policy Copy Chapter

TL;DR: If you're using wgpu, our MSRV is 1.87. If you're running our tests or examples, our MSRV is 1.93.

We will avoid bumping the MSRV of wgpu without good reason, and such a change is considered breaking.

Specific Details

Due to complex dependants, we have three MSRV policies:

• wgpu's MSRV is 1.87
• wgpu-core (and hence wgpu-hal, naga, naga-types and wgpu-types)'s MSRV is 1.87.
• The rest of the workspace has an MSRV of 1.93.

It is enforced on CI (in "/.github/workflows/ci.yml") with the WGPU_MSRV, CORE_MSRV, and REPO_MSRV variables, respectively.
This version can only be upgraded in breaking releases, though we release a breaking version every three months.

The following rules apply:

• The wgpu-core crate should never require an MSRV ahead of Firefox's MSRV for nightly builds, as
  determined by the value of MINIMUM_RUST_VERSION in python/mozboot/mozboot/util.py.
• The wgpu crate should never require an MSRV ahead of Servo's MSRV, as determined by the value of
  their rust-version declaration in Cargo.toml
• The repository MSRV should never require an MSRV higher than stable - 3. For example, if stable is
  at 1.97, the repository MSRV should be no higher than 1.94. This is to allow people who are using a decently-updated
  OS-provided rust to be able to build our repository. Consider cross checking with NixOS, though this
  is not required.

Testing and Environment Variables Copy Chapter

Information about testing, including where tests of various kinds live, and how to run the tests.

Tracking the WebGPU and WGSL draft specifications Copy Chapter

The wgpu crate is meant to be an idiomatic Rust translation of the WebGPU API.
That specification, along with its shading language, WGSL,
are both still in the "Working Draft" phase,
and while the general outlines are stable,
details change frequently.
Until the specification is stabilized, the wgpu crate and the version of WGSL it implements
will likely differ from what is specified,
as the implementation catches up.

Exactly which WGSL features wgpu supports depends on how you are using it:

• When running as native code, wgpu uses Naga
  to translate WGSL code into the shading language of your platform's native GPU API.
  Naga is working on catching up to the WGSL specification,
  with bugs tracking various issues,
  but there is no concise summary of differences from the specification.

• When running in a web browser (by compilation to WebAssembly)
  without the "webgl" feature enabled,
  wgpu relies on the browser's own WebGPU implementation.
  WGSL shaders are simply passed through to the browser,
  so that determines which WGSL features you can use.

• When running in a web browser with wgpu's "webgl" feature enabled,
  wgpu uses Naga to translate WGSL programs into GLSL.
  This uses the same version of Naga as if you were running wgpu as native code.

File: benches/README.md Copy Chapter

Collection of CPU benchmarks for wgpu.

These benchmarks are designed as a first line of defence against performance regressions and generally approximate the performance for users.

Usage Copy Chapter

# Run all benchmarks
cargo bench -p wgpu-benchmark
# Run a specific benchmarks that contains "filter" in its name
cargo bench -p wgpu-benchmark -- "filter"

Use WGPU_BACKEND and WGPU_ADAPTER_NAME to adjust which device the benchmarks use. More info on env vars.

Comparing Against a Baseline Copy Chapter

To compare the current benchmarks against a baseline, you can use the --save-baseline and --baseline flags.

For example, to compare v28 against trunk, you could run the following:

git checkout v28
# Run the baseline benchmarks
cargo bench -p wgpu-benchmark -- --save-baseline "v28"

git checkout trunk
# Run the current benchmarks
cargo bench -p wgpu-benchmark -- --baseline "v28"

The current benchmarking framework was added before v28, so comparisons only work after it was added. Before that the same commands will work, but comparison will be done using criterion.

Integration with Profilers Copy Chapter

The benchmarks can be run with a profiler to get more detailed information about where time is being spent.
Integrations are available for tracy and superluminal.

Tracy

Tracy is available prebuilt for Windows on github.

# Once this is running, you can connect to it with the Tracy Profiler
cargo bench -p wgpu-benchmark --features tracy,profiling/profile-with-tracy

Superluminal

Superluminal is a paid product for windows available here.

# This command will build the benchmarks, and display the path to the executable
cargo bench -p wgpu-benchmark --features profiling/profile-with-superluminal -- -h

# Have Superluminal run the following command (replacing with the path to the executable)
 --bench "filter"

`perf` and others

You can follow the same pattern as above to run the benchmarks with other profilers.
For example, the command line tool perf can be used to profile the benchmarks.

# This command will build the benchmarks, and display the path to the executable
cargo bench -p wgpu-benchmark -- -h

# Run the benchmarks with perf
perf record  --bench "filter"

Benchmarks Copy Chapter

`Renderpass Encoding`

This benchmark measures the performance of recording and submitting a render pass with a large
number of draw calls and resources, emulating an intense, more traditional graphics application.
By default it measures 10k draw calls, with 90k total resources.

Within this benchmark, both single threaded and multi-threaded recording are tested, as well as splitting
the render pass into multiple passes over multiple command buffers.
If available, it also tests a bindless approach, binding all textures at once instead of switching
the bind group for every draw call.

`Computepass Encoding`

This benchmark measures the performance of recording and submitting a compute pass with a large
number of dispatches and resources.
By default it measures 10k dispatch calls, with 60k total resources, emulating an unusually complex and sequential compute workload.

Within this benchmark, both single threaded and multi-threaded recording are tested, as well as splitting
the compute pass into multiple passes over multiple command buffers.
If available, it also tests a bindless approach, binding all resources at once instead of switching
the bind group for every draw call.
TODO(https://github.com/gfx-rs/wgpu/issues/5766): The bindless version uses only 1k dispatches with 6k resources since it would be too slow for a reasonable benchmarking time otherwise.

`Device::create_buffer`

This benchmark measures the performance of creating large buffers.

`Device::create_bind_group`

This benchmark measures the performance of creating large bind groups of 5 to 50,000 resources.

`naga::back`, `naga::compact`, `naga::front`, and `naga::valid`

These benchmark measures the performance of naga parsing, validating, and generating shaders.

File: cts_runner/README.md Copy Chapter

cts_runner

This crate contains infrastructure for running the WebGPU conformance tests on
Deno's wgpu-based implementation of WebGPU.

Instructions for running the tests via the CTS xtask are in
docs/testing.md.
The file revision.txt specifies the version of the CTS that
will be used by default.

cts_runner is somewhat misnamed at this point, in that it is useful for
things other than just running the CTS:

• The tests directory contains a few directed tests for
  Deno's bindings to wgpu.
• Standalone JavaScript snippets that use WebGPU can be run
  with a command like: cargo run -p cts_runner -- test.js.

File: deno_webgpu/README.md Copy Chapter

deno_webgpu

This op crate implements the WebGPU API as defined in
https://gpuweb.github.io/gpuweb/ in Deno. The implementation targets the spec
draft as of March 31, 2024. The spec is still very much in flux. This extension
tries to stay up to date with the spec, but is constrained by the features
implemented in our GPU backend library wgpu.

The spec is still very bare bones, and is still missing many details. As the
spec becomes more concrete, we will implement to follow the spec more closely.

In addition, setting the DENO_WEBGPU_TRACE environmental variable will output
a
wgpu trace
to the specified directory.

This op crate is tested primarily by running the
WebGPU conformance test suite using wgpu's
cts_runner.
cts_runner also has a few
directed tests
to fill in missing coverage.

GPU availability in GitHub CI is limited, so some configurations rely on
software like DX WARP & Vulkan lavapipe.

Links Copy Chapter

Specification: https://gpuweb.github.io/gpuweb/

Design documents: https://github.com/gpuweb/gpuweb/tree/main/design

Conformance tests suite: https://github.com/gpuweb/cts

WebGPU examples for Deno: https://github.com/crowlKats/webgpu-examples

wgpu-users matrix channel: https://matrix.to/#/#wgpu-users:matrix.org

File: docs/api-specs/cooperative_matrix.md Copy Chapter

Cooperative Matrix Extensions

🧪Experimental🧪

wgpu supports an experimental cooperative matrix feature when Features::EXPERIMENTAL_COOPERATIVE_MATRIX is enabled.
This exposes hardware-accelerated matrix multiply-accumulate (MMA) operations (for example, NVIDIA tensor cores,
Metal SIMD-group matrices, and Vulkan VK_KHR_cooperative_matrix).

Note: The features documented here may have bugs and are subject to breaking changes. The API and shader
semantics are expected to evolve. Please refer to the GitHub issue tracker for the latest status and discussions.

Overview Copy Chapter

Cooperative matrices allow a workgroup (or equivalent execution group) to collectively:

• load small matrix tiles from memory,
• perform matrix multiply-accumulate operations on those tiles, and
• store the results back to memory.

Conceptually, this is specialized hardware that evaluates:

C = A * B + C

for relatively small tiles, but at very high throughput compared to composing the same operation from
scalar/vector instructions.

Cooperative matrix operations are most useful in workloads such as:

• machine learning and inference,
• dense linear algebra and scientific computing,
• image processing, filtering, and transforms.

The cooperative nature means that all lanes in the cooperating execution group must participate in
the operations; individual invocations cannot diverge.

Typical example:

• A is an M×K matrix.
• B is a K×N matrix.
• C is an M×N matrix, acting as the accumulator and result.

Querying hardware support (host side) Copy Chapter

Before using cooperative matrices in shaders, you must query what configurations your hardware and backend support.

On the Adapter, wgpu exposes:

• Adapter::cooperative_matrix_properties() -> Vec<CooperativeMatrixProperties>

Each CooperativeMatrixProperties describes a single supported configuration. Fields are:

• m_size: height of matrices A and C (type: naga::CooperativeSize)
• n_size: width of matrices B and C (type: naga::CooperativeSize)
• k_size: shared inner dimension of A and B (type: naga::CooperativeSize)
• ab_type: scalar element type for A and B (type: naga::Scalar)
• cr_type: scalar element type for C and the result (type: naga::Scalar)
• saturating_accumulation: bool indicating whether overflow clamping on accumulation
  is supported for this configuration

Example usage:

let coop_props = adapter.cooperative_matrix_properties();
for prop in &coop_props {
    println!(
        "{:?}x{:?}x{:?} - AB: {:?}, CR: {:?}, saturating: {}",
        prop.m_size, prop.n_size, prop.k_size,
        prop.ab_type, prop.cr_type,
        prop.saturating_accumulation,
    );
}

You must:

1. Enable Features::EXPERIMENTAL_COOPERATIVE_MATRIX on the Device.
2. Query adapter.cooperative_matrix_properties() and ensure that the configuration(s) you intend
   to use in WGSL are actually available on the running adapter/backend.
3. Treat the sizes and types as a contract between your shaders and the underlying hardware implementation.
   Using unsupported configurations is an error.

Feature and backend requirements Copy Chapter

`wgpu` feature Copy Section

• Using cooperative matrices requires enabling:
  • Features::EXPERIMENTAL_COOPERATIVE_MATRIX

This feature may be restricted to certain backends and hardware.

Hardware / backend notes Copy Section

These are general guidelines, not a complete compatibility matrix:

• Metal:

  • Requires Apple7+ (A14) or Mac2+ (M1) GPU with MSL 2.3+.
  • Strong support for 8×8 f32, 8×8 f16, and mixed-precision modes (e.g. f16 A/B and f32 accumulator C).
  • Implementation is based on SIMD-group matrix operations.

• Vulkan:

  • Requires the VK_KHR_cooperative_matrix extension.
  • Many NVIDIA and AMD GPUs support f16 at 16×16 tile sizes and similar.
  • 8×8 f32 support is hardware-dependent.
  • Exact configurations are enumerated by Adapter::cooperative_matrix_properties().

• Other backends:

  • May not support cooperative matrices at all. In that case the feature will not be exposed, and
    adapter.cooperative_matrix_properties() will return an empty list.

Always treat the properties returned at runtime as the source of truth.

`wgpu` API surface Copy Chapter

This section summarizes the host-side API elements related to cooperative matrices.
(For exact signatures and details, refer to the Rust documentation.)

Adapter Copy Section

• Adapter::cooperative_matrix_properties() -> Vec<CooperativeMatrixProperties>

Returns all cooperative matrix configurations supported by the adapter/backend.

Structures Copy Section

• CooperativeMatrixProperties
  • m_size: naga::CooperativeSize
  • n_size: naga::CooperativeSize
  • k_size: naga::CooperativeSize
  • ab_type: naga::Scalar
  • cr_type: naga::Scalar
  • saturating_accumulation: bool

The naga types (CooperativeSize, Scalar) are part of the shader translation layer and
determine the legal WGSL/cooperative matrix combinations.

There are currently no dedicated wgpu buffer or texture types for cooperative matrices; they are
expressed in WGSL as special value types accessed via pointers into ordinary var<storage> /
var<workgroup> / var / etc.

WGSL extension specification Copy Chapter

Cooperative matrices are enabled and accessed via WGSL extensions. The exact extension spelling
may change; the details below describe the intended semantics.

Enabling cooperative matrices in WGSL Copy Section

Any WGSL program using cooperative matrices must declare an extension at the top of the shader, for example:

enable wgpu_cooperative_matrix;

The shader is invalid if any cooperative matrix types or builtins are used without enabling this extension.

Cooperative matrix types Copy Section

A cooperative matrix is a value type parameterized by:

• tile size (M×N),
• scalar element type T, and
• role R indicating how the matrix participates in the multiply-accumulate:
  • A: left operand
  • B: right operand
  • C: accumulator / result

Conceptually:

// A: MxK, B: KxN, C: MxN
type coop_matMxN<T, A>;
type coop_matMxN<T, B>;
type coop_matMxN<T, C>;

Concrete examples (sizes and types must match a supported configuration from
Adapter::cooperative_matrix_properties):

// 8x8 single-precision tiles
alias CoopMatA = coop_mat8x8<f32, A>;
alias CoopMatB = coop_mat8x8<f32, B>;
alias CoopMatC = coop_mat8x8<f32, C>;

// 16x16 half-precision inputs, 16x16 f32 accumulator (mixed precision)
alias CoopMat16x16A = coop_mat16x16<f16, A>;
alias CoopMat16x16B = coop_mat16x16<f16, B>;
alias CoopMat16x16C = coop_mat16x16<f32, C>;

The actual set of legal (M, N, T, R) combinations is defined by the cooperative matrix
properties returned at runtime; shaders must not use arbitrary combinations.

Roles and semantics Copy Section

• A role:

  • Treated as the left operand in the multiplication. Has shape M×K.
  • Participates as A in A * B + C.

• B role:

  • Treated as the right operand in the multiplication. Has shape K×N.
  • Participates as B in A * B + C.

• C role:

  • Treated as accumulator and result. Has shape M×N.
  • Participates as C in A * B + C.

These roles are part of the type; they are not interchangeable.

Cooperative matrix operations Copy Section

WGSL provides built-in functions for operating on cooperative matrices. The exact spelling may
change; the semantics are:

`coopLoad` / `coopLoadT`

Collectively load a tile from memory into a cooperative matrix. Two variants
select the memory layout:

• coopLoad — matrix is stored column-major in memory; stride is the
  number of elements between adjacent columns.
• coopLoadT — matrix is stored row-major in memory (i.e. transposed
  relative to the canonical column-major layout used by coopLoad);
  stride is the number of elements between adjacent rows. This is the
  natural fit for C-style ptr[i * num_cols + j] storage.

fn coopLoad<T, R>(
    ptr: ptr<STORAGE_CLASS, T>, // base pointer to scalar or vector elements
    stride: u32                  // elements between adjacent columns
) -> coop_matMxN<T, R>;

fn coopLoadT<T, R>(
    ptr: ptr<STORAGE_CLASS, T>, // base pointer to scalar or vector elements
    stride: u32                  // elements between adjacent rows
) -> coop_matMxN<T, R>;

• Loads an M×N tile (or M×K / K×N, depending on role and operation) from memory pointed to by ptr.
• All invocations in the cooperative group must call the chosen variant in a converged fashion.
• Memory address range must be valid and properly aligned for the scalar type.

Implementation note: Each lane contributes to filling the tile based on an implementation-defined mapping from
invocation/lane ID to sub-fragment of the matrix.

`coopStore` / `coopStoreT`

Collectively store a cooperative matrix tile back to memory. Variant
selection mirrors the load builtins:

• coopStore — writes column-major; stride between columns.
• coopStoreT — writes row-major; stride between rows.

fn coopStore<T, R>(
    value: coop_matMxN<T, R>,
    ptr: ptr<STORAGE_CLASS, T>,
    stride: u32
);

fn coopStoreT<T, R>(
    value: coop_matMxN<T, R>,
    ptr: ptr<STORAGE_CLASS, T>,
    stride: u32
);

• Stores value into the memory region addressed by ptr with given stride.
• All invocations in the cooperative group must participate.
• The store must not alias overlapping tiles in undefined ways.

`coopMultiplyAdd`

Perform a matrix multiply-accumulate operation on cooperative matrices:

fn coopMultiplyAdd<Tab, Tcr, MA, KA, KB, NB>(
    a: coop_matMAxKA<Tab, A>, // A: MAxKA tile
    b: coop_matKBxNB<Tab, B>, // B: KBxNB tile (KB == KA)
    c: coop_matMAxNB<Tcr, C>  // C: MAxNB accumulator/result
) -> coop_matMAxNB<Tcr, C>;

Semantics:

• Computes C' = A * B + C.
• Returns the resulting accumulator tile C'.
• Implies:
  • KA == KB (inner dimension must match).
  • Types (Tab, Tcr) must be one of the supported AB/CR combinations given by
    CooperativeMatrixProperties.
  • Sizes (MA, NB, KA) must match a supported (m_size, n_size, k_size) triple.

For example, with a supported configuration:

enable wgpu_cooperative_matrix;

alias MatA = coop_mat8x8<f32, A>;
alias MatB = coop_mat8x8<f32, B>;
alias MatC = coop_mat8x8<f32, C>;

// Assumes each tile is stored column-major in memory (the plain `coopLoad`
// / `coopStore` form); use `coopLoadT` / `coopStoreT` for row-major storage.
fn matmul_tile(
    ptr_a: ptr<storage, f32>,
    ptr_b: ptr<storage, f32>,
    ptr_c: ptr<storage, f32>,
    stride: u32,
) {
    let a: MatA = coopLoad<_, A>(ptr_a, stride);
    let b: MatB = coopLoad<_, B>(ptr_b, stride);
    let c: MatC = coopLoad<_, C>(ptr_c, stride);

    let result: MatC = coopMultiplyAdd(a, b, c);
    coopStore(result, ptr_c, stride);
}

If saturating_accumulation is true for the chosen configuration, then overflow during accumulation
is clamped (e.g. saturating arithmetic). If false, overflow behavior for the accumulator follows the
underlying scalar type semantics (e.g. IEEE-754 for floats).

Workgroup cooperation and execution model Copy Section

Cooperative matrix operations are collective:

• All invocations in the relevant execution group must execute each cooperative operation in uniform control flow:

  • Using coopLoad / coopLoadT, coopStore / coopStoreT, or coopMultiplyAdd in divergent control flow
    (e.g. some lanes taking a branch, others not) is undefined behavior.
  • The exact execution group may be a workgroup, a SIMD-group / subgroup, or another backend-specific
    granularity; shaders must treat it abstractly.

• The workgroup (or cooperating group) size is constrained by both:

  • the cooperative matrix configuration, and
  • backend-specific implementation details.

For portable code:

• Choose a workgroup size that is known to be supported efficiently on your target backends, for example:

  • @workgroup_size(8, 8, 1) to operate on an 8×8 tile, or
  • a multiple of the tile size where each subgroup handles a tile.

• Avoid control-flow divergence around cooperative operations.

Example:

enable wgpu_cooperative_matrix;

struct Matrices {
    // Row-major tiles for A, B, C — use the `…T` load/store variants.
    data: array<f32>,
};

@group(0) @binding(0)
var<storage, read>  buf_a: Matrices;
@group(0) @binding(1)
var<storage, read>  buf_b: Matrices;
@group(0) @binding(2)
var<storage, read_write> buf_c: Matrices;

alias MatA = coop_mat8x8<f32, A>;
alias MatB = coop_mat8x8<f32, B>;
alias MatC = coop_mat8x8<f32, C>;

@compute @workgroup_size(8, 8, 1)
fn main(
    @builtin(workgroup_id) wg_id: vec3,
    @builtin(local_invocation_id) lid: vec3,
) {
    // Compute tile offset; this is one of many possible mappings.
    let tile_index = wg_id.x; // 1D tiling in this simple example
    let tile_offset = tile_index * 64u; // 8x8 tile has 64 elements

    // Base pointers for tiles of A, B, C.
    let base_a = &buf_a.data[tile_offset];
    let base_b = &buf_b.data[tile_offset];
    let base_c = &buf_c.data[tile_offset];

    let a: MatA = coopLoadT<f32, A>(base_a, 8u);
    let b: MatB = coopLoadT<f32, B>(base_b, 8u);
    let c: MatC = coopLoadT<f32, C>(base_c, 8u);

    let result: MatC = coopMultiplyAdd(a, b, c);
    coopStoreT(result, base_c, 8u);
}

Validation rules and undefined behavior Copy Chapter

Implementations must validate the following where possible:

• The wgpu_cooperative_matrix WGSL extension is enabled if any cooperative matrix types
  or builtins are used.
• Tile sizes (M, N, K) and scalar types (ab_type, cr_type) match at least one
  CooperativeMatrixProperties entry for the current adapter/backend.
• Workgroup size, shader stage, and other pipeline configuration constraints required
  by the backend are satisfied.

The following are examples of undefined behavior (non-exhaustive):

• Using cooperative matrix operations without enabling the WGSL extension.
• Using a cooperative matrix type (M, N, T, R) not supported by
  Adapter::cooperative_matrix_properties().
• Mismatching sizes or roles in coopMultiplyAdd (e.g. incompatible M/N/K, or incorrect roles).
• Executing coopLoad / coopLoadT, coopStore / coopStoreT, or coopMultiplyAdd in divergent
  control flow within the cooperating execution group.
• Providing invalid, misaligned, or out-of-bounds pointers to any of the load/store builtins.
• Using a load/store variant (coopLoad vs coopLoadT, coopStore vs coopStoreT) whose memory
  layout does not match how the tile is actually stored.
• Overlapping coopStore / coopStoreT targets in a way that creates data races or aliasing that
  the memory model does not allow.

Example: 64×64 matrix multiply using 8×8 tiles Copy Chapter

The example in examples/features/src/cooperative_matrix demonstrates using cooperative matrices to
compute:

• C = A * B + C where:
  • A is 64×64,
  • B is 64×64,
  • C is 64×64.

A high-level tiling strategy:

1. Partition A, B, and C into 8×8 tiles.
2. Launch one workgroup per output tile of C (i.e. 8×8 tiles for a 64×64 matrix = 8×8 = 64 tiles).
3. Within each workgroup:
   • Loop over K-dimension tiles.
   • For each k tile:
     • Load an 8×8 tile of A (MatA).
     • Load an 8×8 tile of B (MatB).
     • Maintain an 8×8 accumulator tile (MatC) and repeatedly apply coopMultiplyAdd.
4. After the K loop, store the final accumulator tile back to C.

Key points from the example:

• Workgroup size is chosen so that all cooperative operations are well-defined and efficient for 8×8 tiles.
• Host-side code:
  • Enables Features::EXPERIMENTAL_COOPERATIVE_MATRIX.
  • Queries cooperative_matrix_properties and verifies that 8×8 f32 or chosen configuration is supported.
  • Dispatches the compute pipeline with appropriate grid dimensions.

Notes and best practices Copy Chapter

• Always query adapter.cooperative_matrix_properties() and check that the configuration your shaders use exists.
  Do not hard-code assumptions about available tile sizes or element types.
• Treat the cooperative execution group as an abstract concept; avoid making assumptions about how
  tiles are mapped to lanes beyond what is guaranteed by the spec.
• Avoid divergent control flow around cooperative operations.
• Consider providing a fallback non-cooperative implementation for devices that do not support the feature.
• This is an experimental extension; API and semantics may change across versions of wgpu and naga.

File: docs/api-specs/mesh_shading.md Copy Chapter

Mesh Shader Extensions

🧪Experimental🧪

wgpu supports an experimental version of mesh shading when Features::EXPERIMENTAL_MESH_SHADER is enabled.
The status of the implementation is documented in the mesh-shading issue.

Note: The features documented here may have major bugs in them and are expected to be subject
to breaking changes. Suggestions for the API exposed by this should be posted on the issue above.

Mesh shaders overview Copy Chapter

What are mesh shaders? Copy Section

Mesh shaders are a new kind of rasterization pipeline intended to address some of the shortfalls with the vertex shader pipeline. The core idea of mesh shaders is that the GPU decides how to render the many small parts of a scene instead of the CPU issuing a draw call for every small part or issuing an inefficient monolithic draw call for a large part of the scene.

Mesh shaders are specifically designed to be used with meshlet rendering, a technique where every object is split into many subobjects called meshlets that are each rendered with their own parameters. With the standard vertex pipeline, each draw call specifies an exact number of primitives to render and the same parameters for all vertex shaders on an entire object (or even multiple objects). This doesn't leave room for different LODs for different parts of an object, for example a closer part having more detail, nor does it allow culling smaller sections (or primitives) of objects. With mesh shaders, each task workgroup might get assigned to a single object. It can then analyze the different meshlets(sections) of that object, determine which are visible and should actually be rendered, and for those meshlets determine what LOD to use based on the distance from the camera. It can then dispatch a mesh workgroup for each meshlet, with each mesh workgroup then reading the data for that specific LOD of its meshlet, determining which and how many vertices and primitives to output, determining which remaining primitives need to be culled, and passing the resulting primitives to the rasterizer.

Mesh shaders are most effective in scenes with many polygons. They can allow skipping processing of entire groups of primitives that are facing away from the camera or otherwise occluded, which reduces the number of primitives that need to be processed by more than half in most cases, and they can reduce the number of primitives that need to be processed for more distant objects. Scenes that are not bottlenecked by geometry (perhaps instead by fragment processing or post processing) will not see much benefit from using them.

Mesh shaders were first shown off in NVIDIA's asteroids demo. Now, they form the basis for Unreal Engine's Nanite.

Mesh shader pipeline Copy Section

With the current pipeline set to a mesh pipeline, a draw command like
render_pass.draw_mesh_tasks(x, y, z) takes the following steps:

• If the pipeline has a task shader stage:

  • Dispatch a grid of task shader workgroups, where x, y, and z give
    the number of workgroups along each axis of the grid. Each task shader
    workgroup produces a mesh shader workgroup grid size (mx, my, mz) and a
    task payload value mp.

  • For each task shader workgroup, dispatch a grid of mesh shader workgroups,
    where mx, my, and mz give the number of workgroups along each axis
    of the grid. Pass mp to each of these workgroup's mesh shader
    invocations.

• Alternatively, if the pipeline does not have a task shader stage:

  • Dispatch a single grid of mesh shader workgroups, where x, y, and z
    give the number of workgroups along each axis of the grid. These mesh
    shaders receive no task payload value.

• Each mesh shader workgroup produces a list of output vertices, and a list of
  primitives built from those vertices. The workgroup can supply per-primitive
  values as well, if needed. Each primitive selects its vertices by index, like
  an indexed draw call, from among the vertices generated by this workgroup.

  Unlike a grid of ordinary compute shader workgroups collaborating to build
  vertex and index data in common storage buffers, the vertices and primitives
  produced by a mesh shader workgroup are entirely private to that workgroup,
  and are not accessible by other workgroups.

• Primitives produced by a mesh shader workgroup can have a culling flag. If a
  primitive's culling flag is false, it is skipped during rasterization.

• The primitives produced by all mesh shader workgroups are then rasterized in
  the usual way, with each fragment shader invocation handling one pixel.

  Attributes from the vertices produced by the mesh shader workgroup are
  provided to the fragment shader with interpolation applied as appropriate.

  If the mesh shader workgroup supplied per-primitive values, these are
  available to each primitive's fragment shader invocations. Per-primitive
  values are never interpolated; fragment shaders simply receive the values
  the mesh shader workgroup associated with their primitive.

`wgpu` API Copy Chapter

New `wgpu` functions Copy Section

Device::create_mesh_pipeline - Creates a mesh shader pipeline. This is very similar to creating a standard render pipeline, except that it takes a mesh shader state and optional task shader state instead of a vertex state. If the task state is omitted, during rendering the number of workgroups is passed directly from the draw call to the mesh shader state, with an empty payload.

RenderPass::draw_mesh_tasks - Dispatches the mesh shader pipeline. This ignores render pipeline specific information, such as vertex buffer bindings and index buffer bindings. The dispatch size must adhere to the limits described below.

RenderPass::draw_mesh_tasks_indirect, RenderPass::multi_draw_mesh_tasks_indirect and RenderPass::multi_draw_mesh_tasks_indirect_count - Dispatches the mesh shader pipeline with dispatch size taken from a buffer. This ignores render pipeline specific information, such as vertex buffer bindings and index buffer bindings. The dispatch size must adhere to the limits described below. Analogous to draw_indirect, multi_draw_indirect and multi_draw_indirect_count. Requires the corresponding indirect feature to be enabled.

An example of using mesh shaders to render a single triangle can be seen here.

Features Copy Section

• Using mesh shaders requires enabling Features::EXPERIMENTAL_MESH_SHADER.
• Using mesh shaders with multiview requires enabling Features::EXPERIMENTAL_MESH_SHADER_MULTIVIEW.
• Using mesh shaders with point primitives requires enabling Features::EXPERIMENTAL_MESH_SHADER_POINTS.
• Queries are unsupported
• Primitive index support will be added once support lands in for them in general.

Limits Copy Section

NOTE: More limits will be added when support is added to naga.

• Limits::max_task_mesh_workgroup_total_count - the maximum total number of workgroups from a draw_mesh_tasks command or similar. The dimensions passed must be less than or equal to this limit when multiplied together.
• Limits::max_task_mesh_workgroups_per_dimension - the maximum for each of the 3 workgroup dimensions in a draw_mesh_tasks command. Each dimension passed must be less than or equal to this limit.
• max_task_invocations_per_workgroup - The maximum total number of threads in a task shader workgroup, given by workgroupSize.x * workgroupSize.y * workgroupSize.z.
• max_task_invocations_per_dimension the maximum value for each of workgroupSize.x, workgroupSize.y and workgroupSize.z in task shader workgroups.
• max_mesh_invocations_per_workgroup - The maximum total number of threads in a mesh shader workgroup, given by workgroupSize.x * workgroupSize.y * workgroupSize.z.
• max_mesh_invocations_per_dimension the maximum value for each of workgroupSize.x, workgroupSize.y and workgroupSize.z in mesh shader workgroups.
• max_task_payload_size - the size of an var<task_payload> variable, in bytes.
• max_mesh_output_vertices - the maximum number of vertices that a single mesh shader workgroup may output.
• max_mesh_output_primitives - the maximum number of primitives that a single mesh shader workgroup may output.
• max_mesh_multiview_count - the maximum number of views used when multiview rendering with a mesh shader pipeline.
• max_mesh_output_layers - the maximum number of output layers for a mesh shader pipeline.

Naga implementation Copy Chapter

Supported frontends Copy Section

• 🛠️ WGSL
• ❌ SPIR-V
• 🚫 GLSL

Supported backends Copy Section

• 🛠️ SPIR-V
• 🛠️ HLSL
• ❌ MSL
• 🚫 GLSL
• 🚫 WGSL

✔️ = Complete
🛠️ = In progress
❌ = Planned
🚫 = Unplanned/impossible

`WGSL` extension specification Copy Chapter

The majority of changes relating to mesh shaders will be in WGSL and naga.

Using any of these features in a wgsl program will require adding the enable wgpu_mesh_shader; directive to the top of a program.

Two new shader stages will be added to WGSL. Fragment shaders are also modified slightly. Both task shaders and mesh shaders are allowed to use any compute-available functionality, including subgroup operations.

Task shader Copy Section

A function with the @task attribute is a task shader entry point. A mesh shader pipeline may optionally specify a task shader entry point, and if it does, mesh draw commands using that pipeline dispatch a task shader grid of workgroups running the task shader entry point. Like compute shader dispatches, the three-component size passed to draw_mesh_tasks, or drawn from the indirect buffer for its indirect variants, specifies the size of the task shader grid as the number of workgroups along each of the grid's three axes.

A task shader entry point must have a @workgroup_size attribute, meeting the same requirements as one appearing on a compute shader entry point.

A task shader entry point must also have a @payload(G) property, where G is the name of a global variable in the task_payload address space. Each task shader workgroup has its own instance of this variable, visible to all invocations in the workgroup. Whatever value the workgroup collectively stores in that global variable becomes the task payload, and is provided to all invocations in the mesh shader grid dispatched for the workgroup. A task payload variable must be at least 4 bytes in size.

A task shader entry point must return a vec3 value decorated with @builtin(mesh_task_size). The return value of each workgroup's first invocation (that is, the one whose local_invocation_index is 0) is taken as the size of a mesh shader grid to dispatch, measured in workgroups. (If the task shader entry point returns vec3(0, 0, 0), then no mesh shaders are dispatched.) Mesh shader grids are described in the next section.

The output of a task shader is set to zero if it violates either of the limits max_task_mesh_workgroup_total_count or max_task_mesh_workgroups_per_dimension.

Each task shader workgroup dispatches an independent mesh shader grid: in mesh shader invocations, @builtin values like workgroup_id and global_invocation_id describe the position of the workgroup and invocation within that grid;
and @builtin(num_workgroups) matches the task shader workgroup's return value. If this output violates any limits, it may be zeroed or cause undefined behavior, depending on the compilation options. Mesh shaders dispatched for other task shader workgroups are not included in the count. If it is necessary for a mesh shader to know which task shader workgroup dispatched it, the task shader can include its own workgroup id in the task payload.

Task shaders can use compute and subgroup builtin inputs, in addition to view_index and draw_id.

Mesh shader Copy Section

A function with the @mesh attribute is a mesh shader entry point. Mesh shaders must not return anything.

Like compute shaders, mesh shaders are invoked in a grid of workgroups, called a mesh shader grid. If the mesh shader pipeline has a task shader, then each task shader workgroup determines the size of a mesh shader grid to be dispatched, as described above. Otherwise, the three-component size passed to draw_mesh_tasks, or drawn from the indirect buffer for its indirect variants, specifies the size of the mesh shader grid directly, as the number of workgroups along each of the grid's three axes.

If the mesh shader pipeline has a task shader entry point, then the pipeline's mesh shader entry point must also have a @payload(G) attribute, and the sizes of the variables must match. Mesh shader invocations can read from, but not write to, this variable, which is initialized to whatever value was written to it by the task shader workgroup that dispatched this mesh shader grid.

If the mesh shader pipeline does not have a task shader entry point, then the mesh shader entry point must not have any @payload attribute.

A mesh shader entry point must have the following attributes:

• @workgroup_size: this has the same meaning as when it appears on a compute shader entry point.

• @mesh(VAR): Here, VAR represents a workgroup variable storing the output information.

All mesh shader outputs are per-workgroup, and taken from the workgroup variable specified above. The type must have exactly 4 fields:

• A field decorated with @builtin(vertex_count), with type u32: this field represents the number of vertices that will be drawn
• A field decorated with @builtin(primitive_count), with type u32: this field represents the number of primitives that will be drawn
• A field decorated with @builtin(vertices), typed as an array of V, where V is the vertex output type as specified below
• A field decorated with @builtin(primitives), typed as an array of P, where P is the primitive output type as specified below

For a vertex count NV, the first NV elements of the vertex array above are outputted. Therefore, NV must be less than or equal to the size of the vertex array. The same is true for primitives with NP.

The vertex output type V must meet the same requirements as a struct type returned by a @vertex entry point: all members must have either @builtin or @location attributes, there must be a @builtin(position), and so on.

The primitive output type P must be a struct type, every member of which either has a @location or @builtin attribute. All members decorated with @location must also be decorated with @per_primitive, as must the corresponding fragment input. The @per_primitive decoration may only be applied to members decorated with @location. The following @builtin attributes are allowed:

• triangle_indices, line_indices, or point_index: The annotated member must be of type vec3, vec2, or u32.

  The member's components are indices (or, its value is an index) into the list of vertices generated by this workgroup, identifying the vertices of the primitive to be drawn. These indices must be less than the value of numVertices passed to setMeshOutputs.

  The type P must contain exactly one member with one of these attributes, determining what sort of primitives the mesh shader generates.

• cull_primitive: The annotated member must be of type bool. If it is true, then the primitive is skipped during rendering.

The @location attributes of P and V must not overlap, since they are merged to produce the user-defined inputs to the fragment shader.

Mesh shaders may write to the primitive_index builtin. This is treated just like a field decorated with @location, so if the mesh shader outputs primitive_index the fragment shader must input it, and if the fragment shader inputs it, the mesh shader must write it (unlike vertex shader pipelines).

Mesh shaders can use compute and mesh shader builtin inputs, in addition to view_index, and if no task shader is present, draw_id.

Fragment shader Copy Section

Fragment shaders can access vertex output data as if it is from a vertex shader. They can also access primitive output data, provided the input is decorated with @per_primitive. The @per_primitive decoration may only be applied to inputs or struct members decorated with @location.

The primitive state is part of the fragment input and must match the output of the mesh shader in the pipeline. Using @per_primitive also requires enabling the mesh shader extension. Additionally, the locations of vertex and primitive input cannot overlap.

Full example Copy Section

The following is a full example of WGSL shaders that could be used to create a mesh shader pipeline, showing off many of the features.

/* Detailed source-code truncated for AI context efficiency. */

File: docs/api-specs/ray_tracing.md Copy Chapter

Ray Tracing Extensions

🧪Experimental🧪

wgpu supports an experimental version of ray tracing which is subject to change. The extensions allow for acceleration structures to be created and built (with
Features::EXPERIMENTAL_RAY_QUERY enabled) and interacted with in shaders. Currently naga only supports ray queries
(accessible with Features::EXPERIMENTAL_RAY_QUERY enabled in wgpu).

Note: The features documented here may have major bugs in them and are expected to be subject
to breaking changes, suggestions for the API exposed by this should be posted on the ray-tracing issue.
Large changes may mean that this documentation may be out of date.

This is not an introduction to raytracing, and assumes basic prior knowledge, to look at the fundamentals look at
an introduction.

`wgpu`'s raytracing API: Copy Chapter

The documentation and specific details of the functions and structures provided
can be found with their definitions.

Acceleration structures do not have a separate feature, instead they are enabled by Features::EXPERIMENTAL_RAY_QUERY, unlike vulkan.
When ray tracing pipelines are added, that feature will also enable acceleration structures.

A Blas can be created with Device::create_blas.
A Tlas can be created with Device::create_tlas.

The Tlas reference can be placed in a bind group to be used in a shader. A reference to a Blas can
be used to create TlasInstance alongside a transformation matrix, custom data
(this can be any data that should be given to the shader on a hit) which only the first 24
bits may be set, and a mask to filter hits in the shader.

A Blas must be built in either the same build as any Tlas it is used to build or an earlier build call.
Before a Tlas is used in a shader it must

• have been built
• have all Blases that it was last built with to have last been built in either the same build as
  this Tlas or an earlier build call.

[`Blas`]es and [`Tlas`]es Copy Section

Both acceleration structures are opaque objects. These objects typically (but are not guaranteed to)
contain a tree-like structure of objects that are quick to be intersected with rays (as of 2026, these
are usually axis-aligned bounding boxes) and contain two or more "branches" (these objects) and "leaves".
For a Tlas these "leaves" are references to the Blas (not copies, hence a Tlas must be rebuild after
any Blas it is storing is modified). For a Blas with triangle geometry the "leaves" are triangles. For
a Blas with AABBs, the "leaves" might not exist or might contain only a small amount of information. This is
because there is no requirement to store the AABBs as this might require a third intersection type. However,
the "branches" will probably be fairly closely approximating the AABBs to maintain fairly good trace performance.

Building

Building an acceleration structure is a slow operation as it is likely to require building a tree-like structure.
This happens on the GPU and can be encoded using CommandEncoder::build_acceleration_structures.

For Blases with triangle geometry, this will copy the triangles out of the vertex buffer (using indexing if
provided), multiply them by the geometry transform matrix (if provided), and store a representation of this
somewhere in the structure to be traced against (and computes surrounding objects for sets of triangles, repeating
the process for objects of the surrounding objects etc. if the implementation uses a tree-like structure).

For Blases with AABB geometry, this will construct object(s) in such a way that all possible rays will generate
a candidate intersection that would have generated a candidate intersection with the AABBs. Note that this may be
(though is very unlikely to be) a volume containing all space. (These may then have objects constructed around them
if the implementation has a tree-like structure.)

For Tlases, these will store references to the Blases within the TlasInstances as well as transforms,
masks, and any other information required. (These Blases then may have objects constructed around their transformed
positions if the implementation uses a tree-like structure.) wgpu will store references to the Blases to keep
them within device memory and alive while the Tlas uses them. However, if you use functions to build it using
the underlying functions (e.g. by using as_hal functions), you are responsible for keeping the Blases alive.

Some memory is allocated when building to be "scratch" data (a temporary buffer used by the GPU to store data during
the build) and instance staging memory (to copy the instances to the GPU). The some of this can be reused for between
the Blas and Tlas builds so in general it is advisable to try and integrate the builds together. Because
building is slow, it should be done as few times as possible. For moving geometry (players, particles, etc.), use
PREFER_FAST_BUILD
to speed up builds. Updating (performing a partial rebuild) is currently unsupported, but may be implemented in the
future. You should compact any geometry which will not change (e.g. static level geometry) once it has been built.

[`Blas`] compaction Copy Section

Once a Blas has been built, it can be compacted. Acceleration structures are allocated conservatively, without
knowing the exact data that is inside them. Once a Blas has been built, the driver can make data specific
optimisations to make the BLAS smaller. To begin compaction call Blas::prepare_compaction_async on it. This
method waits until all builds operating on the Blas are finished, prepares the Blas to be compacted, and runs
the given callback. To check whether the Blas is ready, you can also call Blas::ready_for_compaction instead of
waiting for the callback (useful if you are asynchronously compacting a large number of Blases). Submitting a
rebuild of a Blas terminates any Blas::prepare_compaction_async, preventing the callback from being called, and
making the Blas no longer ready to compact. Once a Blas is ready for compaction, it can be compacted using
Queue::compact_blas this returns the new compacted Blas, which is independent of the Blas passed in. The
other Blas can be used for other things, including being rebuilt without affecting the new Blas. The returned
Blas behaves largely like the Blas it was created from, except that it can be neither rebuilt, nor compacted
again.

An example of compaction being run when Blases are ready, this would be in a situation when memory was not a major
problem, otherwise (e.g. if you get an out of memory error) you should compact immediately (and switching all
non-compacted Blases to compacted ones).

use std::iter;
use wgpu::Blas;

struct BlasToBeCompacted {
    blas: Blas,
    /// The index into the TlasInstance this BLAS is used in.
    tlas_index: usize,
}

fn render(/*whatever args you need to render*/) {
  /* additional code to prepare the renderer */
  //An iterator of whatever BLASes you have called `prepare_compaction_async` on.
  let blas_s_pending_compaction: impl Iterator = iter::empty();
  for blas_to_be_compacted in blas_s_pending_compaction {
    if blas_to_be_compacted.blas.ready_for_compaction() {
        let compacted_blas = queue.compact_blas(&blas_to_be_compacted.blas);
        tlas_instance[blas_to_be_compacted.tlas_index].set_blas(&compacted_blas);
    }
  }
  let mut encoder =
    device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
  /* do other preparations on the TlasInstance.*/
  encoder.build_acceleration_structures(iter::empty(), iter::once(&tlas_package));
  /* more render code */
  queue.submit([encoder.finish()]);
}

`naga`'s raytracing API: Copy Chapter

naga supports ray queries (also known as inline raytracing). To enable basic ray query functions you must add
enable wgpu_ray_query to the shader, ray queries and acceleration structures also support tags which require extra
enable extensions (see Acceleration structure tags for more info). Ray tracing pipelines are currently in
development. Naming is mostly taken from vulkan.

Ray Queries Copy Section

/* Detailed source-code truncated for AI context efficiency. */

Ray Tracing Pipelines Copy Section

Functions

// Begins to check where (if anywhere) the ray defined by `ray_desc` hits in `acceleration_structure` calling through the `any_hit` shaders and `closest_hit` shader if something was hit or the `miss` shader if no hit was found
traceRay<T>(acceleration_structure: acceleration_structure, ray_desc: RayDesc, payload: ptr<ray_payload, T>)

[!CAUTION]

⚠️Undefined behavior ⚠️:

Calling traceRay inside another traceRay more than max_recursion_depth times

*this is only known undefined behaviour, and will be worked around in the future.

New shader stages

// First stage to be called, allowed to call `traceRay`
@ray_generation
fn rg() {}

// Stage called on any hit that is not opaque, not allowed to call `traceRay`
@any_hit
fn ah() {}

// Stage called on the closest hit, allowed to call `traceRay`
@closest_hit
fn ch() {}

// Stage call if there was never a hit, allowed to call `traceRay`
@miss
fn miss() {}

Acceleration structure tags Copy Section

These are tags that can be added to a acceleration structure (acceleration_structure ->
acceleration_structure<... insert tags here! ...>) and to a ray query (ray_query ->
ray_query<... insert tags here! ...>). These require more features.

File: docs/release-checklist.md Copy Chapter

This is an overview of how to run wgpu releases.

Structure Copy Chapter

We do a major breaking release every 12 weeks. This happens no matter the status of various in-flight projects.

We do a patch releases as needed in the weeks between major releases. Once a new major release is cut, we stop doing patch releases for the previous major release unless there is a critical bug or a compilation issue.

People Copy Chapter

Anyone in the @gfx-rs/wgpu team can perform these steps.

Major Release Process Copy Chapter

Approx 1 Week Before:

• Determine if glow (@groves), rspirv (@gfx-rs/wgpu) or any other dependant crates will need a release. If so, coordinate with their maintainers.
• Go through the changelog:
  • Re-categorize miscategorized items.
  • Edit major changes so a user can easily understand what they need to do.
  • Add missing major changes that users need to know about.
  • Copy-edit the changelog for clarity.

Day of Release:

• Update the version number in the root Cargo.toml to the new version, this will update all crates to the new version.
• Bump the wgpu dependency numbers in the following places:
  • Cargo.toml
  • examples/standalone/*
  • examples/bug-repro/*
• Grep for the previous version to ensure various documentation links are updated.
  • For example, if the previous version was v24.0.0, grep for v24 and 24.0
• Ensure glow and rspirv are updated to the latest version if needed.
• Add a new header for the changelog with the release version and date.
• Create a PR with all of the version changes and changelog updates.
• While waiting on the PR, do a dry run of publishing.
  bash Copy

  cargo publish --dry-run --workspace --all-features --exclude deno_webgpu

• Once the PR is CI clean and publish worked, (force) merge it.
• Checkout trunk with the merged PR.
• Publish! These commands can be pasted directly into your terminal in a single command, and they will publish everything.
  bash Copy

  cargo publish --workspace --all-features --exclude deno_webgpu

• If there were any newly published crates, ensure github:gfx-rs:wgpu is added as an owner of that crate.
• Create a new signed tag called vX.Y.Z and push it to the repo.
  • For each crate being released (viz., every publish-able crate that is not deno*), create a new tag of the form {crate_name}-vX.Y.X.
• Create a new release on the wgpu repo with the changelog from this version, targeting that tag
• Create a branch with the with the new version vX and push it to the repo.
  • On this branch, remove the [!NOTE] at the top of [examples/README.md].
• Complete the release's milestone on GitHub.
• Create a new milestone for the next release, in 12 weeks time.
• Update the release checklist with any needed changes.
• Publish the link to the github release in the following places.
  • r/rust.
    • Add an AMA comment.
  • Crosspost to r/rust_gamedev.
    • Add an AMA comment.
  • Include the r/rust post shortlink in the following posts as well:
  • wgpu matrix
  • Rust Gamedev Discord in the #crates and #wgpu channel
  • Bevy Discord in the #rendering-dev channel
  • Graphics Programming Discord in the #webgpu channel
  • Rust Community Discord in the #games-and-graphics channel

Patch Release Process Copy Chapter

• Enumerate all PRs that haven't been backported yet. These use the PR: needs back-porting label. GH Link
• On your own branch based on the latest release branch. Cherry-pick the PRs that need to be backported. When modifying the commits, use --append to retain their original authorship.
• Remove the needs-backport label from the PRs.
• Fix the changelogs items and add a new header for the patch release with the release version and date.
• Create a PR with all of the version changes and changelog updates into the release branch.
• While waiting on the PR, do a dry run of publishing.
  bash Copy

  cargo publish --dry-run --workspace --all-features --exclude deno_webgpu

• Once the PR is CI clean and publish worked, (force) merge it.
• Checkout the version branch with the merged PR.
• Publish! These commands can be pasted directly into your terminal in a single command, and they will publish everything.
  bash Copy

  cargo publish --workspace --all-features --exclude deno_webgpu

• Create a new release on the wgpu repo with the relevant changelog included, based on a new tag called vX.Y.Z in the release branch.
  • For each crate released, also create a tag {crate_name}-vX.Y.Z.
• Backport the changelog and version bumps to the trunk branch.
  • Ensure that any items in the newly-released changelog don't appear in the "unreleased" section of the trunk changelog.
• Update the release checklist with any needed changes.

File: docs/review-checklist.md Copy Chapter

Review Checklist

This is a collection of notes on things to watch out for when
reviewing pull requests submitted to wgpu and Naga.

Ideally, we want to keep items off this list entirely:

• Using Rust effectively can turn some mistakes into compile-time
  errors. For example, in Naga, using exhaustive matching ensures that
  changes to the IR will cause compile-time errors in any code that
  hasn't been updated.

• Refactoring can gather together all the code responsible for
  enforcing some invariant in one place, making it clear whether a
  change preserves it or not. For example, Naga localizes all handle
  validation to naga::valid::Validator::validate_module_handles,
  allowing the rest of the validator to assume that all handles are
  valid.

• Offering custom abstractions can help contributors avoid
  implementing a weaker abstraction by themselves. For example,
  because HandleSet and HandleVec are used throughout Naga,
  contributors are less likely to write code that uses a BitSet or
  Vec on handle indices, which would invite bugs by erasing the
  handle types.

This checklist gathers up the concerns that we haven't found a
satisfying way to address in a more robust way.

Naga Copy Chapter

General Copy Section

• If your change iterates over a collection, did you ensure the
  order of iteration was deterministic? Using HashMap and
  HashSet is fine, as long as you don't iterate over it.
• If you insert elements into a set or map that you expect are not
  already present, did you make an assertion about insert's
  return value?

WGSL Extensions Copy Section

• If you added a new feature to WGSL that is not covered by the
  WebGPU specification:
  • Did you add a Capability flag for it?
  • Did you document the feature fully in that flag's doc comment?
  • Did you ensure the validator rejects programs that use the
    feature unless its capability is enabled?

IR changes Copy Section

If your change adds or removes Handles from the IR:

• Did you update handle validation in valid::handles?
• Did you update the compactor in compact?
• Did you update back::pipeline_constants::adjust_expr?

If your change adds a new operation:

• Did you update the typifier in proc::typifier?
• Did you update the validator in valid::expression?
• If the operation can be used in constant expressions, did you
  update the constant evaluator in proc::constant_evaluator?

Backend changes Copy Section

• If your change introduces any new identifiers to generated code,
  how did you ensure they won't conflict with the users'
  identifiers? (This is usually not relevant to the SPIR-V
  backend.)
  • Did you use the Namer to generate a fresh identifier?
  • Did you register the identifier as a reserved word with the Namer?
  • Did you use a reserved prefix registered with the Namer?

File: docs/testing.md Copy Chapter

Testing in `wgpu` and `naga`

There exist a large variety of tests within the wgpu repository
to make sure we can easily test all the aspects of our libraries.
This document serves as a guide to each class of test, and what
they are used for.

Requirements Copy Chapter

The tests require that the Vulkan SDK
is installed on the system and the bin folder of the SDK is in your PATH.
Without this some tests may fail to run, or report false negatives.

Additionally you require you run the tests with cargo-nextest.
This is what our xtask calls. You can install it with cargo install cargo-nextest.

Run All Tests Copy Chapter

To run all tests, run cargo xtask test from the root of the repository.

Test Breakdown Copy Chapter

This is a table of contents, in the form of the repository's directory structure.

• benches
  • benches
• cts_runner
• examples
  • features
• naga
  • tests
    • example_wgsl
    • snapshot
    • spirv-capabilities
    • validation
    • wgsl_errors
• player
  • tests
• tests
  • compile
  • dependency
  • gpu
  • trace
  • validation

And where applicable unit-tests
are scatteredthroughout the codebase.

Benchmark Tests Copy Chapter

• Located in: benches/benches
• Run with cargo nextest run --bench wgpu-benchmark
• wgpu benchmarks for performance testing.

These are benchmarks that test the performance of wgpu in various
scenarios. When run as part of the test suite, they run a single
iteration of each benchmark to ensure they continue to function.

These tests only run on your system's default GPU.

The benchmarks should be very careful to avoid doing any significant
work (including connecting to a GPU) outside of the various benchmark
criterion functions. If this is done, the benchmarks will take a long
time to list available tests, slowing down the test suite.

To run the benchmarks for benchmarking purposes, use cargo bench.

Example Tests Copy Chapter

• Located in: examples/features
• Run with cargo xtask test --bin wgpu-examples
• Uses a custom #[apply(gpu_test!)] harness.
• wgpu integration tests, with access to wgpu_test helpers.

These tests validate that the examples are functioning correctly
and do not have any regressions. They use the same harness as the
gpu tests, see that section for more information
on the harness.

These tests use nv-flip's image comparison through the wgpu
example framework to validate that the images outputted by the
examples are within tolerance of the expected output.

Examples written in examples/standalone do not have tests, as
they should be easy to copy into a standalone project.

`naga` Example Tests Copy Chapter

• Located in: naga/tests/naga/example_wgsl
• Run with cargo nextest run --test naga example_wgsl

This simple test ensures that all wgsl files in the examples
directory can be parsed by naga's wgsl parser and validate correctly.

`naga` Snapshot Tests Copy Chapter

• Located in: naga/tests/naga/snapshot, naga/tests/in, and naga/tests/out
• Run with cargo nextest run --test naga snapshots
• Data driven snapshot tests for naga's input/output.

These tests are snapshot tests for nagas parsers and code generators.
There are inputs in wgsl, spirv, and glsl. There are outputs for
hlsl, spirv, wgsl, msl, glsl, and naga's internal IR. The tests
can be configured by a sidecar toml file of the same name as the input file.

This is the goto tool for testing all kinds of codegen and parsing features.

To avoid clutter we generally use the following pattern:

• wgsl tests generate output to all backends.
• spirv, glsl tests generate wgsl output

This "butterfly" pattern ensures we don't need to test the
full matrix of possibilities to get full coverage.

While we do not run the results of the code generators, we do
test that the generated code is valid. This is done by running
cargo xtask validate in the naga directory and
will use the respective tool to validate the generated code.

`naga` SPIR-V Capabilities Tests Copy Chapter

• Located in: naga/tests/naga/spirv_capabilities
• Run with cargo nextest run --test naga spirv_capabilities
• Uses the standard #[test] harness.

These tests convert the given wgsl snippet to spirv and
then assert that the spirv has enabled the expected capabilities.

`naga` Validation Tests Copy Chapter

• Located in: naga/tests/naga/validation
• Run with cargo nextest run --test naga validation

These are hand rolled tests against the naga's validator.
If you don't need to test the validator with a custom module,
and can use the wgsl frontend, you should put the test in
the wgsl errors tests.

`naga` WGSL Error Tests Copy Chapter

• Located in: naga/tests/naga/wgsl_errors
• Run with cargo nextest run --test naga wgsl_errors

These are tests for the error messages that the wgsl frontend
produces. Additionally you can check that a given validation error
is produced by the validator from a given wgsl snippet.

`player` Tests Copy Chapter

• Located in: player/tests
• Run with cargo nextest run --test player
• Data driven tests using the player's replay system.
• wgpu integration tests.

These are soft-deprecated tests which are another way to write
API tests. These use captures of the api calls and replay them
to assert on the behavior. They are very difficult to write, and
the trace capturing system is currently broken, so these
tests exist, but you should not write new ones.

These tests only run on your system's default GPU.

`wgpu` Compile Tests Copy Chapter

• Located in: tests/tests/wgpu-compile
• Run with cargo nextest run --test wgpu-compile
• trybuild tests of all rust files in tests/tests/wgpu-compile/fail directory.

These use the trybuild crate to test a few scenarios where
the wgpu crate is expected to fail to compile. This mainly
revolves around ensuring lifetimes are properly handled when
dropping passes, etc.

`wgpu` Dependency Tests Copy Chapter

• Located in: tests/tests/wgpu-dependency
• Run with cargo nextest run --test wgpu-dependency
• Tests against cargo tree.

These tests ensure that the wgpu crate has the correct dependency
tree on all platforms. It's super easy to subtly mess up the dependencies
which can cause issues or extra dependencies to be pulled in.

This provides a way to ensure that our toml files are correct.

`wgpu` GPU Tests Copy Chapter

• Located in: tests/tests/wgpu-gpu
• Run with cargo xtask test --test wgpu-gpu
• Uses a custom #[apply(gpu_test!)] harness.
• wgpu integration tests, with access to wgpu_test helpers.

These tests use a custom harness to run each test on all GPUs
available on the system. They are general integration tests
that write code against the normal wgpu API and assert on the behavior.

These tests are useful to check the runtime behavior of a program,
validate that there are no validation errors coming from the
vulkan/dx12/metal validation layers, and ensure behavior
is the same across GPUs. If the test does not need to run on a
real GPU, it should be in the validation tests instead.

There is a special parameter system that deals with if a GPU
can support the given test, and dealing with expectation
management for tests that are expected to fail due to driver or wgpu bugs.

Normal #[test]s will not be found in this test crate, as we use a custom harness.

See also the example tests for additional GPU tests.

`wgpu` Trace Tests Copy Chapter

• Located in: tests/tests/wgpu_trace.rs
• Run with cargo nextest run --test wgpu_trace
• Use the standard #[test] harness.

These tests are focused on testing the tracing functionality in wgpu. They
use the a special noop backend which does not connect to a real GPU.

`wgpu` Validation Tests Copy Chapter

• Located in: tests/tests/wgpu-validation
• Run with cargo nextest run --test wgpu-validation
• Use the standard #[test] harness.
• wgpu integration tests, with access to wgpu_test helpers.

These tests are focused on testing the validation inside of wgpu-core.
They are written against the wgpu API, but are targeting a special noop
backend which does not connect to a real GPU.

This is significantly faster and simpler than running on real hardware,
and allows any validation logic to be checked, even if real hardware
does not support those features.

Unit Tests Copy Chapter

• Located throughout the codebase.
• Run with cargo nextest test -p
• Standard #[test]s.

Throughout the codebase we have standard #[test]s that test individual
functions or small parts of the codebase. These don't run on the gpu.

WebGPU CTS Copy Chapter

WebGPU includes a Conformance Test Suite to validate that implementations are
working correctly. We run cases from the CTS against wgpu using
Deno. A default list of enabled
tests is automatically run on pull requests in CI.

To run the default set of CTS tests locally, run:

cargo xtask cts

You can also specify a test selector on the command line:

cargo xtask cts 'webgpu:api,operation,command_buffer,basic:*'

Or supply your own test list in a file:

cargo xtask cts -f your_tests.lst

To find the full list of tests, go to the
web-based standalone CTS runner.

The version of the CTS used by cargo xtask cts is specified in
cts_runner/revision.txt.

Memory Initialization Testing Copy Chapter

Simple tests can fail to detect when necessary memory initialization is omitted because
allocations that happen to be satisfied by fresh kernel-provided pages are zero, even though
the actual allocator being invoked does not guarantee that.

To improve our coverage of memory initialization, we set LVP_POISON_MEMORY=true in Linux
(Vulkan) CI. This instructs llvmpipe to fill all newly initialized memory with a non-zero
value, so tests will reliably fail if the memory is not properly initialized.