## File: README.md # wgpu [](https://github.com/gfx-rs/wgpu/actions) [](https://codecov.io/gh/gfx-rs/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][webgpu], but is a fully native Rust library. It serves as the core of the WebGPU integration in Firefox, Servo, and Deno. ## Getting Started See our examples online at . You can see the Rust sources at [examples](examples) and run them directly with `cargo run --bin wgpu-examples `. ### Learning `wgpu` 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. [Learn Wgpu]: https://sotrh.github.io/learn-wgpu/ [WebGPU Fundamentals]: https://webgpufundamentals.org/ ### Wiki We have a [wiki](https://github.com/gfx-rs/wgpu/wiki) which has information on useful architecture patterns, debugging tips, and more getting started information. ### Need Help? Want to Contribute? The wgpu community uses Matrix and Discord to discuss. - [](https://matrix.to/#/#wgpu:matrix.org) - discussion of wgpu's development. - [](https://matrix.to/#/#wgpu-users:matrix.org) - discussion of using the library and the surrounding ecosystem. - [ ](https://discord.gg/X3MYBNXUMJ) - Dedicated support channel on the Rust Gamedev Discord. ### Other Languages To use wgpu in C or dozens of other languages, look at [wgpu-native](https://github.com/gfx-rs/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++. [Learn WebGPU (for C++)]: https://eliemichel.github.io/LearnWebGPU/ [webgpu]: https://gpuweb.github.io/gpuweb/ ## Quick Links | Docs | Examples | Changelog | | :-------------------: | :-----------------------: | :---------------------: | | [v30][rel-docs] | [v30][rel-examples] | [v30][rel-change] | | [`trunk`][trunk-docs] | [`trunk`][trunk-examples] | [`trunk`][trunk-change] | Contributors are welcome! See [CONTRIBUTING.md][contrib] for more information. [rel-docs]: https://docs.rs/wgpu/ [rel-examples]: https://github.com/gfx-rs/wgpu/tree/v30/examples#readme [rel-change]: https://github.com/gfx-rs/wgpu/releases [trunk-docs]: https://wgpu.rs/doc/wgpu/ [trunk-examples]: https://github.com/gfx-rs/wgpu/tree/trunk/examples#readme [trunk-change]: https://github.com/gfx-rs/wgpu/blob/trunk/CHANGELOG.md#unreleased [contrib]: CONTRIBUTING.md ## Supported Platforms | API | Windows | Linux/Android | macOS/iOS | Web (wasm) | | ------ | ------------------- | --------------- | --------- | ----------- | | Vulkan | ✅ | ✅ | 🌋 | | | Metal | | | ✅ | | | DX12 | ✅ | | | | | OpenGL | 🆗 (GL 3.3+), or 📐 | 🆗 (GL ES 3.0+) | 📐 | 🆗 (WebGL2) | | WebGPU | | | | ✅ | ✅ = First Class Support 🆗 = Downlevel/Best Effort Support 📐 = Requires the [ANGLE](https://github.com/gfx-rs/wgpu/wiki/Running-on-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](https://vulkan.lunarg.com/sdk/home#mac) translation layer 🛠️ = Unsupported, though open to contributions ## Environment Variables 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](https://docs.rs/wgpu/latest/wgpu/index.html?search=env) 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 For an overview of all the components in the gfx-rs ecosystem, see [the big picture](./docs/big-picture.png). ## MSRV policy 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`][moz-msrv]. - 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`][servo-msrv] - 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][nixos-msrv], though this is not required. [moz-msrv]: https://searchfox.org/mozilla-central/source/python/mozboot/mozboot/util.py [servo-msrv]: https://github.com/servo/servo/blob/main/Cargo.toml#L23 [nixos-msrv]: https://search.nixos.org/packages?show=rustc ## Testing and Environment Variables [Information about testing](./docs/testing.md), including where tests of various kinds live, and how to run the tests. ## Tracking the WebGPU and WGSL draft specifications The `wgpu` crate is meant to be an idiomatic Rust translation of the [WebGPU API][webgpu spec]. That specification, along with its shading language, [WGSL][wgsl spec], 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][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][naga 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. [webgpu spec]: https://www.w3.org/TR/webgpu/ [wgsl spec]: https://gpuweb.github.io/gpuweb/wgsl/ [naga]: https://github.com/gfx-rs/wgpu/tree/trunk/naga/ [naga bugs]: https://github.com/gfx-rs/wgpu/issues?q=is%3Aissue%20state%3Aopen%20label%3A%22naga%22 --- ## File: benches/README.md 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 ```sh # 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](../README.md#environment-variables). ## Comparing Against a Baseline 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: ```sh 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 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](https://github.com/wolfpld/tracy/releases/latest/). ```sh # 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](https://superluminal.eu/). ```sh # 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. ```sh # 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 #### `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 # 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](https://github.com/gfx-rs/wgpu/blob/trunk/docs/testing.md#webgpu-cts). The file [revision.txt](./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](./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 # 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](https://github.com/gfx-rs/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](https://github.com/gfx-rs/wgpu/wiki/Debugging-wgpu-Applications#tracing-infrastructure) to the specified directory. This op crate is tested primarily by running the [WebGPU conformance test suite](https://github.com/gpuweb/cts) using `wgpu`'s [`cts_runner`](https://github.com/gfx-rs/wgpu/blob/trunk/README.md#webgpu-conformance-test-suite). `cts_runner` also has a few [directed tests](https://github.com/gfx-rs/wgpu/tree/trunk/cts_runner/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 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 # 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 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) 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` 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: ```/dev/null/cooperative-matrix-host.rs#L1-40 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 ### `wgpu` feature - Using cooperative matrices requires enabling: - `Features::EXPERIMENTAL_COOPERATIVE_MATRIX` This feature may be restricted to certain backends and hardware. ### Hardware / backend notes 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 This section summarizes the host-side API elements related to cooperative matrices. (For exact signatures and details, refer to the Rust documentation.) ### Adapter - `Adapter::cooperative_matrix_properties() -> Vec` Returns all cooperative matrix configurations supported by the adapter/backend. ### Structures - `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` / `var` / `var` / etc. --- ## WGSL extension specification 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 Any WGSL program using cooperative matrices must declare an extension at the top of the shader, for example: ```/dev/null/example.wgsl#L1-3 enable wgpu_cooperative_matrix; ``` The shader is invalid if any cooperative matrix types or builtins are used without enabling this extension. ### Cooperative matrix types 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: ```/dev/null/example.wgsl#L1-8 // A: MxK, B: KxN, C: MxN type coop_matMxN; type coop_matMxN; type coop_matMxN; ``` Concrete examples (sizes and types must match a supported configuration from `Adapter::cooperative_matrix_properties`): ```/dev/null/example.wgsl#L10-20 // 8x8 single-precision tiles alias CoopMatA = coop_mat8x8; alias CoopMatB = coop_mat8x8; alias CoopMatC = coop_mat8x8; // 16x16 half-precision inputs, 16x16 f32 accumulator (mixed precision) alias CoopMat16x16A = coop_mat16x16; alias CoopMat16x16B = coop_mat16x16; alias CoopMat16x16C = coop_mat16x16; ``` 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 - `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 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. ```/dev/null/example.wgsl#L1-10 fn coopLoad( ptr: ptr, // base pointer to scalar or vector elements stride: u32 // elements between adjacent columns ) -> coop_matMxN; fn coopLoadT( ptr: ptr, // base pointer to scalar or vector elements stride: u32 // elements between adjacent rows ) -> coop_matMxN; ``` - 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. ```/dev/null/example.wgsl#L12-23 fn coopStore( value: coop_matMxN, ptr: ptr, stride: u32 ); fn coopStoreT( value: coop_matMxN, ptr: ptr, 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: ```/dev/null/example.wgsl#L15-23 fn coopMultiplyAdd( a: coop_matMAxKA, // A: MAxKA tile b: coop_matKBxNB, // B: KBxNB tile (KB == KA) c: coop_matMAxNB // C: MAxNB accumulator/result ) -> coop_matMAxNB; ``` 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: ```/dev/null/example.wgsl#L25-39 enable wgpu_cooperative_matrix; alias MatA = coop_mat8x8; alias MatB = coop_mat8x8; alias MatC = coop_mat8x8; // 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, ptr_b: ptr, ptr_c: ptr, 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 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: ```/dev/null/example.wgsl#L1-42 enable wgpu_cooperative_matrix; struct Matrices { // Row-major tiles for A, B, C — use the `…T` load/store variants. data: array, }; @group(0) @binding(0) var buf_a: Matrices; @group(0) @binding(1) var buf_b: Matrices; @group(0) @binding(2) var buf_c: Matrices; alias MatA = coop_mat8x8; alias MatB = coop_mat8x8; alias MatC = coop_mat8x8; @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(base_a, 8u); let b: MatB = coopLoadT(base_b, 8u); let c: MatC = coopLoadT(base_c, 8u); let result: MatC = coopMultiplyAdd(a, b, c); coopStoreT(result, base_c, 8u); } ``` --- ## Validation rules and undefined behavior 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 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 - 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 # 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](https://github.com/gfx-rs/wgpu/issues/7197). **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 ### What are mesh shaders? 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](https://www.youtube.com/watch?v=CRfZYJ_sk5E). Now, they form the basis for [Unreal Engine's Nanite](https://www.unrealengine.com/en-US/blog/unreal-engine-5-is-now-available-in-preview#Nanite). ### Mesh shader pipeline 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 ### New `wgpu` functions `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](../../examples/features/src/mesh_shader). ### Features - 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 > **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` 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 ### Supported frontends - 🛠️ WGSL - ❌ SPIR-V - 🚫 GLSL ### Supported backends - 🛠️ SPIR-V - 🛠️ HLSL - ❌ MSL - 🚫 GLSL - 🚫 WGSL ✔️ = Complete 🛠️ = In progress ❌ = Planned 🚫 = Unplanned/impossible ## `WGSL` extension specification 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 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 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 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 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 # 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](https://github.com/gfx-rs/wgpu/issues/1040). 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](https://developer.nvidia.com/blog/introduction-nvidia-rtx-directx-ray-tracing/). ## `wgpu`'s raytracing API: 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 [`Blas`]es 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 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 [`Blas`]es 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 [`Blas`]es 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 [`Tlas`]es, these will store references to the [`Blas`]es within the [`TlasInstance`]s as well as transforms, masks, and any other information required. (These [`Blas`]es then may have objects constructed around their transformed positions if the implementation uses a tree-like structure.) `wgpu` will store references to the [`Blas`]es 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 [`Blas`]es 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`](https://wgpu.rs/doc/wgpu_hal/struct.AccelerationStructureBuildFlags.html#associatedconstant.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 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 [`Blas`]es). 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 [`Blas`]es 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 [`Blas`]es to compacted ones). ```rust 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()]); } ``` [`Device::create_blas`]: https://wgpu.rs/doc/wgpu/struct.Device.html#method.create_blas [`CommandEncoder::build_acceleration_structures`]: https://wgpu.rs/doc/wgpu/struct.CommandEncoder.html#method.build_acceleration_structures [`Device::create_tlas`]: https://wgpu.rs/doc/wgpu/struct.Device.html#method.create_tlas [`Tlas`]: https://wgpu.rs/doc/wgpu/struct.Tlas.html [`Blas`]: https://wgpu.rs/doc/wgpu/struct.Blas.html [`TlasInstance`]: https://wgpu.rs/doc/wgpu/struct.TlasInstance.html [`Blas::prepare_compaction_async`]: https://wgpu.rs/doc/wgpu/struct.Blas.html#method.prepare_compaction_async [`Blas::ready_for_compaction`]: https://wgpu.rs/doc/wgpu/struct.Blas.html#method.ready_for_compaction [`Queue::compact_blas`]: https://wgpu.rs/doc/wgpu/struct.Queue.html#method.compact_blas ## `naga`'s raytracing API: `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 ``` /* Detailed source-code truncated for AI context efficiency. */ ``` ### Ray Tracing Pipelines Functions ```wgsl // 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(acceleration_structure: acceleration_structure, ray_desc: RayDesc, payload: ptr) ``` > [!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 ```wgsl // 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 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. | Tag | Requirements | Description | | --------------- | ------------------------------------- | ---------------------------------------------------------------------- | | `vertex_return` | `enable wgpu_ray_query_vertex_return` | Allows getting the vertices of the hit triangle when using ray queries | --- ## File: docs/release-checklist.md This is an overview of how to run wgpu releases. ## Structure 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 Anyone in the @gfx-rs/wgpu team can perform these steps. ## Major Release Process 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 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 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](https://www.reddit.com/r/rust/). - Add an AMA comment. - Crosspost to [r/rust_gamedev](https://www.reddit.com/r/rust_gamedev/). - Add an AMA comment. - Include the r/rust post shortlink in the following posts as well: - [wgpu matrix](https://matrix.to/#/#wgpu:matrix.org) - [Rust Gamedev Discord](https://discord.gg/X3MYBNXUMJ) in the #crates and #wgpu channel - [Bevy Discord](https://discord.com/invite/bevy) in the #rendering-dev channel - [Graphics Programming Discord](https://discord.gg/6mgNGk7) in the #webgpu channel - [Rust Community Discord](https://discord.gg/rust-lang-community) in the #games-and-graphics channel ## Patch Release Process - Enumerate all PRs that haven't been backported yet. These use the `PR: needs back-porting` label. [GH Link](https://github.com/gfx-rs/wgpu/pulls?q=sort%3Aupdated-desc+is%3Apr+label%3A%22PR%3A+needs+back-porting%22) - 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 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 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 # 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 ### General - [ ] 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 - [ ] 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 If your change adds or removes `Handle`s 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 - [ ] 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 # 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 The tests require that the [Vulkan SDK](https://vulkan.lunarg.com/sdk/home) 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 To run all tests, run `cargo xtask test` from the root of the repository. ## Test Breakdown This is a table of contents, in the form of the repository's directory structure. - benches - [benches](#benchmark-tests) - [cts_runner](#webgpu-cts) - examples - [features](#example-tests) - naga - tests - [example_wgsl](#naga-example-tests) - [snapshot](#naga-snapshot-tests) - [spirv-capabilities](#naga-spirv-capabilities-test) - [validation](#naga-validation) - [wgsl_errors](#naga-wgsl-error-tests) - player - [tests](#player-tests) - tests - [compile](#wgpu-compile-tests) - [dependency](#wgpu-dependency-tests) - [gpu](#wgpu-gpu-tests) - [trace](#wgpu-trace-tests) - [validation](#wgpu-validation-tests) And where applicable [unit-tests](#unit-tests) are scatteredthroughout the codebase. ## Benchmark Tests - 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 - 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](#wgpu-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 - 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 - 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 `naga`s 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 - 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 - 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](#naga-wgsl-error-tests) tests. ## `naga` WGSL Error Tests - 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 - 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 - 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 - 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 - 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](#wgpu-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](#example-tests) for additional GPU tests. ## `wgpu` Trace Tests - 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 - 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 - 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 WebGPU includes a Conformance Test Suite to validate that implementations are working correctly. We run cases from the CTS against wgpu using [Deno](https://deno.com/). A [default list of enabled tests](../cts_runner/test.lst) 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](https://gpuweb.github.io/cts/standalone/?runnow=0&worker=0&debug=0&q=webgpu:*). The version of the CTS used by `cargo xtask cts` is specified in [`cts_runner/revision.txt`](../cts_runner/revision.txt). ## Memory Initialization Testing 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.