Shaders and bindings¶

Both GPU traits supply WGSL to the engine as a &'static str, because henad-core depends on nothing, not even wgpu, and therefore cannot refer to wgpu types. Most of what surrounds that string is generated at build time, and this page covers the generated part.

Generated from the WGSL¶

A crate's build.rs runs henad-build over its shaders, and henad::include_shaders!() at the crate root brings the output in as two modules, shader_bindings and binding_decls. henad-build is a build dependency, of the same release as henad. A project made from the template already has the build script and the macro call:

//! Generates bindings for the shaders under `src`, and stamps the build of this crate's models.

fn main() -> Result<(), henad_build::ShaderBuildError> {
    henad_build::stamp_commit();
    henad_build::ShaderBuild::discover("src")?.generate()?;
    Ok(())
}

stamp_commit records the commit the crate was built from, whether its sources differed from that commit, and a hash of its sources, and a sweep's manifest records them beside each model the crate registers. A crate without shaders keeps its build.rs for the stamp, and drops the ShaderBuild line and henad::include_shaders!(). A file that a model reads at compile time with include_bytes! or include_str! belongs under src, where the stamp sees it. The stamp follows no symlink to a folder, and sees no file outside src that a shader imports. A change to a shader in either place reaches the bindings and goes unrecorded. A model crate keeps its shaders under src and passes "src" to discover.

henad::include_shaders!();

ShaderBuild::discover treats every .wgsl file under src without a #define_import_path line as an entry point, and a new shader joins the bindings the next time the crate builds. A symlink to a folder under src is followed, and a folder reached twice is read once. A shader's path below src becomes its names. gpu_sir/step.wgsl is the module shader_bindings::gpu_sir::step and the constant binding_decls::bindings::GPU_SIR_STEP, the path's components upper-cased and joined by _. Each component of a .wgsl path is therefore a Rust identifier and not a keyword, and a folder named gpu-sir fails the build. So do two shaders whose names collide, as gpu/sir_step.wgsl and gpu_sir/step.wgsl both produce GPU_SIR_STEP. Two structs can collide as well. Each struct's layout assertion is named after the struct's module path and name in upper snake case, so TallyParams in gpu_vote/step.wgsl and Params in gpu_vote/step_tally.wgsl both produce GPU_VOTE_STEP_TALLY_PARAMS_ASSERTS. rustc reports that pair inside the generated file, and renaming one of the structs resolves it.

Uniform structs, workgroup sizes and bind group layouts come from the WGSL instead of being retyped in Rust. Your model fills in the generated struct and returns its bytes, and does not declare its own uniform struct. A uniform declared as a bare vector, as GPU Game of Life's step uniform is, has no struct, and the model returns the values themselves. The generated struct always has the shader's layout. An initialiser that lists every field also stops compiling when the WGSL struct Params gains a field, while an initialiser ending in ..bytemuck::Zeroable::zeroed() leaves the new field at zero.

use crate::shader_bindings::gpu_boids::step::Params as StepParams;

The shader source a model declares comes from the same place, as SHADER_STRING.

An imported constant or type reaches the generated bindings exactly when an entry point references it, since naga keeps only what an entry point references. A type or a constant that no shader in the crate uses has no Rust twin.

The generated bindings define their own module named henad, for the shared modules below. Code inside include_shaders! refers to Henad as $crate, and a hand-written module that includes the generated files itself refers to Henad as ::henad. A bare use henad::... there is ambiguous between the crate and the generated module, and fails with E0659.

Shared WGSL¶

The shared modules ship with henad-core, in henad-core/src/authoring/primitives/wgsl/, and a shader imports them with #import henad::<module>, resolved at build time.

#import henad::dispatch::linear_index
#import henad::space::{TORUS, axis_delta, heading_octant, wrap_index}

Most primitives here pair with a Rust function under henad::authoring::primitives, and a parity test pins each pair of pure functions together. Authoring primitives is the index. It lists the WGSL-only primitives and records what is deliberately absent.

Your own module is a file with a #define_import_path line. An import resolves by file path alone, so the import path mirrors the file's path below src, or below the directory of the shader importing it. #define_import_path gpu_ants::state sits in gpu_ants/state.wgsl, and the same line in common/state.wgsl is never found. The root henad is reserved for the shared modules, in any letter case. A file named henad.wgsl, or a folder named henad holding a .wgsl file, shadows one of them, and fails the build. So does a path starting with a name the generated bindings use at their root: wgpu, bytemuck, std, core, alloc, _root, ShaderEntry, layout_asserts or bytemuck_impls. The bindings name an imported module after its import path, and an import path starting with one of those names fails the build too. A module declares no binding, since bindings belong in the entry shader. An import can also refer to a file by its path, in quotes, as #import "../shared/noise" as noise. The bindings name that module after the stem of the file name, here noise, and a stem such as std or henad fails the build as well. A change to such a file reruns the build script and the generator, inside src or outside it.

Bindings¶

henad-build reads the @group(0) lines of every entry point into binding_decls, in @binding order. Every pass of either GPU trait points at one of its constants, as crate::binding_decls::bindings::GPU_SIR_STEP is for gpu_sir/step.wgsl. Each binding sits on one line, as @group(0) @binding(N) var<...> name: Type;, and a line holding @binding or @group in any other form fails the build. Group 0 of an entry point holds storage buffers, uniforms and storage textures, and a sampler or a sampled texture fails the build. Keep a render shader outside src, or list the entry points one by one with ShaderBuild::new. A compile-time assertion checks that each list has the same length as the layout that naga derives from the composed shader. The engine resolves each name itself. Otherwise a slot index could disagree with the shader that owns it.

Seven names are reserved for resources the engine owns, and anything else refers to one of the model's own buffers by its label. The full list is in GPU agent models. A GPU grid model has a resource for four of them, params, dims, output and counters.

The shipped grid models bind each buffer twice in the step shader, <label>_in for reading and <label>_out for writing, and bind params after the last pair. That suffix is a naming convention. The access mode decides which side a name resolves to.

Dispatch¶

An agent pass folds its linear invocation domain onto a 2D workgroup grid, because a hundred million agents overflow one row of workgroups.

@compute @workgroup_size(256)
fn main(@builtin(local_invocation_id) lid: vec3<u32>,
        @builtin(workgroup_id) wid: vec3<u32>) {
    let i = linear_index(lid, wid, params.groups_x);
    if i >= params.num_agents { return; }
    // ...
}

The fold width that the engine picked arrives as groups_x in the uniform block, through PassCtx::groups_x, so a shader that folds must carry that field.

A grid model's shaders dispatch 2D directly and declare a @workgroup_size(N, N) matching WORKGROUP_SIZE.

Reading the composed source¶

A shader is composed from its imports and re-emitted by naga, so the text the engine compiles is not the file as you wrote it, and a WGSL error refers to the composed text rather than your source.

HENAD_DUMP_WGSL=/tmp/wgsl cargo run --release

With that variable set, every shader the engine compiles is written to <dir>/<label>.wgsl, which lets you read a validation error against the composed source as ordinary text. See environment variables.

Next¶

• GPU grid models and GPU agent models are the two traits this machinery sits under.
• Authoring primitives is the index of what a kernel can call.