Introduction

Welcome to the wgsl-rs operator's manual. This book is the canonical, user-facing reference for writing GPU shaders with wgsl-rs.

What is wgsl-rs?

With wgsl-rs you write a subset of Rust code and it automatically generates WGSL shaders and wgpu runtime linkage. Rust code written this way is fully operational — it can be run on the CPU — while the transpiled WGSL is isomorphic and should generate the same results on the GPU.

In short, with wgsl-rs, you can unit test and run your code on the CPU in Rust, and use the generated WGSL on the GPU, while sharing the same type definitions between the two.

Procedural macros are provided by the wgsl-rs-macros crate.

The Two Worlds Problem

A key insight that shapes everything about wgsl-rs is that it maintains two parallel representations of your shader:

1. Rust World: The code must compile as valid Rust that runs on the CPU. This is design decision #1 in the devlog.
2. WGSL World: The proc-macro transpiles to WGSL that runs on the GPU.

These are fundamentally different execution contexts with different memory models, and yet running a wgsl-rs program should produce roughly the same results in both "worlds".

Program setup (or preamble, if you will) and the runtime behavior is expected to be different for each world, but the results should match, within reason. This is why wgsl-rs provides CPU-side implementations of every WGSL builtin in wgsl_rs::std, and why the roundtrip test harness exists — to verify that the two worlds agree.

wgsl-rs vs Rust-GPU

Maybe — it depends on your needs.

Pros of wgsl-rs

• Lower barrier to entry: No custom Rust compiler backend required.
• Works with stable Rust: No need for nightly or custom toolchains.
• Editor support: The #[wgsl] macro makes supported syntax explicit, so your editor (via rust-analyzer) can help you write valid code.
• Immediate WGSL output: Use, inspect, and debug the generated WGSL anywhere WGSL is supported, including browsers and non-Rust projects.
• Human readable WGSL output: The WGSL that wgsl-rs produces is very close in structure to the Rust code you write, including binding names and types.
• Easy interop: Generated WGSL can be used in any WebGPU environment.

Cons of wgsl-rs

• WGSL only: Only works on platforms that support WGSL.
• Limited to WebGPU features: No support for features not present in WGSL (e.g., bindless resources).
• Subset of Rust: Only a strict subset of Rust is supported.
  • No traits
  • No borrowing
  • Very restricted module support

Note: wgsl-rs and Rust-GPU are not mutually exclusive! You can start with wgsl-rs and switch to Rust-GPU when you need more advanced features.

How to Read This Book

• New to wgsl-rs? Start with Installation and Hello, Triangle.
• Want to see working code? Jump to the Examples chapter, which catalogs all 35+ example modules with both Rust source and generated WGSL.
• Writing a real renderer? The wgpu Linkage and Generics & Templates chapters cover building pipelines from generic shaders.
• Building tooling on top of wgsl-rs? The Extensions and Memory Layout chapters are for you.
• Curious about the why? The Design Decisions chapter curates highlights from the devlog; the full DEVLOG.md has every decision recorded with dates and rationale.

Project Structure

The project is split into a few parts:

There's also a devlog that explains some of the decisions and tradeoffs made during the making of this library.

Funding

This project is funded through NGI Zero Commons, a fund established by NLnet with financial support from the European Commission's Next Generation Internet program. Learn more at the 2025 NLnet project page.

Sponsor

This work will always be free and open source. If you use it (outright or for inspiration), please consider donating.

💰 Sponsor 💝

Installation

Prerequisites

• Rust (stable toolchain). Install via rustup if you don't already have it.
• A GPU with a driver supported by wgpu. Required for roundtrip tests and running example renderers. On macOS, Metal works out of the box.

Adding wgsl-rs to your project

Add wgsl-rs to your Cargo.toml. The crate re-exports its proc macros, so you only need the one dependency:

[dependencies]
wgsl-rs = { version = "0.1" }

If you prefer to depend on the macro crate directly, the equivalent is:

[dependencies]
wgsl-rs = { version = "0.1" }
wgsl-rs-macros = { version = "0.1" }

The validation feature is enabled by default and pulls in naga to validate generated WGSL at test time. No extra configuration is required to get it.

Cargo features

See Cargo Features for details on each.

Running the example crate

The repository ships an example crate containing 35+ transpiled modules. List them with:

cargo run -p example -- show

Print the generated WGSL for a specific module with:

cargo run -p example -- source hello_triangle

Verifying your setup

Run the test suite to confirm validation is working:

cargo test -p example

Each #[wgsl] module auto-generates a #[test] fn __validate_wgsl() that feeds the emitted WGSL_SOURCE through naga. If the suite passes, your installation is correct.

Hello, Triangle

This chapter walks through the canonical hello_triangle example end to end. The module is ordinary Rust that the #[wgsl] macro transpiles to WGSL, and it is also a valid Rust module you can compile and test.

The source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod hello_triangle {
    use wgsl_rs::std::*;

    uniform!(group(0), binding(0), FRAME: u32);

    #[vertex]
    pub fn vtx_main(#[builtin(vertex_index)] vertex_index: u32) -> Vec4f {
        const POS: [Vec2f; 3] = [vec2f(0.0, 0.5), vec2f(-0.5, -0.5), vec2f(0.5, -0.5)];
        let position = POS[vertex_index as usize];
        vec4f(position.x, position.y, 0.0, 1.0)
    }

    #[fragment]
    pub fn frag_main() -> Vec4f {
        vec4f(1.0, sin(f32(get!(FRAME)) / 128.0), 0.0, 1.0)
    }
}
}

#[wgsl] pub mod hello_triangle { ... }

The #[wgsl] attribute marks a module for transpilation. The macro consumes the module body, builds an owned IR, and emits a WGSL_SOURCE static containing the generated shader text. The module remains valid Rust: the functions are callable from the CPU side, and the types resolve against wgsl_rs::std.

use wgsl_rs::std::*

The glob import is required. It brings the WGSL type aliases (Vec2f, Vec3f, Vec4f, ...), constructor functions (vec2f, vec3f, vec4f, ...), and the built-in WGSL functions (sin, cos, dot, ...) into scope so the Rust body type-checks and maps one-to-one onto WGSL declarations.

uniform!(...)

#![allow(unused)]
fn main() {
uniform!(group(0), binding(0), FRAME: u32);
}

The uniform! macro declares a uniform binding that is visible in both worlds. It expands to a WGSL var<uniform> declaration in the generated source and to a Rust handle that the runtime can bind and read. Here FRAME is a u32 at group 0, binding 0.

Entry points: #[vertex] and #[fragment]

Functions annotated with #[vertex] and #[fragment] become WGSL entry points tagged with @vertex and @fragment respectively. Other functions in the module without these annotations transpile to plain WGSL functions.

#[builtin(vertex_index)]

#![allow(unused)]
fn main() {
pub fn vtx_main(#[builtin(vertex_index)] vertex_index: u32) -> Vec4f
}

Argument annotations carry WGSL I/O attributes through to the generated signature. #[builtin(vertex_index)] becomes @builtin(vertex_index) in the WGSL output. The same mechanism supports @location(n), @interpolate(...), and other I/O attributes via the corresponding #[...] annotations.

Vector types and constructors

Vec4f and Vec2f are type aliases for vec4<f32> and vec2<f32> exposed by wgsl_rs::std. The lowercase vec2f / vec4f functions are the matching constructors. They mirror WGSL exactly, so Rust expressions like vec4f(1.0, 0.0, 0.0, 1.0) transpile directly to vec4<f32>(1.0, 0.0, 0.0, 1.0).

get!(FRAME)

#![allow(unused)]
fn main() {
sin(f32(get!(FRAME)) / 128.0)
}

get!(...) is the runtime accessor for a declared uniform. On the Rust side it reads the bound value; in the generated WGSL it expands to the bare uniform reference FRAME. This lets the same expression serve both CPU evaluation (e.g. in dispatch-runtime tests) and the shader.

Generated WGSL

@group(0) @binding(0) var<uniform> FRAME: u32;

@vertex
fn vtx_main(@builtin(vertex_index) vertex_index: u32) -> vec4<f32> {
    const POS: array<vec2<f32>, 3> = array<vec2<f32>, 3>(vec2<f32>(0.0, 0.5), vec2<f32>(-0.5, -0.5), vec2<f32>(0.5, -0.5));
    var position: vec2<f32> = POS[vertex_index];
    return vec4<f32>(position.x(), position.y(), 0.0, 1.0);
}

@fragment
fn frag_main() -> vec4<f32> {
    return vec4<f32>(1.0, sin(f32(FRAME) / 128.0), 0.0, 1.0);
}

Note how each Rust construct maps onto WGSL: const to const, let to var, array literals to explicit array<T, N>(...) constructors, and get!(FRAME) to the bare FRAME reference.

Validation

#[wgsl] auto-generates a hidden test:

#![allow(unused)]
fn main() {
#[test]
fn __validate_wgsl() { /* ... */ }
}

For non-template modules this test feeds WGSL_SOURCE through naga and fails on any validation error. Run it with:

cargo test hello_triangle

A passing test means the transpiled WGSL is well-formed according to naga.

Cargo Features

wgsl-rs exposes three cargo features. Only validation is on by default.

validation (default)

Enables naga-based validation of generated WGSL. For every non-template #[wgsl] module, the macro auto-generates a #[test] fn __validate_wgsl() that compiles the emitted WGSL_SOURCE through naga and fails on any validation error.

This is the primary safety net: if your Rust module transpiles but the WGSL is malformed, cargo test catches it.

Disable it with default-features = false:

[dependencies]
wgsl-rs = { version = "0.1", default-features = false }

Use this when you want no naga dependency at all, e.g. in a build that only consumes WGSL_SOURCE text and validates downstream.

dispatch-runtime

Enables the CPU-side fragment dispatch runtime at wgsl_rs::std::runtime. It lets you run a fragment shader on the CPU over a set of inputs and compare the output against a GPU render. This is the mechanism used for roundtrip testing: the same shader code runs on both sides and the results are diffed.

Enable it explicitly:

[dependencies]
wgsl-rs = { version = "0.1", features = ["dispatch-runtime"] }

Use this when writing tests that exercise fragment shaders without spinning up a full wgpu pipeline, or when debugging shader logic on the CPU.

linkage-wgpu

Enables wgpu linkage generation at wgsl_rs::linkage::wgpu. With this feature on, #[wgsl] modules emit the metadata needed to build wgpu render/compute pipelines from the generated shader source, including bind group layouts and entry-point descriptors.

Enable it explicitly:

[dependencies]
wgsl-rs = { version = "0.1", features = ["linkage-wgpu"] }

Use this in application crates that render or compute via wgpu. For pure shader authoring and validation it is not needed.

Combining features

Features compose freely. A typical application crate enables both runtime features:

[dependencies]
wgsl-rs = { version = "0.1", features = ["dispatch-runtime", "linkage-wgpu"] }

A shader-only library crate leaves everything at defaults:

[dependencies]
wgsl-rs = "0.1"

Reference

The #[wgsl] Macro

The #[wgsl] attribute macro is the entry point to wgsl-rs. It is applied to a module and transpiles the Rust inside it to WGSL. The generated source string is stored in a WGSL_SOURCE constant, and the module remains valid Rust that runs on the CPU.

Syntax and Placement

#[wgsl] is placed on a pub mod:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod example {
    use wgsl_rs::std::*;

    pub fn square(x: f32) -> f32 {
        x * x
    }
}
}

The macro produces a pub static WGSL_SOURCE: &str containing the transpiled WGSL. You can read it at runtime:

#![allow(unused)]
fn main() {
println!("{}", example::WGSL_SOURCE);
}

For the module above, the generated WGSL is:

fn square(x: f32) -> f32 {
  return x * x;
}

The code inside the module is ordinary Rust: type-checks on the CPU, runs in cargo test, and transpiles to WGSL for the GPU.

Macro Attributes

Attributes are passed inside the #[wgsl(...)] list.

crate_path

When the macro cannot locate the wgsl_rs crate (for example, from within the crate itself), set the path explicitly:

#![allow(unused)]
fn main() {
#[wgsl(crate_path = "crate")]
pub mod example {
    use wgsl_rs::std::*;
}
}

skip_validation

Disable validation of the generated WGSL:

#![allow(unused)]
fn main() {
#[wgsl(skip_validation)]
pub mod example { use wgsl_rs::std::*; }
}

validate_with_instantiation_types

Validate the module with concrete types for generic/template entry points:

#![allow(unused)]
fn main() {
#[wgsl(validate_with_instantiation_types(f32, u32))]
pub mod example { use wgsl_rs::std::*; }
}

Multiple types may be passed as a comma-separated list.

extensions

Enable WGSL extensions during validation:

#![allow(unused)]
fn main() {
#[wgsl(extensions = [wgsl_rs::WgslExtension::Fxaalp32)] // pseudonymous
pub mod example { use wgsl_rs::std::*; }
}

Extensions are listed inside the extensions = [...] array and must be referenced by their full path in wgsl_rs::WgslExtension.

#[wgsl_ignore]

Items annotated with #[wgsl_ignore] are compiled as Rust but omitted from WGSL generation:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod example {
    use wgsl_rs::std::*;

    pub fn gpu_only(x: f32) -> f32 {
        x * 2.0
    }

    #[wgsl_ignore]
    pub fn cpu_helper(x: f32) -> f32 {
        x.sin() // not transpiled
    }
}
}

#[wgsl_allow(...)]

Suppress transpiler warnings on an expression. Allowed flags:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod example {
    use wgsl_rs::std::*;

    pub fn loopy(n: u32) -> u32 {
        let mut s: u32 = 0;
        #[wgsl_allow(non_literal_loop_bounds)]
        for i in 0..n {
            s += i;
        }
        s
    }
}
}

See Control Flow for related usage of these flags.

Modules and Imports

Every wgsl-rs shader module is a #[wgsl] pub mod name { ... }. The module boundary defines what the macro transpiles and how shaders reference each other.

Module Structure

#![allow(unused)]
fn main() {
#[wgsl]
pub mod lighting {
    use wgsl_rs::std::*;

    pub fn diffuse(n: Vec3f, l: Vec3f) -> f32 {
        max(dot(n, l), 0.0)
    }
}
}

The generated WGSL is emitted into WGSL_SOURCE for that module.

Glob Imports Only

wgsl-rs supports glob imports exclusively. Named imports are not transpiled. The two valid import forms are:

#![allow(unused)]
fn main() {
use wgsl_rs::std::*;        // standard WGSL types and builtins
use super::other_module::*; // another #[wgsl] module in the same parent
}

wgsl_rs::std provides the scalar/vector/matrix types (Vec2f, Vec3f, Vec4f, vec2f, Mat4f, ...), builtins, and texture types.

Importing Other wgsl Modules

A module imported via use super::other_module::* must itself be a #[wgsl] module. This lets you split shaders across files and call functions defined elsewhere:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod math {
    use wgsl_rs::std::*;

    pub fn clampf(x: f32, lo: f32, hi: f32) -> f32 {
        min(max(x, lo), hi)
    }
}

#[wgsl]
pub mod surface {
    use wgsl_rs::std::*;
    use super::math::*;

    pub fn roughness(r: f32) -> f32 {
        clampf(r, 0.0, 1.0)
    }
}
}

Cross-Module Imports and Deduplication

When a module is imported by multiple sibling modules, wgsl-rs deduplicates the generated functions so each WGSL function appears only once in the final output. You do not need to manage inclusion guards.

Doc Comments

Inner doc comments (//!) at the top of a module are preserved in the generated WGSL as comments:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod kernel {
    //! Compute lighting contribution for a single light.
    use wgsl_rs::std::*;
}
}

Outer doc comments (///) on items are not emitted into the WGSL; they stay in the Rust docs.

Functions

Functions are the basic unit of shader logic. They are written as ordinary Rust functions inside a #[wgsl] module.

Syntax

#![allow(unused)]
fn main() {
pub fn name(arg: T, arg2: U) -> R {
    body
}
}

All functions are pub fn. The return type is mandatory unless the function is void (use -> () or omit the arrow for a trailing-statement body).

#![allow(unused)]
fn main() {
#[wgsl]
pub mod funcs {
    use wgsl_rs::std::*;

    pub fn add(a: f32, b: f32) -> f32 {
        a + b
    }

    pub fn noop(x: f32) {
        let _ = x;
    }
}
}

let and let mut

• let x: T = ... transpiles to a WGSL let.
• let mut x: T = ... transpiles to a WGSL var.

#![allow(unused)]
fn main() {
pub fn accumulate(items: Vec4f) -> f32 {
    let mut sum: f32 = 0.0;
    sum += items.x + items.y + items.z + items.w;
    sum
}
}

Early Returns and Implicit Return

Early return is supported:

#![allow(unused)]
fn main() {
pub fn first_nonzero(v: Vec3f) -> f32 {
    if v.x != 0.0 { return v.x; }
    if v.y != 0.0 { return v.y; }
    v.z
}
}

A trailing expression without a semicolon is the implicit return value, matching Rust.

const Inside Functions

Function-scoped const items are supported and transpile to WGSL const:

#![allow(unused)]
fn main() {
pub fn area(r: f32) -> f32 {
    const PI: f32 = 3.14159265;
    PI * r * r
}
}

Function Arguments with IO Annotations

Entry-point and inter-stage functions may carry IO annotations on arguments (#[location(N)], #[builtin(position)], #[interpolate(flat)], etc.). See Entry Points and Inter-Stage IO for the full list of builtins and annotations.

#![allow(unused)]
fn main() {
#[vertex]
pub fn vs_main(
    #[location(0)] pos: Vec3f,
    #[location(1)] uv: Vec2f,
) -> Vec4f {
    vec4f(pos, 1.0)
}
}

Pointer Parameters

WGSL functions can take pointer arguments so the callee can mutate the caller's local or workgroup variable. In wgsl-rs you express this with the ptr! macro — bare &mut T parameters are not supported inside #[wgsl] modules.

#![allow(unused)]
fn main() {
pub fn increment(p: ptr!(function, i32)) {
    *p += 1;
}
}

ptr!(address_space, T) expands to &mut T in Rust (so the code runs on the CPU) and transpiles to ptr<address_space, T> in WGSL. The supported address spaces are:

Dereference with *p, and pass a mutable reference with &mut x:

#![allow(unused)]
fn main() {
pub fn swap(a: ptr!(function, f32), b: ptr!(function, f32)) {
    let tmp = *a;
    *a = *b;
    *b = tmp;
}

pub fn caller() {
    let mut x: f32 = 1.0;
    let mut y: f32 = 2.0;
    swap(&mut x, &mut y);
}
}

Both &x and &mut x transpile to &x in WGSL — mutability is determined by the access mode in the pointer type, not by the reference syntax. The ptr! macro always produces a &mut T on the Rust side so the CPU path can mutate the value.

Structs, Impls, and Enums

Structs

Structs are declared with public fields:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod types {
    use wgsl_rs::std::*;

    pub struct Light {
        pub position: Vec3f,
        pub color: Vec3f,
        pub intensity: f32,
    }
}
}

Transpiles to:

struct Light {
  position: vec3<f32>,
  color: vec3<f32>,
  intensity: f32,
}

#[derive(Wgsl)]

Structs used in storage or uniform buffers should derive Wgsl. This generates the host-side layout and zero-value logic needed for binding setup:

#![allow(unused)]
fn main() {
#[derive(Wgsl)]
pub struct Camera {
    pub view: Mat4f,
    pub proj: Mat4f,
    pub pos: Vec3f,
}
}

Inherent Impls

Methods in impl Type blocks become free WGSL functions named Type_method:

#![allow(unused)]
fn main() {
impl Light {
    pub fn direction(self: Light, target: Vec3f) -> Vec3f {
        normalize(target - self.position)
    }
}
}

In Rust you call Light::direction(light, target). In WGSL this becomes:

fn Light_direction(self_1: Light, target: vec3<f32>) -> vec3<f32> {
  return normalize(target - self_1.position);
}

Associated Constants

const items inside an impl block become associated constants in WGSL:

#![allow(unused)]
fn main() {
impl Light {
    pub const MAX_COUNT: u32 = 64;
}
}

Trait Impls

Trait definitions are Rust-only (the trait is not emitted to WGSL), but the methods in a trait impl are transpiled as if they were inherent methods. This lets you share method syntax between CPU and GPU code:

#![allow(unused)]
fn main() {
pub trait Packed {
    fn pack(self) -> u32;
}

impl Packed for Vec4f {
    pub fn pack(self) -> u32 {
        // bit-packing logic
        0u32
    }
}
}

The pack method transpiles to Vec4f_pack.

Non-pub items in trait impls

Rust forbids pub on any item inside a trait impl (E0449), so wgsl-rs does not require pub on trait-impl methods or associated constants — only inherent impl blocks require pub. This matches Rust's own visibility rules:

#![allow(unused)]
fn main() {
pub trait SlabItem {
    const SLAB_SIZE: usize;
    fn read_at(slab_index: u32) -> Self;
}

impl SlabItem for u32 {
    const SLAB_SIZE: usize = 1;  // no `pub` — correct for a trait impl
    fn read_at(slab_index: u32) -> u32 {
        slab_index
    }
}
}

The associated const is mangled to u32__1SLAB_SIZE in WGSL (the _1 escapes the underscore in SLAB_SIZE per the bijective mangling scheme).

Associated Types in Trait Impls

Trait impls may include associated type definitions (type Array = ...;). The concrete type is resolved at monomorphization time and emitted as a WGSL alias:

#![allow(unused)]
fn main() {
pub trait SlabItem {
    const SLAB_SIZE: usize;
    type Array: Default;
    fn to_array(data: Self) -> Self::Array;
    fn from_array(arr: Self::Array) -> Self;
}

impl SlabItem for u32 {
    const SLAB_SIZE: usize = 1;
    type Array = [u32; 1];
    fn to_array(data: Self) -> Self::Array {
        [data]
    }
    fn from_array(arr: Self::Array) -> Self {
        arr[0]
    }
}
}

This produces:

alias u32_Array = array<u32, 1>;

fn u32__1to_array(data: u32) -> array<u32, 1> {
    return array(data);
}

fn u32__1from_array(arr: array<u32, 1>) -> u32 {
    return arr[0];
}

Self::Array in method signatures is resolved to the concrete type ([u32; 1] → array<u32, 1>) before rendering.

Enums

#[repr(u32)] enums with explicit discriminants transpile to a u32 alias plus const variants:

#![allow(unused)]
fn main() {
#[repr(u32)]
pub enum Mode {
    Add = 0,
    Multiply = 1,
    Screen = 2,
}
}

alias Mode = u32;
const Add: Mode = 0;
const Multiply: Mode = 1;
const Screen: Mode = 2;

Enums are used with match (see Control Flow), which transpiles to a WGSL switch.

Constants

Constants are declared with const. They transpile directly to WGSL const declarations.

Module-Level

#![allow(unused)]
fn main() {
#[wgsl]
pub mod config {
    use wgsl_rs::std::*;

    pub const MAX_LIGHTS: u32 = 64;
    pub const TILE_SIZE: u32 = 16;
    pub const AMBIENT: Vec3f = vec3f(0.1, 0.1, 0.12);
}
}

const MAX_LIGHTS: u32 = 64;
const TILE_SIZE: u32 = 16;
const AMBIENT: vec3<f32> = vec3(0.1, 0.1, 0.12);

Function-Level

const declared inside a function becomes a function-scoped WGSL const:

#![allow(unused)]
fn main() {
pub fn circle_area(r: f32) -> f32 {
    const PI: f32 = 3.14159265;
    PI * r * r
}
}

Uses

Constants are commonly used for:

• Array sizes (const N: u32 = 256; ... var<workgroup> buf: array<f32, N>;)
• Configuration knobs (light counts, tile sizes)
• Fixed colors or directions

Note: WGSL const values are compile-time constants. For values that may vary per dispatch, use a uniform binding (see Binding Macros).

Control Flow

wgsl-rs supports the usual Rust control-flow constructs. Each transpiles to the corresponding WGSL statement.

if / else if / else

#![allow(unused)]
fn main() {
pub fn classify(x: f32) -> u32 {
    if x > 0.0 {
        1u32
    } else if x < 0.0 {
        2u32
    } else {
        0u32
    }
}
}

while

#![allow(unused)]
fn main() {
pub fn gcd(mut a: u32, mut b: u32) -> u32 {
    while b != 0 {
        let t = b;
        b = a % b;
        a = t;
    }
    a
}
}

loop (Infinite Loop)

Rust's bare loop transpiles to a WGSL loop. Use break to exit:

#![allow(unused)]
fn main() {
pub fn first_zero(arr: Vec4f) -> u32 {
    let mut i: u32 = 0;
    loop {
        if i >= 4 { break; }
        if arr[i as usize] == 0.0 { return i; }
        i += 1;
    }
    i
}
}

for

Exclusive range (0..N) and inclusive range (0..=N) are both supported:

#![allow(unused)]
fn main() {
pub fn sum_to(n: u32) -> u32 {
    let mut s: u32 = 0;
    for i in 0..n {
        s += i;
    }
    s
}

pub fn sum_inclusive(n: u32) -> u32 {
    let mut s: u32 = 0;
    for i in 0..=n {
        s += i;
    }
    s
}
}

Loop bounds must be literals or const. For variable bounds, annotate the expression with #[wgsl_allow(non_literal_loop_bounds)]:

#![allow(unused)]
fn main() {
pub fn partial_sum(n: u32) -> u32 {
    let mut s: u32 = 0;
    #[wgsl_allow(non_literal_loop_bounds)]
    for i in 0..n {
        s += i;
    }
    s
}
}

Why does this warning exist? WGSL for loops require explicit, compile-time-known bounds — the spec mandates that loop iteration bounds be literal or const expressions so the shader compiler can reason about termination and resource limits. When wgsl-rs transpiles for i in 0..n, it emits for (var i = 0; i < n; i++). If n is a runtime value, the bound cannot be verified at macro time to be ascending (or even finite), so on stable Rust the macro emits a compile error — proc-macro warnings aren't possible on stable. On nightly it emits a warning instead. The #[wgsl_allow(non_literal_loop_bounds)] attribute suppresses both: you're telling wgsl-rs you've ensured the bound is valid at runtime, taking responsibility the compiler can't.

match (WGSL switch)

match on an integer-typed value transpiles to a WGSL switch:

#![allow(unused)]
fn main() {
#[repr(u32)]
pub enum Op { Add = 0, Sub = 1, Mul = 2 }

pub fn apply(op: Op, a: f32, b: f32) -> f32 {
    match op {
        Op::Add => a + b,
        Op::Sub => a - b,
        Op::Mul => a * b,
        _ => 0.0,
    }
}
}

Or-patterns and non-literal patterns require #[wgsl_allow(non_literal_match_statement_patterns)]:

#![allow(unused)]
fn main() {
pub fn is_zero(op: Op) -> bool {
    #[wgsl_allow(non_literal_match_statement_patterns)]
    match op {
        Op::Add | Op::Sub => false,
        _ => true,
    }
}
}

Why does this warning exist? WGSL switch case selectors must be literal integer constants — the spec doesn't allow arbitrary expressions or enum variants as case labels. When wgsl-rs transpiles a Rust match, it maps each arm to a switch case. Rust enum variants (Op::Add) and const references (LOW) are not WGSL literals — they're names that resolve at the Rust or IR level, not at WGSL compile time. On stable, the macro can't emit a warning (proc-macro diagnostics are nightly-only), so it emits a compile error instead. The #[wgsl_allow(non_literal_match_statement_patterns)] attribute suppresses it: you're asserting the patterns are valid WGSL case selectors once resolved (e.g. #[repr(u32)] enum variants become literal u32 values, consts become literal integers).

break, continue, return

All three are supported inside loops and functions as in Rust. break and continue work in while, loop, and for. return works in any function and supports early returns (see Functions).

Operators

wgsl-rs transpiles Rust operators to their WGSL equivalents. Most have a 1:1 mapping.

Arithmetic

Comparison

Logical

Bitwise

Compound Assignment

+=, -=, *=, /=, %=, &=, |=, ^=, <<=, >>= all transpile directly:

#![allow(unused)]
fn main() {
pub fn bump(x: ptr!(function, f32)) {
    *x += 1.0;
    *x *= 2.0;
}
}

Mutable references in function signatures must use the ptr! macro — bare &mut T parameters are not supported in #[wgsl] modules.

select

select(false_val, true_val, condition) maps to the WGSL select builtin. Argument order matches WGSL:

#![allow(unused)]
fn main() {
pub fn abs_or(x: f32, sign: bool) -> f32 {
    select(-x, x, sign)
}
}

fn abs_or(x: f32, sign: bool) -> f32 {
  return select(-x, x, sign);
}

Unary

as Casts

as casts are transpiled when meaningful in WGSL. The common case is as usize for array indexing — this is stripped in WGSL, which uses the index directly:

#![allow(unused)]
fn main() {
pub fn at(arr: Vec4f, i: u32) -> f32 {
    arr[i as usize]
}
}

fn at(arr: vec4<f32>, i: u32) -> f32 {
  return arr[i];
}

Binding Macros

wgsl-rs provides declarative macros for declaring GPU bindings. Each macro emits both the WGSL binding declaration and a Rust-side static so the same code works on CPU and GPU. To auto-generate wgpu bind group layouts and buffer descriptors from these bindings, see wgpu Linkage.

Overview

Declaration and Access

Binding macros are used at module scope inside a #[wgsl] module. They declare the WGSL binding and the Rust-side static simultaneously:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod shader {
    use wgsl_rs::std::*;

    uniform!(group(0), binding(0), CAMERA: Camera);

    pub fn view() -> Mat4f {
        get!(CAMERA).view
    }
}
}

get! and get_mut!

• get!(VAR) reads a uniform!, storage!, or workgroup! binding. It returns a guard that derefs to the value.
• get!(VAR, T) reads with an explicit type, used inside generic/template entry points.
• get_mut!(VAR) returns a mutable guard for storage! and workgroup! bindings.

#![allow(unused)]
fn main() {
pub fn add_delta() {
    let mut s = get_mut!(COUNTER);
    s.value += 1;
}
}

Slab Helpers

For packed slab buffers, use the slab_copy! macro. It is bidirectional — pass the slab as the source to read from a storage buffer into a local array, or pass the slab as the destination to write from a local array into a storage buffer:

#![allow(unused)]
fn main() {
slab_copy!(src, src_offset, dest, dest_offset, size)
}

Copies size elements from src[src_offset..] into dest[dest_offset..]. On the GPU this emits a WGSL for loop; on the CPU it is a simple element-by-element copy.

#![allow(unused)]
fn main() {
let mut raw = [0u32; 4];
slab_copy!(get!(SLAB), index, raw, 0, 4);
slab_copy!(raw, 0, get_mut!(SLAB), index, 4);
}

uniform!

Declares a uniform buffer binding.

Syntax

#![allow(unused)]
fn main() {
uniform!(group(N), binding(M), NAME: Type);
}

What It Generates

WGSL:

@group(N) @binding(M) var<uniform> NAME: Type;

Rust:

#![allow(unused)]
fn main() {
pub static NAME: Uniform<Type>;
}

Access

Read with get!(NAME). The returned guard dereferences to &Type:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod shader {
    use wgsl_rs::std::*;

    #[derive(Wgsl)]
    pub struct Camera {
        pub view: Mat4f,
        pub proj: Mat4f,
        pub pos: Vec3f,
    }

    uniform!(group(0), binding(0), CAMERA: Camera);

    pub fn world_to_clip(p: Vec3f) -> Vec4f {
        let c = get!(CAMERA);
        c.proj * c.view * vec4f(p, 1.0)
    }
}
}

get!(NAME) returns a guard, so field access uses . directly. For generic entry points, supply the type explicitly: get!(CAMERA, Camera).

Notes

• Uniforms are read-only on the GPU.
• Type should be #[derive(Wgsl)] so the host side can lay out and upload the buffer.
• One uniform binding per (group, binding) pair.

storage!

Declares a storage buffer binding, either read-only or read-write.

Syntax

#![allow(unused)]
fn main() {
// read-only
storage!(group(N), binding(M), NAME: Type);

// read-write
storage!(group(N), binding(M), read_write, NAME: Type);
}

What It Generates

Read-only:

@group(N) @binding(M) var<storage, read> NAME: Type;

Read-write:

@group(N) @binding(M) var<storage, read_write> NAME: Type;

Rust:

#![allow(unused)]
fn main() {
pub static NAME: Storage<Type>;
}

Access

• get!(NAME) reads the buffer.
• get_mut!(NAME) writes to the buffer (only valid for read_write).

Example: Compute Shader

#![allow(unused)]
fn main() {
#[wgsl]
pub mod prefix {
    use wgsl_rs::std::*;

    #[derive(Wgsl)]
    pub struct Data {
        pub value: f32,
    }

    storage!(group(0), binding(0), read_write, INPUT: Data);
    storage!(group(0), binding(1), read_write, OUTPUT: Data);

    #[compute]
    #[workgroup_size(64)]
    pub fn cs_main() {
        let mut src = get_mut!(INPUT);
        let mut dst = get_mut!(OUTPUT);
        dst.value = src.value * 2.0;
    }
}
}

Notes

• Use arrays in Type (e.g. array<f32, N>) for large buffers.
• read_write storage requires the binding to be created with the read_write access flag on the host side.
• The slab_copy! helper operates on storage! bindings (see Binding Macros).

workgroup!

Declares a workgroup-scoped variable shared across invocations in a compute workgroup.

Syntax

#![allow(unused)]
fn main() {
workgroup!(NAME: Type);
}

What It Generates

WGSL:

var<workgroup> NAME: Type;

Rust:

#![allow(unused)]
fn main() {
pub static NAME: Workgroup<Type>;
}

The Rust-side static is backed by a LazyLock<RwLock<T>>, so CPU code can read and write the same value across threads for testing.

Access

• get!(NAME) reads.
• get_mut!(NAME) writes.

Example: Shared Sum

#![allow(unused)]
fn main() {
#[wgsl]
pub mod reduce {
    use wgsl_rs::std::*;

    pub const WG_SIZE: u32 = 64;

    workgroup!(SHARED: array<f32, 64>);

    #[compute]
    #[workgroup_size(64)]
    pub fn cs_main(
        #[builtin(workgroup_id)] wg: Vec3u,
        #[builtin(local_invocation_index)] li: u32,
    ) {
        let mut s = get_mut!(SHARED);
        s[li as usize] = f32(li);

        // barrier omitted for brevity; use workgroupBarrier() via a builtin

        let mut sum: f32 = 0.0;
        for i in 0..WG_SIZE {
            sum += get!(SHARED)[i as usize];
        }

        if li == 0 {
            let mut out = get_mut!(RESULT);
            out.value = sum;
        }
    }
}
}

Notes

• Workgroup variables are visible to all invocations sharing the same workgroup_id.
• Use a workgroup barrier before reading values written by other invocations.
• Lifetime is a single workgroup dispatch; values are not preserved between dispatches.

texture! and sampler!

Declare texture and sampler bindings.

texture!

#![allow(unused)]
fn main() {
texture!(group(N), binding(M), NAME: TextureKind<SampleType>);
}

Generates:

@group(N) @binding(M) var NAME: TextureKind<SampleType>;

Texture Kinds

The sample type for color textures is typically <f32>. Depth textures need no sample type parameter.

Storage Textures

Storage textures (texture_storage_* in WGSL) are declared with TextureStorage types. They take two type parameters: a texel format marker and an access mode marker:

#![allow(unused)]
fn main() {
texture!(group(0), binding(0), OUTPUT: TextureStorage2D<Rgba8unorm, Write>);
}

Texel format markers (unit structs implementing WgslTexelFormat):

Access mode markers:

The enable texture_formats_tier1; directive is auto-hoisted to the start of the assembled WGSL translation unit when storage textures are present — you don't need to add it manually.

sampler!

#![allow(unused)]
fn main() {
sampler!(group(N), binding(M), NAME: Sampler);
sampler!(group(N), binding(M), NAME: SamplerComparison);
}

Generates:

@group(N) @binding(M) var NAME: sampler;
@group(N) @binding(M) var NAME: sampler_comparison;

Two-Level Binding

Both macros produce a hidden __NAME static plus a visible pub const NAME: &'static ... so the binding can be passed by value. You reference NAME directly in functions:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod texturing {
    use wgsl_rs::std::*;

    texture!(group(0), binding(0), ALBEDO: Texture2D<f32>);
    sampler!(group(0), binding(1), LIN: Sampler);

    #[fragment]
    pub fn fs_main(
        #[location(0)] uv: Vec2f,
    ) -> Vec4f {
        texture_sample(ALBEDO, LIN, uv)
    }
}
}

Passing to Functions

Texture and sampler bindings are passed by value (no &) — the visible NAME is already a reference:

#![allow(unused)]
fn main() {
pub fn sample_albedo(uv: Vec2f, tex: Texture2D<f32>, smp: Sampler) -> Vec4f {
    texture_sample(tex, smp, uv)
}
}

Notes

• SamplerComparison is used with textureSampleCompare and textureSampleCompareLevel for shadow maps.
• Pair each texture with its sampler; binding numbers must not collide within a group.

ptr!

Declares a WGSL pointer parameter. Used in function signatures where the function needs to read or write a variable in a specific address space.

Syntax

#![allow(unused)]
fn main() {
fn name(p: ptr!(address_space, Type)) { ... }
}

address_space is one of function, private, or workgroup.

What It Generates

Rust:

#![allow(unused)]
fn main() {
fn name(p: &mut Type) { ... }
}

WGSL:

fn name(p: ptr<address_space, Type>) { ... }

Dereference

Read or write through the pointer with *p:

#![allow(unused)]
fn main() {
pub fn increment(p: ptr!(function, f32)) {
    *p += 1.0;
}
}

fn increment(p: ptr<function, f32>) {
  *p += 1.0;
}

Example: Swap

#![allow(unused)]
fn main() {
#[wgsl]
pub mod utils {
    use wgsl_rs::std::*;

    pub fn swap(a: ptr!(function, f32), b: ptr!(function, f32)) {
        let t = *a;
        *a = *b;
        *b = t;
    }

    pub fn sort_pair(mut x: f32, mut y: f32) -> Vec2f {
        if x > y {
            swap(&mut x, &mut y);
        }
        vec2f(x, y)
    }
}
}

Notes

• Use &mut at the call site in Rust; the macro translates this to a WGSL pointer of the declared address space.
• function is the most common address space for local variables. Use workgroup for pointers to workgroup! variables and private for module-private variables.

discard!

Discards the current fragment. In WGSL this emits discard;; on CPU it sets a thread-local flag.

Syntax

#![allow(unused)]
fn main() {
discard!();
}

Behavior

• In WGSL, transpiles to discard; and stops further output for the fragment.
• On CPU, sets a thread-local flag that dispatch_fragments checks; execution continues after the call, matching WGSL's helper-invocation semantics.

Reachable From

discard!() may be called from any function reachable from a #[fragment] entry point, including helper functions.

Example

#![allow(unused)]
fn main() {
#[wgsl]
pub mod alpha {
    use wgsl_rs::std::*;

    texture!(group(0), binding(0), MASK: Texture2D<f32>);
    sampler!(group(0), binding(1), LIN: Sampler);

    pub fn threshold(uv: Vec2f, min: f32) {
        let a = texture_sample(MASK, LIN, uv).x();
        if a < min {
            discard!();
        }
    }

    #[fragment]
    pub fn fs_main(#[location(0)] uv: Vec2f) -> Vec4f {
        threshold(uv, 0.5);
        vec4f(1.0, 0.0, 0.0, 1.0)
    }
}
}

Notes

• Execution after discard!() continues on the CPU side, so code following the call still runs. Guard any side effects accordingly.
• discard!() is only valid inside fragment shaders. Calling it from a compute or vertex entry point is a validation error.

Scalars & Literals

wgsl-rs shares the four WGSL scalar types with Rust directly. The same names mean the same thing in both worlds.

Because shader code must type-check as ordinary Rust, scalar types are not aliases: they are the literal Rust primitive types. The transpiler maps them onto the matching WGSL keyword.

Literal Suffixes

Rust literal suffixes carry type information into the generated WGSL. Always suffix literals when the surrounding context does not pin the type (e.g. constants, function arguments to generic constructors).

#![allow(unused)]
fn main() {
const WIDTH: u32 = 1024u32;
const EPS: f32 = 1e-5f32;
let count: i32 = 0i32;
}

Unsuffixed integer literals are accepted by Rust and inferred from context, but explicit suffixes make the transpiler's job unambiguous and the generated WGSL easier to read.

as Casts

Rust's as cast operator transpiles to a WGSL conversion expression of the same form.

#![allow(unused)]
fn main() {
let i: i32 = 7;
let u: u32 = i as u32;     // -> u32(i)
let f: f32 = u as f32;     // -> f32(u)
let n: f32 = i as f32;     // -> f32(i)
}

Cross-kind conversions (i32 <-> u32 <-> f32) all generate the corresponding WGSL T(x) conversion. Booleans cannot be cast with as; use select or a manual comparison instead.

Vectors & Swizzles

wgsl-rs exposes all 12 WGSL vector types as Rust aliases and constructor functions in wgsl_rs::std.

Type Aliases

The generic form (Vec2<f32>, Vec3<f32>, Vec4<f32>) is accepted wherever a concrete alias is, including struct fields and function signatures.

Constructors

Constructors are lowercase functions named after the WGSL vecN* builtin:

#![allow(unused)]
fn main() {
let a: Vec2f = vec2f(0.0, 1.0);
let b: Vec3f = vec3f(0.0, 1.0, 2.0);
let c: Vec4f = vec4f(0.0, 1.0, 2.0, 3.0);
let u: Vec3u = vec3u(0u32, 1u32, 2u32);
let i: Vec4i = vec4i(0i32, 1i32, 2i32, 3i32);
}

Single-argument splat constructors are supported, mirroring WGSL:

#![allow(unused)]
fn main() {
let all_ones: Vec4f = vec4f(1.0);   // -> vec4<f32>(1.0)
}

These map directly to WGSL vecN<T>(...) constructor calls.

Swizzles are Method Calls

WGSL lets you access vector components as fields (v.xyz). In wgsl-rs, swizzles are method calls, not field accesses:

#![allow(unused)]
fn main() {
let v: Vec4f = vec4f(1.0, 2.0, 3.0, 4.0);
let xyz: Vec3f = v.xyz();   // -> v.xyz
let x: f32    = v.x();      // -> v.x
let xy: Vec2f = v.xy();     // -> v.xy
let rgb: Vec3f = v.rgb();   // -> v.rgb
}

Why Method Calls?

wgsl-rs maintains a single source file that must compile as unmodified Rust on the CPU and transpile to WGSL for the GPU (the "two worlds" constraint). Rust does not permit .xyz field access on a generic vector alias the way WGSL does, and the canonical Rust vector libraries (glam, etc.) expose swizzles via the Vec4Swizzle-style trait — i.e. as methods. Using method-call syntax keeps the shader source a legal Rust program that mirrors an idiomatic CPU implementation.

Field access (v.xyz) is therefore a parse error in your shader; always call the swizzle as a method.

Supported Swizzle Names

Any combination of the xyzw, rgba, or stpq component sets, length 1-4, single-set only:

The result type matches the source vector's scalar kind: a Vec3u swizzle returns u32 or Vec2u/Vec3u/Vec4u.

Matrices

wgsl-rs covers all nine WGSL matrix shapes. Square matrices have short aliases; non-square matrices use the MatCxRf naming convention where C is column count and R is row count.

Square Aliases

Non-Square Aliases

All matrices are f32 only, matching the WGSL specification. The generic form (Mat4x4<f32>, Mat2x3<f32>, etc.) is accepted anywhere an alias is.

Constructors

Matrix constructors take column vectors. The column vector width must match the row count of the matrix.

#![allow(unused)]
fn main() {
const IDENTITY: Mat4f = mat4x4f(
    vec4f(1.0, 0.0, 0.0, 0.0),
    vec4f(0.0, 1.0, 0.0, 0.0),
    vec4f(0.0, 0.0, 1.0, 0.0),
    vec4f(0.0, 0.0, 0.0, 1.0),
);

const ROTATION_2D: Mat3f = mat3x3f(
    vec3f(0.866, 0.5,  0.0),
    vec3f(-0.5,  0.866, 0.0),
    vec3f(0.0,   0.0,  1.0),
);

const SCALE_2D: Mat2f = mat2x2f(vec2f(2.0, 0.0), vec2f(0.0, 2.0));

const M_3X2: Mat3x2f = mat3x2f(vec2f(1.0, 0.0), vec2f(0.0, 1.0), vec2f(0.0, 0.0));
}

The constructor name mirrors WGSL: matCxRf(col0, col1, ...). Each column argument must be a VecRf (or VecR<i32>/VecR<u32> where the matrix is integer — currently only f32).

Multiplication

Matrix-times-vector and matrix-times-matrix use Rust's * operator and emit WGSL *:

#![allow(unused)]
fn main() {
let m: Mat4f = IDENTITY;
let v: Vec4f = vec4f(1.0, 2.0, 3.0, 1.0);
let transformed: Vec4f = m * v;          // -> m * v

let a: Mat4f = IDENTITY;
let b: Mat4f = IDENTITY;
let composed: Mat4f = a * b;             // -> a * b
}

The result type is inferred by Rust and verified by the transpiler against the WGSL rules: matCxR * vecC yields vecR; matCxR * matRxC2 yields matCx2.

Arrays & RuntimeArray<T>

Fixed-Size Arrays

A Rust fixed-size array [T; N] transpiles to a WGSL array<T, N> with the size baked in.

#![allow(unused)]
fn main() {
const POS: [Vec2f; 3] = [
    vec2f(0.0, 0.5),
    vec2f(-0.5, -0.5),
    vec2f(0.5, -0.5),
];
}

const POS: array<vec2<f32>, 3> = array<vec2<f32>, 3>(
    vec2<f32>(0.0, 0.5),
    vec2<f32>(-0.5, -0.5),
    vec2<f32>(0.5, -0.5),
);

Indexing

WGSL array indexing requires the index to be a u32/i32. Because Rust idioms (and CPU-side data structures) commonly carry usize, wgsl-rs accepts an i as usize cast on the index and emits the inner expression directly:

#![allow(unused)]
fn main() {
let p = POS[vertex_index as usize];   // -> POS[vertex_index]
}

Use arr[i as usize] whenever the index source is a u32 builtin (e.g. vertex_index, global_invocation_id.x()).

Zero-Value Arrays

A Rust zero-value array [0u32; 4] is recognized and turned into an explicit WGSL array constructor of the same length, populated with the zero value:

#![allow(unused)]
fn main() {
let zeros: [u32; 4] = [0u32; 4];
}

var zeros: array<u32, 4> = array<u32, 4>();

The element expression must be a literal 0 of the element type. Non-zero repeated-element arrays are not given this special treatment.

Runtime-Size Arrays

RuntimeArray<T> maps to the unsized WGSL array<T> (no count parameter). Runtime arrays are restricted by the WGSL specification:

• They may only appear in storage buffers.
• They must be the last field of a struct.

On the CPU side RuntimeArray<T> is backed by a Vec<T> so the same struct can be populated and read by host code.

#![allow(unused)]
fn main() {
#[derive(Wgsl)]
pub struct ParticleSystem {
    pub count: u32,
    pub particles: RuntimeArray<Particle>,
}

storage!(group(0), binding(0), read_write, PARTICLES: ParticleSystem);
}

struct ParticleSystem {
  count: u32,
  particles: array<Particle>,
};

@group(0) @binding(0) var<storage, read_write> PARTICLES: ParticleSystem;

array_length

Query the length of a runtime array with array_length(&arr). Pass the array field by reference:

#![allow(unused)]
fn main() {
let n = array_length(&get!(PARTICLES).particles);
}

let n = arrayLength(&PARTICLES.particles);

For fixed-size arrays the length is known statically; use arr.len() or the N from the type where convenient, but prefer array_length only for runtime arrays.

Atomics

Atomic<T> is the wgsl-rs wrapper around the WGSL atomic<T> type. It is the only mechanism for shared mutable state across invocations within a workgroup or across storage-buffer accesses.

Allowed Element Types

WGSL atomics are restricted to i32 and u32. wgsl-rs enforces the same restriction:

On the CPU side, Atomic<T> is backed by the matching std::sync::atomic type so the same shader module can run as ordinary Rust in tests and produce consistent results.

Where Atomics May Appear

Atomics are only valid inside address spaces where multiple invocations can observe each other's writes:

• workgroup variables, declared with the workgroup! macro.
• storage buffers, declared with the storage! macro (typically read_write).

Function-local atomics are not useful (a single invocation has no contention) and are rejected.

use wgsl_rs::std::*;

workgroup!(COUNTER: Atomic<u32>);
workgroup!(FLAGS:  Atomic<i32>);

#[compute]
#[workgroup_size(64)]
pub fn main(#[builtin(local_invocation_index)] local_idx: u32) {
    let _idx = local_idx;
}

var<workgroup> COUNTER: atomic<u32>;
var<workgroup> FLAGS: atomic<i32>;

@compute @workgroup_size(64)
fn main(@builtin(local_invocation_index) local_idx: u32) {
  let _idx = local_idx;
}

Atomic Operations

The standard WGSL atomic builtins (atomicLoad, atomicStore, atomicAdd, atomicSub, atomicMin, atomicMax, atomicAnd, atomicOr, atomicXor, atomicExchange, atomicCompareExchangeWeak) are exposed by wgsl_rs::std as free functions and transpile to the matching WGSL call. Use them through the get! / get_mut! accessors that yield a reference to the underlying Atomic<T>:

#![allow(unused)]
fn main() {
let current: u32 = atomicLoad(&get!(COUNTER));
atomicAdd(&get_mut!(COUNTER), 1u32);
}

let current: u32 = atomicLoad(&COUNTER);
atomicAdd(&COUNTER, 1u32);

Use the get! accessor for read-only atomic loads and get_mut! for mutating atomic operations, mirroring the storage-buffer conventions in Binding Macros.

Generic Functions

Generic free functions let you write a shader helper once and specialize it for several concrete types without preprocessor macros. The macro monomorphizes each call-site instantiation into its own concrete WGSL function with a mangled name.

Defining a Generic Function

A generic function is ordinary Rust with trait bounds. The bounds are required only so Rust can type-check the body; they are stripped from the generated WGSL.

#![allow(unused)]
fn main() {
pub fn double<T: Copy + std::ops::Add<Output = T>>(x: T) -> T {
    x + x
}
}

Turbofish is Required

Because the transpiler must know which concrete type to monomorphize, every call to a generic function must use turbofish:

#![allow(unused)]
fn main() {
pub fn apply_f32(value: f32) -> f32 {
    double::<f32>(value)
}

pub fn apply_i32(value: i32) -> i32 {
    double::<i32>(value)
}
}

Calls without ::<T> are rejected by the macro even when Rust could infer the type.

Monomorphization & Name Mangling

Each unique (function, type-args) pair produces one concrete WGSL function. The name is mangled as <fn>_<type> (with extra type parameters joined):

fn double_f32(x: f32) -> f32 {
  return x + x;
}

fn double_i32(x: i32) -> i32 {
  return x + x;
}

Duplicate instantiations across the module (or transitively through other generic functions) are deduplicated — only one copy of each monomorphized function is emitted.

Transitive Generic Calls

A generic function may call another generic function. The inner turbofish drives its own monomorphization:

#![allow(unused)]
fn main() {
pub fn select_val<T: Copy>(a: T, b: T, cond: bool) -> T {
    if cond { a } else { b }
}

pub fn double_or_keep<T: Copy + std::ops::Add<Output = T>>(x: T, use_double: bool) -> T {
    select_val::<T>(double::<T>(x), x, use_double)
}
}

Calling double_or_keep::<f32>(...) pulls in both double_f32 and select_val_f32 automatically.

Multiple Type Parameters

Functions may take more than one type parameter. Each is monomorphized over the full tuple of concrete type arguments:

#![allow(unused)]
fn main() {
pub fn mix<A: Copy, B: Copy>(a: A, b: B) -> A {
    a
}
}

A call mix::<f32, u32>(x, y) produces mix_f32_u32.

Const Generic Parameters

Functions can also take const generic parameters of type u32 or usize — the only const param types that make sense in WGSL (they're used as array lengths). The const param is substituted with a concrete integer literal at monomorphization time:

#![allow(unused)]
fn main() {
pub fn sum_n<const N: usize>(arr: [u32; N]) -> u32 {
    let mut total: u32 = 0;
    for i in 0..N {
        total += arr[i];
    }
    total
}

pub fn run() -> u32 {
    sum_n::<4>([1, 2, 3, 4])
}
}

The call sum_n::<4> produces a WGSL function sum_n_4 with N replaced by the literal 4 throughout (including the array type and loop bound). Const and type params can coexist on the same function; they're monomorphized over the full tuple of arguments.

Const param references are always bare identifiers (e.g. N), per stable Rust's const generics syntax. They're substituted to Expr::Lit at monomorphization time — no new IR variant is needed.

Trait Bounds are Rust-Only

Copy, Clone, Add, PartialEq, custom traits — all bounds exist solely for the Rust type checker. They generate no WGSL output. This is the "two worlds" split in action: Rust validates the generic body once on the CPU; WGSL receives fully concrete, monomorphized code for the GPU with no notion of traits or generics.

A Worked Example

The generic_functions example module demonstrates the full pipeline:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod generic_functions {
    pub fn double<T: Copy + std::ops::Add<Output = T>>(x: T) -> T {
        x + x
    }

    pub fn select_val<T: Copy>(a: T, b: T, cond: bool) -> T {
        if cond { a } else { b }
    }

    pub fn double_or_keep<T: Copy + std::ops::Add<Output = T>>(x: T, use_double: bool) -> T {
        select_val::<T>(double::<T>(x), x, use_double)
    }

    pub fn apply_f32(value: f32) -> f32 {
        double_or_keep::<f32>(value, true)
    }

    pub fn apply_i32(value: i32) -> i32 {
        double_or_keep::<i32>(value, false)
    }
}
}

apply_f32 and apply_i32 each pull in their own copies of double_or_keep_*, double_*, and select_val_*, with the duplicate instantiations of double_or_keep collapsed as needed.

Generic Structs

Generic structs follow the same monomorphization model as generic functions: you write one Rust definition with type parameters, and the macro emits a separate concrete WGSL struct for each (struct, type-args) pair used in the module.

Defining a Generic Struct

#![allow(unused)]
fn main() {
pub struct Pair<T: Copy> {
    pub a: T,
    pub b: T,
}
}

The Copy (or other) bound is Rust-only and stripped from the WGSL.

Usage & Mangling

At every use site, supply the concrete type either as a turbofish on the path or as a type annotation. Each unique instantiation becomes a mangled WGSL struct:

#![allow(unused)]
fn main() {
pub fn use_pair_f32() -> f32 {
    let p = Pair { a: 1.0, b: 2.0 };
    Pair::<f32>::sum(p)
}

pub fn use_pair_i32() -> i32 {
    let p: Pair<i32> = Pair::<i32> { a: 10, b: 20 };
    Pair::<i32>::first(p)
}
}

struct Pair_f32 {
  a: f32,
  b: f32,
}

struct Pair_i32 {
  a: i32,
  b: i32,
}

Generic Impl Blocks

impl<T> Pair<T> blocks are monomorphized alongside the struct. Each method becomes a mangled WGSL function named <Struct>_<type>_<method>:

#![allow(unused)]
fn main() {
impl<T: Copy + std::ops::Add<Output = T>> Pair<T> {
    pub fn first(p: Pair<T>) -> T {
        p.a
    }

    pub fn sum(p: Pair<T>) -> T {
        p.a + p.b
    }
}
}

For Pair::<f32> this yields:

fn Pair_f32_first(p: Pair_f32) -> f32 {
  return p.a;
}

fn Pair_f32_sum(p: Pair_f32) -> f32 {
  return p.a + p.b;
}

Struct Construction

Construct a generic struct by writing the literal form Pair::<f32> { a, b } or by relying on a type annotation. The macro emits a positional WGSL constructor call with the mangled name:

#![allow(unused)]
fn main() {
let p = Pair::<f32> { a: 1.0, b: 2.0 };
}

let p = Pair_f32(1.0, 2.0);

Fields are emitted in declaration order.

Known Limitation: Struct Constructor Mangling

There is a known bug in which the bare struct-constructor form Pair { a, b } (without a turbofish or annotation that the macro can resolve) is not mangled correctly, producing invalid WGSL. Until this is fixed, the recommended workarounds are:

• Always use the turbofish form Pair::<T> { ... } at construction sites, or
• Annotate the binding: let p: Pair<T> = Pair { ... }.
• For modules that exercise the bug and cannot be restructured, suppress auto-validation with #[wgsl(skip_validation)] (see Disabling Validation) so the failing constructor does not break cargo test.

The generic_structs example currently uses #[wgsl(skip_validation)] for this reason:

#![allow(unused)]
fn main() {
#[wgsl(skip_validation)]
pub mod generic_structs {
    pub struct Pair<T: Copy> {
        pub a: T,
        pub b: T,
    }

    impl<T: Copy + std::ops::Add<Output = T>> Pair<T> {
        pub fn first(p: Pair<T>) -> T { p.a }
        pub fn sum(p: Pair<T>) -> T { p.a + p.b }
    }
}
}

Multiple Type Parameters

A struct may take several type parameters; the mangled name joins all concrete types:

#![allow(unused)]
fn main() {
pub struct Cell<K: Copy, V: Copy> {
    pub key: K,
    pub value: V,
}
}

Cell::<u32, f32> produces Cell_u32_f32.

Const Generic Parameters

Structs can also take const N: usize or const N: u32 parameters, which are substituted with concrete integer literals at monomorphization time. This is the natural way to express arrays whose length varies per instantiation:

#![allow(unused)]
fn main() {
pub struct Grid<const N: usize> {
    pub cells: [u32; N],
}

impl<const N: usize> Grid<N> {
    pub fn first(cells: [u32; N]) -> u32 {
        cells[0]
    }
}

pub fn run() -> u32 {
    let g = Grid::<4> { cells: [0, 0, 0, 0] };
    g.cells[0]
}
}

Grid::<4> produces a WGSL struct Grid_4 with cells: array<u32, 4>, and Grid::<4>::first becomes Grid_4_first.

Generic Trait Impls on Array Types

Generic impl blocks on array self types (impl<T: Trait> Trait for [T; N]) are supported. The monomorphizer substitutes the concrete element type and mangles the methods:

#![allow(unused)]
fn main() {
pub trait Zeroable {
    fn zero() -> Self;
}

impl<T: Zeroable> Zeroable for [T; 4] {
    fn zero() -> [T; 4] {
        [T::zero(), T::zero(), T::zero(), T::zero()]
    }
}

pub fn caller_u32_array() -> [u32; 4] {
    Zeroable::zero::<[u32; 4]>()
}
}

The call with [u32; 4] produces a WGSL function _2array_u32_4_zero (the _2 prefix is the bijective mangled encoding of array_u32_4). Similarly, [f32; 4] produces _2array_f32_4_zero.

Limitation: Direct <[u32; 4]>::method() call syntax (QSelf paths) is not yet supported — only T::method() resolved via monomorphization. Tracked in GitHub issue #131.

PhantomData<T> Marker Fields

A generic struct may carry PhantomData<T> fields as type-parameter markers (e.g. for slab-id tags or type-level metadata). PhantomData is re-exported from wgsl_rs::std so the glob import brings it into scope. The proc-macro recognizes PhantomData<_> fields specially: they are retained in the IR (so extensions can observe which type parameter each phantom slot binds) but omitted from the rendered WGSL:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod phantom_example {
    use wgsl_rs::std::*;

    pub struct Id<T> {
        pub index: u32,
        pub phantom: PhantomData<T>,
    }

    pub struct Tagged<T, A> {
        pub x: f32,
        pub t: PhantomData<T>,
        pub a: PhantomData<A>,
    }

    pub fn make_id() -> Id<f32> {
        Id { index: 0u32, phantom: PhantomData }
    }

    pub fn make_tagged() -> Tagged<f32, u32> {
        Tagged { x: 1.0, t: PhantomData, a: PhantomData }
    }
}
}

The rendered WGSL drops phantom fields entirely — Id_f32 has only index: u32, and Tagged_f32_u32 has only x: f32:

struct Id_f32 {
    index: u32
}

struct Tagged_f32_u32 {
    x: f32
}

Construction expressions use the bare PhantomData value (no turbofish). The macro strips PhantomData from the positional constructor call so the rendered arity matches the non-phantom field count: Id { index: 0u32, phantom: PhantomData } becomes Id_f32(0u).

Why retain phantom fields in the IR? Extensions consuming the IR via WgslExtension::modify_ir must be able to see the full type-parameter binding structure of a generic struct. If phantom fields were skipped at parse time, an extension inspecting struct Tagged<T, A> would see type_params: ["T", "A"] but only { x: f32 }, with no way to recover which phantom slot bound which parameter. Keeping Type::Phantom { elem } in the IR preserves the T↔field provenance.

Template Modules & Instantiation

Generic functions and structs are monomorphized to concrete WGSL at macro time — every type parameter is resolved before the module is emitted. Template modules are the complementary mechanism for deferring type parameters until runtime: the macro emits template WGSL carrying TypeParam placeholders, and you instantiate it with concrete types at runtime to obtain a valid shader.

When to Use a Template

Use a template module when the type of an entry point, a linkage binding, or a struct field cannot be pinned down at macro time — for example, a renderer that wants to swap f32 precision for f16, or a uniform whose host-side type is chosen per pipeline.

A module becomes a template when any of the following appear:

• An entry-point function with type parameters.
• An entry-point function with const parameters (const N: usize or const N: u32).
• A linkage macro (uniform!, storage!, ...) whose declared type uses impl Trait.
• A get!(VAR, T) accessor that introduces a fresh type variable bound to a linkage variable.

Defining a Template Module

The hello_triangle_generic example shows the pattern:

#![allow(unused)]
fn main() {
#[wgsl(validate_with_instantiation_types(f32, f32))]
pub mod hello_triangle_generic {
    use wgsl_rs::std::*;

    // The `impl Convert<f32>` syntax declares that FRAME's concrete type
    // is chosen at instantiation time.
    uniform!(group(0), binding(0), FRAME: impl Convert<f32>);

    #[vertex]
    pub fn vtx_main(#[builtin(vertex_index)] vertex_index: u32) -> Vec4f {
        const POS: [Vec2f; 3] = [
            vec2f(0.0, 0.5),
            vec2f(-0.5, -0.5),
            vec2f(0.5, -0.5),
        ];
        let position = POS[vertex_index as usize];
        vec4f(position.x, position.y, 0.0, 1.0)
    }

    #[fragment]
    pub fn frag_main<T: Convert<f32> + Wgsl + Clone>() -> Vec4f {
        let frame_t = get!(FRAME, T);
        vec4f(1.0, sin(f32(frame_t) / 128.0), 0.0, 1.0)
    }
}
}

Two type parameters are in play here: the impl Convert<f32> on FRAME and the T on frag_main. They are linked through get!(FRAME, T).

Generic Linkage: impl Trait

A linkage macro may declare its type as impl SomeTrait:

#![allow(unused)]
fn main() {
uniform!(group(0), binding(0), FRAME: impl Convert<f32>);
}

This says: FRAME has some concrete type, chosen at instantiation time, that implements Convert<f32>. The trait bounds are replayed onto the typestate builder's set_frame method, so a host caller can only bind FRAME to a type satisfying those bounds.

get!(VAR, T) Constraints

Inside a generic entry point, get!(VAR, T) reads the linkage variable VAR and introduces a fresh type variable T connected to VAR's declared type. This generates a constraint of the form

VAR: linkage::Type<Is = T>

on the module's instantiate function. The transpiler threads that constraint so that the same concrete type is used for both the binding and the body of the entry point.

Instantiation at Runtime

A template module exposes an instantiate function you call with concrete turbofish type arguments. It returns a concrete ir::Module whose TypeParam placeholders have been substituted:

#![allow(unused)]
fn main() {
use example::hello_triangle_generic as tmpl;

let module: ir::Module = tmpl::instantiate::<f32, f32>();
let source: String = module.to_wgsl();
}

The number and order of type arguments match the template's declared type parameters. The resulting ir::Module is a fully concrete, validatable WGSL module. To build wgpu pipelines from an instantiated template, see Template Linkage.

Validating Templates

Template modules are not auto-validated by #[wgsl], because the raw TypeParam placeholders are not valid WGSL. There are two ways to validate:

1. At test time with the validate_with_instantiation_types(T1, T2, ...) attribute. The auto-generated test instantiates the template with the given types and validates the result through naga:

   #![allow(unused)]
   fn main() {
   #[wgsl(validate_with_instantiation_types(f32, f32))]
   pub mod hello_triangle_generic { /* ... */ }
   }

2. At runtime by calling module.validate() on the instantiated ir::Module (requires the validation feature). See Runtime Validation.

If you omit both, the template is never validated by cargo test.

Multiple Type Parameters & Transitive Use

Templates support multiple type parameters and transitive generic calls just like monomorphized generics. Each instantiate call substitutes the full tuple of type arguments through the module's IR; the runtime performs deduplication of any shared monomorphized pieces inside the resulting module.

Const Parameters on Entry Points

Entry points can also take const N: usize (or const N: u32) parameters. The module becomes a template and is instantiated with a concrete integer:

#[wgsl(skip_validation)]
pub mod entry_point {
    use wgsl_rs::std::*;

    #[compute]
    #[workgroup_size(1)]
    pub fn main<const N: usize>() -> u32 {
        let arr: [u32; N] = [0u32; N];
        arr[0]
    }
}

Instantiate with the concrete const value:

#![allow(unused)]
fn main() {
let module: ir::Module = entry_point::instantiate::<4>();
}

Const params use a separate positional namespace ({fn}_c{i}) so type and const params on the same entry point don't collide. The instantiate::<...>() turbofish accepts both type and const arguments in the order they're declared on the entry point.

Vertex / Fragment / Compute

WGSL has three shader stages, each with its own entry-point attribute. wgsl-rs exposes them as Rust attributes that the macro translates to @vertex, @fragment, and @compute.

A module may contain any combination of entry points. Functions without these attributes transpile to ordinary WGSL functions.

Vertex

A vertex entry point takes per-vertex/per-instance inputs (builtins and location-tagged values) and returns a position. Returning a bare Vec4f is automatically annotated with @builtin(position); see Default Annotations.

#![allow(unused)]
fn main() {
#[vertex]
pub fn vtx_main(#[builtin(vertex_index)] vertex_index: u32) -> Vec4f {
    const POS: [Vec2f; 3] = [
        vec2f(0.0, 0.5),
        vec2f(-0.5, -0.5),
        vec2f(0.5, -0.5),
    ];
    let position = POS[vertex_index as usize];
    vec4f(position.x, position.y, 0.0, 1.0)
}
}

@vertex
fn vtx_main(@builtin(vertex_index) vertex_index: u32) -> vec4<f32> {
  /* ... */
}

Fragment

A fragment entry point takes inter-stage inputs (locations and builtins such as @builtin(front_facing)) and returns a color. Returning a bare Vec4f is automatically annotated with @location(0).

#![allow(unused)]
fn main() {
#[fragment]
pub fn frag_main() -> Vec4f {
    vec4f(1.0, 0.0, 0.0, 1.0)
}
}

@fragment
fn frag_main() -> vec4<f32> {
  return vec4<f32>(1.0, 0.0, 0.0, 1.0);
}

Compute

A compute entry point must declare a workgroup size. Use a single integer for a 1D dispatch or three integers for a 3D dispatch:

#[compute]
#[workgroup_size(64)]
pub fn main(#[builtin(global_invocation_id)] global_id: Vec3u) {
    let idx = global_id.x() as usize;
    /* ... */
}

#[compute]
#[workgroup_size(8, 8, 1)]
pub fn tiled(#[builtin(global_invocation_id)] global_id: Vec3u) {
    let x = global_id.x();
    let y = global_id.y();
    /* ... */
}

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) { /* ... */ }

@compute @workgroup_size(8, 8, 1)
fn tiled(@builtin(global_invocation_id) global_id: vec3<u32>) { /* ... */ }

Compute entry points frequently take no return value (a () return type) and access storage or workgroup resources via the binding macros (see Binding Macros).

Inputs

Entry-point inputs are declared as ordinary function parameters. Each parameter may carry an I/O attribute:

• #[builtin(NAME)] — a WGSL builtin value.
• #[location(N)] — a per-vertex attribute (vertex stage) or inter-stage value (fragment stage).

For complex inter-stage IO, use a struct input/output; see Inter-stage IO.

All Three Stages Together

A single module may declare all three stages. The example module below shows vertex, fragment, and compute entry points coexisting (see Binding Macros for storage! and get_mut!):

#![allow(unused)]
fn main() {
#[wgsl]
pub mod pipeline {
    use wgsl_rs::std::*;

    storage!(group(0), binding(0), read_write, COUNTER: u32);

    #[vertex]
    pub fn vs_main(#[builtin(vertex_index)] vi: u32) -> Vec4f {
        vec4f(0.0, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn fs_main() -> Vec4f {
        vec4f(1.0, 1.0, 1.0, 1.0)
    }

    #[compute]
    #[workgroup_size(64)]
    pub fn cs_main(#[builtin(global_invocation_id)] gid: Vec3u) {
        let i = gid.x() as usize;
        get_mut!(COUNTER);
    }
}
}

Inter-stage IO

Vertex outputs and fragment inputs are connected by passing data through a struct whose fields carry WGSL IO attributes. wgsl-rs mirrors the WGSL pattern directly: attributes go on struct fields, and the same struct can serve as both a vertex return type and a fragment parameter.

IO Attributes

Interpolation

#[interpolate(...)] accepts a type and an optional sampling qualifier:

#![allow(unused)]
fn main() {
#[interpolate(flat)]
#[interpolate(linear)]
#[interpolate(perspective)]
#[interpolate(perspective, centroid)]
#[interpolate(perspective, sample)]
}

The default when #[interpolate] is omitted is @interpolate(perspective) with the default sampling, matching WGSL.

Shared Inter-stage Struct

The idiomatic pattern is a single struct used as both the vertex output and the fragment input — the shared_inter_stage example:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod shared_inter_stage {
    use wgsl_rs::std::*;

    pub struct VertexOutput {
        #[builtin(position)]
        pub clip_position: Vec4f,
        #[location(0)]
        pub color: Vec4f,
    }

    #[vertex]
    pub fn vs_main(#[builtin(vertex_index)] vertex_index: u32) -> VertexOutput {
        const POS: [Vec2f; 3] = [
            vec2f(0.0, 0.5),
            vec2f(-0.5, -0.5),
            vec2f(0.5, -0.5),
        ];
        let position = POS[vertex_index as usize];
        VertexOutput {
            clip_position: vec4f(position.x, position.y, 0.0, 1.0),
            color: vec4f(1.0, 0.0, 0.0, 1.0),
        }
    }

    #[fragment]
    pub fn fs_main(input: VertexOutput) -> Vec4f {
        input.color
    }
}
}

struct VertexOutput {
  @builtin(position) clip_position: vec4<f32>,
  @location(0) color: vec4<f32>,
}

@vertex
fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
  /* ... */
  return VertexOutput(/* ... */);
}

@fragment
fn fs_main(input: VertexOutput) -> vec4<f32> {
  return input.color;
}

There is no separate attribute on the struct itself — the field-level attributes do all the work, exactly as in WGSL.

IO Attributes are Stripped from Rust

The #[wgsl] macro strips #[builtin], #[location], #[interpolate], #[blend_src], and #[invariant] from the emitted Rust so the module remains valid Rust without needing wrapper attributes. You do not need to gate these annotations behind a cfg or feature; the macro removes them before the Rust compiler sees the post-expansion module.

This means VertexOutput is a plain #[derive(Wgsl)] struct on the CPU side, and the same field list becomes a fully attributed WGSL struct on the GPU side.

Supported Builtins

wgsl-rs recognizes the following builtin names inside #[builtin(...)]:

position may additionally carry #[invariant] on the vertex output to force invariant interpolation.

Mixing Builtins and Locations

A struct may mix builtins and locations freely:

#![allow(unused)]
fn main() {
pub struct VertexOutput {
    #[builtin(position)]
    #[invariant]
    pub clip_position: Vec4f,
    #[location(0)]
    pub color: Vec4f,
    #[location(1)]
    #[interpolate(flat)]
    pub material_id: u32,
}
}

Default Annotations

To keep simple shaders short, wgsl-rs applies a few default WGSL I/O annotations when an entry point returns a bare vector without an explicit struct.

Vertex Returning Vec4f

A vertex entry point that returns Vec4f directly (rather than a struct) is automatically annotated with @builtin(position) on the return value:

#![allow(unused)]
fn main() {
#[vertex]
pub fn vtx_main(#[builtin(vertex_index)] vertex_index: u32) -> Vec4f {
    vec4f(0.0, 0.0, 0.0, 1.0)
}
}

@vertex
fn vtx_main(@builtin(vertex_index) vertex_index: u32) -> @builtin(position) vec4<f32> {
  return vec4<f32>(0.0, 0.0, 0.0, 1.0);
}

This matches the most common vertex shader shape: produce a clip-space position and nothing else.

Fragment Returning Vec4f

A fragment entry point that returns Vec4f directly is automatically annotated with @location(0):

#![allow(unused)]
fn main() {
#[fragment]
pub fn frag_main() -> Vec4f {
    vec4f(1.0, 0.0, 0.0, 1.0)
}
}

@fragment
fn frag_main() -> @location(0) vec4<f32> {
  return vec4<f32>(1.0, 0.0, 0.0, 1.0);
}

This is the single-render-target case.

When to Use Explicit Annotations

Use a struct return type instead of the bare defaults whenever you need to emit more than one output value:

• A vertex shader that also writes inter-stage varyings (color, UV, normal, ...).
• A fragment shader writing multiple render targets (MRT) — use #[location(N)] per field.
• A vertex shader whose position should be @invariant — put #[builtin(position)] and #[invariant] on the field.
• Dual-source blending — use #[blend_src(0)] and #[blend_src(1)] on two fields.

For example, switching from the default to a struct:

#![allow(unused)]
fn main() {
pub struct VertexOutput {
    #[builtin(position)]
    pub clip_position: Vec4f,
    #[location(0)]
    pub color: Vec4f,
}

#[vertex]
pub fn vs_main(#[builtin(vertex_index)] vi: u32) -> VertexOutput {
    VertexOutput {
        clip_position: vec4f(0.0, 0.0, 0.0, 1.0),
        color: vec4f(1.0, 0.0, 0.0, 1.0),
    }
}
}

See Inter-stage IO for the full set of field attributes.

Compute Entry Points

Compute entry points return () (no value) and have no default annotations; their I/O is entirely through builtins on parameters and through binding macros. See Vertex / Fragment / Compute.

Custom IO via Struct Fields

For fragment inputs, accept a struct parameter whose fields mirror the vertex output struct. wgsl-rs allows the same struct to be used on both sides, which is the recommended shared-inter-stage pattern:

#![allow(unused)]
fn main() {
#[fragment]
pub fn fs_main(input: VertexOutput) -> Vec4f {
    input.color
}
}

The defaults apply only to bare Vec4f returns; once you return or accept a struct, every IO attribute must be explicit on its fields.

Auto-generated Tests

For every non-template #[wgsl] module, the macro automatically generates a hidden test that validates the transpiled WGSL through naga. Running cargo test therefore validates every shader in the crate.

What Gets Generated

Given:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod example {
    use wgsl_rs::std::*;
}
}

the macro emits (approximately):

#![allow(unused)]
fn main() {
#[test]
fn __validate_wgsl() {
    example::WGSL_SOURCE.validate().expect("WGSL validation failed");
}
}

The test name is fixed; there is one per #[wgsl] module. A failure prints the naga diagnostic and the failing module path.

Running the Tests

cargo test                          # validate every #[wgsl] module
cargo test example                  # narrow to one module
cargo test __validate_wgsl          # run only the auto-generated validation tests

A passing test means the generated WGSL parses and passes naga's type-check for the module's declared bindings, entry points, and function bodies.

Template Modules are Excluded

A template module (one with type-parameterized entry points or impl Trait linkages) emits WGSL containing TypeParam placeholders. That text is not valid WGSL on its own, so the macro does not generate an auto-test for templates.

To validate a template, supply concrete types via:

#![allow(unused)]
fn main() {
#[wgsl(validate_with_instantiation_types(f32, f32))]
pub mod hello_triangle_generic { /* ... */ }
}

The attribute makes the auto-test instantiate the template with the given types and validate the resulting concrete module. The number and order of types must match the template's declared type parameters.

See Runtime Validation for validating instantiated templates from your own code, and Templates for the broader template story.

Skipping a Single Module

If a specific module must opt out of auto-validation (e.g. it intentionally exercises a known transpiler bug, or it depends on an extension naga does not yet support), annotate it:

#![allow(unused)]
fn main() {
#[wgsl(skip_validation)]
pub mod example { /* ... */ }
}

No __validate_wgsl test is generated for that module. Other modules are unaffected. See Disabling Validation.

Runtime Validation

In addition to the auto-generated cargo test checks, wgsl-rs exposes runtime validation APIs that let you validate shader source on demand — useful for instantiated templates, dynamically composed pipelines, and CI scripts that want to fail fast on bad WGSL.

WGSL_SOURCE.validate()

Every #[wgsl] module exposes a pub static WGSL_SOURCE: &str containing the transpiled WGSL. When the validation feature is enabled, calling .validate() on it runs naga and returns a Result:

#![allow(unused)]
fn main() {
use example::hello_triangle::WGSL_SOURCE;

WGSL_SOURCE.validate().expect("hello_triangle failed validation");
}

This is exactly what the auto-generated __validate_wgsl test calls. See Auto-generated Tests.

module.validate()

For template modules you instantiate at runtime, the returned ir::Module also has a .validate() method (with the validation feature):

#![allow(unused)]
fn main() {
use example::hello_triangle_generic as tmpl;
use wgsl_rs::ir;

let module: ir::Module = tmpl::instantiate::<f32, f32>();
module.validate().expect("instantiated template failed validation");
let source: String = module.to_wgsl();
}

module.validate() runs naga over the substituted, concrete WGSL — the same path the validate_with_instantiation_types attribute uses at test time, but driven from your own code.

validate_with_instantiation_types at Runtime

The validate_with_instantiation_types(T1, T2, ...) helper is also callable as a runtime function on a template module:

#![allow(unused)]
fn main() {
use example::hello_triangle_generic as tmpl;

tmpl::validate_with_instantiation_types(f32, f32)
    .expect("template validation failed");
}

This is convenient when you want to validate several instantiations of the same template from a single test or application entry point. The argument list mirrors the template's declared type parameters.

The validation Feature

Both WGSL_SOURCE.validate() and module.validate() require the validation feature on the wgsl-rs crate. It is enabled by default. If you turn it off (see Disabling Validation), those methods are removed and any call site will fail to compile — there is no stub that silently returns Ok.

Error Reporting

Validation errors surface naga's diagnostics directly. A typical failure looks like:

WGSL validation failed: SomeWrongSnafu { ... }
  in module `hello_triangle`
  at @vertex fn vtx_main(...)

The error includes the offending module name (or the instantiated template's type arguments) and the naga span where available. For template modules, validate after instantiation so the error points at concrete WGSL rather than TypeParam placeholders.

Disabling Validation

Validation is on by default and runs through naga. There are two granularities at which you can disable it: per module, or globally via Cargo features.

Per Module: #[wgsl(skip_validation)]

Annotate a single module to suppress its auto-generated __validate_wgsl test:

#![allow(unused)]
fn main() {
#[wgsl(skip_validation)]
pub mod generic_structs { /* ... */ }
}

Effects:

• No #[test] fn __validate_wgsl() is emitted for that module.
• WGSL_SOURCE is still produced and usable at runtime.
• Other modules in the crate are still validated normally.

Use this when a module intentionally cannot pass naga (e.g. it exercises a known transpiler bug such as the generic struct constructor mangling issue, or it relies on a WGSL extension naga does not yet support). The generic_structs example uses it for exactly this reason:

#![allow(unused)]
fn main() {
#[wgsl(skip_validation)]
pub mod generic_structs {
    pub struct Pair<T: Copy> { pub a: T, pub b: T }
}
}

Globally: default-features = false

Validation pulls in naga, which is a non-trivial dependency. To remove it entirely from your build, disable the default features of the wgsl-rs crate:

[dependencies]
wgsl-rs = { version = "...", default-features = false }

Effects:

• The validation feature is off.
• WGSL_SOURCE.validate() and ir::Module::validate() are removed (not stubbed). Any call site will fail to compile, so you must also remove your own calls to these methods.
• The auto-generated __validate_wgsl tests are not emitted for any module, so cargo test no longer validates shaders.
• WGSL_SOURCE and ir::Module::to_wgsl() are still available — generation is unaffected, only validation is removed.

If you later want validation back in a specific build (e.g. CI), enable the feature explicitly:

cargo test --features wgsl-rs/validation

Choosing a Granularity

Disabling validation does not change the generated WGSL text — only whether it is checked. Always run validation somewhere in your pipeline (CI, dev builds, or an explicit test) before shipping shaders.

The Standard Library

wgsl_rs::std is the prelude for every #[wgsl] module. It provides WGSL types, builtin functions, binding macros, entry-point attributes, and runtime macros that bridge the Rust and WGSL worlds.

The glob import

Every #[wgsl] module begins with a glob import:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod my_shader {
    use wgsl_rs::std::*;
    // ...
}
}

The glob import is required. Only use wgsl_rs::std::* is recognized by the transpiler; named imports from std are not supported inside #[wgsl] modules.

What it provides

The Wgsl trait

Wgsl marks a type usable in a #[wgsl] module. Any type passed to a WGSL function, stored in a uniform!/storage!, or returned from an entry point must implement Wgsl. The macro and the runtime both rely on this trait to marshal values between Rust and WGSL.

Related traits:

WgslTextureScalar is a stricter subset of WgslScalar: only types that have a corresponding WGSL texture format qualify, so bool is excluded.

CPU and WGSL agreement

Every builtin in wgsl_rs::std has a CPU implementation that mirrors WGSL semantics. When you run a #[wgsl] module as ordinary Rust (e.g. under cargo test), the builtins execute on the CPU; when the transpiler emits WGSL, the same names map to native WGSL builtins. The roundtrip tests verify the two worlds agree.

Numeric Builtins

The numeric builtins mirror WGSL's numeric functions. Each is a free function exported by wgsl_rs::std::*. Many are defined per concrete type via a one-trait-per-builtin strategy: the transpiler resolves the right WGSL builtin based on argument types.

Trigonometric

Exponential, logarithmic, and root

Rounding and floating-point decomposition

modf

modf(x) returns a struct with two fields:

#![allow(unused)]
fn main() {
let r = modf(-1.5);
let frac: f32 = r.fract; //  0.5
let whole: f32 = r.whole; // -1.0
}

frexp

frexp(x) returns a struct with:

• .fract — significand in [0.5, 1.0)
• .exp — exponent such that x = fract * 2^exp

Interpolation and clamping

Geometric

Per-type dispatch

Some builtins are implemented as traits with one method per concrete scalar type (the one-trait-per-builtin strategy). This keeps the CPU implementation type-correct and lets the transpiler emit the exact WGSL overload. You call them as ordinary free functions; the correct specialization is inferred from argument types.

Example

#![allow(unused)]
fn main() {
#[wgsl]
pub mod numeric_example {
    use wgsl_rs::std::*;

    pub fn to_srgb(linear: f32) -> f32 {
        if linear <= 0.0031308 {
            linear * 12.92
        } else {
            1.055 * pow(linear, 1.0 / 2.4) - 0.055
        }
    }

    pub fn smoothstep(edge0: f32, edge1: f32, x: f32) -> f32 {
        let t = clamp((x - edge0) / (edge1 - edge0), 0.0, 1.0);
        t * t * (3.0 - 2.0 * t)
    }
}
}

smoothstep shown above is a user-authored helper using clamp and pow. If a WGSL smoothStep builtin is supported by your target, prefer it.

Matrix & Vector Functions

wgsl_rs::std provides vector and matrix types, constructors, and the WGSL builtins that operate on them.

Types

Vector and matrix types are aliases over the generic Vec and Mat structs:

See Scalars & Literals, Vectors & Swizzles, and Matrices for full coverage.

Constructors

Constructors are free functions named like the WGSL type:

#![allow(unused)]
fn main() {
let a = vec2f(1.0, 2.0);
let b = vec3f(1.0, 2.0, 3.0);
let c = vec4f(0.0);                 // splat
let m = mat4x4f(
    1.0, 0.0, 0.0, 0.0,
    0.0, 1.0, 0.0, 0.0,
    0.0, 0.0, 1.0, 0.0,
    0.0, 0.0, 0.0, 1.0,
);
}

Constructors accept scalars, smaller vectors, and combinations thereof, just like WGSL.

Vector operations

Matrix builtins

#![allow(unused)]
fn main() {
#[wgsl]
pub mod matrix_example {
    use wgsl_rs::std::*;

    pub fn normal_matrix(model: Mat4f) -> Mat3f {
        let upper = mat3x3f(
            model[0].xyz(),
            model[1].xyz(),
            model[2].xyz(),
        );
        let det = determinant(upper);
        if abs(det) < 1e-8 {
            return mat3x3f(1.0, 0.0, 0.0,
                           0.0, 1.0, 0.0,
                           0.0, 0.0, 1.0);
        }
        transpose(upper) * (1.0 / det)
    }
}
}

inverse is not a WGSL builtin. Compute it from the adjugate and determinant, or use transpose of the cofactor matrix for the common 3×3 normal-matrix case.

Component access

Vector components are accessed with .x(), .y(), .z(), .w() or via swizzle methods like .xyz(), .xy(), .xx(). See Vectors & Swizzles.

Matrix columns are indexed with m[i] (returns a vector) and individual entries with m[i][j], matching WGSL semantics.

Texture & Sampler Functions

wgsl_rs::std provides the WGSL texture and sampler builtins plus the sampler! and texture! binding macros used to declare them.

Sampler types

See Binding Macros: texture! & sampler! for declaration syntax.

Sampling functions

Each function has multiple overloads (2D, 2DArray, 3D, Cube, CubeArray, etc.), implemented as separate Rust functions that map to the same WGSL builtin. The transpiler picks the right overload from argument types.

Functions suffixed ..._level are usable from any stage; plain texture_sample and texture_sample_bias/texture_sample_compare are restricted to the fragment stage in WGSL.

Load / store / query

Storage texture access

texture_load_storage and texture_store are gated by the ReadableStorageAccess and WritableStorageAccess traits respectively, so the access mode of the storage texture is enforced at compile time:

A separate texture_load_storage function (rather than overloading texture_load) is used because the WGSL storage overload of textureLoad takes no level parameter, unlike the sampled texture overload.

Example

#![allow(unused)]
fn main() {
#[wgsl]
pub mod texture_example {
    use wgsl_rs::std::*;

    texture!(group(0), binding(0), COLOR_TEX: Texture2D<f32>);
    sampler!(group(0), binding(1), COLOR_SMP: Sampler);

    pub struct FragmentInput {
        #[location(0)]
        pub uv: Vec2f,
    }

    #[fragment]
    pub fn fs_main(input: FragmentInput) -> Vec4f {
        let uv = input.uv;
        let color = texture_sample(COLOR_TEX, COLOR_SMP, uv);
        let dim = texture_dimensions(COLOR_TEX);
        color
    }
}
}

Overload resolution

Because Rust has no WGSL-style overload sets, each texture builtin is implemented as a distinct Rust function per texture kind, e.g. there is one texture_sample for Texture2D, another for Texture2DArray, another for Texture3D, another for TextureCube, and so on. The transpiler emits the WGSL textureSample builtin regardless of which Rust overload you called — the overload exists only to type-check on the CPU side.

Derivatives

Derivatives compute per-pixel rate-of-change of a value with respect to screen-space coordinates. They are only valid inside fragment shaders; calling them from any other stage is a WGSL error.

Fine / coarse variants

WGSL provides three precision tiers. wgsl-rs exposes all of them:

The bare dpdx/dpdy/fwidth map to WGSL's default-precision builtins, which WGSL defines as coarse. Prefer the explicit _fine / _coarse variants when the choice matters for your application.

Common uses

• Mip selection in non-fragment-aware sampling: pass gradients to textureSampleGrad from dpdx/dpdy of the texture coordinates.
• Edge detection / anti-aliasing: fwidth to compute pixel-local width.
• Screen-space dependent branching: compare fwidth against a threshold.

Example

#![allow(unused)]
fn main() {
#[wgsl]
pub mod derivatives_example {
    use wgsl_rs::std::*;

    uniform!(FRAME, Frame);
    texture!(COLOR_TEX, Texture2d);
    sampler!(COLOR_SMP, Sampler);

    pub struct Frame {
        pub time: f32,
    }

    #[fragment]
    pub fn fs(in: VertexOutput) -> Vec4f {
        let uv = in.uv;
        let dx = dpdx_fine(uv);
        let dy = dpdy_fine(uv);
        let w = fwidth_fine(uv);
        let color = textureSampleGrad(COLOR_TEX, COLOR_SMP, uv, dx, dy);
        color
    }
}
}

CPU behavior

On the CPU, derivatives return zero for dpdx/dpdy and the input value's magnitude for fwidth — enough to keep CPU tests running without panic. The GPU is the source of truth for derivative accuracy.

Bitcast

bitcast reinterprets the bit pattern of a value as a different type without changing any bits — unlike f32 as u32, which performs a numeric conversion. WGSL's bitcast builtin takes a single argument and the result type is inferred from context; wgsl-rs instead provides one named function per target type so Rust type inference is unambiguous.

Functions

Each function is named bitcast_<targettype>:

The set of accepted input types per target follows WGSL §17: the source and target must have the same bit width, and only numeric scalar/vector types are allowed (no bool).

Why per-target functions

WGSL resolves bitcast overloading from the surrounding expression context, which Rust cannot do without type annotations. Naming each target type makes the intent explicit on the CPU side and keeps type inference deterministic.

Example

#![allow(unused)]
fn main() {
#[wgsl]
pub mod bitcast_example {
    use wgsl_rs::std::*;

    pub fn pack_normal_as_u32(n: Vec3f) -> u32 {
        let q = vec4f(n.x() * 0.5 + 0.5,
                      n.y() * 0.5 + 0.5,
                      n.z() * 0.5 + 0.5,
                      0.0);
        bitcast_vec4u(q).x()
    }

    pub fn unpack_normal_from_u32(packed: u32) -> Vec3f {
        let q = bitcast_vec4f(vec4u(packed, 0, 0, 0));
        q.xyz() * 2.0 - 1.0
    }
}
}

CPU behavior

On the CPU, these map to f32::from_bits / f32::to_bits (and the Vec equivalents), so the bit pattern is preserved exactly. This is what makes bitcast safe to use in roundtrip tests.

Packing

WGSL provides builtins to pack and unpack vectors of normalized or floating-point values into a single u32. wgsl-rs exposes all of them as free functions in wgsl_rs::std.

Pack / unpack pairs

Inputs are Vec4f / Vec2f for the pack functions; outputs are u32. The unpack functions take u32 and return the corresponding vector.

Rounding behavior

The pack functions round normalized floats to the nearest representable integer using round-to-nearest-even, matching WGSL. Values outside the normalized range are clamped.

Example

#![allow(unused)]
fn main() {
#[wgsl]
pub mod packing_example {
    use wgsl_rs::std::*;

    pub fn encode_tangent(t: Vec4f) -> u32 {
        pack4x8snorm(t)
    }

    pub fn decode_tangent(packed: u32) -> Vec4f {
        unpack4x8snorm(packed)
    }
}
}

CPU behavior

The CPU implementations use the same rounding and clamping as WGSL, so pack followed by unpack returns a value within one ULP of the original. This makes packing safe to exercise in unit tests.

Synchronization

Synchronization builtins order memory accesses between invocations in a compute workgroup. They are only valid inside compute shaders; calling them from vertex or fragment stages is a WGSL error.

Barriers

All three take no arguments and return unit. On the CPU they are no-ops — a single-threaded CPU dispatch has nothing to synchronize — but they still compile and execute, so the same shader can run in both worlds.

workgroupUniformLoad

workgroupUniformLoad<T>(&var: &T) -> T reads a workgroup variable such that all invocations in the workgroup observe the same value. WGSL requires the loaded address to be uniform across the workgroup.

wgsl-rs models this via the WorkgroupUniformLoad trait, implemented for the types that are safe to load this way.

#![allow(unused)]
fn main() {
#[wgsl]
pub mod sync_example {
    use wgsl_rs::std::*;

    workgroup!(SHARED: [u32; 64]);

    #[compute]
    #[workgroup_size(64)]
    pub fn cs(#[builtin(local_invocation_id)] lid: Vec3u) {
        let idx = lid.x() as usize;
        get_mut!(SHARED)[idx] = lid.x();
        workgroup_barrier();
        let partner = get!(SHARED)[((lid.x() + 1u32) % 64u32) as usize];
        workgroup_barrier();
        let uniform_first = workgroup_uniform_load(&SHARED);
    }
}
}

When to use which

• workgroup_barrier — gate control flow: ensure every invocation has reached a point before any proceeds.
• storage_barrier — gate storage! reads/writes across invocations.
• texture_barrier — gate storage-texture writes before subsequent reads.
• workgroup_uniform_load — broadcast one uniform value to the whole workgroup (useful for divergent control flow convergence).

discard!()

discard!() aborts the current fragment's output. In WGSL it emits the discard; statement; on the CPU it sets a thread-local flag.

Syntax

#![allow(unused)]
fn main() {
discard!();
}

It is a macro (not a function) because it must be recognized by the transpiler as a control-flow side effect.

Semantics

• WGSL: execution of the rest of the fragment invocation has undefined behavior; outputs (color, depth) are not committed. Helper-invocation semantics apply — derivatives and similar may still run.
• CPU: sets a thread-local "discarded" flag and returns. The CPU dispatch runtime checks this flag after the entry point returns and skips committing the fragment's output. Code after discard!() continues to execute unless you explicitly return; this mirrors WGSL's helper-invocation model, where the shader is not abruptly terminated.

Reachability

discard!() can appear in any function reachable from a fragment entry point. The transpiler tracks this through the call graph. Using it outside fragment-reachable code is a compile error.

Example

#![allow(unused)]
fn main() {
#[wgsl]
pub mod discard_example {
    use wgsl_rs::std::*;

    pub struct Material {
        pub alpha_cutoff: f32,
    }
    uniform!(MATERIAL, Material);
    texture!(ALBEDO_TEX, Texture2d);
    sampler!(ALBEDO_SMP, Sampler);

    #[fragment]
    pub fn fs(in: VertexOutput) -> Vec4f {
        let albedo = textureSample(ALBEDO_TEX, ALBEDO_SMP, in.uv);
        if albedo.w() < MATERIAL.alpha_cutoff {
            discard!();
        }
        albedo
    }
}
}

Difference from an early return

discard!() is not a return. It marks the fragment for rejection but lets subsequent statements run. If you want to stop executing immediately, combine it with an explicit return:

#![allow(unused)]
fn main() {
if albedo.w() < cutoff {
    discard!();
    return vec4f(0.0);
}
}

On the CPU this keeps the thread-local flag set and the dispatch runtime will ignore the returned value. On the GPU the WGSL discard; ensures the output is not committed regardless of what the return produces.

wgpu Linkage Overview

wgpu linkage is the bridge between a #[wgsl] module and a running wgpu pipeline. Given a shader module, the linkage analyzer produces everything wgpu needs to create buffers, bind groups, and pipelines, with sizes computed per WGSL §14.4.1 (not Rust's sizeof).

Enabling linkage

Linkage is gated behind the linkage-wgpu cargo feature:

[dependencies]
wgsl-rs = { version = "0.1", features = ["linkage-wgpu"] }

See Cargo Features.

WgpuLinkage

WgpuLinkage is the main type. It owns the concrete IR it was built from and renders the WGSL source via ir::render_module. You get one of two ways:

#![allow(unused)]
fn main() {
use wgsl_rs::linkage::wgpu::{analyze_wgsl_module, analyze_ir_module, WgpuLinkage};

// From a concrete (non-template) Source:
let linkage: WgpuLinkage = analyze_wgsl_module(&my_shader::SOURCE)?;

// From an instantiated template's IR (see [Template Modules & Instantiation](../generics/templates.md)):
let concrete_ir: ir::Module = template.instantiate::<f32, u32>()?;
let linkage: WgpuLinkage = analyze_ir_module(concrete_ir);
}

Source and ir::Module both expose generate_linkage() — an extension trait method (IrModuleExt) re-exported via wgsl_rs::*:

#![allow(unused)]
fn main() {
use wgsl_rs::*;
let linkage = my_shader::SOURCE.generate_linkage()?;
}

What WgpuLinkage provides

See Bind Groups & Buffers and Pipeline Layouts for detail.

High-level workflow

1. Write your shader in a #[wgsl] module.
2. Build the WgpuLinkage (via generate_linkage() or analyze_*).
3. Create the wgpu::ShaderModule from the linkage.
4. Create the buffers and bind groups you need.
5. Create the pipeline layout (cached) and the render/compute pipeline.

#![allow(unused)]
fn main() {
use wgsl_rs::*;

let mut linkage = my_shader::SOURCE.generate_linkage()?;
let shader = linkage.shader_module(&device);

let frame_bg = linkage.create_bind_group_named(0, &device, &[
    ("FRAME", frame_buffer.as_entire_binding()),
])?;

let layout = linkage.pipeline_layout(&mut self, &device, Some("main_layout"));
let pipeline = device.create_render_pipeline(wgpu::RenderPipelineDescriptor {
    label: Some("main"),
    layout: Some(&layout),
    vertex: linkage.entry_point("vs").unwrap().vertex_state(&device, &[], &[]),
    fragment: Some(wgpu::FragmentState {
        module: &shader,
        entry_point: linkage.entry_point("fs").unwrap().name(),
        targets: &[Some(wgpu::TextureFormat::Bgra8Unorm.into())],
    }),
    primitive: wgpu::PrimitiveState::default(),
    depth_stencil: None,
    multisample: wgpu::MultisampleState::default(),
    multiview: None,
});
}

Sizing

Buffer and binding sizes are computed from the IR using WGSL §14.4.1 alignment and size rules, not Rust's size_of. This means a Rust struct with padding or a different layout from its WGSL equivalent still gets the right wgpu binding size. See Memory Layout for the underlying trait machinery.

Layout caching

WgpuLinkage lazily caches wgpu::BindGroupLayouts (per group index) and the wgpu::PipelineLayout. Methods that create these take &mut self; the returned wgpu types are Arc-backed, so cloning is cheap. You typically keep one WgpuLinkage per shader and call the &mut self methods once at pipeline-construction time.

Bind Groups & Buffers

WgpuLinkage exposes the bind groups, bindings, and buffers declared by a shader via uniform!, storage!, workgroup!, texture!, and sampler!.

Bind groups

Each @group(n) in WGSL corresponds to a BindGroupInfo. Access them by index or iterate:

#![allow(unused)]
fn main() {
let bg0: &BindGroupInfo = linkage.bind_group(0).expect("group 0 exists");
for (n, bg) in linkage.bind_groups() {
    println!("group {n}: {} bindings", bg.bindings.len());
}
}

BindGroupInfo carries:

• The group index.
• The list of bindings (@binding(i) entries) with their names, types, and visibility (see Per-binding Shader Stages).
• A method to create the wgpu::BindGroupLayout.

Creating a bind group

The simplest path: WgpuLinkage::create_bind_group_named, which builds the bind group and caches the bind group layout for that group index:

#![allow(unused)]
fn main() {
let frame_bg = linkage.create_bind_group_named(0, &device, &[
    ("FRAME", frame_buffer.as_entire_binding()),
    ("COLOR_TEX", &color_texture_view),
    ("COLOR_SMP", &color_sampler),
])?;
}

Each entry is ("NAME", resource) where NAME is the binding name declared in the shader (the first argument to uniform!/storage!/texture!/ sampler!) and resource is a wgpu::BindingResource.

For more control, use BindGroupInfo::create directly:

#![allow(unused)]
fn main() {
let layout = bg0.create_layout(&device);
let bg = bg0.create(&device, &layout, &[
    frame_buffer.as_entire_binding(),
    &color_texture_view,
    &color_sampler,
]);
}

Resources must be supplied in binding-index order, matching the @binding(i) numbering in the generated WGSL.

create_named lets you supply resources by name regardless of order:

#![allow(unused)]
fn main() {
let bg = bg0.create_named("renderer", &device, &layout, &[
    ("FRAME", frame_buffer.as_entire_binding()),
    ("COLOR_TEX", &color_texture_view),
]);
}

Buffers

uniform! and storage! bindings produce a BufferDescriptorInfo entry that knows the buffer's WGSL size and usage. Find a buffer by its declared name:

#![allow(unused)]
fn main() {
let frame_buf_info = linkage.buffer("FRAME").expect("FRAME buffer exists");
let frame_buffer = frame_buf_info.create_buffer(&device);
}

BufferDescriptorInfo::create_buffer(device) returns a wgpu::Buffer sized per WGSL §14.4.1, with wgpu::BufferUsages derived from how the shader declares the binding (uniform vs storage, read-only vs read-write).

Workgroup variables

workgroup! bindings do not appear in bind groups — they are workgroup-scoped storage. They do not need wgpu host-side resources.

Example: full bind group setup

#![allow(unused)]
fn main() {
let frame_info = linkage.buffer("FRAME").unwrap();
let frame_buffer = frame_info.create_buffer(&device);

let frame_bg = linkage.create_bind_group_named(0, &device, &[
    ("FRAME", frame_buffer.as_entire_binding()),
    ("ALBEDO", &albedo_view),
    ("SMP", &linear_sampler),
])?;
}

Layout caching

create_bind_group_named caches the wgpu::BindGroupLayout inside the WgpuLinkage so subsequent calls for the same group index return the cached layout. BindGroupInfo::create/create_named take an explicit layout, so you control whether to cache yourself.

Pipeline Layouts

WgpuLinkage builds and caches the wgpu::PipelineLayout from the bind group layouts it derives from the shader. Combined with EntryPointInfo, you can create a complete render or compute pipeline without manually writing layout descriptors.

pipeline_layout

#![allow(unused)]
fn main() {
let layout: wgpu::PipelineLayout = linkage.pipeline_layout(&mut linkage, &device, Some("main"));
}

• Takes &mut self because the result is cached.
• Returns an owned wgpu::PipelineLayout (internally Arc-backed, so cloning is cheap).
• The label argument is passed straight through to wgpu.

Calling pipeline_layout again returns a clone of the cached layout.

Entry points

linkage.entry_point(name) returns Option<&EntryPointInfo>. The name is the Rust function name of your #[vertex] / #[fragment] / #[compute] entry point.

EntryPointInfo exposes the wgpu descriptor builders:

The *_state builders take the same trailing arguments as the corresponding wgpu descriptor fields (vertex buffer layouts, color targets, pipeline constants), so you keep full control over the pipeline descriptor while the linkage supplies the module and entry_point.

Full pipeline creation

#![allow(unused)]
fn main() {
use wgsl_rs::*;

let mut linkage = my_shader::SOURCE.generate_linkage()?;
let shader = linkage.shader_module(&device);

let frame_bg = linkage.create_bind_group_named(0, &device, &[
    ("FRAME", frame_buffer.as_entire_binding()),
    ("ALBEDO", &albedo_view),
    ("SMP", &linear_sampler),
])?;

let pipeline_layout = linkage.pipeline_layout(&mut linkage, &device, Some("main_layout"));

let vs_info = linkage.entry_point("vs").expect("vertex entry point");
let fs_info = linkage.entry_point("fs").expect("fragment entry point");

let pipeline = device.create_render_pipeline(wgpu::RenderPipelineDescriptor {
    label: Some("main_pipeline"),
    layout: Some(&pipeline_layout),
    vertex: vs_info.vertex_state(&device, &vertex_buffer_layouts, &[]),
    fragment: Some(fs_info.fragment_state(&device, &[Some(wgpu::TextureFormat::Bgra8Unorm.into())], &[])),
    primitive: wgpu::PrimitiveState::default(),
    depth_stencil: None,
    multisample: wgpu::MultisampleState::default(),
    multiview: None,
});
}

Compute pipelines

For compute, use EntryPointInfo::compute_state:

#![allow(unused)]
fn main() {
let cs_info = linkage.entry_point("main").expect("compute entry point");
let pipeline = device.create_compute_pipeline(wgpu::ComputePipelineDescriptor {
    label: Some("compute"),
    layout: Some(&pipeline_layout),
    module: &shader,
    entry_point: cs_info.name(),
    compilation_options: Default::default(),
    cache: None,
});
}

compute_state returns a small struct with module and entry_point fields so you can spread them into the descriptor yourself; name() is available when you want just the string.

Template Linkage

A template module is generic: it cannot be analyzed for wgpu linkage until it is instantiated with concrete type arguments. This chapter covers how to go from a generic shader to a usable WgpuLinkage.

Templates cannot be analyzed directly

Calling analyze_wgsl_module on the Source of a template fails:

#![allow(unused)]
fn main() {
let linkage = analyze_wgsl_module(&template::SOURCE);
// Err(Error::TemplateResolution)
}

The analyzer needs concrete types to compute binding sizes and entry-point signatures. A template's Source still has unsubstituted generic type parameters, so there is nothing concrete to link.

Instantiate, then analyze

The correct flow is:

1. Get the template's IR (or Source).
2. Call instantiate::<T1, T2, ...>() to produce a concrete ir::Module.
3. Pass the concrete IR to analyze_ir_module.

#![allow(unused)]
fn main() {
use wgsl_rs::linkage::wgpu::analyze_ir_module;
use wgsl_rs::*;

let template: &Source = &renderer::SOURCE;
let concrete_ir: ir::Module = template.instantiate::<f32, Vec4f>()?;
let linkage = analyze_ir_module(concrete_ir);
}

ir::Module::generate_linkage() (the IrModuleExt re-export) does the second step in one call:

#![allow(unused)]
fn main() {
let linkage = concrete_ir.generate_linkage()?;
}

WgpuLinkage owns the concrete IR

WgpuLinkage holds the ir::Module it was given. This matters: the WGSL source it returns via wgsl_source() is rendered from that exact concrete IR, so the source always matches what was analyzed. There is no risk of re-rendering a different template instantiation.

For a non-template Source, analyze_wgsl_module builds the IR internally and owns it the same way.

Example: instantiate a generic shader

#![allow(unused)]
fn main() {
#[wgsl]
pub mod generic_blit {
    use wgsl_rs::std::*;

    pub struct Params<T> { pub scale: T, pub bias: T }

    uniform!(PARAMS, Params<Vec4f>);     // concrete after instantiation
    texture!(SRC, Texture2d);
    sampler!(SMP, Sampler);

    #[vertex]
    pub fn vs(...) -> VertexOutput { /* ... */ }

    #[fragment]
    pub fn fs(in: VertexOutput) -> Vec4f {
        let c = textureSample(SRC, SMP, in.uv);
        c * PARAMS.scale + PARAMS.bias
    }
}

// In application code:
use wgsl_rs::*;

let concrete = generic_blit::SOURCE.instantiate::<Vec4f>()?;
let mut linkage = concrete.generate_linkage()?;
let shader = linkage.shader_module(&device);
let layout = linkage.pipeline_layout(&mut linkage, &device, Some("blit"));
// ... build pipeline as usual
}

Summary

Per-binding Shader Stages

Each wgpu binding has a wgpu::ShaderStages visibility flag. The linkage analyzer computes this per binding by walking the bodies of the shader's entry-point functions and collecting which bindings each references.

How visibility is computed

1. For each entry point (#[vertex], #[fragment], #[compute]), the analyzer walks the function body (and any called functions) and records every binding name that is read or written.
2. A binding's ShaderStages is the union of the stages of the entry points that reference it.
3. Bindings not referenced by any entry point default to wgpu::ShaderStages::COMPUTE.

The default reflects WGSL's own default stage for unreferenced bindings and keeps the common compute-only case working without annotation.

Example

Given:

#![allow(unused)]
fn main() {
#[wgsl]
pub mod stages_example {
    use wgsl_rs::std::*;

    uniform!(FRAME, Frame);
    storage!(VERTICES, [Vertex; 1024]);
    storage!(COUNTERS, [u32; 4]);

    #[vertex]
    pub fn vs(...) -> VertexOutput {
        let v = VERTICES[i];
        // ...
    }

    #[fragment]
    pub fn fs(in: VertexOutput) -> Vec4f {
        let f = FRAME.time;
        // ...
    }

    #[compute]
    pub fn cs(...) {
        COUNTERS[0] += 1;
    }
}
}

Resulting visibility:

If a binding is referenced from both vertex and fragment stages, the visibility is VERTEX | FRAGMENT.

Why this matters: read-write storage in vertex stage

storage! bindings default to read_write access in WGSL unless restricted. A read_write storage buffer visible to the vertex stage requires the wgpu::Features::BUFFER_BINDING_ARRAY_NON_UNIFORM_INDEXING family / VERTEX_WRITABLE_STORAGE feature — specifically wgpu::Features::VERTEX_WRITABLE_STORAGE — because not all GPUs support writes from the vertex stage.

Because the analyzer derives visibility from actual references, a storage binding touched only by a compute entry point will not force the vertex-writable-storage feature on. If you instead set ShaderStages::all() globally, wgpu may reject the pipeline on hardware that lacks vertex-writable-storage.

The per-binding computation keeps the requested feature set as small as possible: only bindings actually referenced by a vertex-stage entry point acquire vertex visibility, and only read_write storage among those requires the feature.

Overriding visibility

The analyzer's per-binding result is what gets passed to wgpu::BindGroupLayoutEntry::visibility. If you need to widen or narrow visibility (for example, to bind the same layout across multiple shaders), build the BindGroupLayout yourself with explicit ShaderStages and pass it to BindGroupInfo::create/create_named. See Bind Groups & Buffers.

The WgslExtension Trait

The WgslExtension trait lets downstream crates inspect and modify a shader's WGSL IR after transpilation but before type instantiation. It is the primary extension point for post-transpile code generation and analysis.

Definition

The trait lives in wgsl_rs::extension and is re-exported at the crate root.

#![allow(unused)]
fn main() {
pub trait WgslExtension {
    const MACROS: &'static [&'static str] = &[];
    fn modify_ir(module: &mut crate::ir::Module);
}
}

Import it directly from the crate root:

#![allow(unused)]
fn main() {
use wgsl_rs::WgslExtension;
}

Purpose

modify_ir receives a mutable reference to the IR module, giving an extension full read/write access to every item, field, function argument, and attribute. Extensions can:

• Inject helper functions derived from #[derive(...)] attributes.
• Rewrite or remove items.
• Inspect attributes that are preserved on IR nodes but never rendered to WGSL.
• Lower custom statement macros (see Statement Macro Lowering below).

Anything an extension adds or rewrites carries through every subsequent instantiate() call (see Template Modules & Instantiation). TypeParam nodes in injected code are substituted automatically.

Wiring

Extensions are activated via the extensions argument on the wgsl attribute:

#![allow(unused)]
fn main() {
#[wgsl(extensions = [my_crate::NoopExt, my_crate::SlabItemExt])]
mod shader {
    // ...
}
}

Extensions run in declaration order on every wgsl_source() and instantiate() call. There is no priority mechanism — order is determined solely by the list order in the attribute.

Minimal Example

A no-op extension useful as a smoke test:

#![allow(unused)]
fn main() {
use wgsl_rs::WgslExtension;
use wgsl_rs::ir;

pub struct NoopExt;

impl WgslExtension for NoopExt {
    fn modify_ir(_module: &mut ir::Module) {}
}
}

For details on walking the IR and the SlabItemExt worked example, see Modifying the IR.

Statement Macro Lowering

When the #[wgsl] parser encounters a statement macro that is not one of its builtins (slab_copy!, discard!), it passes it through as an ir::Stmt::Macro { name, args } variant instead of rejecting it. An extension can then recognize the macro by name in modify_ir and replace the Stmt::Macro with lowered IR statements.

Extensions declare which macro names they handle via the MACROS associated const:

#![allow(unused)]
fn main() {
pub struct SlabItemExt;

impl WgslExtension for SlabItemExt {
    const MACROS: &'static [&'static str] = &["slab_read", "slab_write"];

    fn modify_ir(module: &mut ir::Module) {
        for item in &mut module.items {
            if let ir::Item::Fn(f) = item {
                lower_in_block(&mut f.block);
            }
        }
    }
}
}

The #[wgsl] macro emits a compile-time const check ensuring every Stmt::Macro in the module is claimed by at least one listed extension. If a macro name is used in the module but no listed extension declares it in MACROS, the result is a compile error (E0080), not a runtime error.

Extensions that do not lower statement macros can leave MACROS as the empty default.

Why statement-position only

#[wgsl] runs before macro_rules! expand, so expression-position macros are untyped black boxes — the parser can only accept ones it recognizes by name and knows the return type of (like get!/get_mut!). Statement-position macros don't need return types, so they can be passed through as Stmt::Macro and lowered by an extension. This is why downstream statement macros (e.g. crabslab's slab_read!/slab_write!) are statement macros rather than expressions.

For a full worked example, see Worked Examples.

Modifying the IR

modify_ir is the single entry point an extension implements. Understanding when and how it runs is essential for writing correct extensions.

When It Runs

modify_ir is invoked by the IR constructor:

1. IR items are built from the transpiled source.
2. modify_ir runs for each extension, in declaration order.
3. Type instantiation occurs (for templates), substituting TypeParam nodes.

This means extensions see the fully constructed IR but operate before any type-parameter substitution. Anything an extension injects containing TypeParam nodes will be substituted automatically during instantiation.

modify_ir runs on every wgsl_source() and instantiate() call — not once at definition time.

Walking the Module

ir::Module is:

#![allow(unused)]
fn main() {
pub struct Module {
    pub name: String,
    pub items: Vec<Item>,
    pub attrs: Vec<Attribute>,
}
}

ir::Item is an enum with variants: Struct, Fn, Const, Uniform, Storage, Workgroup, Sampler, Texture, Impl, Enum.

Inside function bodies, ir::Stmt includes a Macro { name, args } variant for unrecognized statement macros (see Statement Macro Lowering). Extensions that claim macros via MACROS walk function blocks and replace Stmt::Macro nodes with lowered IR.

To iterate and mutate items, match on the variant:

#![allow(unused)]
fn main() {
use wgsl_rs::WgslExtension;
use wgsl_rs::ir::{Item, Module};

pub struct SlabItemExt;

impl WgslExtension for SlabItemExt {
    fn modify_ir(module: &mut Module) {
        let slab_structs: Vec<String> = module
            .items
            .iter()
            .filter_map(|item| match item {
                Item::Struct(s) => {
                    let derives_slab = s.attrs.iter().any(|a| {
                        a.path == "derive" && a.args.iter().any(|arg| arg == "SlabItem")
                    });
                    derives_slab.then(|| s.name.clone())
                }
                _ => None,
            })
            .collect();

        for name in slab_structs {
            module.items.push(slab_read_fn(&name));
            module.items.push(slab_write_fn(&name));
        }
    }
}

fn slab_read_fn(struct_name: &str) -> Item {
    // Build an ir::Item::Fn that reads a slab item at a given offset.
    // ...
}

fn slab_write_fn(struct_name: &str) -> Item {
    // Build an ir::Item::Fn that writes a slab item at a given offset.
    // ...
}
}

Type Substitution

Extensions do not need to handle type parameters themselves. If an injected function references a TypeParam node, the instantiation pass substitutes it with the concrete type from each instantiate() call. Do not attempt to outsmart this by string-replacing type names — operate on IR nodes and let substitution handle generics.

See IR Attributes for how to filter items by attribute, and Examples for the full SlabItemExt and other worked examples.

IR Attributes

The IR preserves Rust attributes on every node where they appear, making them available for extension inspection.

The ir::Attribute Struct

#![allow(unused)]
fn main() {
pub struct Attribute {
    pub path: String,
    pub args: Vec<String>,
}
}

Each attribute is decomposed into a path and a list of argument strings:

Where Attributes Live

Attributes are preserved on:

• Every ir::Item (struct, fn, const, etc.)
• Each ir::Field within a struct
• Each ir::FnArg within a function
• The ir::Module itself

Never Rendered to WGSL

Attributes exist solely for extension inspection. They are never emitted into the final WGSL source. This means an extension can stash metadata on items via attributes and trust that it will not leak into shader output.

Filtering on Attributes

The common pattern is to find items carrying a specific derive:

#![allow(unused)]
fn main() {
let slab_structs: Vec<&ir::Item> = module
    .items
    .iter()
    .filter(|item| matches!(item, ir::Item::Struct(s) if s.attrs.iter().any(|a| {
        a.path == "derive" && a.args.iter().any(|arg| arg == "SlabItem")
    })))
    .collect();
}

Intentional Duplication

Some attribute information also appears in dedicated IR fields such as FnAttrs and InterStageIo. This duplication is intentional: those dedicated fields drive WGSL rendering of entry-point decorators (@vertex, @location, etc.), while the raw Attribute list is preserved verbatim for extensions that want the unfiltered Rust-level view.

Extension Examples

NoopExt

The smallest possible extension, useful as a smoke test that the wiring works:

#![allow(unused)]
fn main() {
use wgsl_rs::WgslExtension;
use wgsl_rs::ir;

pub struct NoopExt;

impl WgslExtension for NoopExt {
    fn modify_ir(_module: &mut ir::Module) {}
}
}

#![allow(unused)]
fn main() {
#[wgsl(extensions = [my_crate::NoopExt])]
mod shader {
    // ...
}
}

SlabItemExt

A derive-driven code generator. It finds structs annotated with #[derive(SlabItem)] and injects slab_read and slab_write helper functions for each.

#![allow(unused)]
fn main() {
use wgsl_rs::WgslExtension;
use wgsl_rs::ir::{Item, Module};

pub struct SlabItemExt;

impl WgslExtension for SlabItemExt {
    fn modify_ir(module: &mut Module) {
        let slab_structs: Vec<String> = module
            .items
            .iter()
            .filter_map(|item| match item {
                Item::Struct(s) => {
                    let derives_slab = s.attrs.iter().any(|a| {
                        a.path == "derive" && a.args.iter().any(|arg| arg == "SlabItem")
                    });
                    derives_slab.then(|| s.name.clone())
                }
                _ => None,
            })
            .collect();

        for name in slab_structs {
            module.items.push(build_slab_read(&name));
            module.items.push(build_slab_write(&name));
        }
    }
}

fn build_slab_read(struct_name: &str) -> Item {
    // Construct an ir::Item::Fn that reads `struct_name` from a slab buffer
    // at a given element index.
    todo!()
}

fn build_slab_write(struct_name: &str) -> Item {
    // Construct an ir::Item::Fn that writes a `struct_name` value into a slab
    // buffer at a given element index.
    todo!()
}
}

wgsl-rs-layout

wgsl-rs-layout is the first real-world extension crate. It is a standalone crate that depends on wgsl-rs for its types and implements WgslExtension to compute WGSL memory layout for Rust structs. It demonstrates that the extension mechanism is sufficient to build a non-trivial, redistributable tool on top of wgsl-rs without forking the transpiler.

See the Memory Layout section for full coverage of wgsl-rs-layout.

Statement Macro Lowering: LowerMyMacro

An extension can claim a custom statement macro via MACROS and replace Stmt::Macro nodes with lowered IR. This example from the trybuild test suite defines my_macro!() and an extension that lowers it to let result: u32 = 42;:

#![allow(unused)]
fn main() {
use wgsl_rs::{ir, wgsl, WgslExtension};

pub struct LowerMyMacro;

impl WgslExtension for LowerMyMacro {
    const MACROS: &'static [&'static str] = &["my_macro"];

    fn modify_ir(module: &mut ir::Module) {
        for item in &mut module.items {
            if let ir::Item::Fn(f) = item {
                lower_in_block(&mut f.block);
            }
        }
    }
}

fn lower_in_block(block: &mut ir::Block) {
    for i in 0..block.stmts.len() {
        if let ir::Stmt::Macro { name, .. } = &block.stmts[i] {
            if name == "my_macro" {
                block.stmts[i] = ir::Stmt::Local(ir::Local {
                    mutable: false,
                    name: "result".to_string(),
                    ty: Some(ir::Type::Scalar(ir::ScalarType::U32)),
                    init: Some(ir::Expr::Lit(ir::Lit::Int {
                        digits: "42".to_string(),
                        suffix: "u32".to_string(),
                    })),
                });
            }
        }
    }
}
}

The shader module uses my_macro!() in statement position, and the extension's modify_ir replaces it before rendering:

#[wgsl(extensions = [super::LowerMyMacro])]
mod ext_macro_shader {
    pub fn main() -> u32 {
        my_macro!();
        42u32
    }
}

After modify_ir runs, the Stmt::Macro is replaced by Stmt::Local, and the rendered WGSL contains no trace of my_macro.

Pitfalls and Constraints

Run Order Is Declaration Order

Extensions run in the order listed in #[wgsl(extensions = [...])]. There is no priority field, no topological sort, no guaranteed order beyond the list. If two extensions conflict, order them explicitly in the attribute.

Type Substitution Happens After modify_ir

modify_ir runs before type instantiation. Do not attempt to outsmart substitution by string-replacing type names or pre-resolving TypeParam nodes. Operate on IR nodes and let the instantiation pass substitute TypeParam nodes automatically into anything you inject.

Attributes Are Not in WGSL Output

ir::Attribute values are preserved on IR nodes for extension inspection only. They are never rendered into WGSL. Do not rely on them appearing in the final shader source, and do not try to emit WGSL by stuffing text into attribute args.

Extension Types Must Be Visible at the Call Site

The paths in #[wgsl(extensions = [path::ExtA, path::ExtB])] must resolve at the location of the wgsl attribute. Import the extension types or use fully-qualified paths.

Non-WgslExtension Types Are Rejected at Compile Time

The macro generates code that calls Ext::modify_ir. If a listed type does not implement WgslExtension, the error surfaces as a compile-time trait bound failure, not a runtime error.

Don't Rely on WGSL_MODULE Being Mutable

Extensions receive &mut ir::Module, not &mut wgsl_rs::Source. The surrounding Source is not mutable from within an extension. Do all your work through the Module reference.

Trait Is Re-exported at the Crate Root

Import the trait from wgsl_rs, not wgsl_rs::extension:

#![allow(unused)]
fn main() {
use wgsl_rs::WgslExtension;
}

Both paths work, but the crate-root re-export is the documented public path.

wgsl-rs-layout Overview

wgsl-rs-layout computes WGSL memory layout for Rust types, implementing the rules in WGSL spec section 14.4.1 ("Alignment and Size"). It answers a single practical question: where do bytes go in the GPU buffer?

Crates

Quick Start

Annotate a struct with #[derive(Layout)] and assert its WGSL layout constants:

use wgsl_rs_layout::{Layout, WgslLayout};

#[derive(Layout)]
struct Particle {
    pos: [f32; 3],
    velocity: [f32; 3],
    charge: f32,
}

fn main() {
    assert_eq!(Particle::SIZE, 32);
    assert_eq!(Particle::ALIGN, 16);
}

SIZE and ALIGN are the WGSL-spec size and alignment of the struct, not the Rust layout. Use FIELDS to find each field's offset:

#![allow(unused)]
fn main() {
for field in Particle::FIELDS {
    println!("{:>10} offset={:<3} size={:<3} align={}",
        field.name, field.offset, field.size, field.alignment);
}
}

What It Is Not

The derive computes the WGSL memory layout only. It does not align the Rust CPU-side representation. For example, a struct with #[repr(C)] and a Vec3f field has Rust alignment 4, but WGSL vec3<f32> has alignment 16. To use a struct as a staging buffer that matches WGSL layout, you must separately ensure the Rust layout matches (for example by padding fields manually) or use a serialization step that writes bytes per FieldLayout.

See Traits, Derive, and Field Layout for the details.

Layout Traits

WgslLayout

WgslLayout is implemented for every built-in WGSL type: scalars, vectors, matrices, arrays, and atomics.

#![allow(unused)]
fn main() {
pub trait WgslLayout {
    const SIZE: usize;
    const ALIGN: usize;
}
}

These constants are the source of truth for all downstream layout computation.

Layout

Layout extends WgslLayout with per-field metadata for composite types:

#![allow(unused)]
fn main() {
pub trait Layout: WgslLayout {
    const FIELDS: &'static [FieldLayout];
}
}

FieldLayout is described in Field Layout.

Generic Structs

Generic structs are supported. Each type parameter receives a T: WgslLayout bound in the generated impl, so SIZE, ALIGN, and FIELDS are computed in terms of the substituted type's constants:

#![allow(unused)]
fn main() {
#[derive(Layout)]
struct Cell<T> {
    value: T,
    next: u32,
}
}

The generated impl is roughly:

#![allow(unused)]
fn main() {
impl<T: WgslLayout> Layout for Cell<T> {
    const FIELDS: &'static [FieldLayout] = &[ /* computed from T::SIZE, T::ALIGN */ ];
}
}

Because the bounds propagate WgslLayout, generic structs compose freely with other layout-annotated types and built-in WGSL types.

The #[derive(Layout)] Macro

#[derive(Layout)] (from wgsl-rs-layout-macros) generates the WgslLayout and Layout impls for a struct by computing offsets, sizes, and alignments per WGSL spec 14.4.1.

Generated Inherent Associated Constants

In addition to the trait impls, the derive emits inherent associated constants on the struct itself:

#![allow(unused)]
fn main() {
impl Particle {
    const __OFFSET_0: usize = /* field 0 offset */;
    const __OFFSET_1: usize = /* field 1 offset */;
    const __SIZE_0: usize   = /* field 0 size   */;
    const __SIZE_1: usize   = /* field 1 size   */;
    const __ALIGN_0: usize  = /* field 0 align  */;
    const __ALIGN_1: usize  = /* field 1 align  */;
    // ... one triple per field
}
}

These are accessible from both the WgslLayout and Layout impls, as well as from user code that wants a specific field's layout without indexing into FIELDS.

Computation

Each field's offset is computed via roundUp(current_offset, field_align), matching the WGSL spec. The struct alignment is the maximum of all field alignments. The struct size is roundUp(last_field_end, struct_align).

Why Inherent Constants

The derive emits inherent constants rather than inline const expressions inside the trait impl because the Rust const evaluator has complexity limits when evaluating deeply nested roundUp expressions inside associated const bodies. Moving the computed values to inherent constants keeps the trait impls thin and avoids hitting those limits on large structs.

The constant values are identical whether accessed via the inherent constants or via FIELDS — they are generated from the same computation in the proc-macro.

Field Layout

FieldLayout

#![allow(unused)]
fn main() {
pub struct FieldLayout {
    pub name: &'static str,
    pub offset: usize,
    pub size: usize,
    pub alignment: usize,
    pub pad_after: usize,
}
}

Each entry in Layout::FIELDS describes one field of the struct.

pad_after Semantics

pad_after is the number of padding bytes between the end of this field's data and the start of the next field (or the end of the struct). When writing a struct to a buffer byte-by-byte:

1. Write the field's size bytes.
2. Write pad_after zero bytes.
3. Proceed to the next field.

The final field's pad_after accounts for struct-end padding so that the total equals SIZE.

RuntimeArray<T>

A runtime-sized array has no statically knowable size:

#![allow(unused)]
fn main() {
assert_eq!(<RuntimeArray<f32> as WgslLayout>::SIZE, 0);
}

SIZE is 0 because the array length is runtime-dependent. Such arrays may only appear as the last field of a storage-buffer struct.

Empty Structs

An empty struct is the identity element for layout:

#![allow(unused)]
fn main() {
assert_eq!(Empty::SIZE, 0);
assert_eq!(Empty::ALIGN, 1);
assert!(Empty::FIELDS.is_empty());
}

The WGSL spec does not define the empty-struct case; wgsl-rs-layout defines ALIGN = 1 so that roundUp(offset, 1) is a no-op and empty structs compose without disturbing surrounding layout.

SVG Byte-Layout Diagrams

Behind the doc-diagrams cargo feature, wgsl-rs-layout can generate self-contained SVG diagrams visualizing a type's byte layout.

Enabling the Feature

[dependencies]
wgsl-rs-layout = { version = "0.1", features = ["doc-diagrams"] }

generate_svg

#![allow(unused)]
fn main() {
use wgsl_rs_layout::diagram::{generate_svg, DiagramConfig};

let svg: String = generate_svg::<Particle>(&DiagramConfig::default());
std::fs::write("particle_layout.svg", svg).unwrap();
}

The returned string is a complete SVG document — no external assets, CSS, or fonts required.

Style

The diagram style mirrors webgpufundamentals.org: each field is a labeled box sized to its size, padding cells are shaded, and rows are laid out left-to-right.

Row width is T::ALIGN bytes. This keeps each row exactly one alignment unit wide, so padding to alignment is visually obvious as a partial row.

All dimensions are sourced from the WgslLayout and Layout trait constants — the diagrams are a pure rendering of the same data exposed by FIELDS.

Wiring into cargo doc

To embed diagrams in rustdoc, generate the SVG files and place them in a directory passed via --resource-files:

cargo doc --resource-files --resources-path ./doc-resources

Reference the image from doc-comments using relative paths:

#![allow(unused)]
fn main() {
/// # Layout
///
/// ![Particle layout](particle_layout.svg)
pub struct Particle { /* ... */ }
}

Examples

This section contains a catalog of example wgsl-rs modules. Each example demonstrates a specific feature of the transpiler, showing the Rust source and the generated WGSL output.

Running examples

The example crate provides two subcommands for inspecting examples:

• cargo run -p example -- show — list all available example names
• cargo run -p example -- source {name} — print the generated WGSL for the named example

Example catalog

Hello Triangle

A "hello world" shader that renders a triangle with a color that changes over time. It demonstrates vertex and fragment entry points, the uniform! macro, glob imports, and builtins like vertex_index.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod hello_triangle {
    //! This is a "hello world" shader that shows a triangle with changing
    //! color. Original source is [here](https://google.github.io/tour-of-wgsl/).

    // Only glob-imports are supported, but hey, imports work!
    use wgsl_rs::std::*;

    // Define a uniform in both Rust and WGSL using the uniform! macro.
    uniform!(group(0), binding(0), FRAME: u32);

    #[vertex]
    pub fn vtx_main(#[builtin(vertex_index)] vertex_index: u32) -> Vec4f {
        const POS: [Vec2f; 3] = [vec2f(0.0, 0.5), vec2f(-0.5, -0.5), vec2f(0.5, -0.5)];

        let position = POS[vertex_index as usize];
        vec4f(position.x, position.y, 0.0, 1.0)
    }

    #[fragment]
    pub fn frag_main() -> Vec4f {
        vec4f(1.0, sin(f32(get!(FRAME)) / 128.0), 0.0, 1.0)
    }
}
}

Generated WGSL

@group(0) @binding(0) var<uniform> FRAME: u32;

@vertex fn vtx_main(@builtin(vertex_index) vertex_index: u32) -> @builtin(position) vec4f {
    const POS: array<vec2f, 3> = array(vec2f(0.0, 0.5), vec2f(-0.5, -0.5), vec2f(0.5, -0.5));
    let position = POS[u32(vertex_index)];
    return vec4f(position.x, position.y, 0.0, 1.0);
}

@fragment fn frag_main() -> @location(0) vec4f {
    return vec4f(1.0, sin(f32(FRAME) / 128.0), 0.0, 1.0);
}

Notes

• The uniform! macro declares a uniform variable available in both Rust and WGSL.
• get!(FRAME) reads the uniform; it is a no-op in WGSL and is stripped during parsing.
• #[vertex] and #[fragment] mark entry points.

Structs

Demonstrates user-defined structs used as fragment shader inputs and outputs, mixing #[location] and #[builtin] attributes, plus #[interpolate].

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod structs {
    use wgsl_rs::std::*;

    // Mixed builtins and user-defined inputs.
    pub struct MyInputs {
        #[location(0)]
        pub x: Vec4<f32>,

        #[builtin(front_facing)]
        pub y: bool,

        #[location(1)]
        #[interpolate(flat)]
        pub z: u32,

        #[location(2)]
        pub other: f32,
    }

    pub struct MyOutputs {
        #[location(0)]
        pub x: f32,

        #[location(1)]
        pub y: Vec4<f32>,
    }

    #[fragment]
    pub fn frag_shader(in1: MyInputs) -> MyOutputs {
        MyOutputs { x: 0.0, y: in1.x }
    }
}
}

Generated WGSL

struct MyInputs {
    @location(0) x: vec4f,
    @builtin(front_facing) y: bool,
    @location(1) @interpolate(flat) z: u32,
    @location(2) other: f32
}

struct MyOutputs {
    @location(0) x: f32,
    @location(1) y: vec4f
}

@fragment fn frag_shader(in1: MyInputs) -> MyOutputs {
    return MyOutputs(0.0, in1.x);
}

Notes

• Struct fields may carry #[location(...)], #[builtin(...)], and #[interpolate(...)] attributes that map directly to WGSL decorations.
• Struct literals in Rust (MyOutputs { x: 0.0, y: in1.x }) become positional constructor calls in WGSL.

Compute Shader

A simple compute shader that demonstrates defining and accessing storage buffers with the storage!, get!, and get_mut! macros, plus the #[derive(Wgsl)] macro for user-defined storage types.

Rust Source

#[wgsl]
pub mod compute_shader {
    //! A simple compute shader that demonstrates defining and accessing storage
    //! buffers.
    //!
    //! Storage buffers are special on the Rust side and require locking,
    //! so they are accessed with the `get!` and `get_mut!` macros, which
    //! do the heavy lifting for you. These macros are a noop in WGSL and are
    //! stripped during parsing.
    use wgsl_rs::std::*;

    // Read-only input buffer
    storage!(group(0), binding(0), INPUT: [f32; 256]);

    #[derive(Wgsl)]
    pub struct Output {
        pub inner: f32,
    }

    // Read-write output buffer
    storage!(group(0), binding(1), read_write, OUTPUT: Output);

    #[compute]
    #[workgroup_size(64)]
    pub fn main(#[builtin(global_invocation_id)] global_id: Vec3u) {
        // Compute the index from global invocation ID
        let idx = global_id.x() as usize;
        // Use the `get!` macro to access the storage
        let input = get!(INPUT)[idx];
        // Use the `get_mut!` macro to access the storage mutably
        get_mut!(OUTPUT).inner = input;
    }
}

Generated WGSL

@group(0) @binding(0) var<storage, read> INPUT: array<f32, 256>;

struct Output {
    inner: f32
}
@group(0) @binding(1) var<storage, read_write> OUTPUT: Output;

@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) global_id: vec3u) {
    let idx = u32(global_id.x);
    let input = INPUT[idx];
    OUTPUT.inner = input;
}

Notes

• storage!(group(0), binding(0), INPUT: [f32; 256]) declares a read-only storage buffer; read_write makes it read-write.
• get! and get_mut! access storage buffers on the Rust side (performing locking). They are no-ops in WGSL and are stripped during parsing.
• #[derive(Wgsl)] makes a struct eligible for use in storage buffers.
• #[compute] and #[workgroup_size(64)] mark the entry point.

Matrix

Demonstrates matrix types (Mat2f, Mat3f, Mat4f) and their constructors (mat2x2f, mat3x3f, mat4x4f) used in module-level constants.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[expect(dead_code, reason = "demonstration")]
pub mod matrix_example {
    //! Demonstrates matrix types and constructors.
    use wgsl_rs::std::*;

    // 4x4 identity matrix constant
    const IDENTITY: Mat4f = mat4x4f(
        vec4f(1.0, 0.0, 0.0, 0.0),
        vec4f(0.0, 1.0, 0.0, 0.0),
        vec4f(0.0, 0.0, 1.0, 0.0),
        vec4f(0.0, 0.0, 0.0, 1.0),
    );

    // 3x3 2D rotation matrix (30 degrees)
    // cos(30°) ≈ 0.866, sin(30°) = 0.5
    const ROTATION_2D: Mat3f = mat3x3f(
        vec3f(0.866, 0.5, 0.0),
        vec3f(-0.5, 0.866, 0.0),
        vec3f(0.0, 0.0, 1.0),
    );

    // 2x2 matrix constant
    const SCALE_2D: Mat2f = mat2x2f(vec2f(2.0, 0.0), vec2f(0.0, 2.0));

    #[vertex]
    pub fn matrix_vertex() -> Vec4f {
        vec4f(0.0, 0.0, 0.0, 1.0)
    }
}
}

Generated WGSL

const IDENTITY: mat4x4f = mat4x4f(vec4f(1.0, 0.0, 0.0, 0.0), vec4f(0.0, 1.0, 0.0, 0.0), vec4f(0.0, 0.0, 1.0, 0.0), vec4f(0.0, 0.0, 0.0, 1.0));
const ROTATION_2D: mat3x3f = mat3x3f(vec3f(0.866, 0.5, 0.0), vec3f(-0.5, 0.866, 0.0), vec3f(0.0, 0.0, 1.0));
const SCALE_2D: mat2x2f = mat2x2f(vec2f(2.0, 0.0), vec2f(0.0, 2.0));

@vertex fn matrix_vertex() -> @builtin(position) vec4f {
    return vec4f(0.0, 0.0, 0.0, 1.0);
}

Notes

• Mat2f/Mat3f/Mat4f alias to the WGSL mat2x2f/mat3x3f/mat4x4f types.
• Module-level const items become WGSL module-scope constants.

Impl

Demonstrates struct impl blocks with associated constants and methods. Methods are called with explicit Type::method(receiver, args) syntax and constants with Type::CONSTANT; both translate to Type_member in WGSL.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod impl_example {
    //! Demonstrates struct impl blocks with explicit receiver syntax.
    //!
    //! Methods and constants are defined in impl blocks.
    //! - Methods are called using `Type::method(receiver, args)` syntax
    //! - Constants are accessed using `Type::CONSTANT` syntax
    //!
    //! Both translate to `Type_member` in WGSL output.
    use wgsl_rs::std::*;

    pub struct Light {
        pub position: Vec3f,
        pub intensity: f32,
    }

    impl Light {
        // Associated constants
        pub const DEFAULT_INTENSITY: f32 = 1.0;
        pub const DEFAULT_RANGE: f32 = 10.0;

        // Create a new light at the given position with the given intensity.
        pub fn new(position: Vec3f, intensity: f32) -> Light {
            Light {
                position,
                intensity,
            }
        }

        // Calculate light attenuation based on distance.
        // Uses inverse-square falloff.
        pub fn attenuate(light: Light, distance: f32) -> f32 {
            light.intensity / (distance * distance)
        }

        // Get the light's position.
        pub fn get_position(light: Light) -> Vec3f {
            light.position
        }
    }

    #[fragment]
    pub fn frag_main() -> Vec4f {
        // Create a light using the explicit receiver syntax
        let light = Light::new(vec3f(0.0, 5.0, 0.0), Light::DEFAULT_INTENSITY);

        // Call a method using explicit path syntax: Type::method(receiver, args)
        let attenuation = Light::attenuate(light, Light::DEFAULT_RANGE / 5.0);

        // Return a color based on attenuation
        vec4f(attenuation, attenuation, attenuation, 1.0)
    }
}
}

Generated WGSL

struct Light {
    position: vec3f,
    intensity: f32
}
const Light__1DEFAULT_INTENSITY: f32 = 1.0;
const Light__1DEFAULT_RANGE: f32 = 10.0;

fn Light_new(position: vec3f, intensity: f32) -> Light {
    return Light(position, intensity);
}

fn Light_attenuate(light: Light, distance: f32) -> f32 {
    return light.intensity / (distance * distance);
}

fn Light__1get_position(light: Light) -> vec3f {
    return light.position;
}

@fragment fn frag_main() -> @location(0) vec4f {
    let light = Light_new(vec3f(0.0, 5.0, 0.0), Light__1DEFAULT_INTENSITY);
    let attenuation = Light_attenuate(light, Light__1DEFAULT_RANGE / 5.0);
    return vec4f(attenuation, attenuation, attenuation, 1.0);
}

Notes

• Associated constants and methods are mangled into Type_member (with _1 separating the type name from the member name when needed for uniqueness).
• Type::method(receiver, args) calls become free functions Type_method(receiver, args) in WGSL.

Enum

Demonstrates limited enum support. #[repr(u32)] enums are translated to a u32 alias and a set of u32 constants, one per variant. match on enum variants becomes a WGSL switch.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod enum_example {
    //! Limited support for enums.
    use wgsl_rs::std::*;

    /// Analytical lighting types.
    #[repr(u32)]
    pub enum LightType {
        Directional = 1337,
        Spot = 420,
        Point = 666,
    }

    #[repr(u32)]
    #[derive(Wgsl)]
    pub enum Holidays {
        // Syntax error!
        // Halloween = -23,
        AprilFoolsDay,
        WaitangiDay,
    }

    storage!(group(0), binding(0), read_write, INPUT: [Holidays; 256]);

    #[compute]
    #[workgroup_size(16)]
    pub fn compute_holidays(#[builtin(global_invocation_id)] global_id: Vec3u) {
        let index = global_id.x();

        let holiday = &mut get_mut!(INPUT)[index as usize];

        #[wgsl_allow(non_literal_match_statement_patterns)]
        match *holiday {
            Holidays::AprilFoolsDay => {
                *holiday = Holidays::WaitangiDay;
            }
            Holidays::WaitangiDay => {
                *holiday = Holidays::AprilFoolsDay;
            }
        }
    }
}
}

Generated WGSL

alias LightType = u32;
const LightType_Directional: u32 = 1337u;
const LightType_Spot: u32 = 420u;
const LightType_Point: u32 = 666u;
alias Holidays = u32;
const Holidays_AprilFoolsDay: u32 = 0u;
const Holidays_WaitangiDay: u32 = 1u;
@group(0) @binding(0) var<storage, read_write> INPUT: array<Holidays, 256>;

@compute @workgroup_size(16) fn compute_holidays(@builtin(global_invocation_id) global_id: vec3u) {
    let index = global_id.x;
    let holiday = &INPUT[u32(index)];
    switch *holiday {
        case Holidays_AprilFoolsDay: {
            *holiday = Holidays_WaitangiDay;
        }
        case Holidays_WaitangiDay: {
            *holiday = Holidays_AprilFoolsDay;
        }
        default: { }
    }
}

Notes

• Enums must be #[repr(u32)]; variants become u32 constants (auto-numbered from 0 if not explicitly assigned).
• #[wgsl_allow(non_literal_match_statement_patterns)] is required when match arms use enum variant paths rather than literal patterns, because WGSL switch cases must be literal — the transpiler substitutes the variant constants.
• #[derive(Wgsl)] enables enum use in storage buffers.

Binary Operators

Demonstrates all supported binary operators in four fragment shaders: arithmetic (+ - * / %), comparison (== != < <= > >=), logical (&& ||), and bitwise (& | ^ << >>).

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod binary_ops_example {
    //! Demonstrates all supported binary operators including:
    //! - Arithmetic: + - * / %
    //! - Comparison: == != < <= > >=
    //! - Logical: && ||
    //! - Bitwise: & | ^ << >>

    use wgsl_rs::std::*;

    // Demonstrates arithmetic operators including remainder.
    #[fragment]
    pub fn test_arithmetic() -> Vec4f {
        let a = vec3f(10.0, 11.0, 12.0);
        let b = 3.0;
        let add = a + b;
        let sub = a - b;
        let mul = a * b;
        let div = a / b;
        let rem = a % b;
        vec4f(add.x(), sub.y(), (mul * div).z(), rem.z())
    }

    // Demonstrates comparison operators.
    // All comparison operators return bool (or vecN<bool> for vectors).
    #[fragment]
    pub fn test_comparison() -> Vec4f {
        let a = 5;
        let b = 10;

        // Comparison operators
        let lt = a < b;
        let le = a <= b;
        let _gt = a > b;
        let ge = a >= b;
        let eq = a == b;
        let ne = a != b;

        // Use the booleans in a calculation
        // In WGSL, we use select() to convert bool to numeric
        let lt_val = select(0.0, 1.0, lt);
        let eq_val = select(0.0, 1.0, eq);
        let ne_val = select(0.0, 1.0, ne);
        let combined = select(0.0, 1.0, le && ge);

        vec4f(lt_val, eq_val, ne_val, combined)
    }

    // Demonstrates logical operators (short-circuit and/or).
    #[fragment]
    pub fn test_logical() -> Vec4f {
        let a = true;
        let b = false;

        // Logical operators (short-circuit evaluation)
        let and_result = a && b;
        let or_result = a || b;
        let complex = (a && b) || (!a && !b);

        let and_val = select(0.0, 1.0, and_result);
        let or_val = select(0.0, 1.0, or_result);
        let complex_val = select(0.0, 1.0, complex);

        vec4f(and_val, or_val, complex_val, 1.0)
    }

    // Demonstrates bitwise operators.
    #[fragment]
    pub fn test_bitwise() -> Vec4f {
        let a: u32 = 0xFF00;
        let b: u32 = 0x0F0F;

        // Bitwise operators
        let and_result = a & b;
        let or_result = a | b;
        let xor_result = a ^ b;

        // Shift operators
        let shl_result = a << 4u32;
        let shr_result = a >> 4u32;

        // Convert to floats for output (normalized)
        let and_f = f32(and_result) / 65535.0;
        let or_f = f32(or_result) / 65535.0;
        let xor_f = f32(xor_result) / 65535.0;
        let shift_f = f32(shl_result ^ shr_result) / 65535.0;

        vec4f(and_f, or_f, xor_f, shift_f)
    }
}
}

Generated WGSL

@fragment fn test_arithmetic() -> @location(0) vec4f {
    let a = vec3f(10.0, 11.0, 12.0);
    let b = 3.0;
    let add = a + b;
    let sub = a - b;
    let mul = a * b;
    let div = a / b;
    let rem = a % b;
    return vec4f(add.x, sub.y, (mul * div).z, rem.z);
}

@fragment fn test_comparison() -> @location(0) vec4f {
    let a = 5;
    let b = 10;
    let lt = a < b;
    let le = a <= b;
    let _gt = a > b;
    let ge = a >= b;
    let eq = a == b;
    let ne = a != b;
    let lt_val = select(0.0, 1.0, lt);
    let eq_val = select(0.0, 1.0, eq);
    let ne_val = select(0.0, 1.0, ne);
    let combined = select(0.0, 1.0, le && ge);
    return vec4f(lt_val, eq_val, ne_val, combined);
}

@fragment fn test_logical() -> @location(0) vec4f {
    let a = true;
    let b = false;
    let and_result = a && b;
    let or_result = a || b;
    let complex = (a && b) || (!a && !b);
    let and_val = select(0.0, 1.0, and_result);
    let or_val = select(0.0, 1.0, or_result);
    let complex_val = select(0.0, 1.0, complex);
    return vec4f(and_val, or_val, complex_val, 1.0);
}

@fragment fn test_bitwise() -> @location(0) vec4f {
    let a: u32 = 65280;
    let b: u32 = 3855;
    let and_result = a & b;
    let or_result = a | b;
    let xor_result = a ^ b;
    let shl_result = a << 4u;
    let shr_result = a >> 4u;
    let and_f = f32(and_result) / 65535.0;
    let or_f = f32(or_result) / 65535.0;
    let xor_f = f32(xor_result) / 65535.0;
    let shift_f = f32(shl_result ^ shr_result) / 65535.0;
    return vec4f(and_f, or_f, xor_f, shift_f);
}

Notes

• Hex literals like 0xFF00 are emitted as decimal WGSL literals.
• Short-circuit &&/|| are preserved in the raw output; the naga-validated form lowers them to if/else because WGSL has no short-circuit logical operators.

For Loop

Demonstrates for-loop support with range expressions: exclusive (0..n), inclusive (0..=n), and literal bounds. Variable bounds require #[wgsl_allow(non_literal_loop_bounds)].

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod for_loop_example {
    //! Demonstrates for-loop support with range expressions.
    //! - Exclusive ranges: `for i in 0..10 { ... }`
    //! - Inclusive ranges: `for i in 0..=9 { ... }`
    //! - Variable bounds: `for i in start..end { ... }` (requires
    //!   `#[wgsl_allow]`)
    use wgsl_rs::std::*;

    // Sum values from 0 to n-1 using exclusive range.
    // Uses #[wgsl_allow] on for-loop because `n` is a variable bound.
    pub fn sum_exclusive(n: i32) -> i32 {
        let mut total = 0;
        #[wgsl_allow(non_literal_loop_bounds)]
        for i in 0..n {
            total += i;
        }
        total
    }

    // Sum values from 0 to n (inclusive) using inclusive range.
    // Uses #[wgsl_allow] on for-loop because `n` is a variable bound.
    pub fn sum_inclusive(n: i32) -> i32 {
        let mut total = 0;
        #[wgsl_allow(non_literal_loop_bounds)]
        for i in 0..=n {
            total += i;
        }
        total
    }

    // Compute dot product of two arrays using for-loop.
    // No #[wgsl_allow] needed because bounds are literals.
    pub fn dot_product(a: [f32; 4], b: [f32; 4]) -> f32 {
        let mut result = 0.0;
        for i in 0..4 {
            result += a[i as usize] * b[i as usize];
        }
        result
    }

    // Nested for-loops: initialize a 2D-like structure.
    // No #[wgsl_allow] needed because bounds are literals.
    pub fn nested_loops() -> i32 {
        let mut sum = 0;
        for i in 0..3 {
            for j in 0..4 {
                sum += i * 4 + j;
            }
        }
        sum
    }

    #[fragment]
    pub fn for_loop_fragment() -> Vec4f {
        // Test sum_exclusive: sum of 0..10 = 0+1+2+...+9 = 45
        let exclusive_sum = sum_exclusive(10);

        // Test sum_inclusive: sum of 0..=9 = 0+1+2+...+9 = 45
        let inclusive_sum = sum_inclusive(9);

        // Test dot_product
        let a: [f32; 4] = [1.0, 2.0, 3.0, 4.0];
        let b: [f32; 4] = [1.0, 1.0, 1.0, 1.0];
        let dot = dot_product(a, b); // 1+2+3+4 = 10

        // Test nested loops: sum of (i*4+j) for i in 0..3, j in 0..4
        // = (0,1,2,3) + (4,5,6,7) + (8,9,10,11) = 6 + 22 + 38 = 66
        let nested = nested_loops();

        vec4f(
            f32(exclusive_sum) / 100.0,
            f32(inclusive_sum) / 100.0,
            dot / 10.0,
            f32(nested) / 100.0,
        )
    }
}
}

Generated WGSL

fn sum_exclusive(n: i32) -> i32 {
    var total = 0;
    for (var i = 0; i < n; i++) {
        total += i;
    }
    return total;
}

fn sum_inclusive(n: i32) -> i32 {
    var total = 0;
    for (var i = 0; i <= n; i++) {
        total += i;
    }
    return total;
}

fn dot_product(a: array<f32, 4>, b: array<f32, 4>) -> f32 {
    var result = 0.0;
    for (var i = 0; i < 4; i++) {
        result += a[u32(i)] * b[u32(i)];
    }
    return result;
}

fn nested_loops() -> i32 {
    var sum = 0;
    for (var i = 0; i < 3; i++) {
        for (var j = 0; j < 4; j++) {
            sum += i * 4 + j;
        }
    }
    return sum;
}

@fragment fn for_loop_fragment() -> @location(0) vec4f {
    let exclusive_sum = sum_exclusive(10);
    let inclusive_sum = sum_inclusive(9);
    let a: array<f32, 4> = array(1.0, 2.0, 3.0, 4.0);
    let b: array<f32, 4> = array(1.0, 1.0, 1.0, 1.0);
    let dot = dot_product(a, b);
    let nested = nested_loops();
    return vec4f(f32(exclusive_sum) / 100.0, f32(inclusive_sum) / 100.0, dot / 10.0, f32(nested) / 100.0);
}

Notes

• Rust for i in 0..n becomes a WGSL C-style for (var i = 0; i < n; i++).
• Inclusive ranges (0..=n) translate to i <= n.
• #[wgsl_allow(non_literal_loop_bounds)] is required when loop bounds are not literal constants, because WGSL's for requires constant bounds — the transpiler emits them anyway and naga lowers the loop to a loop with a break condition.

While & Loop

Demonstrates while loops and WGSL loop (infinite loop) statements. Includes continue, compound conditions, and nested loops.

while_loop_example

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[allow(dead_code, unused_assignments)]
pub mod while_loop_example {
    use wgsl_rs::std::*;

    #[fragment]
    pub fn test_simple_while() -> Vec4f {
        let mut i = 0;
        let mut sum = 0.0;

        while i < 10 {
            i += 1;
            // Skip even numbers using continue
            if i % 2 == 0 {
                continue;
            }
            sum += f32(i);
        }

        vec4f(sum / 10.0, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_while_with_condition() -> Vec4f {
        let mut value = 1.0;
        let mut iterations = 0;

        while value < 100.0 && iterations < 20 {
            value *= 1.5;
            iterations += 1;
        }

        vec4f(value / 100.0, f32(iterations) / 20.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_nested_while() -> Vec4f {
        let mut i = 0;
        let mut j = 0;
        let mut count = 0;

        while i < 5 {
            j = 0;
            while j < 5 {
                count += 1;
                j += 1;
            }
            i += 1;
        }

        vec4f(f32(count) / 25.0, 0.0, 0.0, 1.0)
    }
}
}

Generated WGSL

@fragment fn test_simple_while() -> @location(0) vec4f {
    var i = 0;
    var sum = 0.0;
    while i < 10 {
        i += 1;
        if i % 2 == 0 {
            continue;
        }
        sum += f32(i);
    }
    return vec4f(sum / 10.0, 0.0, 0.0, 1.0);
}

@fragment fn test_while_with_condition() -> @location(0) vec4f {
    var value = 1.0;
    var iterations = 0;
    while value < 100.0 && iterations < 20 {
        value *= 1.5;
        iterations += 1;
    }
    return vec4f(value / 100.0, f32(iterations) / 20.0, 0.0, 1.0);
}

@fragment fn test_nested_while() -> @location(0) vec4f {
    var i = 0;
    var j = 0;
    var count = 0;
    while i < 5 {
        j = 0;
        while j < 5 {
            count += 1;
            j += 1;
        }
        i += 1;
    }
    return vec4f(f32(count) / 25.0, 0.0, 0.0, 1.0);
}

loop_example

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[allow(dead_code, unused_assignments)]
pub mod loop_example {
    //! Demonstrates WGSL loop statements (infinite loops).
    //! Note: These are demonstration examples only.

    use wgsl_rs::std::*;

    #[fragment]
    pub fn test_simple_loop() -> Vec4f {
        let mut counter: u32 = 0;
        let mut sum: f32 = 0.0;

        loop {
            sum += f32(counter);
            counter += 1;
            if counter >= 10 {
                break;
            }
        }

        vec4f(sum / 10.0, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_nested_loop() -> Vec4f {
        let mut i: u32 = 0;
        let mut j: u32 = 0;
        let mut result: f32 = 0.0;

        loop {
            j = 0;
            loop {
                result += 1.0;
                j += 1;
                if j >= 5 {
                    break;
                }
            }
            i += 1;
            if i >= 5 {
                break;
            }
        }

        vec4f(result / 25.0, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_loop_with_operations() -> Vec4f {
        let mut value: f32 = 1.0;
        let mut iterations: u32 = 0;

        loop {
            value *= 1.5;
            iterations += 1;
            if value >= 100.0 || iterations >= 20 {
                break;
            }
        }

        vec4f(value / 100.0, f32(iterations) / 20.0, 0.0, 1.0)
    }
}
}

Generated WGSL

@fragment fn test_simple_loop() -> @location(0) vec4f {
    var counter: u32 = 0;
    var sum: f32 = 0.0;
    loop {
        sum += f32(counter);
        counter += 1;
        if counter >= 10 {
            break;
        }
    }
    return vec4f(sum / 10.0, 0.0, 0.0, 1.0);
}

@fragment fn test_nested_loop() -> @location(0) vec4f {
    var i: u32 = 0;
    var j: u32 = 0;
    var result: f32 = 0.0;
    loop {
        j = 0;
        loop {
            result += 1.0;
            j += 1;
            if j >= 5 {
                break;
            }
        }
        i += 1;
        if i >= 5 {
            break;
        }
    }
    return vec4f(result / 25.0, 0.0, 0.0, 1.0);
}

@fragment fn test_loop_with_operations() -> @location(0) vec4f {
    var value: f32 = 1.0;
    var iterations: u32 = 0;
    loop {
        value *= 1.5;
        iterations += 1;
        if value >= 100.0 || iterations >= 20 {
            break;
        }
    }
    return vec4f(value / 100.0, f32(iterations) / 20.0, 0.0, 1.0);
}

Notes

• Rust while maps directly to WGSL while.
• Rust loop { ... } maps to WGSL loop { ... } (infinite loop with explicit break).
• Compound conditions and nested loops are supported.

Control Flow

Demonstrates control-flow constructs: if/else if/else, break (including nested), explicit return statements, and match/switch with literal, or-patterns, and const patterns.

if_example

Demonstrates if statements: simple if, if/else, if/else if/else chains, and nested if.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[allow(dead_code)]
pub mod if_example {
    //! Demonstrates if statements including:
    //! - Simple if
    //! - if/else
    //! - if/else if/else chains
    //! - Nested if statements

    use wgsl_rs::std::*;

    #[fragment]
    pub fn test_simple_if() -> Vec4f {
        let mut result = 0.0;
        if true {
            result = 1.0;
        }
        vec4f(result, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_if_else() -> Vec4f {
        let mut result = 0.0;
        if result < 1.0 {
            result = 1.0;
        } else {
            result = 2.0;
        }
        vec4f(result, 0.0, 0.0, 1.0)
    }

    #[fragment]
    #[allow(unused_assignments)]
    pub fn test_if_else_if_else() -> Vec4f {
        let x = 5;
        let mut result = 0.0;
        if x < 3 {
            result = 1.0;
        } else if x < 7 {
            result = 2.0;
        } else {
            result = 3.0;
        }
        vec4f(result, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_nested_if() -> Vec4f {
        let x = 5;
        let y = 10;
        let mut result = 0.0;
        if x > 0 {
            if y > 5 {
                result = 1.0;
            } else {
                result = 0.5;
            }
        }
        vec4f(result, 0.0, 0.0, 1.0)
    }
}
}

Generated WGSL

@fragment fn test_simple_if() -> @location(0) vec4f {
    var result = 0.0;
    if true {
        result = 1.0;
    }
    return vec4f(result, 0.0, 0.0, 1.0);
}

@fragment fn test_if_else() -> @location(0) vec4f {
    var result = 0.0;
    if result < 1.0 {
        result = 1.0;
    } else {
        result = 2.0;
    }
    return vec4f(result, 0.0, 0.0, 1.0);
}

@fragment fn test_if_else_if_else() -> @location(0) vec4f {
    let x = 5;
    var result = 0.0;
    if x < 3 {
        result = 1.0;
    } else if x < 7 {
        result = 2.0;
    } else {
        result = 3.0;
    }
    return vec4f(result, 0.0, 0.0, 1.0);
}

@fragment fn test_nested_if() -> @location(0) vec4f {
    let x = 5;
    let y = 10;
    var result = 0.0;
    if x > 0 {
        if y > 5 {
            result = 1.0;
        } else {
            result = 0.5;
        }
    }
    return vec4f(result, 0.0, 0.0, 1.0);
}

break_example

Demonstrates break statements inside while loops, including conditional breaks and nested break.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[allow(dead_code, unused_assignments)]
pub mod break_example {
    use wgsl_rs::std::*;

    #[fragment]
    pub fn test_break_in_while() -> Vec4f {
        let mut i = 0;
        let mut sum = 0.0;

        while i < 100 {
            if i >= 10 {
                break;
            }
            sum += f32(i);
            i += 1;
        }

        vec4f(sum / 100.0, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_break_with_condition() -> Vec4f {
        let mut value = 1.0;
        let mut iterations = 0;

        while iterations < 100 {
            value *= 1.1;
            iterations += 1;

            if value > 50.0 {
                break;
            }
        }

        vec4f(value / 100.0, f32(iterations) / 100.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_nested_break() -> Vec4f {
        let mut i = 0;
        let mut j = 0;
        let mut found = 0;

        while i < 10 {
            j = 0;
            while j < 10 {
                if i * 10 + j == 55 {
                    found = 1;
                    break;
                }
                j += 1;
            }
            if found == 1 {
                break;
            }
            i += 1;
        }

        vec4f(f32(i) / 10.0, f32(j) / 10.0, f32(found), 1.0)
    }
}
}

Generated WGSL

@fragment fn test_break_in_while() -> @location(0) vec4f {
    var i = 0;
    var sum = 0.0;
    while i < 100 {
        if i >= 10 {
            break;
        }
        sum += f32(i);
        i += 1;
    }
    return vec4f(sum / 100.0, 0.0, 0.0, 1.0);
}

@fragment fn test_break_with_condition() -> @location(0) vec4f {
    var value = 1.0;
    var iterations = 0;
    while iterations < 100 {
        value *= 1.1;
        iterations += 1;
        if value > 50.0 {
            break;
        }
    }
    return vec4f(value / 100.0, f32(iterations) / 100.0, 0.0, 1.0);
}

@fragment fn test_nested_break() -> @location(0) vec4f {
    var i = 0;
    var j = 0;
    var found = 0;
    while i < 10 {
        j = 0;
        while j < 10 {
            if i * 10 + j == 55 {
                found = 1;
                break;
            }
            j += 1;
        }
        if found == 1 {
            break;
        }
        i += 1;
    }
    return vec4f(f32(i) / 10.0, f32(j) / 10.0, f32(found), 1.0);
}

return_example

Demonstrates explicit return statements: early returns, return with expressions, and mixed explicit/implicit returns.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[allow(dead_code, clippy::needless_return, clippy::mixed_attributes_style)]
pub mod return_example {
    //! Demonstrates explicit return statements including:
    //! - Early returns from functions
    //! - Return with expressions
    //! - Mixed explicit and implicit returns
    use wgsl_rs::std::*;

    // Helper function with early return
    pub fn clamp_positive(x: f32) -> f32 {
        if x < 0.0 {
            return 0.0;
        }
        return x;
    }

    // Function with multiple return paths
    pub fn sign(x: f32) -> f32 {
        if x > 0.0 {
            return 1.0;
        }
        if x < 0.0 {
            return -1.0;
        }
        return 0.0;
    }

    // Mixed explicit and implicit return
    pub fn abs_or_zero(x: f32, threshold: f32) -> f32 {
        if abs(x) < threshold {
            return 0.0;
        }
        abs(x)
    }

    #[fragment]
    pub fn test_explicit_returns() -> Vec4f {
        let pos = clamp_positive(-5.0); // 0.0
        let neg = clamp_positive(3.0); // 3.0
        let s1 = sign(5.0); // 1.0
        let s2 = sign(-2.0); // -1.0
        let a1 = abs_or_zero(0.1, 0.5); // 0.0
        let a2 = abs_or_zero(2.0, 0.5); // 2.0

        vec4f(pos + neg / 10.0, s1 + s2, a1 + a2 / 10.0, 1.0)
    }
}
}

Generated WGSL

fn clamp_positive(x: f32) -> f32 {
    if x < 0.0 {
        return 0.0;
    }
    return x;
}

fn sign(x: f32) -> f32 {
    if x > 0.0 {
        return 1.0;
    }
    if x < 0.0 {
        return -1.0;
    }
    return 0.0;
}

fn abs_or_zero(x: f32, threshold: f32) -> f32 {
    if abs(x) < threshold {
        return 0.0;
    }
    return abs(x);
}

@fragment fn test_explicit_returns() -> @location(0) vec4f {
    let pos = clamp_positive(-5.0);
    let neg = clamp_positive(3.0);
    let s1 = sign(5.0);
    let s2 = sign(-2.0);
    let a1 = abs_or_zero(0.1, 0.5);
    let a2 = abs_or_zero(2.0, 0.5);
    return vec4f(pos + neg / 10.0, s1 + s2, a1 + a2 / 10.0, 1.0);
}

switch_example

Demonstrates match/switch support: simple integer matching, or-patterns (multiple cases), default cases, auto-generated default when missing, and const patterns (with warning suppression).

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[allow(dead_code, unused_assignments)]
pub mod switch_example {
    //! Demonstrates switch/match statement support including:
    //! - Simple integer matching
    //! - Or-patterns (multiple cases)
    //! - Default cases
    //! - Auto-generated default when missing
    //! - Const patterns (with warning suppression)

    use wgsl_rs::std::*;

    const LOW: i32 = 0;
    const MID: i32 = 1;
    const HIGH: i32 = 2;

    #[fragment]
    pub fn test_simple_switch() -> Vec4f {
        let x: i32 = 2;
        let mut result = 0.0;
        match x {
            0 => {
                result = 0.0;
            }
            1 => {
                result = 0.25;
            }
            2 => {
                result = 0.5;
            }
            _ => {
                result = 1.0;
            }
        }
        vec4f(result, 0.0, 0.0, 1.0)
    }

    #[fragment]
    #[allow(clippy::manual_range_patterns)]
    pub fn test_or_patterns() -> Vec4f {
        let x: u32 = 5;
        let mut result = 0.0;
        match x {
            1 | 2 | 3 => {
                result = 0.25;
            }
            4 | 5 | 6 => {
                result = 0.5;
            }
            _ => {
                result = 1.0;
            }
        }
        vec4f(result, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_missing_default() -> Vec4f {
        let x: i32 = 1;
        let mut result = 0.0;
        // No default arm - WGSL will get auto-generated `default: {}`
        // But Rust requires exhaustive matching, so we use a catch-all underscore
        // that will be optimized out in the test below
        match x {
            0 => {
                result = 0.0;
            }
            1 => {
                result = 1.0;
            }
            _ => {}
        }
        vec4f(result, 0.0, 0.0, 1.0)
    }

    #[fragment]
    pub fn test_const_patterns() -> Vec4f {
        let level: i32 = 1;
        let mut brightness = 0.0;
        #[wgsl_allow(non_literal_match_statement_patterns)]
        match level {
            LOW => {
                brightness = 0.0;
            }
            MID => {
                brightness = 0.5;
            }
            HIGH => {
                brightness = 1.0;
            }
            _ => {
                brightness = 0.0;
            }
        }
        vec4f(brightness, 0.0, 0.0, 1.0)
    }
}
}

Generated WGSL

const LOW: i32 = 0;
const MID: i32 = 1;
const HIGH: i32 = 2;

@fragment fn test_simple_switch() -> @location(0) vec4f {
    let x: i32 = 2;
    var result = 0.0;
    switch x {
        case 0: {
            result = 0.0;
        }
        case 1: {
            result = 0.25;
        }
        case 2: {
            result = 0.5;
        }
        default: {
            result = 1.0;
        }
    }
    return vec4f(result, 0.0, 0.0, 1.0);
}

@fragment fn test_or_patterns() -> @location(0) vec4f {
    let x: u32 = 5;
    var result = 0.0;
    switch x {
        case 1, 2, 3: {
            result = 0.25;
        }
        case 4, 5, 6: {
            result = 0.5;
        }
        default: {
            result = 1.0;
        }
    }
    return vec4f(result, 0.0, 0.0, 1.0);
}

@fragment fn test_missing_default() -> @location(0) vec4f {
    let x: i32 = 1;
    var result = 0.0;
    switch x {
        case 0: {
            result = 0.0;
        }
        case 1: {
            result = 1.0;
        }
        default: {
        }
    }
    return vec4f(result, 0.0, 0.0, 1.0);
}

@fragment fn test_const_patterns() -> @location(0) vec4f {
    let level: i32 = 1;
    var brightness = 0.0;
    switch level {
        case LOW: {
            brightness = 0.0;
        }
        case MID: {
            brightness = 0.5;
        }
        case HIGH: {
            brightness = 1.0;
        }
        default: {
            brightness = 0.0;
        }
    }
    return vec4f(brightness, 0.0, 0.0, 1.0);
}

Notes

• Rust match becomes WGSL switch. The _ arm maps to default.
• Or-patterns (1 | 2 | 3) become comma-separated case selectors (case 1, 2, 3:).
• #[wgsl_allow(non_literal_match_statement_patterns)] is required when match arms reference const values rather than literals, since WGSL cases must be literals — the transpiler emits the const names directly (naga substitutes them).
• A missing default arm in Rust (with a _ => {} catch-all) emits an empty default: {} in WGSL.

Runtime Array

Demonstrates runtime-sized arrays (RuntimeArray<T>) in storage buffers. Runtime arrays transpile to array<T> in WGSL (no size) and must be the last field of a struct in a storage buffer.

Rust Source

#[wgsl]
#[allow(dead_code)]
pub mod runtime_array_example {
    //! Demonstrates runtime-sized arrays (`RuntimeArray<T>`).
    //!
    //! Runtime-sized arrays transpile to `array<T>` in WGSL (no size
    //! parameter). They can only be used in storage buffers, typically as
    //! the last field of a struct.
    use wgsl_rs::std::*;

    #[derive(Wgsl)]
    pub struct Particle {
        pub position: Vec3f,
        pub velocity: Vec3f,
    }

    #[derive(Wgsl)]
    pub struct ParticleSystem {
        pub count: u32,
        pub particles: RuntimeArray<Particle>,
    }

    storage!(group(0), binding(0), read_write, PARTICLES: ParticleSystem);

    #[compute]
    #[workgroup_size(16, 16, 1)]
    pub fn main(#[builtin(global_invocation_id)] global_id: Vec3u) {
        let num_particles = array_length(&get!(PARTICLES).particles);
        let index = global_id.y() * 16 + global_id.x();
        if num_particles < index {
            let velocity = get!(PARTICLES).particles[index].velocity;
            let position = &mut get_mut!(PARTICLES).particles[index].position;
            *position = *position + velocity;
        }
    }
}

Generated WGSL

struct Particle {
    position: vec3f,
    velocity: vec3f
}

struct ParticleSystem {
    count: u32,
    particles: array<Particle>
}
@group(0) @binding(0) var<storage, read_write> PARTICLES: ParticleSystem;

@compute @workgroup_size(16, 16, 1) fn main(@builtin(global_invocation_id) global_id: vec3u) {
    let num_particles = arrayLength(&PARTICLES.particles);
    let index = global_id.y * 16 + global_id.x;
    if num_particles < index {
        let velocity = PARTICLES.particles[index].velocity;
        let position = &PARTICLES.particles[index].position;
        *position = *position + velocity;
    }
}

Notes

• RuntimeArray<T> transpiles to array<T> (unsized) in WGSL.
• array_length(&...) maps to the WGSL arrayLength builtin.
• Runtime arrays may only appear in storage buffers, typically as the last field.
• The bounds check in the example (if num_particles < index) is intentionally verbatim from the source — note the condition is reversed from what you'd typically want (index < num_particles). This is a known quirk of the example, preserved for roundtrip-test compatibility.

Pointer

Demonstrates pointer types in function parameters via the ptr! macro. Pointers translate to WGSL ptr<function, T> parameters.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
#[allow(dead_code, clippy::manual_swap, clippy::assign_op_pattern)]
pub mod ptr_example {
    //! Demonstrates pointer types in function parameters.
    use wgsl_rs::std::*;

    // Increment a value through a pointer.
    pub fn increment(p: ptr!(function, i32)) {
        *p += 1;
    }

    // Swap two values through pointers.
    // Note: We use manual swap because std::mem::swap is not available in WGSL.
    pub fn swap(a: ptr!(function, f32), b: ptr!(function, f32)) {
        let tmp = *a;
        *a = *b;
        *b = tmp;
    }

    // Double a value in-place through a pointer.
    // Note: We use *p = *p * 2.0 instead of *p *= 2.0 to demonstrate dereference.
    pub fn double_value(p: ptr!(function, f32)) {
        *p = *p * 2.0;
    }

    #[fragment]
    pub fn test_ptr() -> Vec4f {
        let mut x: i32 = 5;
        increment(&mut x);
        // x is now 6

        let mut a: f32 = 1.0;
        let mut b: f32 = 2.0;
        swap(&mut a, &mut b);
        // a is now 2.0, b is now 1.0

        let mut c: f32 = 3.0;
        double_value(&mut c);
        // c is now 6.0

        vec4f(f32(x), a, b, c / 10.0)
    }
}
}

Generated WGSL

fn increment(p: ptr<function, i32>) {
    *p += 1;
}

fn swap(a: ptr<function, f32>, b: ptr<function, f32>) {
    let tmp = *a;
    *a = *b;
    *b = tmp;
}

fn double_value(p: ptr<function, f32>) {
    *p = *p * 2.0;
}

@fragment fn test_ptr() -> @location(0) vec4f {
    var x: i32 = 5;
    increment(&x);
    var a: f32 = 1.0;
    var b: f32 = 2.0;
    swap(&a, &b);
    var c: f32 = 3.0;
    double_value(&c);
    return vec4f(f32(x), a, b, c / 10.0);
}

Notes

• ptr!(function, i32) expands to the WGSL pointer type ptr<function, i32>.
• &mut x at the call site becomes &x in WGSL (WGSL pointers do not distinguish mutability in the reference syntax).

Atomics

Demonstrates atomic types and workgroup variables. Atomic types provide thread-safe operations for concurrent access in compute shaders and may only hold i32 or u32. Workgroup variables are shared between all invocations in a workgroup.

Rust Source

#[wgsl]
pub mod atomic_example {
    //! Demonstrates atomic types and workgroup variables.
    //!
    //! Atomic types provide thread-safe operations for concurrent access in
    //! compute shaders. They can only hold `i32` or `u32` values.
    //!
    //! Workgroup variables are shared between all invocations in a workgroup
    //! and can only be used in compute shaders.
    use wgsl_rs::std::*;

    // Workgroup variable with atomic counter - shared between all invocations
    workgroup!(COUNTER: Atomic<u32>);

    // Workgroup variable with atomic flags
    workgroup!(FLAGS: Atomic<i32>);

    #[compute]
    #[workgroup_size(64)]
    pub fn main(#[builtin(local_invocation_index)] local_idx: u32) {
        // Each invocation can access the shared atomic counter
        // Note: atomicLoad/atomicStore builtins will be added in a future update
        // For now, this demonstrates the type parsing and code generation
        let _idx = local_idx;
    }
}

Generated WGSL

var<workgroup> COUNTER: atomic<u32>;
var<workgroup> FLAGS: atomic<i32>;

@compute @workgroup_size(64) fn main(@builtin(local_invocation_index) local_idx: u32) {
    let _idx = local_idx;
}

Notes

• Atomic<T> maps to WGSL atomic<T>; T must be i32 or u32.
• workgroup!(NAME: T) declares a var<workgroup> variable. Workgroup variables are only valid in compute shaders.

Texture

Demonstrates textures, samplers, and the textureSample builtin. Covers the texture! and sampler! macros and a fragment shader that samples a 2D texture.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod texture_example {
    //! Demonstrates using textures and texture builtin functions.
    //!
    //! WGSL provides several categories of texture operations:
    //! - **Query functions**: `textureDimensions`, `textureNumLayers`, etc.
    //! - **Load functions**: `textureLoad` - direct texel access without
    //!   filtering
    //! - **Sample functions**: `textureSample` - filtered sampling with a
    //!   sampler
    //! - **Depth comparison**: `textureSampleCompare` - for shadow mapping
    use wgsl_rs::std::*;

    // A 2D texture for color/albedo
    texture!(group(0), binding(0), DIFFUSE_TEX: Texture2D<f32>);
    // A sampler for filtering the texture
    sampler!(group(0), binding(1), TEX_SAMPLER: Sampler);

    // Fragment input with texture coordinates
    pub struct FragmentInput {
        #[location(0)]
        pub uv: Vec2f,
    }

    // Output struct
    pub struct FragmentOutput {
        #[location(0)]
        pub color: Vec4f,
    }

    // Main fragment shader demonstrating texture operations.
    #[fragment]
    pub fn frag_main(input: FragmentInput) -> FragmentOutput {
        // Sample the diffuse texture
        let albedo = texture_sample(DIFFUSE_TEX, TEX_SAMPLER, input.uv);

        FragmentOutput { color: albedo }
    }
}
}

Generated WGSL

@group(0) @binding(0) var DIFFUSE_TEX: texture_2d<f32>;
@group(0) @binding(1) var TEX_SAMPLER: sampler;

struct FragmentInput {
    @location(0) uv: vec2f
}

struct FragmentOutput {
    @location(0) color: vec4f
}

@fragment fn frag_main(input: FragmentInput) -> FragmentOutput {
    let albedo = textureSample(DIFFUSE_TEX, TEX_SAMPLER, input.uv);
    return FragmentOutput(albedo);
}

Notes

• texture!(group(0), binding(0), DIFFUSE_TEX: Texture2D<f32>) declares a texture_2d<f32> binding.
• sampler!(...) declares a sampler binding.
• texture_sample(...) maps to the WGSL textureSample builtin.

Bitcast

Demonstrates bitcast builtin functions for reinterpreting the bit pattern of a value as another type. In wgsl-rs, each target type has a dedicated function (e.g. bitcast_f32, bitcast_u32, bitcast_vec4i).

Rust Source

#[wgsl]
#[expect(dead_code, reason = "demonstration")]
pub mod bitcast_example {
    //! Demonstrates using `bitcast` to reinterpret the bits of a value as
    //! another type.
    //!
    //! WGSL `bitcast<T>(e)` reinterprets the bit pattern of `e` as type `T`
    //! without changing any bits. This is useful for packing/unpacking data,
    //! interpreting raw buffer contents, and working with IEEE 754
    //! representations.
    //!
    //! In `wgsl-rs`, each target type has a dedicated function:
    //!   - `bitcast_f32(e)` → `bitcast<f32>(e)`
    //!   - `bitcast_u32(e)` → `bitcast<u32>(e)`
    //!   - `bitcast_i32(e)` → `bitcast<i32>(e)`
    //!   - `bitcast_vec2f(e)` → `bitcast<vec2<f32>>(e)`, etc.
    use wgsl_rs::std::*;

    // Input: raw u32 data representing packed floats
    storage!(group(0), binding(0), INPUT: [u32; 256]);

    // Output: reinterpreted as floats
    storage!(group(0), binding(1), read_write, OUTPUT: [f32; 256]);

    // Reinterpret a u32 bit pattern as an f32 value.
    pub fn reinterpret_as_float(bits: u32) -> f32 {
        bitcast_f32(bits)
    }

    // Reinterpret an f32 value as its u32 bit pattern.
    pub fn float_to_bits(value: f32) -> u32 {
        bitcast_u32(value)
    }

    // Reinterpret a u32 vector as an i32 vector.
    pub fn reinterpret_vec_as_signed(v: Vec4u) -> Vec4i {
        bitcast_vec4i(v)
    }

    #[compute]
    #[workgroup_size(64)]
    pub fn main(#[builtin(global_invocation_id)] global_id: Vec3u) {
        let idx = global_id.x() as usize;
        // Read raw u32 bits from input and reinterpret as f32
        let raw_bits = get!(INPUT)[idx];
        get_mut!(OUTPUT)[idx] = bitcast_f32(raw_bits);
    }
}

Generated WGSL

@group(0) @binding(0) var<storage, read> INPUT: array<u32, 256>;
@group(0) @binding(1) var<storage, read_write> OUTPUT: array<f32, 256>;

fn reinterpret_as_float(bits: u32) -> f32 {
    return bitcast<f32>(bits);
}

fn float_to_bits(value: f32) -> u32 {
    return bitcast<u32>(value);
}

fn reinterpret_vec_as_signed(v: vec4u) -> vec4i {
    return bitcast<vec4<i32>>(v);
}

@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) global_id: vec3u) {
    let idx = u32(global_id.x);
    let raw_bits = INPUT[idx];
    OUTPUT[idx] = bitcast<f32>(raw_bits);
}

Notes

• bitcast_<ty>(e) maps to bitcast<<ty>>(e) in WGSL.
• Useful for packing/unpacking data, interpreting raw buffer contents, and working with IEEE 754 representations.

Packing

Demonstrates WGSL data packing and unpacking builtin functions, which convert between vector types and packed u32 representations. Useful for vertex attribute compression and storage optimization.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod packing_example {
    //! Demonstrates WGSL data packing and unpacking builtin functions.
    //!
    //! These functions convert between vector types and packed `u32`
    //! representations, useful for vertex attribute compression and storage
    //! optimization.
    use wgsl_rs::std::*;

    pub fn demo_pack4x8snorm(v: Vec4f) -> u32 {
        pack4x8snorm(v)
    }

    pub fn demo_unpack4x8snorm(e: u32) -> Vec4f {
        unpack4x8snorm(e)
    }

    pub fn demo_pack4x8unorm(v: Vec4f) -> u32 {
        pack4x8unorm(v)
    }

    pub fn demo_unpack4x8unorm(e: u32) -> Vec4f {
        unpack4x8unorm(e)
    }

    pub fn demo_pack2x16snorm(v: Vec2f) -> u32 {
        pack2x16snorm(v)
    }

    pub fn demo_unpack2x16snorm(e: u32) -> Vec2f {
        unpack2x16snorm(e)
    }

    pub fn demo_pack2x16unorm(v: Vec2f) -> u32 {
        pack2x16unorm(v)
    }

    pub fn demo_unpack2x16unorm(e: u32) -> Vec2f {
        unpack2x16unorm(e)
    }

    pub fn demo_pack2x16float(v: Vec2f) -> u32 {
        pack2x16float(v)
    }

    pub fn demo_unpack2x16float(e: u32) -> Vec2f {
        unpack2x16float(e)
    }
}
}

Generated WGSL

fn demo_pack4x8snorm(v: vec4f) -> u32 {
    return pack4x8snorm(v);
}

fn demo_unpack4x8snorm(e: u32) -> vec4f {
    return unpack4x8snorm(e);
}

fn demo_pack4x8unorm(v: vec4f) -> u32 {
    return pack4x8unorm(v);
}

fn demo_unpack4x8unorm(e: u32) -> vec4f {
    return unpack4x8unorm(e);
}

fn demo_pack2x16snorm(v: vec2f) -> u32 {
    return pack2x16snorm(v);
}

fn demo_unpack2x16snorm(e: u32) -> vec2f {
    return unpack2x16snorm(e);
}

fn demo_pack2x16unorm(v: vec2f) -> u32 {
    return pack2x16unorm(v);
}

fn demo_unpack2x16unorm(e: u32) -> vec2f {
    return unpack2x16unorm(e);
}

fn demo_pack2x16float(v: vec2f) -> u32 {
    return pack2x16float(v);
}

fn demo_unpack2x16float(e: u32) -> vec2f {
    return unpack2x16float(e);
}

Notes

• The Rust snake_case names (pack4x8snorm, etc.) map directly to the WGSL builtin names.

Advanced Numeric

Demonstrates the advanced numeric builtin functions: modf, frexp, and ldexp. modf and frexp return structs with named fields (fract, whole, exp) that map to WGSL struct member access.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod advanced_numeric_example {
    //! Demonstrates the advanced numeric builtin functions: `modf`, `frexp`,
    //! and `ldexp`.
    use wgsl_rs::std::*;

    pub fn demo_modf_fract(e: f32) -> f32 {
        let result = modf(e);
        result.fract
    }

    pub fn demo_modf_whole(e: f32) -> f32 {
        modf(e).whole
    }

    pub fn demo_frexp_fract(e: f32) -> f32 {
        frexp(e).fract
    }

    pub fn demo_frexp_exp(e: f32) -> i32 {
        frexp(e).exp
    }

    pub fn demo_ldexp(significand: f32, exponent: i32) -> f32 {
        ldexp(significand, exponent)
    }
}
}

Generated WGSL

fn demo_modf_fract(e: f32) -> f32 {
    let result = modf(e);
    return result.fract;
}

fn demo_modf_whole(e: f32) -> f32 {
    return modf(e).whole;
}

fn demo_frexp_fract(e: f32) -> f32 {
    return frexp(e).fract;
}

fn demo_frexp_exp(e: f32) -> i32 {
    return frexp(e).exp;
}

fn demo_ldexp(significand: f32, exponent: i32) -> f32 {
    return ldexp(significand, exponent);
}

Notes

• modf(e) and frexp(e) return structs with fract/whole and fract/exp fields respectively, accessed via field access syntax that maps directly to WGSL.

Matrix Builtin

Demonstrates the matrix builtin functions determinant and transpose for 2x2, 3x3, and 4x4 matrices.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod matrix_builtin_example {
    //! Demonstrates matrix builtin functions: `determinant` and `transpose`.
    use wgsl_rs::std::*;

    pub fn demo_determinant_2x2(m: Mat2f) -> f32 {
        determinant(m)
    }

    pub fn demo_determinant_3x3(m: Mat3f) -> f32 {
        determinant(m)
    }

    pub fn demo_determinant_4x4(m: Mat4f) -> f32 {
        determinant(m)
    }

    pub fn demo_transpose_4x4(m: Mat4f) -> Mat4f {
        transpose(m)
    }
}
}

Generated WGSL

fn demo_determinant_2x2(m: mat2x2f) -> f32 {
    return determinant(m);
}

fn demo_determinant_3x3(m: mat3x3f) -> f32 {
    return determinant(m);
}

fn demo_determinant_4x4(m: mat4x4f) -> f32 {
    return determinant(m);
}

fn demo_transpose_4x4(m: mat4x4f) -> mat4x4f {
    return transpose(m);
}

Notes

• determinant and transpose map directly to the WGSL builtins of the same names.

Synchronization

Demonstrates synchronization builtin functions for compute shaders: workgroup_barrier, storage_barrier, and workgroup_uniform_load. These coordinate memory visibility and execution ordering across invocations within a workgroup and must only be called from compute entry points in uniform control flow.

Rust Source

#[wgsl]
pub mod synchronization_example {
    //! Demonstrates synchronization builtin functions for compute shaders.
    //!
    //! These functions coordinate memory visibility and execution ordering
    //! across invocations within a workgroup. They must only be called from
    //! compute shader entry points in uniform control flow.
    use wgsl_rs::std::*;

    workgroup!(SCRATCH: [u32; 64]);

    storage!(group(0), binding(0), INPUT: [u32; 64]);
    storage!(group(0), binding(1), read_write, OUTPUT: [u32; 64]);

    #[compute]
    #[workgroup_size(64)]
    pub fn main(#[builtin(local_invocation_index)] local_idx: u32) {
        // Copy input data into workgroup-shared memory.
        get_mut!(SCRATCH)[local_idx as usize] = get!(INPUT)[local_idx as usize];

        // Ensure all workgroup memory writes are visible to every invocation.
        workgroup_barrier();

        // Read from a neighbor's slot (with wrap-around) to demonstrate
        // that the barrier made all writes visible.
        let neighbor_idx: u32 = (local_idx + 1u32) % 64u32;
        let neighbor_val: u32 = get!(SCRATCH)[neighbor_idx as usize];

        // Ensure all storage writes from the workgroup are complete
        // before writing results.
        storage_barrier();

        get_mut!(OUTPUT)[local_idx as usize] = neighbor_val;
    }

    #[compute]
    #[workgroup_size(64)]
    pub fn uniform_load_example(#[builtin(local_invocation_index)] local_idx: u32) {
        // Each invocation writes its index into shared memory.
        get_mut!(SCRATCH)[local_idx as usize] = local_idx;

        // Ensure all writes are visible before uniform load.
        workgroup_barrier();

        // Uniformly load the first element across the entire workgroup.
        // All invocations receive the same value.
        let first: [u32; 64] = workgroup_uniform_load(&SCRATCH);
        get_mut!(OUTPUT)[local_idx as usize] = first[0];
    }
}

Generated WGSL

var<workgroup> SCRATCH: array<u32, 64>;
@group(0) @binding(0) var<storage, read> INPUT: array<u32, 64>;
@group(0) @binding(1) var<storage, read_write> OUTPUT: array<u32, 64>;

@compute @workgroup_size(64) fn main(@builtin(local_invocation_index) local_idx: u32) {
    SCRATCH[u32(local_idx)] = INPUT[u32(local_idx)];
    workgroupBarrier();
    let neighbor_idx: u32 = (local_idx + 1u) % 64u;
    let neighbor_val: u32 = SCRATCH[u32(neighbor_idx)];
    storageBarrier();
    OUTPUT[u32(local_idx)] = neighbor_val;
}

@compute @workgroup_size(64) fn uniform_load_example(@builtin(local_invocation_index) local_idx: u32) {
    SCRATCH[u32(local_idx)] = local_idx;
    workgroupBarrier();
    let first: array<u32, 64> = workgroupUniformLoad(&SCRATCH);
    OUTPUT[u32(local_idx)] = first[0];
}

Notes

• workgroup_barrier() maps to workgroupBarrier; storage_barrier() maps to storageBarrier.
• workgroup_uniform_load(&SCRATCH) maps to workgroupUniformLoad(&SCRATCH).
• These functions must only be called from compute entry points in uniform control flow.

Macro Rules

Demonstrates that macro_rules! definitions and derive macros inside a #[wgsl] module are stripped from WGSL code generation but remain available in the Rust source. The struct below carries #[derive(Debug, Clone, Copy)], which produces no WGSL output.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod macro_rules_definitions {
    //! It is possible to define `macro_rules!` within a WGSL module.
    //!
    //! Macros defined this way **will not generate WGSL code**, but will pass
    //! through to Rust code.
    //!
    //! Said another way - `macro_rules!` definitions will be stripped from WGSL
    //! code generation but will remain in your Rust source.

    #[expect(unused_macros)]
    macro_rules! my_macro {
        ($id:ident) => {
            id
        };
    }

    // It's also possible to use derive macros.
    //
    // Derive macros pass through without generating any extra WGSL.
    #[derive(Debug, Clone, Copy)]
    pub struct Data {
        pub inner: f32,
    }
}
}

Generated WGSL

struct Data {
    inner: f32
}

Notes

• macro_rules! definitions are stripped from WGSL output entirely but remain usable in Rust.
• Derive macros (e.g. #[derive(Debug, Clone, Copy)]) pass through without generating any extra WGSL; only the struct definition itself is emitted.

Slab Read/Write

Demonstrates wgsl-rs macros for reading from and writing to u32 "slabs". The slab can be any indexable item such as an array, RuntimeArray, or storage pointer. The slab_copy! macro expands to an element-wise copy loop in WGSL.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod slab_read_write {
    //! `wgsl-rs` includes macros for reading to and from u32 "slabs".
    //!
    //! The slab can be any indexable item such as an array, RuntimeArray,
    //! storage pointer, etc.

    use wgsl_rs::std::*;

    pub struct Data {
        pub one: f32,
        pub two: u32,
        pub three_four: Vec2f,
    }

    impl Data {
        /// `Data`'s slab size.
        ///
        /// This is the number of u32 slots it occupies in a u32 slab.
        pub const SLAB_SIZE: usize = 4;

        /// Returns an array container to hold ephemeral data read from a slab.
        pub fn array_container() -> [u32; Self::SLAB_SIZE] {
            [0, 0, 0, 0]
        }

        /// Convert an array into `Data`.
        pub fn from_array(arr: [u32; Self::SLAB_SIZE]) -> Self {
            Self {
                one: bitcast_f32(arr[0]),
                two: arr[1],
                three_four: vec2f(bitcast_f32(arr[2]), bitcast_f32(arr[3])),
            }
        }

        /// Convert `Data` into an array.
        pub fn to_array(data: Self) -> [u32; Self::SLAB_SIZE] {
            [
                bitcast_u32(data.one),
                data.two,
                bitcast_u32(data.three_four.x()),
                bitcast_u32(data.three_four.y()),
            ]
        }
    }

    storage!(group(0), binding(0), read_write, SLAB: RuntimeArray<u32>);

    #[compute]
    #[workgroup_size(8)]
    pub fn slab_example(#[builtin(local_invocation_index)] local_idx: u32) {
        let index = local_idx;

        // Create our `Data` struct from extracted data from the slab
        let mut data: Data;
        {
            // Extract the u32 data from the slab
            let mut array_data = Data::array_container();
            slab_copy!(get!(SLAB), index, array_data, 0, Data::SLAB_SIZE);
            data = Data::from_array(array_data);
        }

        // Modify it
        data.three_four.x = 123.0;

        // Write the modified `Data` struct back to the slab
        let out_array = Data::to_array(data);
        slab_copy!(out_array, 0, get_mut!(SLAB), index, Data::SLAB_SIZE);
    }
}
}

Generated WGSL

struct Data {
    one: f32,
    two: u32,
    three_four: vec2f
}
const Data__1SLAB_SIZE: u32 = 4;

fn Data__1array_container() -> array<u32, Data__1SLAB_SIZE> {
    return array(0, 0, 0, 0);
}

fn Data__1from_array(arr: array<u32, Data__1SLAB_SIZE>) -> Data {
    return Data(bitcast<f32>(arr[0]), arr[1], vec2f(bitcast<f32>(arr[2]), bitcast<f32>(arr[3])));
}

fn Data__1to_array(data: Data) -> array<u32, Data__1SLAB_SIZE> {
    return array(bitcast<u32>(data.one), data.two, bitcast<u32>(data.three_four.x), bitcast<u32>(data.three_four.y));
}
@group(0) @binding(0) var<storage, read_write> SLAB: array<u32>;

@compute @workgroup_size(8) fn slab_example(@builtin(local_invocation_index) local_idx: u32) {
    let index = local_idx;
    var data: Data;
    {
        var array_data = Data__1array_container();
        for (var _i: u32 = 0u; _i < Data__1SLAB_SIZE; _i++) {
            array_data[0 + _i] = SLAB[index + _i];
        }
        data = Data__1from_array(array_data);
    }
    data.three_four.x = 123.0;
    let out_array = Data__1to_array(data);
    for (var _i: u32 = 0u; _i < Data__1SLAB_SIZE; _i++) {
        SLAB[index + _i] = out_array[0 + _i];
    }
}

Notes

• slab_copy!(src, src_offset, dest, dest_offset, size) is bidirectional: pass the slab as src to read, or as dest to write.
• Self::SLAB_SIZE (an associated const) is mangled to Data__1SLAB_SIZE in WGSL (the _1 is the bijective mangle escaping the underscore in SLAB_SIZE).

Derivatives

Demonstrates all 9 WGSL derivative builtin functions used in a fragment shader: dpdx, dpdy, fwidth, plus their _fine and _coarse variants, applied to both scalars and vectors.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod derivative_example {
    //! Demonstrates all 9 WGSL derivative builtin functions used in a fragment
    //! shader.

    use wgsl_rs::std::*;

    pub struct FragInput {
        #[builtin(position)]
        pub position: Vec4f,
    }

    pub struct DerivativeOutputs {
        #[location(0)]
        pub dx: Vec4f,
        #[location(1)]
        pub dy: Vec4f,
        #[location(2)]
        pub fw: Vec4f,
    }

    #[fragment]
    pub fn frag_main(input: FragInput) -> DerivativeOutputs {
        let position = input.position;

        // Scalar derivatives.
        let dx_scalar = dpdx(position.x());
        let dy_scalar = dpdy(position.y());
        let fw_scalar = fwidth(position.x());

        // Fine variants on a Vec2f.
        let pos_xy = vec2f(position.x(), position.y());
        let dx_fine = dpdx_fine(pos_xy);
        let dy_fine = dpdy_fine(pos_xy);
        let fw_fine = fwidth_fine(pos_xy);

        // Coarse variants on a scalar.
        let dx_coarse = dpdx_coarse(position.x());
        let dy_coarse = dpdy_coarse(position.y());
        let fw_coarse = fwidth_coarse(position.x());

        DerivativeOutputs {
            dx: vec4f(dx_scalar, dx_fine.x(), dx_fine.y(), dx_coarse),
            dy: vec4f(dy_scalar, dy_fine.x(), dy_fine.y(), dy_coarse),
            fw: vec4f(fw_scalar, fw_fine.x(), fw_fine.y(), fw_coarse),
        }
    }
}
}

Generated WGSL

struct FragInput {
    @builtin(position) position: vec4f
}

struct DerivativeOutputs {
    @location(0) dx: vec4f,
    @location(1) dy: vec4f,
    @location(2) fw: vec4f
}

@fragment fn frag_main(input: FragInput) -> DerivativeOutputs {
    let position = input.position;
    let dx_scalar = dpdx(position.x);
    let dy_scalar = dpdy(position.y);
    let fw_scalar = fwidth(position.x);
    let pos_xy = vec2f(position.x, position.y);
    let dx_fine = dpdxFine(pos_xy);
    let dy_fine = dpdyFine(pos_xy);
    let fw_fine = fwidthFine(pos_xy);
    let dx_coarse = dpdxCoarse(position.x);
    let dy_coarse = dpdyCoarse(position.y);
    let fw_coarse = fwidthCoarse(position.x);
    return DerivativeOutputs(vec4f(dx_scalar, dx_fine.x, dx_fine.y, dx_coarse), vec4f(dy_scalar, dy_fine.x, dy_fine.y, dy_coarse), vec4f(fw_scalar, fw_fine.x, fw_fine.y, fw_coarse));
}

Notes

• The nine functions are: dpdx, dpdy, fwidth, dpdx_fine, dpdy_fine, fwidth_fine, dpdx_coarse, dpdy_coarse, fwidth_coarse.
• Snake_case Rust names map to camelCase WGSL builtins (e.g. dpdx_fine -> dpdxFine).

Discard

Demonstrates the discard!() statement for discarding fragments. The macro expands to the WGSL discard statement.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod discard_example {
    //! Demonstrates the `discard!()` statement for discarding fragments.

    use wgsl_rs::std::*;

    /// Discard fragments with shallow depth (close to the near plane).
    pub fn discard_if_shallow(pos: Vec4f) {
        if pos.z < 0.001 {
            discard!();
        }
    }

    pub struct FragInput {
        #[builtin(position)]
        pub position: Vec4f,
    }

    #[fragment]
    pub fn frag_main(input: FragInput) -> Vec4f {
        discard_if_shallow(input.position);
        vec4f(1.0, 0.0, 0.0, 1.0)
    }
}
}

Generated WGSL

fn discard_if_shallow(pos: vec4f) {
    if pos.z < 0.001 {
        discard;
    }
}

struct FragInput {
    @builtin(position) position: vec4f
}

@fragment fn frag_main(input: FragInput) -> @location(0) vec4f {
    discard_if_shallow(input.position);
    return vec4f(1.0, 0.0, 0.0, 1.0);
}

Notes

• discard!() is a macro that expands to the WGSL discard statement. It must be used inside a fragment shader.

Generic Functions

Demonstrates generic functions with monomorphization. Trait bounds (Copy + std::ops::Add) are required for Rust type-checking but produce no WGSL output. Each concrete call site triggers generation of a specialized WGSL function.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod generic_functions {
    /// A generic function that doubles a value via addition.
    ///
    /// Trait bounds (`Copy + std::ops::Add`) are required for Rust to
    /// type-check the generic body. They produce no WGSL output.
    pub fn double<T: Copy + std::ops::Add<Output = T>>(x: T) -> T {
        x + x
    }

    /// A generic "select" function: returns `a` if `cond` is true, else `b`.
    pub fn select_val<T: Copy>(a: T, b: T, cond: bool) -> T {
        if cond { a } else { b }
    }

    /// A generic function calling another generic function (transitive
    /// monomorphization). Demonstrates nested turbofish: `double::<T>(x)`.
    pub fn double_or_keep<T: Copy + std::ops::Add<Output = T>>(x: T, use_double: bool) -> T {
        select_val::<T>(double::<T>(x), x, use_double)
    }

    /// Concrete function that calls the generic helpers with `f32`.
    pub fn apply_f32(value: f32) -> f32 {
        double_or_keep::<f32>(value, true)
    }

    /// Concrete function that calls the generic helpers with `i32`.
    pub fn apply_i32(value: i32) -> i32 {
        double_or_keep::<i32>(value, false)
    }
}
}

Generated WGSL

fn apply_f32(value: f32) -> f32 {
    return _2double_or_keep_f32(value, true);
}

fn apply_i32(value: i32) -> i32 {
    return _2double_or_keep_i32(value, false);
}

fn _2double_or_keep_f32(x: f32, use_double: bool) -> f32 {
    return _1select_val_f32(double_f32(x), x, use_double);
}

fn _2double_or_keep_i32(x: i32, use_double: bool) -> i32 {
    return _1select_val_i32(double_i32(x), x, use_double);
}

fn _1select_val_f32(a: f32, b: f32, cond: bool) -> f32 {
    if cond {
        return a;
    } else {
        return b;
    }
}

fn double_f32(x: f32) -> f32 {
    return x + x;
}

fn _1select_val_i32(a: i32, b: i32, cond: bool) -> i32 {
    if cond {
        return a;
    } else {
        return b;
    }
}

fn double_i32(x: i32) -> i32 {
    return x + x;
}

Notes

• Each generic function is monomorphized per concrete type used, producing specialized WGSL functions suffixed with the type (e.g. double_f32, double_i32).
• Trait bounds are erased; only the monomorphized bodies appear in WGSL.
• Leading underscores in mangled names (_1, _2) are used to avoid collisions with user-defined names.

Trait Impls

Demonstrates trait definitions and impl blocks resolved via monomorphization. The trait definition itself produces no WGSL output; each impl method becomes a free function named Type_method, and generic calls resolve to the concrete function after monomorphization.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod trait_impl_example {
    /// A trait for types that support an "add" operation.
    /// This definition is Rust-only — it produces no WGSL output.
    pub trait Addable {
        fn add(a: Self, b: Self) -> Self;
    }

    impl Addable for f32 {
        fn add(a: f32, b: f32) -> f32 {
            a + b
        }
    }

    impl Addable for i32 {
        fn add(a: i32, b: i32) -> i32 {
            a + b
        }
    }

    /// Generic function that sums three values using the trait method.
    /// `T::add(a, b)` resolves to `f32_add(a, b)` or `i32_add(a, b)`
    /// after monomorphization.
    pub fn sum_three<T: Addable>(a: T, b: T, c: T) -> T {
        let ab = T::add(a, b);
        T::add(ab, c)
    }

    /// Concrete caller — triggers monomorphization of `sum_three::<f32>`.
    pub fn sum_f32(a: f32, b: f32, c: f32) -> f32 {
        sum_three::<f32>(a, b, c)
    }

    /// Concrete caller — triggers monomorphization of `sum_three::<i32>`.
    pub fn sum_i32(a: i32, b: i32, c: i32) -> i32 {
        sum_three::<i32>(a, b, c)
    }
}
}

Generated WGSL

fn f32_add(a: f32, b: f32) -> f32 {
    return a + b;
}

fn i32_add(a: i32, b: i32) -> i32 {
    return a + b;
}

fn sum_f32(a: f32, b: f32, c: f32) -> f32 {
    return _1sum_three_f32(a, b, c);
}

fn sum_i32(a: i32, b: i32, c: i32) -> i32 {
    return _1sum_three_i32(a, b, c);
}

fn _1sum_three_f32(a: f32, b: f32, c: f32) -> f32 {
    let ab = f32_add(a, b);
    return f32_add(ab, c);
}

fn _1sum_three_i32(a: i32, b: i32, c: i32) -> i32 {
    let ab = i32_add(a, b);
    return i32_add(ab, c);
}

Notes

• Trait definitions (trait Addable { ... }) produce no WGSL output.
• impl Addable for f32 methods become free functions f32_add, i32_add, etc.
• T::add(a, b) inside a generic function resolves to the concrete mangled function after monomorphization.

Renderer Specialization

A full renderer example specialized via traits and turbofish. Three trait axes (Material, LightModel, NormalSource) are composed in a single generic shade_fragment function. Each concrete configuration (selected via turbofish like shade_fragment::<Gradient, BlinnPhong, PerturbNormal>) monomorphizes into a fully-inlined, specialized WGSL function with zero overhead.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod renderer_specialization {
    use wgsl_rs::std::*;

    // ===== Traits: each axis of rendering variation =====

    /// How a material produces a surface color at a given UV coordinate.
    pub trait Material {
        fn surface_color(uv: Vec2f) -> Vec4f;
    }

    /// How lighting is computed for a given surface color and geometry.
    pub trait LightModel {
        fn apply_lighting(
            surface: Vec4f,
            normal: Vec3f,
            light_dir: Vec3f,
            view_dir: Vec3f,
        ) -> Vec4f;
    }

    /// How the surface normal is determined.
    pub trait NormalSource {
        fn get_normal(uv: Vec2f, geom_normal: Vec3f) -> Vec3f;
    }

    // ===== Strategy structs =====
    //
    // Each struct represents a concrete rendering strategy. In a real
    // renderer these might hold configuration data; here they serve as
    // type-level tags that select which code path to monomorphize.

    /// Simple checkerboard material — procedural, no textures needed.
    pub struct Checker {
        pub _tag: u32,
    }

    /// Vertical gradient material — warm-to-cool procedural color.
    pub struct Gradient {
        pub _tag: u32,
    }

    /// Lambert diffuse lighting model.
    pub struct Lambert {
        pub _tag: u32,
    }

    /// Blinn-Phong lighting with specular highlights.
    pub struct BlinnPhong {
        pub _tag: u32,
    }

    /// Use the raw geometric normal as-is.
    pub struct GeomNormal {
        pub _tag: u32,
    }

    /// Perturb the geometric normal (simulates a normal map).
    pub struct PerturbNormal {
        pub _tag: u32,
    }

    // ===== Trait implementations =====

    impl Material for Checker {
        fn surface_color(uv: Vec2f) -> Vec4f {
            let checker: f32 = floor(uv.x() * 4.0) + floor(uv.y() * 4.0);
            let c: f32 = (checker % 2.0) * 0.5 + 0.25;
            vec4f(c, c, c, 1.0)
        }
    }

    impl Material for Gradient {
        fn surface_color(uv: Vec2f) -> Vec4f {
            vec4f(uv.y() * 0.8, 0.3, (1.0 - uv.y()) * 0.9, 1.0)
        }
    }

    impl LightModel for Lambert {
        fn apply_lighting(
            surface: Vec4f,
            normal: Vec3f,
            light_dir: Vec3f,
            _view_dir: Vec3f,
        ) -> Vec4f {
            let ndotl: f32 = max(dot(normal, light_dir), 0.0);
            surface * ndotl
        }
    }

    impl LightModel for BlinnPhong {
        fn apply_lighting(
            surface: Vec4f,
            normal: Vec3f,
            light_dir: Vec3f,
            view_dir: Vec3f,
        ) -> Vec4f {
            let ndotl: f32 = max(dot(normal, light_dir), 0.0);
            let diffuse: Vec4f = surface * ndotl;
            let half_vec: Vec3f = normalize(light_dir + view_dir);
            let spec: f32 = pow(max(dot(normal, half_vec), 0.0), 32.0);
            diffuse + vec4f(spec, spec, spec, 0.0)
        }
    }

    impl NormalSource for GeomNormal {
        fn get_normal(_uv: Vec2f, geom_normal: Vec3f) -> Vec3f {
            normalize(geom_normal)
        }
    }

    impl NormalSource for PerturbNormal {
        fn get_normal(uv: Vec2f, geom_normal: Vec3f) -> Vec3f {
            let perturb: Vec3f = vec3f(uv.x() * 0.1 - 0.05, uv.y() * 0.1 - 0.05, 1.0);
            normalize(geom_normal + perturb)
        }
    }

    // ===== Generic shader pipeline =====

    /// The single generic fragment shading function. It composes material,
    /// lighting, and normal sourcing through trait bounds. After
    /// monomorphization, each configuration produces a fully-inlined,
    /// specialized WGSL function with zero overhead.
    pub fn shade_fragment<M: Material, L: LightModel, N: NormalSource>(
        uv: Vec2f,
        geom_normal: Vec3f,
        light_dir: Vec3f,
        view_dir: Vec3f,
    ) -> Vec4f {
        let normal: Vec3f = N::get_normal(uv, geom_normal);
        let surface: Vec4f = M::surface_color(uv);
        L::apply_lighting(surface, normal, light_dir, view_dir)
    }

    // ===== Concrete shader variants (each is one turbofish line) =====

    /// Fancy renderer: gradient + Blinn-Phong + perturbed normals.
    pub fn shade_fancy(uv: Vec2f, normal: Vec3f, light_dir: Vec3f, view_dir: Vec3f) -> Vec4f {
        shade_fragment::<Gradient, BlinnPhong, PerturbNormal>(uv, normal, light_dir, view_dir)
    }

    /// Mix-and-match: checkerboard + Blinn-Phong + perturbed normals.
    /// Demonstrates that each axis of variation is independent.
    pub fn shade_hybrid(uv: Vec2f, normal: Vec3f, light_dir: Vec3f, view_dir: Vec3f) -> Vec4f {
        shade_fragment::<Checker, BlinnPhong, PerturbNormal>(uv, normal, light_dir, view_dir)
    }
}
}

Generated WGSL

struct Checker {
    _tag: u32
}

struct Gradient {
    _tag: u32
}

struct Lambert {
    _tag: u32
}

struct BlinnPhong {
    _tag: u32
}

struct GeomNormal {
    _tag: u32
}

struct PerturbNormal {
    _tag: u32
}

fn Checker__1surface_color(uv: vec2f) -> vec4f {
    let checker: f32 = floor(uv.x * 4.0) + floor(uv.y * 4.0);
    let c: f32 = (checker % 2.0) * 0.5 + 0.25;
    return vec4f(c, c, c, 1.0);
}

fn Gradient__1surface_color(uv: vec2f) -> vec4f {
    return vec4f(uv.y * 0.8, 0.3, (1.0 - uv.y) * 0.9, 1.0);
}

fn Lambert__1apply_lighting(surface: vec4f, normal: vec3f, light_dir: vec3f, _view_dir: vec3f) -> vec4f {
    let ndotl: f32 = max(dot(normal, light_dir), 0.0);
    return surface * ndotl;
}

fn BlinnPhong__1apply_lighting(surface: vec4f, normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    let ndotl: f32 = max(dot(normal, light_dir), 0.0);
    let diffuse: vec4f = surface * ndotl;
    let half_vec: vec3f = normalize(light_dir + view_dir);
    let spec: f32 = pow(max(dot(normal, half_vec), 0.0), 32.0);
    return diffuse + vec4f(spec, spec, spec, 0.0);
}

fn GeomNormal__1get_normal(_uv: vec2f, geom_normal: vec3f) -> vec3f {
    return normalize(geom_normal);
}

fn PerturbNormal__1get_normal(uv: vec2f, geom_normal: vec3f) -> vec3f {
    let perturb: vec3f = vec3f(uv.x * 0.1 - 0.05, uv.y * 0.1 - 0.05, 1.0);
    return normalize(geom_normal + perturb);
}

fn shade_fancy(uv: vec2f, normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    return _1shade_fragment_Gradient_BlinnPhong_PerturbNormal(uv, normal, light_dir, view_dir);
}

fn shade_hybrid(uv: vec2f, normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    return _1shade_fragment_Checker_BlinnPhong_PerturbNormal(uv, normal, light_dir, view_dir);
}

fn _1shade_fragment_Gradient_BlinnPhong_PerturbNormal(uv: vec2f, geom_normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    let normal: vec3f = PerturbNormal__1get_normal(uv, geom_normal);
    let surface: vec4f = Gradient__1surface_color(uv);
    return BlinnPhong__1apply_lighting(surface, normal, light_dir, view_dir);
}

fn _1shade_fragment_Checker_BlinnPhong_PerturbNormal(uv: vec2f, geom_normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    let normal: vec3f = PerturbNormal__1get_normal(uv, geom_normal);
    let surface: vec4f = Checker__1surface_color(uv);
    return BlinnPhong__1apply_lighting(surface, normal, light_dir, view_dir);
}

Notes

• Each trait method becomes a free function mangled as Type__1method (e.g. Checker__1surface_color).
• The generic shade_fragment<M, L, N> is monomorphized into one function per turbofish configuration, named _1shade_fragment_<M>_<L>_<N>.
• Strategy structs become empty-tagged WGSL structs; they exist only to drive monomorphization.

Renderer Specialization (Simple)

A second specialization of the shared renderer pipeline from renderer-specialization. This module selects Checker + Lambert + GeomNormal via turbofish, producing a distinct specialized WGSL function. It re-exports the trait implementations from the parent module via use super::renderer_specialization::*.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod renderer_specialization_simple {
    use super::renderer_specialization::*;
    use wgsl_rs::std::*;

    /// Simple renderer: checkerboard + Lambert + geometric normals.
    pub fn shade_simple(uv: Vec2f, normal: Vec3f, light_dir: Vec3f, view_dir: Vec3f) -> Vec4f {
        shade_fragment::<Checker, Lambert, GeomNormal>(uv, normal, light_dir, view_dir)
    }
}
}

Generated WGSL

struct Checker {
    _tag: u32
}

struct Gradient {
    _tag: u32
}

struct Lambert {
    _tag: u32
}

struct BlinnPhong {
    _tag: u32
}

struct GeomNormal {
    _tag: u32
}

struct PerturbNormal {
    _tag: u32
}

fn Checker__1surface_color(uv: vec2f) -> vec4f {
    let checker: f32 = floor(uv.x * 4.0) + floor(uv.y * 4.0);
    let c: f32 = (checker % 2.0) * 0.5 + 0.25;
    return vec4f(c, c, c, 1.0);
}

fn Gradient__1surface_color(uv: vec2f) -> vec4f {
    return vec4f(uv.y * 0.8, 0.3, (1.0 - uv.y) * 0.9, 1.0);
}

fn Lambert__1apply_lighting(surface: vec4f, normal: vec3f, light_dir: vec3f, _view_dir: vec3f) -> vec4f {
    let ndotl: f32 = max(dot(normal, light_dir), 0.0);
    return surface * ndotl;
}

fn BlinnPhong__1apply_lighting(surface: vec4f, normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    let ndotl: f32 = max(dot(normal, light_dir), 0.0);
    let diffuse: vec4f = surface * ndotl;
    let half_vec: vec3f = normalize(light_dir + view_dir);
    let spec: f32 = pow(max(dot(normal, half_vec), 0.0), 32.0);
    return diffuse + vec4f(spec, spec, spec, 0.0);
}

fn GeomNormal__1get_normal(_uv: vec2f, geom_normal: vec3f) -> vec3f {
    return normalize(geom_normal);
}

fn PerturbNormal__1get_normal(uv: vec2f, geom_normal: vec3f) -> vec3f {
    let perturb: vec3f = vec3f(uv.x * 0.1 - 0.05, uv.y * 0.1 - 0.05, 1.0);
    return normalize(geom_normal + perturb);
}

fn shade_fancy(uv: vec2f, normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    return _1shade_fragment_Gradient_BlinnPhong_PerturbNormal(uv, normal, light_dir, view_dir);
}

fn shade_hybrid(uv: vec2f, normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    return _1shade_fragment_Checker_BlinnPhong_PerturbNormal(uv, normal, light_dir, view_dir);
}

fn _1shade_fragment_Gradient_BlinnPhong_PerturbNormal(uv: vec2f, geom_normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    let normal: vec3f = PerturbNormal__1get_normal(uv, geom_normal);
    let surface: vec4f = Gradient__1surface_color(uv);
    return BlinnPhong__1apply_lighting(surface, normal, light_dir, view_dir);
}

fn _1shade_fragment_Checker_BlinnPhong_PerturbNormal(uv: vec2f, geom_normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    let normal: vec3f = PerturbNormal__1get_normal(uv, geom_normal);
    let surface: vec4f = Checker__1surface_color(uv);
    return BlinnPhong__1apply_lighting(surface, normal, light_dir, view_dir);
}
fn shade_simple(uv: vec2f, normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    return _1shade_fragment_Checker_Lambert_GeomNormal(uv, normal, light_dir, view_dir);
}
fn _1shade_fragment_Checker_Lambert_GeomNormal(uv: vec2f, geom_normal: vec3f, light_dir: vec3f, view_dir: vec3f) -> vec4f {
    let normal: vec3f = GeomNormal__1get_normal(uv, geom_normal);
    let surface: vec4f = Checker__1surface_color(uv);
    return Lambert__1apply_lighting(surface, normal, light_dir, view_dir);
}

Notes

• use super::renderer_specialization::* pulls in all the strategy structs, trait impls, and the generic shade_fragment from the sibling module, causing them to be re-emitted in this module's WGSL.
• The new shade_simple turbofish line (shade_fragment::<Checker, Lambert, GeomNormal>) produces a new specialized function _1shade_fragment_Checker_Lambert_GeomNormal.
• This demonstrates that each axis of variation is independently combinable.

Generic Structs

Demonstrates generic structs and impl blocks with monomorphization. A Pair<T> struct and its methods are specialized for each concrete type used.

Note: This example uses #[wgsl(skip_validation)] because of a known monomorphization bug: the Pair struct constructor is not mangled correctly when monomorphized. WGSL generation partially works but the constructor call is not produced correctly. There is no validated WGSL output for this example.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl(skip_validation)] // TODO: monomorphization bug — `Pair` struct constructor not mangled
pub mod generic_structs {
    /// A generic pair of values.
    pub struct Pair<T: Copy> {
        pub a: T,
        pub b: T,
    }

    /// Methods on the generic Pair struct.
    impl<T: Copy + std::ops::Add<Output = T>> Pair<T> {
        /// Extract the first element.
        pub fn first(p: Pair<T>) -> T {
            p.a
        }

        /// Sum both elements.
        pub fn sum(p: Pair<T>) -> T {
            p.a + p.b
        }
    }

    pub fn generic_pair_sum<T: Copy + std::ops::Add<Output = T>>(a: T, b: T) -> T {
        let p = Pair { a, b };
        Pair::sum(p)
    }

    /// Uses `Pair<f32>`.
    pub fn use_pair_f32() -> f32 {
        let p = Pair { a: 1.0, b: 2.0 };
        Pair::<f32>::sum(p)
    }

    /// Uses `Pair<i32>`.
    pub fn use_pair_i32() -> i32 {
        let p: Pair<i32> = Pair::<i32> { a: 10, b: 20 };
        Pair::<i32>::first(p)
    }
}
}

Notes

• #[wgsl(skip_validation)] disables naga validation of the generated WGSL. This is required here due to a known bug where the Pair struct constructor is not mangled during monomorphization, so the generated WGSL does not pass validation.
• Generic struct methods (Pair::sum, Pair::first) are expected to monomorphize to Pair_f32_sum, Pair_i32_first, etc., but the constructor mangling is currently broken.
• Track the fix for this bug to remove the skip_validation attribute.

Shared Inter-Stage

Demonstrates a single struct shared between vertex and fragment stages. VertexOutput serves as both the vertex return type and the fragment input, with #[builtin(position)] and #[location] decorations.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl]
pub mod shared_inter_stage {
    use wgsl_rs::std::*;

    /// Vertex output / fragment input — a single struct shared across stages.
    pub struct VertexOutput {
        #[builtin(position)]
        pub clip_position: Vec4f,
        #[location(0)]
        pub color: Vec4f,
    }

    #[vertex]
    pub fn vs_main(#[builtin(vertex_index)] vertex_index: u32) -> VertexOutput {
        const POS: [Vec2f; 3] = [vec2f(0.0, 0.5), vec2f(-0.5, -0.5), vec2f(0.5, -0.5)];
        let position = POS[vertex_index as usize];
        VertexOutput {
            clip_position: vec4f(position.x(), position.y(), 0.0, 1.0),
            color: vec4f(1.0, 0.0, 0.0, 1.0),
        }
    }

    #[fragment]
    pub fn fs_main(input: VertexOutput) -> Vec4f {
        input.color
    }
}
}

Generated WGSL

struct VertexOutput {
    @builtin(position) clip_position: vec4f,
    @location(0) color: vec4f
}

@vertex fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
    const POS: array<vec2f, 3> = array(vec2f(0.0, 0.5), vec2f(-0.5, -0.5), vec2f(0.5, -0.5));
    let position = POS[u32(vertex_index)];
    return VertexOutput(vec4f(position.x, position.y, 0.0, 1.0), vec4f(1.0, 0.0, 0.0, 1.0));
}

@fragment fn fs_main(input: VertexOutput) -> @location(0) vec4f {
    return input.color;
}

Notes

• One struct can serve as both the vertex shader's output and the fragment shader's input, matching the WGSL convention where vertex output structs feed fragment input structs.
• Struct literals become positional constructor calls in WGSL.

PhantomData

Demonstrates PhantomData<T> marker fields on #[wgsl] structs. PhantomData is re-exported from wgsl_rs::std. The proc-macro recognizes PhantomData<_> fields specially: they are retained in the IR (so extensions can observe which type parameter each phantom slot binds) but omitted from the rendered WGSL. Construction expressions using the bare PhantomData value are likewise stripped so the rendered positional constructor has the correct arity.

Rust Source

#![allow(unused)]
fn main() {
#[wgsl(skip_validation)]
pub mod phantom_data {
    use wgsl_rs::std::*;

    /// A typed identifier carrying a phantom type tag. The `phantom`
    /// field is dropped from the WGSL output, leaving only `index`.
    pub struct Id<T> {
        pub index: u32,
        pub phantom: PhantomData<T>,
    }

    /// A struct binding two type parameters to two phantom slots. An
    /// extension inspecting the IR sees `t: PhantomData<T>` and
    /// `a: PhantomData<A>` and can reconstruct which field binds which
    /// parameter.
    pub struct Tagged<T, A> {
        pub x: f32,
        pub t: PhantomData<T>,
        pub a: PhantomData<A>,
    }

    pub fn make_id() -> Id<f32> {
        Id {
            index: 0u32,
            phantom: PhantomData,
        }
    }

    pub fn make_tagged() -> Tagged<f32, u32> {
        Tagged {
            x: 1.0,
            t: PhantomData,
            a: PhantomData,
        }
    }

    pub fn read_id(i: Id<f32>) -> u32 {
        i.index
    }

    pub fn read_tagged(t: Tagged<f32, u32>) -> f32 {
        t.x
    }
}
}

Generated WGSL

fn make_id() -> Id_f32 {
    return Id(0u);
}

fn make_tagged() -> Tagged_f32_u32 {
    return Tagged(1.0);
}

fn read_id(i: Id_f32) -> u32 {
    return i.index;
}

fn read_tagged(t: Tagged_f32_u32) -> f32 {
    return t.x;
}

struct Id_f32 {
    index: u32
}

struct Tagged_f32_u32 {
    x: f32
}

Notes

• PhantomData<T> fields are dropped from the WGSL struct definition — Id_f32 has only index: u32, Tagged_f32_u32 has only x: f32.
• Construction expressions use the bare PhantomData value (no turbofish). The macro strips it from the positional constructor call so the rendered arity matches the non-phantom field count.
• The IR retains phantom fields as Type::Phantom { elem } so extensions can see the full type-parameter binding structure.
• This example uses #[wgsl(skip_validation)] because the auto-validation test doesn't cover phantom field stripping.

Supported Rust Subset

wgsl-rs transpiles a deliberately constrained subset of Rust to WGSL. The macro is additive: it never translates or rewrites user code. Constructs that cannot map cleanly to WGSL are rejected at compile time rather than approximated.

Supported

• Structs (including generic and const-generic structs)
• Enums with #[repr(u32)]
• impl blocks (free functions in WGSL), including trait impls and generic trait impls on arrays
• Free functions
• const items
• let and let mut bindings
• if / else, while, loop, for, match
• All binary, unary, and compound assignment operators
• Arrays
• Generic functions and generic structs (monomorphized)
• Const generic parameters (const N: usize / const N: u32) on functions, structs, impl blocks, and template entry points

Not Supported

Feature Table

Macro Attributes

wgsl Container Attributes

Field / Item Attributes

Entry-Point Attributes

I/O Decorator Attributes

Derive Attributes

Cargo Features

wgsl-rs

wgsl-rs-layout

Usage

[dependencies]
wgsl-rs = { version = "0.1", default-features = false, features = ["validation", "linkage-wgpu"] }
wgsl-rs-layout = { version = "0.1", features = ["doc-diagrams"] }

IR Types

The wgsl_rs::ir module exposes the intermediate representation that extensions and the linker operate on.

Top-Level Types

ir::Item Variants

ir::Type Variants