Source code
Revision control
Copy as Markdown
Other Tools
/*! Universal shader translator.
The central structure of the crate is [`Module`]. A `Module` contains:
- [`Function`]s, which have arguments, a return type, local variables, and a body,
- [`EntryPoint`]s, which are specialized functions that can serve as the entry
point for pipeline stages like vertex shading or fragment shading,
- [`Constant`]s and [`GlobalVariable`]s used by `EntryPoint`s and `Function`s, and
- [`Type`]s used by the above.
The body of an `EntryPoint` or `Function` is represented using two types:
- An [`Expression`] produces a value, but has no side effects or control flow.
`Expressions` include variable references, unary and binary operators, and so
on.
- A [`Statement`] can have side effects and structured control flow.
`Statement`s do not produce a value, other than by storing one in some
designated place. `Statements` include blocks, conditionals, and loops, but also
operations that have side effects, like stores and function calls.
`Statement`s form a tree, with pointers into the DAG of `Expression`s.
Restricting side effects to statements simplifies analysis and code generation.
A Naga backend can generate code to evaluate an `Expression` however and
whenever it pleases, as long as it is certain to observe the side effects of all
previously executed `Statement`s.
Many `Statement` variants use the [`Block`] type, which is `Vec<Statement>`,
with optional span info, representing a series of statements executed in order. The body of an
`EntryPoint`s or `Function` is a `Block`, and `Statement` has a
[`Block`][Statement::Block] variant.
## Arenas
To improve translator performance and reduce memory usage, most structures are
stored in an [`Arena`]. An `Arena<T>` stores a series of `T` values, indexed by
[`Handle<T>`](Handle) values, which are just wrappers around integer indexes.
For example, a `Function`'s expressions are stored in an `Arena<Expression>`,
and compound expressions refer to their sub-expressions via `Handle<Expression>`
values. (When examining the serialized form of a `Module`, note that the first
element of an `Arena` has an index of 1, not 0.)
A [`UniqueArena`] is just like an `Arena`, except that it stores only a single
instance of each value. The value type must implement `Eq` and `Hash`. Like an
`Arena`, inserting a value into a `UniqueArena` returns a `Handle` which can be
used to efficiently access the value, without a hash lookup. Inserting a value
multiple times returns the same `Handle`.
If the `span` feature is enabled, both `Arena` and `UniqueArena` can associate a
source code span with each element.
## Function Calls
Naga's representation of function calls is unusual. Most languages treat
function calls as expressions, but because calls may have side effects, Naga
represents them as a kind of statement, [`Statement::Call`]. If the function
returns a value, a call statement designates a particular [`Expression::CallResult`]
expression to represent its return value, for use by subsequent statements and
expressions.
## `Expression` evaluation time
It is essential to know when an [`Expression`] should be evaluated, because its
value may depend on previous [`Statement`]s' effects. But whereas the order of
execution for a tree of `Statement`s is apparent from its structure, it is not
so clear for `Expressions`, since an expression may be referred to by any number
of `Statement`s and other `Expression`s.
Naga's rules for when `Expression`s are evaluated are as follows:
- [`Literal`], [`Constant`], and [`ZeroValue`] expressions are
considered to be implicitly evaluated before execution begins.
- [`FunctionArgument`] and [`LocalVariable`] expressions are considered
implicitly evaluated upon entry to the function to which they belong.
Function arguments cannot be assigned to, and `LocalVariable` expressions
produce a *pointer to* the variable's value (for use with [`Load`] and
[`Store`]). Neither varies while the function executes, so it suffices to
consider these expressions evaluated once on entry.
- Similarly, [`GlobalVariable`] expressions are considered implicitly
evaluated before execution begins, since their value does not change while
code executes, for one of two reasons:
- Most `GlobalVariable` expressions produce a pointer to the variable's
value, for use with [`Load`] and [`Store`], as `LocalVariable`
expressions do. Although the variable's value may change, its address
does not.
- A `GlobalVariable` expression referring to a global in the
[`AddressSpace::Handle`] address space produces the value directly, not
a pointer. Such global variables hold opaque types like shaders or
images, and cannot be assigned to.
- A [`CallResult`] expression that is the `result` of a [`Statement::Call`],
representing the call's return value, is evaluated when the `Call` statement
is executed.
- Similarly, an [`AtomicResult`] expression that is the `result` of an
[`Atomic`] statement, representing the result of the atomic operation, is
evaluated when the `Atomic` statement is executed.
- A [`RayQueryProceedResult`] expression, which is a boolean
indicating if the ray query is finished, is evaluated when the
[`RayQuery`] statement whose [`Proceed::result`] points to it is
executed.
- All other expressions are evaluated when the (unique) [`Statement::Emit`]
statement that covers them is executed.
Now, strictly speaking, not all `Expression` variants actually care when they're
evaluated. For example, you can evaluate a [`BinaryOperator::Add`] expression
any time you like, as long as you give it the right operands. It's really only a
very small set of expressions that are affected by timing:
- [`Load`], [`ImageSample`], and [`ImageLoad`] expressions are influenced by
stores to the variables or images they access, and must execute at the
proper time relative to them.
- [`Derivative`] expressions are sensitive to control flow uniformity: they
must not be moved out of an area of uniform control flow into a non-uniform
area.
- More generally, any expression that's used by more than one other expression
or statement should probably be evaluated only once, and then stored in a
variable to be cited at each point of use.
Naga tries to help back ends handle all these cases correctly in a somewhat
circuitous way. The [`ModuleInfo`] structure returned by [`Validator::validate`]
provides a reference count for each expression in each function in the module.
Naturally, any expression with a reference count of two or more deserves to be
evaluated and stored in a temporary variable at the point that the `Emit`
statement covering it is executed. But if we selectively lower the reference
count threshold to _one_ for the sensitive expression types listed above, so
that we _always_ generate a temporary variable and save their value, then the
same code that manages multiply referenced expressions will take care of
introducing temporaries for time-sensitive expressions as well. The
`Expression::bake_ref_count` method (private to the back ends) is meant to help
with this.
## `Expression` scope
Each `Expression` has a *scope*, which is the region of the function within
which it can be used by `Statement`s and other `Expression`s. It is a validation
error to use an `Expression` outside its scope.
An expression's scope is defined as follows:
- The scope of a [`Constant`], [`GlobalVariable`], [`FunctionArgument`] or
[`LocalVariable`] expression covers the entire `Function` in which it
occurs.
- The scope of an expression evaluated by an [`Emit`] statement covers the
subsequent expressions in that `Emit`, the subsequent statements in the `Block`
to which that `Emit` belongs (if any) and their sub-statements (if any).
- The `result` expression of a [`Call`] or [`Atomic`] statement has a scope
covering the subsequent statements in the `Block` in which the statement
occurs (if any) and their sub-statements (if any).
For example, this implies that an expression evaluated by some statement in a
nested `Block` is not available in the `Block`'s parents. Such a value would
need to be stored in a local variable to be carried upwards in the statement
tree.
## Constant expressions
A Naga *constant expression* is one of the following [`Expression`]
variants, whose operands (if any) are also constant expressions:
- [`Literal`]
- [`Constant`], for [`Constant`]s
- [`ZeroValue`], for fixed-size types
- [`Compose`]
- [`Access`]
- [`AccessIndex`]
- [`Splat`]
- [`Swizzle`]
- [`Unary`]
- [`Binary`]
- [`Select`]
- [`Relational`]
- [`Math`]
- [`As`]
A constant expression can be evaluated at module translation time.
## Override expressions
A Naga *override expression* is the same as a [constant expression],
except that it is also allowed to reference other [`Override`]s.
An override expression can be evaluated at pipeline creation time.
[`AtomicResult`]: Expression::AtomicResult
[`RayQueryProceedResult`]: Expression::RayQueryProceedResult
[`CallResult`]: Expression::CallResult
[`Constant`]: Expression::Constant
[`ZeroValue`]: Expression::ZeroValue
[`Literal`]: Expression::Literal
[`Derivative`]: Expression::Derivative
[`FunctionArgument`]: Expression::FunctionArgument
[`GlobalVariable`]: Expression::GlobalVariable
[`ImageLoad`]: Expression::ImageLoad
[`ImageSample`]: Expression::ImageSample
[`Load`]: Expression::Load
[`LocalVariable`]: Expression::LocalVariable
[`Atomic`]: Statement::Atomic
[`Call`]: Statement::Call
[`Emit`]: Statement::Emit
[`Store`]: Statement::Store
[`RayQuery`]: Statement::RayQuery
[`Proceed::result`]: RayQueryFunction::Proceed::result
[`Validator::validate`]: valid::Validator::validate
[`ModuleInfo`]: valid::ModuleInfo
[`Literal`]: Expression::Literal
[`ZeroValue`]: Expression::ZeroValue
[`Compose`]: Expression::Compose
[`Access`]: Expression::Access
[`AccessIndex`]: Expression::AccessIndex
[`Splat`]: Expression::Splat
[`Swizzle`]: Expression::Swizzle
[`Unary`]: Expression::Unary
[`Binary`]: Expression::Binary
[`Select`]: Expression::Select
[`Relational`]: Expression::Relational
[`Math`]: Expression::Math
[`As`]: Expression::As
[constant expression]: index.html#constant-expressions
*/
#![allow(
clippy::new_without_default,
clippy::unneeded_field_pattern,
clippy::match_like_matches_macro,
clippy::collapsible_if,
clippy::derive_partial_eq_without_eq,
clippy::needless_borrowed_reference,
clippy::single_match,
clippy::enum_variant_names
)]
#![warn(
trivial_casts,
trivial_numeric_casts,
unused_extern_crates,
unused_qualifications,
clippy::pattern_type_mismatch,
clippy::missing_const_for_fn,
clippy::rest_pat_in_fully_bound_structs,
clippy::match_wildcard_for_single_variants
)]
#![deny(clippy::exit)]
#![cfg_attr(
not(test),
warn(
clippy::dbg_macro,
clippy::panic,
clippy::print_stderr,
clippy::print_stdout,
clippy::todo
)
)]
mod arena;
pub mod back;
mod block;
#[cfg(feature = "compact")]
pub mod compact;
pub mod error;
pub mod front;
pub mod keywords;
pub mod proc;
mod span;
pub mod valid;
pub use crate::arena::{Arena, Handle, Range, UniqueArena};
pub use crate::span::{SourceLocation, Span, SpanContext, WithSpan};
#[cfg(feature = "arbitrary")]
use arbitrary::Arbitrary;
#[cfg(feature = "deserialize")]
use serde::Deserialize;
#[cfg(feature = "serialize")]
use serde::Serialize;
/// Width of a boolean type, in bytes.
pub const BOOL_WIDTH: Bytes = 1;
/// Width of abstract types, in bytes.
pub const ABSTRACT_WIDTH: Bytes = 8;
/// Hash map that is faster but not resilient to DoS attacks.
pub type FastHashMap<K, T> = rustc_hash::FxHashMap<K, T>;
/// Hash set that is faster but not resilient to DoS attacks.
pub type FastHashSet<K> = rustc_hash::FxHashSet<K>;
/// Insertion-order-preserving hash set (`IndexSet<K>`), but with the same
/// hasher as `FastHashSet<K>` (faster but not resilient to DoS attacks).
pub type FastIndexSet<K> =
indexmap::IndexSet<K, std::hash::BuildHasherDefault<rustc_hash::FxHasher>>;
/// Insertion-order-preserving hash map (`IndexMap<K, V>`), but with the same
/// hasher as `FastHashMap<K, V>` (faster but not resilient to DoS attacks).
pub type FastIndexMap<K, V> =
indexmap::IndexMap<K, V, std::hash::BuildHasherDefault<rustc_hash::FxHasher>>;
/// Map of expressions that have associated variable names
pub(crate) type NamedExpressions = FastIndexMap<Handle<Expression>, String>;
/// Early fragment tests.
///
/// In a standard situation, if a driver determines that it is possible to switch on early depth test, it will.
///
/// Typical situations when early depth test is switched off:
/// - Calling `discard` in a shader.
/// - Writing to the depth buffer, unless ConservativeDepth is enabled.
///
/// To use in a shader:
/// - GLSL: `layout(early_fragment_tests) in;`
/// - HLSL: `Attribute earlydepthstencil`
/// - SPIR-V: `ExecutionMode EarlyFragmentTests`
/// - WGSL: `@early_depth_test`
///
/// For more, see:
/// - <https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/sm5-attributes-earlydepthstencil>
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct EarlyDepthTest {
pub conservative: Option<ConservativeDepth>,
}
/// Enables adjusting depth without disabling early Z.
///
/// To use in a shader:
/// - GLSL: `layout (depth_<greater/less/unchanged/any>) out float gl_FragDepth;`
/// - `depth_any` option behaves as if the layout qualifier was not present.
/// - HLSL: `SV_DepthGreaterEqual`/`SV_DepthLessEqual`/`SV_Depth`
/// - SPIR-V: `ExecutionMode Depth<Greater/Less/Unchanged>`
/// - WGSL: `@early_depth_test(greater_equal/less_equal/unchanged)`
///
/// For more, see:
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ConservativeDepth {
/// Shader may rewrite depth only with a value greater than calculated.
GreaterEqual,
/// Shader may rewrite depth smaller than one that would have been written without the modification.
LessEqual,
/// Shader may not rewrite depth value.
Unchanged,
}
/// Stage of the programmable pipeline.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
#[allow(missing_docs)] // The names are self evident
pub enum ShaderStage {
Vertex,
Fragment,
Compute,
}
/// Addressing space of variables.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum AddressSpace {
/// Function locals.
Function,
/// Private data, per invocation, mutable.
Private,
/// Workgroup shared data, mutable.
WorkGroup,
/// Uniform buffer data.
Uniform,
/// Storage buffer data, potentially mutable.
Storage { access: StorageAccess },
/// Opaque handles, such as samplers and images.
Handle,
/// Push constants.
PushConstant,
}
/// Built-in inputs and outputs.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum BuiltIn {
Position { invariant: bool },
ViewIndex,
// vertex
BaseInstance,
BaseVertex,
ClipDistance,
CullDistance,
InstanceIndex,
PointSize,
VertexIndex,
// fragment
FragDepth,
PointCoord,
FrontFacing,
PrimitiveIndex,
SampleIndex,
SampleMask,
// compute
GlobalInvocationId,
LocalInvocationId,
LocalInvocationIndex,
WorkGroupId,
WorkGroupSize,
NumWorkGroups,
// subgroup
NumSubgroups,
SubgroupId,
SubgroupSize,
SubgroupInvocationId,
}
/// Number of bytes per scalar.
pub type Bytes = u8;
/// Number of components in a vector.
#[repr(u8)]
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum VectorSize {
/// 2D vector
Bi = 2,
/// 3D vector
Tri = 3,
/// 4D vector
Quad = 4,
}
impl VectorSize {
const MAX: usize = Self::Quad as u8 as usize;
}
/// Primitive type for a scalar.
#[repr(u8)]
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ScalarKind {
/// Signed integer type.
Sint,
/// Unsigned integer type.
Uint,
/// Floating point type.
Float,
/// Boolean type.
Bool,
/// WGSL abstract integer type.
///
/// These are forbidden by validation, and should never reach backends.
AbstractInt,
/// Abstract floating-point type.
///
/// These are forbidden by validation, and should never reach backends.
AbstractFloat,
}
/// Characteristics of a scalar type.
#[derive(Clone, Copy, Debug, PartialEq, Eq, PartialOrd, Ord, Hash)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct Scalar {
/// How the value's bits are to be interpreted.
pub kind: ScalarKind,
/// This size of the value in bytes.
pub width: Bytes,
}
/// Size of an array.
#[repr(u8)]
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ArraySize {
/// The array size is constant.
Constant(std::num::NonZeroU32),
/// The array size can change at runtime.
Dynamic,
}
/// The interpolation qualifier of a binding or struct field.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum Interpolation {
/// The value will be interpolated in a perspective-correct fashion.
/// Also known as "smooth" in glsl.
Perspective,
/// Indicates that linear, non-perspective, correct
/// interpolation must be used.
/// Also known as "no_perspective" in glsl.
Linear,
/// Indicates that no interpolation will be performed.
Flat,
}
/// The sampling qualifiers of a binding or struct field.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum Sampling {
/// Interpolate the value at the center of the pixel.
Center,
/// Interpolate the value at a point that lies within all samples covered by
/// the fragment within the current primitive. In multisampling, use a
/// single value for all samples in the primitive.
Centroid,
/// Interpolate the value at each sample location. In multisampling, invoke
/// the fragment shader once per sample.
Sample,
}
/// Member of a user-defined structure.
// Clone is used only for error reporting and is not intended for end users
#[derive(Clone, Debug, Eq, Hash, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct StructMember {
pub name: Option<String>,
/// Type of the field.
pub ty: Handle<Type>,
/// For I/O structs, defines the binding.
pub binding: Option<Binding>,
/// Offset from the beginning from the struct.
pub offset: u32,
}
/// The number of dimensions an image has.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ImageDimension {
/// 1D image
D1,
/// 2D image
D2,
/// 3D image
D3,
/// Cube map
Cube,
}
bitflags::bitflags! {
/// Flags describing an image.
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
#[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)]
pub struct StorageAccess: u32 {
/// Storage can be used as a source for load ops.
const LOAD = 0x1;
/// Storage can be used as a target for store ops.
const STORE = 0x2;
}
}
/// Image storage format.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum StorageFormat {
// 8-bit formats
R8Unorm,
R8Snorm,
R8Uint,
R8Sint,
// 16-bit formats
R16Uint,
R16Sint,
R16Float,
Rg8Unorm,
Rg8Snorm,
Rg8Uint,
Rg8Sint,
// 32-bit formats
R32Uint,
R32Sint,
R32Float,
Rg16Uint,
Rg16Sint,
Rg16Float,
Rgba8Unorm,
Rgba8Snorm,
Rgba8Uint,
Rgba8Sint,
Bgra8Unorm,
// Packed 32-bit formats
Rgb10a2Uint,
Rgb10a2Unorm,
Rg11b10Float,
// 64-bit formats
Rg32Uint,
Rg32Sint,
Rg32Float,
Rgba16Uint,
Rgba16Sint,
Rgba16Float,
// 128-bit formats
Rgba32Uint,
Rgba32Sint,
Rgba32Float,
// Normalized 16-bit per channel formats
R16Unorm,
R16Snorm,
Rg16Unorm,
Rg16Snorm,
Rgba16Unorm,
Rgba16Snorm,
}
/// Sub-class of the image type.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ImageClass {
/// Regular sampled image.
Sampled {
/// Kind of values to sample.
kind: ScalarKind,
/// Multi-sampled image.
///
/// A multi-sampled image holds several samples per texel. Multi-sampled
/// images cannot have mipmaps.
multi: bool,
},
/// Depth comparison image.
Depth {
/// Multi-sampled depth image.
multi: bool,
},
/// Storage image.
Storage {
format: StorageFormat,
access: StorageAccess,
},
}
/// A data type declared in the module.
#[derive(Clone, Debug, Eq, Hash, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct Type {
/// The name of the type, if any.
pub name: Option<String>,
/// Inner structure that depends on the kind of the type.
pub inner: TypeInner,
}
/// Enum with additional information, depending on the kind of type.
#[derive(Clone, Debug, Eq, Hash, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum TypeInner {
/// Number of integral or floating-point kind.
Scalar(Scalar),
/// Vector of numbers.
Vector { size: VectorSize, scalar: Scalar },
/// Matrix of numbers.
Matrix {
columns: VectorSize,
rows: VectorSize,
scalar: Scalar,
},
/// Atomic scalar.
Atomic(Scalar),
/// Pointer to another type.
///
/// Pointers to scalars and vectors should be treated as equivalent to
/// [`ValuePointer`] types. Use the [`TypeInner::equivalent`] method to
/// compare types in a way that treats pointers correctly.
///
/// ## Pointers to non-`SIZED` types
///
/// The `base` type of a pointer may be a non-[`SIZED`] type like a
/// dynamically-sized [`Array`], or a [`Struct`] whose last member is a
/// dynamically sized array. Such pointers occur as the types of
/// [`GlobalVariable`] or [`AccessIndex`] expressions referring to
/// dynamically-sized arrays.
///
/// However, among pointers to non-`SIZED` types, only pointers to `Struct`s
/// are [`DATA`]. Pointers to dynamically sized `Array`s cannot be passed as
/// arguments, stored in variables, or held in arrays or structures. Their
/// only use is as the types of `AccessIndex` expressions.
///
/// [`SIZED`]: valid::TypeFlags::SIZED
/// [`DATA`]: valid::TypeFlags::DATA
/// [`Array`]: TypeInner::Array
/// [`Struct`]: TypeInner::Struct
/// [`ValuePointer`]: TypeInner::ValuePointer
/// [`GlobalVariable`]: Expression::GlobalVariable
/// [`AccessIndex`]: Expression::AccessIndex
Pointer {
base: Handle<Type>,
space: AddressSpace,
},
/// Pointer to a scalar or vector.
///
/// A `ValuePointer` type is equivalent to a `Pointer` whose `base` is a
/// `Scalar` or `Vector` type. This is for use in [`TypeResolution::Value`]
/// variants; see the documentation for [`TypeResolution`] for details.
///
/// Use the [`TypeInner::equivalent`] method to compare types that could be
/// pointers, to ensure that `Pointer` and `ValuePointer` types are
/// recognized as equivalent.
///
/// [`TypeResolution`]: proc::TypeResolution
/// [`TypeResolution::Value`]: proc::TypeResolution::Value
ValuePointer {
size: Option<VectorSize>,
scalar: Scalar,
space: AddressSpace,
},
/// Homogeneous list of elements.
///
/// The `base` type must be a [`SIZED`], [`DATA`] type.
///
/// ## Dynamically sized arrays
///
/// An `Array` is [`SIZED`] unless its `size` is [`Dynamic`].
/// Dynamically-sized arrays may only appear in a few situations:
///
/// - They may appear as the type of a [`GlobalVariable`], or as the last
/// member of a [`Struct`].
///
/// - They may appear as the base type of a [`Pointer`]. An
/// [`AccessIndex`] expression referring to a struct's final
/// unsized array member would have such a pointer type. However, such
/// pointer types may only appear as the types of such intermediate
/// expressions. They are not [`DATA`], and cannot be stored in
/// variables, held in arrays or structs, or passed as parameters.
///
/// [`SIZED`]: crate::valid::TypeFlags::SIZED
/// [`DATA`]: crate::valid::TypeFlags::DATA
/// [`Dynamic`]: ArraySize::Dynamic
/// [`Struct`]: TypeInner::Struct
/// [`Pointer`]: TypeInner::Pointer
/// [`AccessIndex`]: Expression::AccessIndex
Array {
base: Handle<Type>,
size: ArraySize,
stride: u32,
},
/// User-defined structure.
///
/// There must always be at least one member.
///
/// A `Struct` type is [`DATA`], and the types of its members must be
/// `DATA` as well.
///
/// Member types must be [`SIZED`], except for the final member of a
/// struct, which may be a dynamically sized [`Array`]. The
/// `Struct` type itself is `SIZED` when all its members are `SIZED`.
///
/// [`DATA`]: crate::valid::TypeFlags::DATA
/// [`SIZED`]: crate::valid::TypeFlags::SIZED
/// [`Array`]: TypeInner::Array
Struct {
members: Vec<StructMember>,
//TODO: should this be unaligned?
span: u32,
},
/// Possibly multidimensional array of texels.
Image {
dim: ImageDimension,
arrayed: bool,
//TODO: consider moving `multisampled: bool` out
class: ImageClass,
},
/// Can be used to sample values from images.
Sampler { comparison: bool },
/// Opaque object representing an acceleration structure of geometry.
AccelerationStructure,
/// Locally used handle for ray queries.
RayQuery,
/// Array of bindings.
///
/// A `BindingArray` represents an array where each element draws its value
/// from a separate bound resource. The array's element type `base` may be
/// [`Image`], [`Sampler`], or any type that would be permitted for a global
/// in the [`Uniform`] or [`Storage`] address spaces. Only global variables
/// may be binding arrays; on the host side, their values are provided by
/// [`TextureViewArray`], [`SamplerArray`], or [`BufferArray`]
/// bindings.
///
/// Since each element comes from a distinct resource, a binding array of
/// images could have images of varying sizes (but not varying dimensions;
/// they must all have the same `Image` type). Or, a binding array of
/// buffers could have elements that are dynamically sized arrays, each with
/// a different length.
///
/// Binding arrays are in the same address spaces as their underlying type.
/// As such, referring to an array of images produces an [`Image`] value
/// directly (as opposed to a pointer). The only operation permitted on
/// `BindingArray` values is indexing, which works transparently: indexing
/// a binding array of samplers yields a [`Sampler`], indexing a pointer to the
/// binding array of storage buffers produces a pointer to the storage struct.
///
/// Unlike textures and samplers, binding arrays are not [`ARGUMENT`], so
/// they cannot be passed as arguments to functions.
///
/// Naga's WGSL front end supports binding arrays with the type syntax
/// `binding_array<T, N>`.
///
/// [`Image`]: TypeInner::Image
/// [`Sampler`]: TypeInner::Sampler
/// [`Uniform`]: AddressSpace::Uniform
/// [`Storage`]: AddressSpace::Storage
/// [`TextureViewArray`]: https://docs.rs/wgpu/latest/wgpu/enum.BindingResource.html#variant.TextureViewArray
/// [`SamplerArray`]: https://docs.rs/wgpu/latest/wgpu/enum.BindingResource.html#variant.SamplerArray
/// [`DATA`]: crate::valid::TypeFlags::DATA
/// [`ARGUMENT`]: crate::valid::TypeFlags::ARGUMENT
BindingArray { base: Handle<Type>, size: ArraySize },
}
#[derive(Debug, Clone, Copy, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum Literal {
/// May not be NaN or infinity.
F64(f64),
/// May not be NaN or infinity.
F32(f32),
U32(u32),
I32(i32),
U64(u64),
I64(i64),
Bool(bool),
AbstractInt(i64),
AbstractFloat(f64),
}
/// Pipeline-overridable constant.
#[derive(Debug, PartialEq, Clone)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct Override {
pub name: Option<String>,
/// Pipeline Constant ID.
pub id: Option<u16>,
pub ty: Handle<Type>,
/// The default value of the pipeline-overridable constant.
///
/// This [`Handle`] refers to [`Module::global_expressions`], not
/// any [`Function::expressions`] arena.
pub init: Option<Handle<Expression>>,
}
/// Constant value.
#[derive(Debug, PartialEq, Clone)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct Constant {
pub name: Option<String>,
pub ty: Handle<Type>,
/// The value of the constant.
///
/// This [`Handle`] refers to [`Module::global_expressions`], not
/// any [`Function::expressions`] arena.
pub init: Handle<Expression>,
}
/// Describes how an input/output variable is to be bound.
#[derive(Clone, Debug, Eq, PartialEq, Hash)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum Binding {
/// Built-in shader variable.
BuiltIn(BuiltIn),
/// Indexed location.
///
/// Values passed from the [`Vertex`] stage to the [`Fragment`] stage must
/// have their `interpolation` defaulted (i.e. not `None`) by the front end
/// as appropriate for that language.
///
/// For other stages, we permit interpolations even though they're ignored.
/// When a front end is parsing a struct type, it usually doesn't know what
/// stages will be using it for IO, so it's easiest if it can apply the
/// defaults to anything with a `Location` binding, just in case.
///
/// For anything other than floating-point scalars and vectors, the
/// interpolation must be `Flat`.
///
/// [`Vertex`]: crate::ShaderStage::Vertex
/// [`Fragment`]: crate::ShaderStage::Fragment
Location {
location: u32,
/// Indicates the 2nd input to the blender when dual-source blending.
second_blend_source: bool,
interpolation: Option<Interpolation>,
sampling: Option<Sampling>,
},
}
/// Pipeline binding information for global resources.
#[derive(Clone, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct ResourceBinding {
/// The bind group index.
pub group: u32,
/// Binding number within the group.
pub binding: u32,
}
/// Variable defined at module level.
#[derive(Clone, Debug, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct GlobalVariable {
/// Name of the variable, if any.
pub name: Option<String>,
/// How this variable is to be stored.
pub space: AddressSpace,
/// For resources, defines the binding point.
pub binding: Option<ResourceBinding>,
/// The type of this variable.
pub ty: Handle<Type>,
/// Initial value for this variable.
///
/// Expression handle lives in global_expressions
pub init: Option<Handle<Expression>>,
}
/// Variable defined at function level.
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct LocalVariable {
/// Name of the variable, if any.
pub name: Option<String>,
/// The type of this variable.
pub ty: Handle<Type>,
/// Initial value for this variable.
///
/// This handle refers to this `LocalVariable`'s function's
/// [`expressions`] arena, but it is required to be an evaluated
/// override expression.
///
/// [`expressions`]: Function::expressions
pub init: Option<Handle<Expression>>,
}
/// Operation that can be applied on a single value.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum UnaryOperator {
Negate,
LogicalNot,
BitwiseNot,
}
/// Operation that can be applied on two values.
///
/// ## Arithmetic type rules
///
/// The arithmetic operations `Add`, `Subtract`, `Multiply`, `Divide`, and
/// `Modulo` can all be applied to [`Scalar`] types other than [`Bool`], or
/// [`Vector`]s thereof. Both operands must have the same type.
///
/// `Add` and `Subtract` can also be applied to [`Matrix`] values. Both operands
/// must have the same type.
///
/// `Multiply` supports additional cases:
///
/// - A [`Matrix`] or [`Vector`] can be multiplied by a scalar [`Float`],
/// either on the left or the right.
///
/// - A [`Matrix`] on the left can be multiplied by a [`Vector`] on the right
/// if the matrix has as many columns as the vector has components (`matCxR
/// * VecC`).
///
/// - A [`Vector`] on the left can be multiplied by a [`Matrix`] on the right
/// if the matrix has as many rows as the vector has components (`VecR *
/// matCxR`).
///
/// - Two matrices can be multiplied if the left operand has as many columns
/// as the right operand has rows (`matNxR * matCxN`).
///
/// In all the above `Multiply` cases, the byte widths of the underlying scalar
/// types of both operands must be the same.
///
/// Note that `Multiply` supports mixed vector and scalar operations directly,
/// whereas the other arithmetic operations require an explicit [`Splat`] for
/// mixed-type use.
///
/// [`Scalar`]: TypeInner::Scalar
/// [`Vector`]: TypeInner::Vector
/// [`Matrix`]: TypeInner::Matrix
/// [`Float`]: ScalarKind::Float
/// [`Bool`]: ScalarKind::Bool
/// [`Splat`]: Expression::Splat
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum BinaryOperator {
Add,
Subtract,
Multiply,
Divide,
/// Equivalent of the WGSL's `%` operator or SPIR-V's `OpFRem`
Modulo,
Equal,
NotEqual,
Less,
LessEqual,
Greater,
GreaterEqual,
And,
ExclusiveOr,
InclusiveOr,
LogicalAnd,
LogicalOr,
ShiftLeft,
/// Right shift carries the sign of signed integers only.
ShiftRight,
}
/// Function on an atomic value.
///
/// Note: these do not include load/store, which use the existing
/// [`Expression::Load`] and [`Statement::Store`].
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum AtomicFunction {
Add,
Subtract,
And,
ExclusiveOr,
InclusiveOr,
Min,
Max,
Exchange { compare: Option<Handle<Expression>> },
}
/// Hint at which precision to compute a derivative.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum DerivativeControl {
Coarse,
Fine,
None,
}
/// Axis on which to compute a derivative.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum DerivativeAxis {
X,
Y,
Width,
}
/// Built-in shader function for testing relation between values.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum RelationalFunction {
All,
Any,
IsNan,
IsInf,
}
/// Built-in shader function for math.
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum MathFunction {
// comparison
Abs,
Min,
Max,
Clamp,
Saturate,
// trigonometry
Cos,
Cosh,
Sin,
Sinh,
Tan,
Tanh,
Acos,
Asin,
Atan,
Atan2,
Asinh,
Acosh,
Atanh,
Radians,
Degrees,
// decomposition
Ceil,
Floor,
Round,
Fract,
Trunc,
Modf,
Frexp,
Ldexp,
// exponent
Exp,
Exp2,
Log,
Log2,
Pow,
// geometry
Dot,
Outer,
Cross,
Distance,
Length,
Normalize,
FaceForward,
Reflect,
Refract,
// computational
Sign,
Fma,
Mix,
Step,
SmoothStep,
Sqrt,
InverseSqrt,
Inverse,
Transpose,
Determinant,
// bits
CountTrailingZeros,
CountLeadingZeros,
CountOneBits,
ReverseBits,
ExtractBits,
InsertBits,
FindLsb,
FindMsb,
// data packing
Pack4x8snorm,
Pack4x8unorm,
Pack2x16snorm,
Pack2x16unorm,
Pack2x16float,
// data unpacking
Unpack4x8snorm,
Unpack4x8unorm,
Unpack2x16snorm,
Unpack2x16unorm,
Unpack2x16float,
}
/// Sampling modifier to control the level of detail.
#[derive(Clone, Copy, Debug, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum SampleLevel {
Auto,
Zero,
Exact(Handle<Expression>),
Bias(Handle<Expression>),
Gradient {
x: Handle<Expression>,
y: Handle<Expression>,
},
}
/// Type of an image query.
#[derive(Clone, Copy, Debug, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum ImageQuery {
/// Get the size at the specified level.
///
/// The return value is a `u32` for 1D images, and a `vecN<u32>`
/// for an image with dimensions N > 2.
Size {
/// If `None`, the base level is considered.
level: Option<Handle<Expression>>,
},
/// Get the number of mipmap levels, a `u32`.
NumLevels,
/// Get the number of array layers, a `u32`.
NumLayers,
/// Get the number of samples, a `u32`.
NumSamples,
}
/// Component selection for a vector swizzle.
#[repr(u8)]
#[derive(Clone, Copy, Debug, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum SwizzleComponent {
///
X = 0,
///
Y = 1,
///
Z = 2,
///
W = 3,
}
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum GatherMode {
/// All gather from the active lane with the smallest index
BroadcastFirst,
/// All gather from the same lane at the index given by the expression
Broadcast(Handle<Expression>),
/// Each gathers from a different lane at the index given by the expression
Shuffle(Handle<Expression>),
/// Each gathers from their lane plus the shift given by the expression
ShuffleDown(Handle<Expression>),
/// Each gathers from their lane minus the shift given by the expression
ShuffleUp(Handle<Expression>),
/// Each gathers from their lane xored with the given by the expression
ShuffleXor(Handle<Expression>),
}
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum SubgroupOperation {
All = 0,
Any = 1,
Add = 2,
Mul = 3,
Min = 4,
Max = 5,
And = 6,
Or = 7,
Xor = 8,
}
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum CollectiveOperation {
Reduce = 0,
InclusiveScan = 1,
ExclusiveScan = 2,
}
bitflags::bitflags! {
/// Memory barrier flags.
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
#[derive(Clone, Copy, Debug, Default, Eq, PartialEq)]
pub struct Barrier: u32 {
/// Barrier affects all `AddressSpace::Storage` accesses.
const STORAGE = 1 << 0;
/// Barrier affects all `AddressSpace::WorkGroup` accesses.
const WORK_GROUP = 1 << 1;
/// Barrier synchronizes execution across all invocations within a subgroup that exectue this instruction.
const SUB_GROUP = 1 << 2;
}
}
/// An expression that can be evaluated to obtain a value.
///
/// This is a Single Static Assignment (SSA) scheme similar to SPIR-V.
#[derive(Clone, Debug, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum Expression {
/// Literal.
Literal(Literal),
/// Constant value.
Constant(Handle<Constant>),
/// Pipeline-overridable constant.
Override(Handle<Override>),
/// Zero value of a type.
ZeroValue(Handle<Type>),
/// Composite expression.
Compose {
ty: Handle<Type>,
components: Vec<Handle<Expression>>,
},
/// Array access with a computed index.
///
/// ## Typing rules
///
/// The `base` operand must be some composite type: [`Vector`], [`Matrix`],
/// [`Array`], a [`Pointer`] to one of those, or a [`ValuePointer`] with a
/// `size`.
///
/// The `index` operand must be an integer, signed or unsigned.
///
/// Indexing a [`Vector`] or [`Array`] produces a value of its element type.
/// Indexing a [`Matrix`] produces a [`Vector`].
///
/// Indexing a [`Pointer`] to any of the above produces a pointer to the
/// element/component type, in the same [`space`]. In the case of [`Array`],
/// the result is an actual [`Pointer`], but for vectors and matrices, there
/// may not be any type in the arena representing the component's type, so
/// those produce [`ValuePointer`] types equivalent to the appropriate
/// [`Pointer`].
///
/// ## Dynamic indexing restrictions
///
/// To accommodate restrictions in some of the shader languages that Naga
/// targets, it is not permitted to subscript a matrix or array with a
/// dynamically computed index unless that matrix or array appears behind a
/// pointer. In other words, if the inner type of `base` is [`Array`] or
/// [`Matrix`], then `index` must be a constant. But if the type of `base`
/// is a [`Pointer`] to an array or matrix or a [`ValuePointer`] with a
/// `size`, then the index may be any expression of integer type.
///
/// You can use the [`Expression::is_dynamic_index`] method to determine
/// whether a given index expression requires matrix or array base operands
/// to be behind a pointer.
///
/// (It would be simpler to always require the use of `AccessIndex` when
/// subscripting arrays and matrices that are not behind pointers, but to
/// accommodate existing front ends, Naga also permits `Access`, with a
/// restricted `index`.)
///
/// [`Vector`]: TypeInner::Vector
/// [`Matrix`]: TypeInner::Matrix
/// [`Array`]: TypeInner::Array
/// [`Pointer`]: TypeInner::Pointer
/// [`space`]: TypeInner::Pointer::space
/// [`ValuePointer`]: TypeInner::ValuePointer
/// [`Float`]: ScalarKind::Float
Access {
base: Handle<Expression>,
index: Handle<Expression>,
},
/// Access the same types as [`Access`], plus [`Struct`] with a known index.
///
/// [`Access`]: Expression::Access
/// [`Struct`]: TypeInner::Struct
AccessIndex {
base: Handle<Expression>,
index: u32,
},
/// Splat scalar into a vector.
Splat {
size: VectorSize,
value: Handle<Expression>,
},
/// Vector swizzle.
Swizzle {
size: VectorSize,
vector: Handle<Expression>,
pattern: [SwizzleComponent; 4],
},
/// Reference a function parameter, by its index.
///
/// A `FunctionArgument` expression evaluates to a pointer to the argument's
/// value. You must use a [`Load`] expression to retrieve its value, or a
/// [`Store`] statement to assign it a new value.
///
/// [`Load`]: Expression::Load
/// [`Store`]: Statement::Store
FunctionArgument(u32),
/// Reference a global variable.
///
/// If the given `GlobalVariable`'s [`space`] is [`AddressSpace::Handle`],
/// then the variable stores some opaque type like a sampler or an image,
/// and a `GlobalVariable` expression referring to it produces the
/// variable's value directly.
///
/// For any other address space, a `GlobalVariable` expression produces a
/// pointer to the variable's value. You must use a [`Load`] expression to
/// retrieve its value, or a [`Store`] statement to assign it a new value.
///
/// [`space`]: GlobalVariable::space
/// [`Load`]: Expression::Load
/// [`Store`]: Statement::Store
GlobalVariable(Handle<GlobalVariable>),
/// Reference a local variable.
///
/// A `LocalVariable` expression evaluates to a pointer to the variable's value.
/// You must use a [`Load`](Expression::Load) expression to retrieve its value,
/// or a [`Store`](Statement::Store) statement to assign it a new value.
LocalVariable(Handle<LocalVariable>),
/// Load a value indirectly.
///
/// For [`TypeInner::Atomic`] the result is a corresponding scalar.
/// For other types behind the `pointer<T>`, the result is `T`.
Load { pointer: Handle<Expression> },
/// Sample a point from a sampled or a depth image.
ImageSample {
image: Handle<Expression>,
sampler: Handle<Expression>,
/// If Some(), this operation is a gather operation
/// on the selected component.
gather: Option<SwizzleComponent>,
coordinate: Handle<Expression>,
array_index: Option<Handle<Expression>>,
/// Expression handle lives in global_expressions
offset: Option<Handle<Expression>>,
level: SampleLevel,
depth_ref: Option<Handle<Expression>>,
},
/// Load a texel from an image.
///
/// For most images, this returns a four-element vector of the same
/// [`ScalarKind`] as the image. If the format of the image does not have
/// four components, default values are provided: the first three components
/// (typically R, G, and B) default to zero, and the final component
/// (typically alpha) defaults to one.
///
/// However, if the image's [`class`] is [`Depth`], then this returns a
/// [`Float`] scalar value.
///
/// [`ScalarKind`]: ScalarKind
/// [`class`]: TypeInner::Image::class
/// [`Depth`]: ImageClass::Depth
/// [`Float`]: ScalarKind::Float
ImageLoad {
/// The image to load a texel from. This must have type [`Image`]. (This
/// will necessarily be a [`GlobalVariable`] or [`FunctionArgument`]
/// expression, since no other expressions are allowed to have that
/// type.)
///
/// [`Image`]: TypeInner::Image
/// [`GlobalVariable`]: Expression::GlobalVariable
/// [`FunctionArgument`]: Expression::FunctionArgument
image: Handle<Expression>,
/// The coordinate of the texel we wish to load. This must be a scalar
/// for [`D1`] images, a [`Bi`] vector for [`D2`] images, and a [`Tri`]
/// vector for [`D3`] images. (Array indices, sample indices, and
/// explicit level-of-detail values are supplied separately.) Its
/// component type must be [`Sint`].
///
/// [`D1`]: ImageDimension::D1
/// [`D2`]: ImageDimension::D2
/// [`D3`]: ImageDimension::D3
/// [`Bi`]: VectorSize::Bi
/// [`Tri`]: VectorSize::Tri
/// [`Sint`]: ScalarKind::Sint
coordinate: Handle<Expression>,
/// The index into an arrayed image. If the [`arrayed`] flag in
/// `image`'s type is `true`, then this must be `Some(expr)`, where
/// `expr` is a [`Sint`] scalar. Otherwise, it must be `None`.
///
/// [`arrayed`]: TypeInner::Image::arrayed
/// [`Sint`]: ScalarKind::Sint
array_index: Option<Handle<Expression>>,
/// A sample index, for multisampled [`Sampled`] and [`Depth`] images.
///
/// [`Sampled`]: ImageClass::Sampled
/// [`Depth`]: ImageClass::Depth
sample: Option<Handle<Expression>>,
/// A level of detail, for mipmapped images.
///
/// This must be present when accessing non-multisampled
/// [`Sampled`] and [`Depth`] images, even if only the
/// full-resolution level is present (in which case the only
/// valid level is zero).
///
/// [`Sampled`]: ImageClass::Sampled
/// [`Depth`]: ImageClass::Depth
level: Option<Handle<Expression>>,
},
/// Query information from an image.
ImageQuery {
image: Handle<Expression>,
query: ImageQuery,
},
/// Apply an unary operator.
Unary {
op: UnaryOperator,
expr: Handle<Expression>,
},
/// Apply a binary operator.
Binary {
op: BinaryOperator,
left: Handle<Expression>,
right: Handle<Expression>,
},
/// Select between two values based on a condition.
///
/// Note that, because expressions have no side effects, it is unobservable
/// whether the non-selected branch is evaluated.
Select {
/// Boolean expression
condition: Handle<Expression>,
accept: Handle<Expression>,
reject: Handle<Expression>,
},
/// Compute the derivative on an axis.
Derivative {
axis: DerivativeAxis,
ctrl: DerivativeControl,
expr: Handle<Expression>,
},
/// Call a relational function.
Relational {
fun: RelationalFunction,
argument: Handle<Expression>,
},
/// Call a math function
Math {
fun: MathFunction,
arg: Handle<Expression>,
arg1: Option<Handle<Expression>>,
arg2: Option<Handle<Expression>>,
arg3: Option<Handle<Expression>>,
},
/// Cast a simple type to another kind.
As {
/// Source expression, which can only be a scalar or a vector.
expr: Handle<Expression>,
/// Target scalar kind.
kind: ScalarKind,
/// If provided, converts to the specified byte width.
/// Otherwise, bitcast.
convert: Option<Bytes>,
},
/// Result of calling another function.
CallResult(Handle<Function>),
/// Result of an atomic operation.
AtomicResult { ty: Handle<Type>, comparison: bool },
/// Result of a [`WorkGroupUniformLoad`] statement.
///
/// [`WorkGroupUniformLoad`]: Statement::WorkGroupUniformLoad
WorkGroupUniformLoadResult {
/// The type of the result
ty: Handle<Type>,
},
/// Get the length of an array.
/// The expression must resolve to a pointer to an array with a dynamic size.
///
/// This doesn't match the semantics of spirv's `OpArrayLength`, which must be passed
/// a pointer to a structure containing a runtime array in its' last field.
ArrayLength(Handle<Expression>),
/// Result of a [`Proceed`] [`RayQuery`] statement.
///
/// [`Proceed`]: RayQueryFunction::Proceed
/// [`RayQuery`]: Statement::RayQuery
RayQueryProceedResult,
/// Return an intersection found by `query`.
///
/// If `committed` is true, return the committed result available when
RayQueryGetIntersection {
query: Handle<Expression>,
committed: bool,
},
/// Result of a [`SubgroupBallot`] statement.
///
/// [`SubgroupBallot`]: Statement::SubgroupBallot
SubgroupBallotResult,
/// Result of a [`SubgroupCollectiveOperation`] or [`SubgroupGather`] statement.
///
/// [`SubgroupCollectiveOperation`]: Statement::SubgroupCollectiveOperation
/// [`SubgroupGather`]: Statement::SubgroupGather
SubgroupOperationResult { ty: Handle<Type> },
}
pub use block::Block;
/// The value of the switch case.
#[derive(Clone, Copy, Debug, Eq, Hash, PartialEq)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum SwitchValue {
I32(i32),
U32(u32),
Default,
}
/// A case for a switch statement.
// Clone is used only for error reporting and is not intended for end users
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct SwitchCase {
/// Value, upon which the case is considered true.
pub value: SwitchValue,
/// Body of the case.
pub body: Block,
/// If true, the control flow continues to the next case in the list,
/// or default.
pub fall_through: bool,
}
/// An operation that a [`RayQuery` statement] applies to its [`query`] operand.
///
/// [`RayQuery` statement]: Statement::RayQuery
/// [`query`]: Statement::RayQuery::query
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum RayQueryFunction {
/// Initialize the `RayQuery` object.
Initialize {
/// The acceleration structure within which this query should search for hits.
///
/// The expression must be an [`AccelerationStructure`].
///
/// [`AccelerationStructure`]: TypeInner::AccelerationStructure
acceleration_structure: Handle<Expression>,
#[allow(rustdoc::private_intra_doc_links)]
/// A struct of detailed parameters for the ray query.
///
/// This expression should have the struct type given in
/// [`SpecialTypes::ray_desc`]. This is available in the WGSL
/// front end as the `RayDesc` type.
descriptor: Handle<Expression>,
},
/// Start or continue the query given by the statement's [`query`] operand.
///
/// After executing this statement, the `result` expression is a
/// [`Bool`] scalar indicating whether there are more intersection
/// candidates to consider.
///
/// [`query`]: Statement::RayQuery::query
/// [`Bool`]: ScalarKind::Bool
Proceed {
result: Handle<Expression>,
},
Terminate,
}
//TODO: consider removing `Clone`. It's not valid to clone `Statement::Emit` anyway.
/// Instructions which make up an executable block.
// Clone is used only for error reporting and is not intended for end users
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum Statement {
/// Emit a range of expressions, visible to all statements that follow in this block.
///
/// See the [module-level documentation][emit] for details.
///
/// [emit]: index.html#expression-evaluation-time
Emit(Range<Expression>),
/// A block containing more statements, to be executed sequentially.
Block(Block),
/// Conditionally executes one of two blocks, based on the value of the condition.
///
/// Naga IR does not have "phi" instructions. If you need to use
/// values computed in an `accept` or `reject` block after the `If`,
/// store them in a [`LocalVariable`].
If {
condition: Handle<Expression>, //bool
accept: Block,
reject: Block,
},
/// Conditionally executes one of multiple blocks, based on the value of the selector.
///
/// Each case must have a distinct [`value`], exactly one of which must be
/// [`Default`]. The `Default` may appear at any position, and covers all
/// values not explicitly appearing in other cases. A `Default` appearing in
/// the midst of the list of cases does not shadow the cases that follow.
///
/// Some backend languages don't support fallthrough (HLSL due to FXC,
/// WGSL), and may translate fallthrough cases in the IR by duplicating
/// code. However, all backend languages do support cases selected by
/// multiple values, like `case 1: case 2: case 3: { ... }`. This is
/// represented in the IR as a series of fallthrough cases with empty
/// bodies, except for the last.
///
/// Naga IR does not have "phi" instructions. If you need to use
/// values computed in a [`SwitchCase::body`] block after the `Switch`,
/// store them in a [`LocalVariable`].
///
/// [`value`]: SwitchCase::value
/// [`body`]: SwitchCase::body
/// [`Default`]: SwitchValue::Default
Switch {
selector: Handle<Expression>,
cases: Vec<SwitchCase>,
},
/// Executes a block repeatedly.
///
/// Each iteration of the loop executes the `body` block, followed by the
/// `continuing` block.
///
/// Executing a [`Break`], [`Return`] or [`Kill`] statement exits the loop.
///
/// A [`Continue`] statement in `body` jumps to the `continuing` block. The
/// `continuing` block is meant to be used to represent structures like the
/// third expression of a C-style `for` loop head, to which `continue`
/// statements in the loop's body jump.
///
/// The `continuing` block and its substatements must not contain `Return`
/// or `Kill` statements, or any `Break` or `Continue` statements targeting
/// this loop. (It may have `Break` and `Continue` statements targeting
/// loops or switches nested within the `continuing` block.) Expressions
/// emitted in `body` are in scope in `continuing`.
///
/// If present, `break_if` is an expression which is evaluated after the
/// continuing block. Expressions emitted in `body` or `continuing` are
/// considered to be in scope. If the expression's value is true, control
/// continues after the `Loop` statement, rather than branching back to the
/// top of body as usual. The `break_if` expression corresponds to a "break
/// if" statement in WGSL, or a loop whose back edge is an
/// `OpBranchConditional` instruction in SPIR-V.
///
/// Naga IR does not have "phi" instructions. If you need to use
/// values computed in a `body` or `continuing` block after the
/// `Loop`, store them in a [`LocalVariable`].
///
/// [`Break`]: Statement::Break
/// [`Continue`]: Statement::Continue
/// [`Kill`]: Statement::Kill
/// [`Return`]: Statement::Return
/// [`break if`]: Self::Loop::break_if
Loop {
body: Block,
continuing: Block,
break_if: Option<Handle<Expression>>,
},
/// Exits the innermost enclosing [`Loop`] or [`Switch`].
///
/// A `Break` statement may only appear within a [`Loop`] or [`Switch`]
/// statement. It may not break out of a [`Loop`] from within the loop's
/// `continuing` block.
///
/// [`Loop`]: Statement::Loop
/// [`Switch`]: Statement::Switch
Break,
/// Skips to the `continuing` block of the innermost enclosing [`Loop`].
///
/// A `Continue` statement may only appear within the `body` block of the
/// innermost enclosing [`Loop`] statement. It must not appear within that
/// loop's `continuing` block.
///
/// [`Loop`]: Statement::Loop
Continue,
/// Returns from the function (possibly with a value).
///
/// `Return` statements are forbidden within the `continuing` block of a
/// [`Loop`] statement.
///
/// [`Loop`]: Statement::Loop
Return { value: Option<Handle<Expression>> },
/// Aborts the current shader execution.
///
/// `Kill` statements are forbidden within the `continuing` block of a
/// [`Loop`] statement.
///
/// [`Loop`]: Statement::Loop
Kill,
/// Synchronize invocations within the work group.
/// The `Barrier` flags control which memory accesses should be synchronized.
/// If empty, this becomes purely an execution barrier.
Barrier(Barrier),
/// Stores a value at an address.
///
/// For [`TypeInner::Atomic`] type behind the pointer, the value
/// has to be a corresponding scalar.
/// For other types behind the `pointer<T>`, the value is `T`.
///
/// This statement is a barrier for any operations on the
/// `Expression::LocalVariable` or `Expression::GlobalVariable`
/// that is the destination of an access chain, started
/// from the `pointer`.
Store {
pointer: Handle<Expression>,
value: Handle<Expression>,
},
/// Stores a texel value to an image.
///
/// The `image`, `coordinate`, and `array_index` fields have the same
/// meanings as the corresponding operands of an [`ImageLoad`] expression;
/// see that documentation for details. Storing into multisampled images or
/// images with mipmaps is not supported, so there are no `level` or
/// `sample` operands.
///
/// This statement is a barrier for any operations on the corresponding
/// [`Expression::GlobalVariable`] for this image.
///
/// [`ImageLoad`]: Expression::ImageLoad
ImageStore {
image: Handle<Expression>,
coordinate: Handle<Expression>,
array_index: Option<Handle<Expression>>,
value: Handle<Expression>,
},
/// Atomic function.
Atomic {
/// Pointer to an atomic value.
pointer: Handle<Expression>,
/// Function to run on the atomic.
fun: AtomicFunction,
/// Value to use in the function.
value: Handle<Expression>,
/// [`AtomicResult`] expression representing this function's result.
///
/// [`AtomicResult`]: crate::Expression::AtomicResult
result: Handle<Expression>,
},
/// Load uniformly from a uniform pointer in the workgroup address space.
///
/// Corresponds to the [`workgroupUniformLoad`](https://www.w3.org/TR/WGSL/#workgroupUniformLoad-builtin)
/// built-in function of wgsl, and has the same barrier semantics
WorkGroupUniformLoad {
/// This must be of type [`Pointer`] in the [`WorkGroup`] address space
///
/// [`Pointer`]: TypeInner::Pointer
/// [`WorkGroup`]: AddressSpace::WorkGroup
pointer: Handle<Expression>,
/// The [`WorkGroupUniformLoadResult`] expression representing this load's result.
///
/// [`WorkGroupUniformLoadResult`]: Expression::WorkGroupUniformLoadResult
result: Handle<Expression>,
},
/// Calls a function.
///
/// If the `result` is `Some`, the corresponding expression has to be
/// `Expression::CallResult`, and this statement serves as a barrier for any
/// operations on that expression.
Call {
function: Handle<Function>,
arguments: Vec<Handle<Expression>>,
result: Option<Handle<Expression>>,
},
RayQuery {
/// The [`RayQuery`] object this statement operates on.
///
/// [`RayQuery`]: TypeInner::RayQuery
query: Handle<Expression>,
/// The specific operation we're performing on `query`.
fun: RayQueryFunction,
},
/// Calculate a bitmask using a boolean from each active thread in the subgroup
SubgroupBallot {
/// The [`SubgroupBallotResult`] expression representing this load's result.
///
/// [`SubgroupBallotResult`]: Expression::SubgroupBallotResult
result: Handle<Expression>,
/// The value from this thread to store in the ballot
predicate: Option<Handle<Expression>>,
},
/// Gather a value from another active thread in the subgroup
SubgroupGather {
/// Specifies which thread to gather from
mode: GatherMode,
/// The value to broadcast over
argument: Handle<Expression>,
/// The [`SubgroupOperationResult`] expression representing this load's result.
///
/// [`SubgroupOperationResult`]: Expression::SubgroupOperationResult
result: Handle<Expression>,
},
/// Compute a collective operation across all active threads in the subgroup
SubgroupCollectiveOperation {
/// What operation to compute
op: SubgroupOperation,
/// How to combine the results
collective_op: CollectiveOperation,
/// The value to compute over
argument: Handle<Expression>,
/// The [`SubgroupOperationResult`] expression representing this load's result.
///
/// [`SubgroupOperationResult`]: Expression::SubgroupOperationResult
result: Handle<Expression>,
},
}
/// A function argument.
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct FunctionArgument {
/// Name of the argument, if any.
pub name: Option<String>,
/// Type of the argument.
pub ty: Handle<Type>,
/// For entry points, an argument has to have a binding
/// unless it's a structure.
pub binding: Option<Binding>,
}
/// A function result.
#[derive(Clone, Debug)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct FunctionResult {
/// Type of the result.
pub ty: Handle<Type>,
/// For entry points, the result has to have a binding
/// unless it's a structure.
pub binding: Option<Binding>,
}
/// A function defined in the module.
#[derive(Debug, Default, Clone)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct Function {
/// Name of the function, if any.
pub name: Option<String>,
/// Information about function argument.
pub arguments: Vec<FunctionArgument>,
/// The result of this function, if any.
pub result: Option<FunctionResult>,
/// Local variables defined and used in the function.
pub local_variables: Arena<LocalVariable>,
/// Expressions used inside this function.
///
/// An `Expression` must occur before all other `Expression`s that use its
/// value.
pub expressions: Arena<Expression>,
/// Map of expressions that have associated variable names
pub named_expressions: NamedExpressions,
/// Block of instructions comprising the body of the function.
pub body: Block,
}
/// The main function for a pipeline stage.
///
/// An [`EntryPoint`] is a [`Function`] that serves as the main function for a
/// graphics or compute pipeline stage. For example, an `EntryPoint` whose
/// [`stage`] is [`ShaderStage::Vertex`] can serve as a graphics pipeline's
/// vertex shader.
///
/// Since an entry point is called directly by the graphics or compute pipeline,
/// not by other WGSL functions, you must specify what the pipeline should pass
/// as the entry point's arguments, and what values it will return. For example,
/// a vertex shader needs a vertex's attributes as its arguments, but if it's
/// used for instanced draw calls, it will also want to know the instance id.
/// The vertex shader's return value will usually include an output vertex
/// position, and possibly other attributes to be interpolated and passed along
/// to a fragment shader.
///
/// To specify this, the arguments and result of an `EntryPoint`'s [`function`]
/// must each have a [`Binding`], or be structs whose members all have
/// `Binding`s. This associates every value passed to or returned from the entry
/// point with either a [`BuiltIn`] or a [`Location`]:
///
/// - A [`BuiltIn`] has special semantics, usually specific to its pipeline
/// stage. For example, the result of a vertex shader can include a
/// [`BuiltIn::Position`] value, which determines the position of a vertex
/// of a rendered primitive. Or, a compute shader might take an argument
/// whose binding is [`BuiltIn::WorkGroupSize`], through which the compute
/// pipeline would pass the number of invocations in your workgroup.
///
/// - A [`Location`] indicates user-defined IO to be passed from one pipeline
/// stage to the next. For example, a vertex shader might also produce a
/// `uv` texture location as a user-defined IO value.
///
/// In other words, the pipeline stage's input and output interface are
/// determined by the bindings of the arguments and result of the `EntryPoint`'s
/// [`function`].
///
/// [`Function`]: crate::Function
/// [`Location`]: Binding::Location
/// [`function`]: EntryPoint::function
/// [`stage`]: EntryPoint::stage
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct EntryPoint {
/// Name of this entry point, visible externally.
///
/// Entry point names for a given `stage` must be distinct within a module.
pub name: String,
/// Shader stage.
pub stage: ShaderStage,
/// Early depth test for fragment stages.
pub early_depth_test: Option<EarlyDepthTest>,
/// Workgroup size for compute stages
pub workgroup_size: [u32; 3],
/// The entrance function.
pub function: Function,
}
/// Return types predeclared for the frexp, modf, and atomicCompareExchangeWeak built-in functions.
///
/// These cannot be spelled in WGSL source.
///
/// Stored in [`SpecialTypes::predeclared_types`] and created by [`Module::generate_predeclared_type`].
#[derive(Debug, PartialEq, Eq, Hash, Clone)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum PredeclaredType {
AtomicCompareExchangeWeakResult(Scalar),
ModfResult {
size: Option<VectorSize>,
width: Bytes,
},
FrexpResult {
size: Option<VectorSize>,
width: Bytes,
},
}
/// Set of special types that can be optionally generated by the frontends.
#[derive(Debug, Default, Clone)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct SpecialTypes {
/// Type for `RayDesc`.
///
/// Call [`Module::generate_ray_desc_type`] to populate this if
/// needed and return the handle.
pub ray_desc: Option<Handle<Type>>,
/// Type for `RayIntersection`.
///
/// Call [`Module::generate_ray_intersection_type`] to populate
/// this if needed and return the handle.
pub ray_intersection: Option<Handle<Type>>,
/// Types for predeclared wgsl types instantiated on demand.
///
/// Call [`Module::generate_predeclared_type`] to populate this if
/// needed and return the handle.
pub predeclared_types: FastIndexMap<PredeclaredType, Handle<Type>>,
}
/// Shader module.
///
/// A module is a set of constants, global variables and functions, as well as
/// the types required to define them.
///
/// Some functions are marked as entry points, to be used in a certain shader stage.
///
/// To create a new module, use the `Default` implementation.
/// Alternatively, you can load an existing shader using one of the [available front ends][front].
///
/// When finished, you can export modules using one of the [available backends][back].
#[derive(Debug, Default, Clone)]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub struct Module {
/// Arena for the types defined in this module.
pub types: UniqueArena<Type>,
/// Dictionary of special type handles.
pub special_types: SpecialTypes,
/// Arena for the constants defined in this module.
pub constants: Arena<Constant>,
/// Arena for the pipeline-overridable constants defined in this module.
pub overrides: Arena<Override>,
/// Arena for the global variables defined in this module.
pub global_variables: Arena<GlobalVariable>,
/// [Constant expressions] and [override expressions] used by this module.
///
/// Each `Expression` must occur in the arena before any
/// `Expression` that uses its value.
///
/// [Constant expressions]: index.html#constant-expressions
/// [override expressions]: index.html#override-expressions
pub global_expressions: Arena<Expression>,
/// Arena for the functions defined in this module.
///
/// Each function must appear in this arena strictly before all its callers.
/// Recursion is not supported.
pub functions: Arena<Function>,
/// Entry points.
pub entry_points: Vec<EntryPoint>,
}