/*! 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`, 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` stores a series of `T` values, indexed by [`Handle`](Handle) values, which are just wrappers around integer indexes. For example, a `Function`'s expressions are stored in an `Arena`, and compound expressions refer to their sub-expressions via `Handle` 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 = rustc_hash::FxHashMap; /// Hash set that is faster but not resilient to DoS attacks. pub type FastHashSet = rustc_hash::FxHashSet; /// Insertion-order-preserving hash set (`IndexSet`), but with the same /// hasher as `FastHashSet` (faster but not resilient to DoS attacks). pub type FastIndexSet = indexmap::IndexSet>; /// Insertion-order-preserving hash map (`IndexMap`), but with the same /// hasher as `FastHashMap` (faster but not resilient to DoS attacks). pub type FastIndexMap = indexmap::IndexMap>; /// Map of expressions that have associated variable names pub(crate) type NamedExpressions = FastIndexMap, 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: /// - /// - /// - #[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, } /// Enables adjusting depth without disabling early Z. /// /// To use in a shader: /// - GLSL: `layout (depth_) 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` /// - 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, /// Type of the field. pub ty: Handle, /// For I/O structs, defines the binding. pub binding: Option, /// 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, /// 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, 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, 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, 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, //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`. /// /// [`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 /// [`BufferArray`]: https://docs.rs/wgpu/latest/wgpu/enum.BindingResource.html#variant.BufferArray /// [`DATA`]: crate::valid::TypeFlags::DATA /// [`ARGUMENT`]: crate::valid::TypeFlags::ARGUMENT /// [naga#1864]: https://github.com/gfx-rs/naga/issues/1864 BindingArray { base: Handle, 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, /// Pipeline Constant ID. pub id: Option, pub ty: Handle, /// The default value of the pipeline-overridable constant. /// /// This [`Handle`] refers to [`Module::global_expressions`], not /// any [`Function::expressions`] arena. pub init: Option>, } /// 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, pub ty: Handle, /// The value of the constant. /// /// This [`Handle`] refers to [`Module::global_expressions`], not /// any [`Function::expressions`] arena. pub init: Handle, } /// 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, sampling: Option, }, } /// 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, /// How this variable is to be stored. pub space: AddressSpace, /// For resources, defines the binding point. pub binding: Option, /// The type of this variable. pub ty: Handle, /// Initial value for this variable. /// /// Expression handle lives in global_expressions pub init: Option>, } /// 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, /// The type of this variable. pub ty: Handle, /// 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>, } /// 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> }, } /// 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), Bias(Handle), Gradient { x: Handle, y: Handle, }, } /// 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` /// for an image with dimensions N > 2. Size { /// If `None`, the base level is considered. level: Option>, }, /// 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), /// Each gathers from a different lane at the index given by the expression Shuffle(Handle), /// Each gathers from their lane plus the shift given by the expression ShuffleDown(Handle), /// Each gathers from their lane minus the shift given by the expression ShuffleUp(Handle), /// Each gathers from their lane xored with the given by the expression ShuffleXor(Handle), } #[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), /// Pipeline-overridable constant. Override(Handle), /// Zero value of a type. ZeroValue(Handle), /// Composite expression. Compose { ty: Handle, components: Vec>, }, /// 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, index: Handle, }, /// Access the same types as [`Access`], plus [`Struct`] with a known index. /// /// [`Access`]: Expression::Access /// [`Struct`]: TypeInner::Struct AccessIndex { base: Handle, index: u32, }, /// Splat scalar into a vector. Splat { size: VectorSize, value: Handle, }, /// Vector swizzle. Swizzle { size: VectorSize, vector: Handle, 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), /// 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), /// Load a value indirectly. /// /// For [`TypeInner::Atomic`] the result is a corresponding scalar. /// For other types behind the `pointer`, the result is `T`. Load { pointer: Handle }, /// Sample a point from a sampled or a depth image. ImageSample { image: Handle, sampler: Handle, /// If Some(), this operation is a gather operation /// on the selected component. gather: Option, coordinate: Handle, array_index: Option>, /// Expression handle lives in global_expressions offset: Option>, level: SampleLevel, depth_ref: Option>, }, /// 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, /// 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, /// 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>, /// A sample index, for multisampled [`Sampled`] and [`Depth`] images. /// /// [`Sampled`]: ImageClass::Sampled /// [`Depth`]: ImageClass::Depth sample: Option>, /// 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>, }, /// Query information from an image. ImageQuery { image: Handle, query: ImageQuery, }, /// Apply an unary operator. Unary { op: UnaryOperator, expr: Handle, }, /// Apply a binary operator. Binary { op: BinaryOperator, left: Handle, right: Handle, }, /// 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, accept: Handle, reject: Handle, }, /// Compute the derivative on an axis. Derivative { axis: DerivativeAxis, ctrl: DerivativeControl, expr: Handle, }, /// Call a relational function. Relational { fun: RelationalFunction, argument: Handle, }, /// Call a math function Math { fun: MathFunction, arg: Handle, arg1: Option>, arg2: Option>, arg3: Option>, }, /// Cast a simple type to another kind. As { /// Source expression, which can only be a scalar or a vector. expr: Handle, /// Target scalar kind. kind: ScalarKind, /// If provided, converts to the specified byte width. /// Otherwise, bitcast. convert: Option, }, /// Result of calling another function. CallResult(Handle), /// Result of an atomic operation. AtomicResult { ty: Handle, comparison: bool }, /// Result of a [`WorkGroupUniformLoad`] statement. /// /// [`WorkGroupUniformLoad`]: Statement::WorkGroupUniformLoad WorkGroupUniformLoadResult { /// The type of the result ty: Handle, }, /// 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), /// 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, 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 }, } 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, #[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, }, /// 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, }, 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), /// 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, //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, cases: Vec, }, /// 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>, }, /// 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> }, /// 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`, 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, value: Handle, }, /// 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, coordinate: Handle, array_index: Option>, value: Handle, }, /// Atomic function. Atomic { /// Pointer to an atomic value. pointer: Handle, /// Function to run on the atomic. fun: AtomicFunction, /// Value to use in the function. value: Handle, /// [`AtomicResult`] expression representing this function's result. /// /// [`AtomicResult`]: crate::Expression::AtomicResult result: Handle, }, /// 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, /// The [`WorkGroupUniformLoadResult`] expression representing this load's result. /// /// [`WorkGroupUniformLoadResult`]: Expression::WorkGroupUniformLoadResult result: Handle, }, /// 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, arguments: Vec>, result: Option>, }, RayQuery { /// The [`RayQuery`] object this statement operates on. /// /// [`RayQuery`]: TypeInner::RayQuery query: Handle, /// 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, /// The value from this thread to store in the ballot predicate: Option>, }, /// 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, /// The [`SubgroupOperationResult`] expression representing this load's result. /// /// [`SubgroupOperationResult`]: Expression::SubgroupOperationResult result: Handle, }, /// 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, /// The [`SubgroupOperationResult`] expression representing this load's result. /// /// [`SubgroupOperationResult`]: Expression::SubgroupOperationResult result: Handle, }, } /// 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, /// Type of the argument. pub ty: Handle, /// For entry points, an argument has to have a binding /// unless it's a structure. pub binding: Option, } /// 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, /// For entry points, the result has to have a binding /// unless it's a structure. pub binding: Option, } /// 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, /// Information about function argument. pub arguments: Vec, /// The result of this function, if any. pub result: Option, /// Local variables defined and used in the function. pub local_variables: Arena, /// Expressions used inside this function. /// /// An `Expression` must occur before all other `Expression`s that use its /// value. pub expressions: Arena, /// 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, /// 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, width: Bytes, }, FrexpResult { size: Option, 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>, /// Type for `RayIntersection`. /// /// Call [`Module::generate_ray_intersection_type`] to populate /// this if needed and return the handle. pub ray_intersection: Option>, /// 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>, } /// 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, /// Dictionary of special type handles. pub special_types: SpecialTypes, /// Arena for the constants defined in this module. pub constants: Arena, /// Arena for the pipeline-overridable constants defined in this module. pub overrides: Arena, /// Arena for the global variables defined in this module. pub global_variables: Arena, /// [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, /// 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, /// Entry points. pub entry_points: Vec, }