Struct Builder 

pub struct Builder {
    pub module: Module,
    pub function: Function,
    pub blocks: Vec<Block>,
    pub workgroup_size: [u32; 3],
    /* private fields */
}

A module and the entry function being built, with a stack of the blocks open for structured control flow.

What a lifter reads decides the shapes of the values it hands the builder, so an operation asked of values that cannot do it (a vector of 2 lanes added to one of 3) is not a bug of the builder and does not panic: it records the first such fault and returns a value of the shape asked for. The contract is the lifter’s: it calls Builder::check where it can name what it was reading (after each instruction or statement), and turns a fault into a refusal. finish, finish_mesh, finish_functions and end_function give back a fault nobody checked, or control flow left open, as an error: a lifter that follows the contract does not see one from its input, but a lifter’s own bug is an error to refuse the program with, not a panic.

Fields

module: Module

The module being built.

function: Function

The entry function being built (its body is assembled from Builder::blocks when the module is finished).

blocks: Vec<Block>

The blocks being filled, innermost last; statements go to the last.

workgroup_size: [u32; 3]

A compute entry point’s thread group size (set before Builder::finish; 0 for the other stages).

Implementations

impl Builder

pub fn atomic_kind(scalar: Scalar) -> Kind

The kind a word of the atomic integer type scalar is read as.

pub fn atomic( &mut self, pointer: Handle<Expression>, fun: AtomicFunction, value: Val, scalar: Scalar, ) -> Val

fun applied to the atomic word pointer names with value (as the word’s own type): the old value.

pub fn image_atomic_result(&mut self, kind: Kind) -> Val

The result expression of an image atomic operation on texels of kind: the texel’s previous value, which the Statement::ImageAtomic that names it writes.

pub fn compare_exchange( &mut self, pointer: Handle<Expression>, compare: Val, value: Val, scalar: Scalar, ) -> (Val, Val)

A compare-and-swap on the atomic word pointer names: the old value and whether the swap happened.

pub fn buffer_word( &mut self, words: Handle<Expression>, word: Val, typed: bool, ) -> Handle<Expression>

The pointer to word word of the buffer whose words words points to (a typed buffer’s element is four words: its first).

pub fn atomic_inside( &mut self, inside: Val, kind: Kind, body: impl FnOnce(&mut Self) -> Val, ) -> Val

An atomic operation body makes only where inside holds, its old value (of kind) zero where it does not: D3D’s answer for an operation past the end of its buffer, which does nothing.

impl Builder

pub fn is_nan(&mut self, v: Val) -> Val

Whether a float is NaN, by its bits: the exponent all ones and a nonzero mantissa.

pub fn d3d_saturate(&mut self, v: Val) -> Val

D3D saturate: clamp to 0 to 1, and a NaN becomes 0.

pub fn d3d_min_max(&mut self, fun: Mf, a: Val, b: Val) -> Val

D3D min and max: one NaN operand yields the other operand.

pub fn d3d_float_to_int(&mut self, x: Val) -> Val

D3D ftoi: truncate toward zero, saturate to the i32 range, NaN becomes 0.

pub fn d3d_float_to_uint(&mut self, x: Val) -> Val

D3D ftou: truncate toward zero, negatives and NaN become 0, saturate at u32::MAX.

pub fn d3d_unsigned_divide(&mut self, a: Val, b: Val) -> (Val, Val)

D3D udiv: the quotient and the remainder of unsigned values; a division by zero gives all ones for both.

pub fn d3d_shift_count(&mut self, count: Val) -> Val

A shift or bit-field count: D3D uses the low five bits.

pub fn mul_high_unsigned(&mut self, a: Val, b: Val) -> Val

The high 32 bits of the 64-bit product of two unsigned values, from 16-bit halves.

impl Builder

pub fn io_array(&mut self, v: Value, n: u32) -> Result<Handle<Type>, String>

The type of an array of n values of v, as an interface array has it (see IoType::Array).

pub fn declare_entry(&mut self, io: &EntryIo) -> Result<EntryDecl, String>

Declare the arguments (one per EntryIo::arguments) and the result structure Out (one member per EntryIo::results) with their bindings, member offsets by the alignment of their types. An argument of an array type is not carried by any stage and is an Err.

pub fn declare_function( &mut self, sig: &FunctionSig, ) -> Result<FunctionDecl, String>

Declare a stage function’s arguments and result structure (no bindings anywhere).

pub fn declare_function_with( &mut self, sig: &FunctionSig, structures: &[Handle<Type>], ) -> Result<FunctionDecl, String>

Self::declare_function for a function with pointer arguments (IoType::Memory): each is a pointer in function memory to the next of structures.

impl Builder

pub fn declare_mesh_output( &mut self, facts: &MeshFacts, vertex: &[IoItem], ) -> Result<MeshOutput, String>

Declare the output variable of a mesh program: vertex are the members of its vertex type (see crate::mesh_vertex_io), the primitive type holds the indices of a primitive.

pub fn finish_mesh( self, arguments: Vec<FunctionArgument>, facts: &MeshFacts, output: &MeshOutput, ) -> Result<Module, String>

Finish the module with main as the mesh entry point.

impl Builder

pub fn new() -> Self

A builder over an empty module with one open block.

pub fn check(&mut self) -> Result<(), String>

The first fault of the operations since the last check, which a lifter turns into a refusal of the program at the place it is reading. A builder with a fault unchecked cannot be finished: that is the lifter’s bug.

pub fn ty(&mut self, inner: TypeInner) -> Handle<Type>

The type inner in the module (types equal to an earlier one are shared).

pub fn ty_of(&mut self, kind: Kind, width: Width) -> Handle<Type>

The type of a value of kind and width lanes: a scalar for one.

pub fn app(&mut self, e: Expression) -> Handle<Expression>

Append an expression to the function; one that is not a pre-emitted kind (a literal, an argument, a pointer) is emitted in the current block.

pub fn atomic_quad(&mut self, scalar: Scalar) -> Handle<Type>

The type of an element of a typed buffer the program reaches with atomic operations: four atomic words.

pub fn result_expr(&mut self, e: Expression) -> Handle<Expression>

A result expression (an atomic’s old value): made before the statement that fills it, covered by no emit.

pub fn when(&mut self, condition: Val, then: impl FnOnce(&mut Self))

if condition { then }: what then appends goes in the accept block.

pub fn begin_repeat(&mut self, count: u32) -> (Val, Handle<LocalVariable>)

Open for i in 0..count { .. } as a loop (not unrolled): what is appended until Self::end_repeat is its body, and the value returned is i, read at the top of each pass.

pub fn end_repeat(&mut self, counter: Handle<LocalVariable>)

Close the loop Self::begin_repeat opened: count the pass and go round.

pub fn push(&mut self, s: Statement)

Append s to the innermost open block.

pub fn local(&mut self, name: &str, ty: Handle<Type>) -> Handle<LocalVariable>

A new local variable of the function, named name and uninitialised.

pub fn local_ptr(&mut self, l: Handle<LocalVariable>) -> Handle<Expression>

A pointer expression to the local variable l.

pub fn global_ptr(&mut self, g: Handle<GlobalVariable>) -> Handle<Expression>

A pointer expression to the global variable g.

pub fn load(&mut self, ptr: Handle<Expression>, kind: Kind, width: Width) -> Val

A load through ptr, typed by the kind and width the caller states the pointee has.

pub fn store(&mut self, ptr: Handle<Expression>, v: Val)

Store v through ptr.

pub fn global(&mut self, g: GlobalVariable) -> Handle<GlobalVariable>

Add the global variable g to the module.

pub fn global_expr(&mut self, e: Expression) -> Handle<Expression>

An expression of the module’s constant arena (an initialiser).

pub fn lit_u(&mut self, x: u32) -> Val

An unsigned integer literal.

pub fn lit_i(&mut self, x: i32) -> Val

A signed integer literal.

pub fn lit_f(&mut self, x: f32) -> Val

A float literal. naga holds only finite ones, so a NaN or an infinity is the bits of one, reinterpreted.

pub fn lit_b(&mut self, x: bool) -> Val

A boolean literal.

pub fn lit_bits(&mut self, kind: Kind, bits: u32) -> Val

The literal of kind whose 32 bits are bits (a bool is whether they are nonzero).

pub fn splat(&mut self, v: Val, width: Width) -> Val

v repeated across width lanes; v itself when it already has that many. A v that is not a scalar records a fault (Builder::check) and gives v back.

pub fn compose(&mut self, parts: &[Val]) -> Val

A vector (or the scalar itself, for one part) from scalars of one kind.

pub fn lane(&mut self, v: Val, i: u8) -> Val

Lane i of a vector.

pub fn swizzle(&mut self, v: Val, lanes: &[u8]) -> Val

Lanes lanes of a vector, in that order, repeats allowed; a scalar is splat.

pub fn bitcast(&mut self, v: Val, kind: Kind) -> Val

Reinterpret the bits as another kind (u32, i32 or f32), width unchanged.

pub fn bitcast_to(&mut self, v: Val, kind: Kind) -> Val

The value as kind: a bitcast between u32, i32 and f32; a bool becomes a word of all ones or zero (D3D’s true and false) and a word a bool by being nonzero.

pub fn convert(&mut self, v: Val, kind: Kind) -> Val

Convert the value to another kind (itof, utof, float to integer) with naga’s own semantics; D3D’s differ for float to integer and are behaviour.

pub fn unary(&mut self, op: UnaryOperator, v: Val) -> Val

A unary operation; the result has v’s kind and width.

pub fn binary(&mut self, op: BinaryOperator, a: Val, b: Val) -> Val

A binary operation; a scalar operand is splat to the other’s width, a comparison gives a bool.

pub fn math(&mut self, fun: MathFunction, args: &[Val]) -> Val

A math function of args (a scalar among vectors is splat to their width). The result is a scalar for Dot, a u32 for Pack2x16float, two floats for Unpack2x16float and otherwise has the first argument’s kind and the widest width. No arguments records a fault (Builder::check).

pub fn select(&mut self, cond: Val, accept: Val, reject: Val) -> Val

select(reject, accept, cond): a scalar condition is splat to the operands’ width.

pub fn mask_of(&mut self, cond: Val) -> Val

A comparison result as D3D’s mask: all ones where true, else zero.

pub fn end_function( &mut self, name: &str, decl: FunctionDecl, ) -> Result<Handle<Function>, String>

Close the function being built as the module’s function name with these arguments and result, and start an empty one: a program of a stage with no entry point is several of them (see crate::stage_functions).

pub fn declare_operation( &mut self, name: &str, arguments: Vec<FunctionArgument>, returns_bool: bool, ) -> Handle<Function>

Declare a function with no body, the module’s name for an operation the pipeline does (see crate::rt); one that returns a value returns false.

pub fn finish_functions(self) -> Result<Module, String>

Finish a module that has functions and no entry point; an error where the builder cannot be finished (a fault nobody checked, control flow left open).

pub fn finish( self, name: &str, stage: ShaderStage, arguments: Vec<FunctionArgument>, result: Option<Handle<Type>>, early_depth_test: bool, ) -> Result<Module, String>

Finish a module whose one entry point name runs in stage with these arguments and result; early_depth_test forces early depth testing. The entry’s workgroup size is Builder::workgroup_size. An error where the builder cannot be finished: a fault nobody checked (Builder::check) or control flow left open, both a lifter’s own bugs.

Trait Implementations

impl Debug for Builder

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Default for Builder

fn default() -> Self

Returns the “default value” for a type.