naga/ir/
mod.rs

1/*!
2The Intermediate Representation shared by all frontends and backends.
3
4The central structure of the IR, and the crate, is [`Module`]. A `Module` contains:
5
6- [`Function`]s, which have arguments, a return type, local variables, and a body,
7
8- [`EntryPoint`]s, which are specialized functions that can serve as the entry
9  point for pipeline stages like vertex shading or fragment shading,
10
11- [`Constant`]s and [`GlobalVariable`]s used by `EntryPoint`s and `Function`s, and
12
13- [`Type`]s used by the above.
14
15The body of an `EntryPoint` or `Function` is represented using two types:
16
17- An [`Expression`] produces a value, but has no side effects or control flow.
18  `Expressions` include variable references, unary and binary operators, and so
19  on.
20
21- A [`Statement`] can have side effects and structured control flow.
22  `Statement`s do not produce a value, other than by storing one in some
23  designated place. `Statements` include blocks, conditionals, and loops, but also
24  operations that have side effects, like stores and function calls.
25
26`Statement`s form a tree, with pointers into the DAG of `Expression`s.
27
28Restricting side effects to statements simplifies analysis and code generation.
29A Naga backend can generate code to evaluate an `Expression` however and
30whenever it pleases, as long as it is certain to observe the side effects of all
31previously executed `Statement`s.
32
33Many `Statement` variants use the [`Block`] type, which is `Vec<Statement>`,
34with optional span info, representing a series of statements executed in order. The body of an
35`EntryPoint`s or `Function` is a `Block`, and `Statement` has a
36[`Block`][Statement::Block] variant.
37
38## Function Calls
39
40Naga's representation of function calls is unusual. Most languages treat
41function calls as expressions, but because calls may have side effects, Naga
42represents them as a kind of statement, [`Statement::Call`]. If the function
43returns a value, a call statement designates a particular [`Expression::CallResult`]
44expression to represent its return value, for use by subsequent statements and
45expressions.
46
47## `Expression` evaluation time
48
49It is essential to know when an [`Expression`] should be evaluated, because its
50value may depend on previous [`Statement`]s' effects. But whereas the order of
51execution for a tree of `Statement`s is apparent from its structure, it is not
52so clear for `Expressions`, since an expression may be referred to by any number
53of `Statement`s and other `Expression`s.
54
55Naga's rules for when `Expression`s are evaluated are as follows:
56
57-   [`Literal`], [`Constant`], and [`ZeroValue`] expressions are
58    considered to be implicitly evaluated before execution begins.
59
60-   [`FunctionArgument`] and [`LocalVariable`] expressions are considered
61    implicitly evaluated upon entry to the function to which they belong.
62    Function arguments cannot be assigned to, and `LocalVariable` expressions
63    produce a *pointer to* the variable's value (for use with [`Load`] and
64    [`Store`]). Neither varies while the function executes, so it suffices to
65    consider these expressions evaluated once on entry.
66
67-   Similarly, [`GlobalVariable`] expressions are considered implicitly
68    evaluated before execution begins, since their value does not change while
69    code executes, for one of two reasons:
70
71    -   Most `GlobalVariable` expressions produce a pointer to the variable's
72        value, for use with [`Load`] and [`Store`], as `LocalVariable`
73        expressions do. Although the variable's value may change, its address
74        does not.
75
76    -   A `GlobalVariable` expression referring to a global in the
77        [`AddressSpace::Handle`] address space produces the value directly, not
78        a pointer. Such global variables hold opaque types like shaders or
79        images, and cannot be assigned to.
80
81-   A [`CallResult`] expression that is the `result` of a [`Statement::Call`],
82    representing the call's return value, is evaluated when the `Call` statement
83    is executed.
84
85-   Similarly, an [`AtomicResult`] expression that is the `result` of an
86    [`Atomic`] statement, representing the result of the atomic operation, is
87    evaluated when the `Atomic` statement is executed.
88
89-   A [`RayQueryProceedResult`] expression, which is a boolean
90    indicating if the ray query is finished, is evaluated when the
91    [`RayQuery`] statement whose [`Proceed::result`] points to it is
92    executed.
93
94-   All other expressions are evaluated when the (unique) [`Statement::Emit`]
95    statement that covers them is executed.
96
97Now, strictly speaking, not all `Expression` variants actually care when they're
98evaluated. For example, you can evaluate a [`BinaryOperator::Add`] expression
99any time you like, as long as you give it the right operands. It's really only a
100very small set of expressions that are affected by timing:
101
102-   [`Load`], [`ImageSample`], and [`ImageLoad`] expressions are influenced by
103    stores to the variables or images they access, and must execute at the
104    proper time relative to them.
105
106-   [`Derivative`] expressions are sensitive to control flow uniformity: they
107    must not be moved out of an area of uniform control flow into a non-uniform
108    area.
109
110-   More generally, any expression that's used by more than one other expression
111    or statement should probably be evaluated only once, and then stored in a
112    variable to be cited at each point of use.
113
114Naga tries to help back ends handle all these cases correctly in a somewhat
115circuitous way. The [`ModuleInfo`] structure returned by [`Validator::validate`]
116provides a reference count for each expression in each function in the module.
117Naturally, any expression with a reference count of two or more deserves to be
118evaluated and stored in a temporary variable at the point that the `Emit`
119statement covering it is executed. But if we selectively lower the reference
120count threshold to _one_ for the sensitive expression types listed above, so
121that we _always_ generate a temporary variable and save their value, then the
122same code that manages multiply referenced expressions will take care of
123introducing temporaries for time-sensitive expressions as well. The
124`Expression::bake_ref_count` method (private to the back ends) is meant to help
125with this.
126
127## `Expression` scope
128
129Each `Expression` has a *scope*, which is the region of the function within
130which it can be used by `Statement`s and other `Expression`s. It is a validation
131error to use an `Expression` outside its scope.
132
133An expression's scope is defined as follows:
134
135-   The scope of a [`Constant`], [`GlobalVariable`], [`FunctionArgument`] or
136    [`LocalVariable`] expression covers the entire `Function` in which it
137    occurs.
138
139-   The scope of an expression evaluated by an [`Emit`] statement covers the
140    subsequent expressions in that `Emit`, the subsequent statements in the `Block`
141    to which that `Emit` belongs (if any) and their sub-statements (if any).
142
143-   The `result` expression of a [`Call`] or [`Atomic`] statement has a scope
144    covering the subsequent statements in the `Block` in which the statement
145    occurs (if any) and their sub-statements (if any).
146
147For example, this implies that an expression evaluated by some statement in a
148nested `Block` is not available in the `Block`'s parents. Such a value would
149need to be stored in a local variable to be carried upwards in the statement
150tree.
151
152## Constant expressions
153
154A Naga *constant expression* is one of the following [`Expression`]
155variants, whose operands (if any) are also constant expressions:
156- [`Literal`]
157- [`Constant`], for [`Constant`]s
158- [`ZeroValue`], for fixed-size types
159- [`Compose`]
160- [`Access`]
161- [`AccessIndex`]
162- [`Splat`]
163- [`Swizzle`]
164- [`Unary`]
165- [`Binary`]
166- [`Select`]
167- [`Relational`]
168- [`Math`]
169- [`As`]
170
171A constant expression can be evaluated at module translation time.
172
173## Override expressions
174
175A Naga *override expression* is the same as a [constant expression],
176except that it is also allowed to reference other [`Override`]s.
177
178An override expression can be evaluated at pipeline creation time.
179
180[`AtomicResult`]: Expression::AtomicResult
181[`RayQueryProceedResult`]: Expression::RayQueryProceedResult
182[`CallResult`]: Expression::CallResult
183[`Constant`]: Expression::Constant
184[`ZeroValue`]: Expression::ZeroValue
185[`Literal`]: Expression::Literal
186[`Derivative`]: Expression::Derivative
187[`FunctionArgument`]: Expression::FunctionArgument
188[`GlobalVariable`]: Expression::GlobalVariable
189[`ImageLoad`]: Expression::ImageLoad
190[`ImageSample`]: Expression::ImageSample
191[`Load`]: Expression::Load
192[`LocalVariable`]: Expression::LocalVariable
193
194[`Atomic`]: Statement::Atomic
195[`Call`]: Statement::Call
196[`Emit`]: Statement::Emit
197[`Store`]: Statement::Store
198[`RayQuery`]: Statement::RayQuery
199
200[`Proceed::result`]: RayQueryFunction::Proceed::result
201
202[`Validator::validate`]: crate::valid::Validator::validate
203[`ModuleInfo`]: crate::valid::ModuleInfo
204
205[`Literal`]: Expression::Literal
206[`ZeroValue`]: Expression::ZeroValue
207[`Compose`]: Expression::Compose
208[`Access`]: Expression::Access
209[`AccessIndex`]: Expression::AccessIndex
210[`Splat`]: Expression::Splat
211[`Swizzle`]: Expression::Swizzle
212[`Unary`]: Expression::Unary
213[`Binary`]: Expression::Binary
214[`Select`]: Expression::Select
215[`Relational`]: Expression::Relational
216[`Math`]: Expression::Math
217[`As`]: Expression::As
218
219[constant expression]: #constant-expressions
220*/
221
222mod block;
223
224use alloc::{boxed::Box, string::String, vec::Vec};
225
226#[cfg(feature = "arbitrary")]
227use arbitrary::Arbitrary;
228use half::f16;
229#[cfg(feature = "deserialize")]
230use serde::Deserialize;
231#[cfg(feature = "serialize")]
232use serde::Serialize;
233
234use crate::arena::{Arena, Handle, Range, UniqueArena};
235use crate::diagnostic_filter::DiagnosticFilterNode;
236use crate::{FastIndexMap, NamedExpressions};
237
238pub use block::Block;
239
240/// Explicitly allows early depth/stencil tests.
241///
242/// Normally, depth/stencil tests are performed after fragment shading. However, as an optimization,
243/// most drivers will move the depth/stencil tests before fragment shading if this does not
244/// have any observable consequences. This optimization is disabled under the following
245/// circumstances:
246///   - `discard` is called in the fragment shader.
247///   - The fragment shader writes to the depth buffer.
248///   - The fragment shader writes to any storage bindings.
249///
250/// When `EarlyDepthTest` is set, it is allowed to perform an early depth/stencil test even if the
251/// above conditions are not met. When [`EarlyDepthTest::Force`] is used, depth/stencil tests
252/// **must** be performed before fragment shading.
253///
254/// To force early depth/stencil tests in a shader:
255///   - GLSL: `layout(early_fragment_tests) in;`
256///   - HLSL: `Attribute earlydepthstencil`
257///   - SPIR-V: `ExecutionMode EarlyFragmentTests`
258///   - WGSL: `@early_depth_test(force)`
259///
260/// This may also be enabled in a shader by specifying a [`ConservativeDepth`].
261///
262/// For more, see:
263///   - <https://www.khronos.org/opengl/wiki/Early_Fragment_Test#Explicit_specification>
264///   - <https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/sm5-attributes-earlydepthstencil>
265///   - <https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#Execution_Mode>
266#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
267#[cfg_attr(feature = "serialize", derive(Serialize))]
268#[cfg_attr(feature = "deserialize", derive(Deserialize))]
269#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
270pub enum EarlyDepthTest {
271    /// Requires depth/stencil tests to be performed before fragment shading.
272    ///
273    /// This will disable depth/stencil tests after fragment shading, so discarding the fragment
274    /// or overwriting the fragment depth will have no effect.
275    Force,
276
277    /// Allows an additional depth/stencil test to be performed before fragment shading.
278    ///
279    /// It is up to the driver to decide whether early tests are performed. Unlike `Force`, this
280    /// does not disable depth/stencil tests after fragment shading.
281    Allow {
282        /// Specifies restrictions on how the depth value can be modified within the fragment
283        /// shader.
284        ///
285        /// This may be taken into account when deciding whether to perform early tests.
286        conservative: ConservativeDepth,
287    },
288}
289
290/// Enables adjusting depth without disabling early Z.
291///
292/// To use in a shader:
293///   - GLSL: `layout (depth_<greater/less/unchanged/any>) out float gl_FragDepth;`
294///     - `depth_any` option behaves as if the layout qualifier was not present.
295///   - HLSL: `SV_DepthGreaterEqual`/`SV_DepthLessEqual`/`SV_Depth`
296///   - SPIR-V: `ExecutionMode Depth<Greater/Less/Unchanged>`
297///   - WGSL: `@early_depth_test(greater_equal/less_equal/unchanged)`
298///
299/// For more, see:
300///   - <https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_conservative_depth.txt>
301///   - <https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-semantics#system-value-semantics>
302///   - <https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#Execution_Mode>
303#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
304#[cfg_attr(feature = "serialize", derive(Serialize))]
305#[cfg_attr(feature = "deserialize", derive(Deserialize))]
306#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
307pub enum ConservativeDepth {
308    /// Shader may rewrite depth only with a value greater than calculated.
309    GreaterEqual,
310
311    /// Shader may rewrite depth smaller than one that would have been written without the modification.
312    LessEqual,
313
314    /// Shader may not rewrite depth value.
315    Unchanged,
316}
317
318/// Stage of the programmable pipeline.
319#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
320#[cfg_attr(feature = "serialize", derive(Serialize))]
321#[cfg_attr(feature = "deserialize", derive(Deserialize))]
322#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
323#[allow(missing_docs)] // The names are self evident
324pub enum ShaderStage {
325    Vertex,
326    Fragment,
327    Compute,
328    Task,
329    Mesh,
330}
331
332/// Addressing space of variables.
333#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
334#[cfg_attr(feature = "serialize", derive(Serialize))]
335#[cfg_attr(feature = "deserialize", derive(Deserialize))]
336#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
337pub enum AddressSpace {
338    /// Function locals.
339    Function,
340    /// Private data, per invocation, mutable.
341    Private,
342    /// Workgroup shared data, mutable.
343    WorkGroup,
344    /// Uniform buffer data.
345    Uniform,
346    /// Storage buffer data, potentially mutable.
347    Storage { access: StorageAccess },
348    /// Opaque handles, such as samplers and images.
349    Handle,
350
351    /// Push constants.
352    ///
353    /// A [`Module`] may contain at most one [`GlobalVariable`] in
354    /// this address space. Its contents are provided not by a buffer
355    /// but by `SetPushConstant` pass commands, allowing the CPU to
356    /// establish different values for each draw/dispatch.
357    ///
358    /// `PushConstant` variables may not contain `f16` values, even if
359    /// the [`SHADER_FLOAT16`] capability is enabled.
360    ///
361    /// Backends generally place tight limits on the size of
362    /// `PushConstant` variables.
363    ///
364    /// [`SHADER_FLOAT16`]: crate::valid::Capabilities::SHADER_FLOAT16
365    PushConstant,
366}
367
368/// Built-in inputs and outputs.
369#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
370#[cfg_attr(feature = "serialize", derive(Serialize))]
371#[cfg_attr(feature = "deserialize", derive(Deserialize))]
372#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
373pub enum BuiltIn {
374    Position { invariant: bool },
375    ViewIndex,
376    // vertex
377    BaseInstance,
378    BaseVertex,
379    ClipDistance,
380    CullDistance,
381    InstanceIndex,
382    PointSize,
383    VertexIndex,
384    DrawID,
385    // fragment
386    FragDepth,
387    PointCoord,
388    FrontFacing,
389    PrimitiveIndex,
390    SampleIndex,
391    SampleMask,
392    // compute
393    GlobalInvocationId,
394    LocalInvocationId,
395    LocalInvocationIndex,
396    WorkGroupId,
397    WorkGroupSize,
398    NumWorkGroups,
399    // subgroup
400    NumSubgroups,
401    SubgroupId,
402    SubgroupSize,
403    SubgroupInvocationId,
404}
405
406/// Number of bytes per scalar.
407pub type Bytes = u8;
408
409/// Number of components in a vector.
410#[repr(u8)]
411#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
412#[cfg_attr(feature = "serialize", derive(Serialize))]
413#[cfg_attr(feature = "deserialize", derive(Deserialize))]
414#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
415pub enum VectorSize {
416    /// 2D vector
417    Bi = 2,
418    /// 3D vector
419    Tri = 3,
420    /// 4D vector
421    Quad = 4,
422}
423
424impl VectorSize {
425    pub const MAX: usize = Self::Quad as usize;
426}
427
428impl From<VectorSize> for u8 {
429    fn from(size: VectorSize) -> u8 {
430        size as u8
431    }
432}
433
434impl From<VectorSize> for u32 {
435    fn from(size: VectorSize) -> u32 {
436        size as u32
437    }
438}
439
440/// Primitive type for a scalar.
441#[repr(u8)]
442#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
443#[cfg_attr(feature = "serialize", derive(Serialize))]
444#[cfg_attr(feature = "deserialize", derive(Deserialize))]
445#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
446pub enum ScalarKind {
447    /// Signed integer type.
448    Sint,
449    /// Unsigned integer type.
450    Uint,
451    /// Floating point type.
452    Float,
453    /// Boolean type.
454    Bool,
455
456    /// WGSL abstract integer type.
457    ///
458    /// These are forbidden by validation, and should never reach backends.
459    AbstractInt,
460
461    /// Abstract floating-point type.
462    ///
463    /// These are forbidden by validation, and should never reach backends.
464    AbstractFloat,
465}
466
467/// Characteristics of a scalar type.
468#[derive(Clone, Copy, Debug, PartialEq, Eq, PartialOrd, Ord, Hash)]
469#[cfg_attr(feature = "serialize", derive(Serialize))]
470#[cfg_attr(feature = "deserialize", derive(Deserialize))]
471#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
472pub struct Scalar {
473    /// How the value's bits are to be interpreted.
474    pub kind: ScalarKind,
475
476    /// This size of the value in bytes.
477    pub width: Bytes,
478}
479
480/// Size of an array.
481#[repr(u8)]
482#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
483#[cfg_attr(feature = "serialize", derive(Serialize))]
484#[cfg_attr(feature = "deserialize", derive(Deserialize))]
485#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
486pub enum ArraySize {
487    /// The array size is constant.
488    Constant(core::num::NonZeroU32),
489    /// The array size is an override-expression.
490    Pending(Handle<Override>),
491    /// The array size can change at runtime.
492    Dynamic,
493}
494
495/// The interpolation qualifier of a binding or struct field.
496#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
497#[cfg_attr(feature = "serialize", derive(Serialize))]
498#[cfg_attr(feature = "deserialize", derive(Deserialize))]
499#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
500pub enum Interpolation {
501    /// The value will be interpolated in a perspective-correct fashion.
502    /// Also known as "smooth" in glsl.
503    Perspective,
504    /// Indicates that linear, non-perspective, correct
505    /// interpolation must be used.
506    /// Also known as "no_perspective" in glsl.
507    Linear,
508    /// Indicates that no interpolation will be performed.
509    Flat,
510}
511
512/// The sampling qualifiers of a binding or struct field.
513#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
514#[cfg_attr(feature = "serialize", derive(Serialize))]
515#[cfg_attr(feature = "deserialize", derive(Deserialize))]
516#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
517pub enum Sampling {
518    /// Interpolate the value at the center of the pixel.
519    Center,
520
521    /// Interpolate the value at a point that lies within all samples covered by
522    /// the fragment within the current primitive. In multisampling, use a
523    /// single value for all samples in the primitive.
524    Centroid,
525
526    /// Interpolate the value at each sample location. In multisampling, invoke
527    /// the fragment shader once per sample.
528    Sample,
529
530    /// Use the value provided by the first vertex of the current primitive.
531    First,
532
533    /// Use the value provided by the first or last vertex of the current primitive. The exact
534    /// choice is implementation-dependent.
535    Either,
536}
537
538/// Member of a user-defined structure.
539// Clone is used only for error reporting and is not intended for end users
540#[derive(Clone, Debug, Eq, Hash, PartialEq)]
541#[cfg_attr(feature = "serialize", derive(Serialize))]
542#[cfg_attr(feature = "deserialize", derive(Deserialize))]
543#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
544pub struct StructMember {
545    pub name: Option<String>,
546    /// Type of the field.
547    pub ty: Handle<Type>,
548    /// For I/O structs, defines the binding.
549    pub binding: Option<Binding>,
550    /// Offset from the beginning from the struct.
551    pub offset: u32,
552}
553
554/// The number of dimensions an image has.
555#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
556#[cfg_attr(feature = "serialize", derive(Serialize))]
557#[cfg_attr(feature = "deserialize", derive(Deserialize))]
558#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
559pub enum ImageDimension {
560    /// 1D image
561    D1,
562    /// 2D image
563    D2,
564    /// 3D image
565    D3,
566    /// Cube map
567    Cube,
568}
569
570bitflags::bitflags! {
571    /// Flags describing an image.
572    #[cfg_attr(feature = "serialize", derive(Serialize))]
573    #[cfg_attr(feature = "deserialize", derive(Deserialize))]
574    #[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
575    #[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)]
576    pub struct StorageAccess: u32 {
577        /// Storage can be used as a source for load ops.
578        const LOAD = 0x1;
579        /// Storage can be used as a target for store ops.
580        const STORE = 0x2;
581        /// Storage can be used as a target for atomic ops.
582        const ATOMIC = 0x4;
583    }
584}
585
586/// Image storage format.
587#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
588#[cfg_attr(feature = "serialize", derive(Serialize))]
589#[cfg_attr(feature = "deserialize", derive(Deserialize))]
590#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
591pub enum StorageFormat {
592    // 8-bit formats
593    R8Unorm,
594    R8Snorm,
595    R8Uint,
596    R8Sint,
597
598    // 16-bit formats
599    R16Uint,
600    R16Sint,
601    R16Float,
602    Rg8Unorm,
603    Rg8Snorm,
604    Rg8Uint,
605    Rg8Sint,
606
607    // 32-bit formats
608    R32Uint,
609    R32Sint,
610    R32Float,
611    Rg16Uint,
612    Rg16Sint,
613    Rg16Float,
614    Rgba8Unorm,
615    Rgba8Snorm,
616    Rgba8Uint,
617    Rgba8Sint,
618    Bgra8Unorm,
619
620    // Packed 32-bit formats
621    Rgb10a2Uint,
622    Rgb10a2Unorm,
623    Rg11b10Ufloat,
624
625    // 64-bit formats
626    R64Uint,
627    Rg32Uint,
628    Rg32Sint,
629    Rg32Float,
630    Rgba16Uint,
631    Rgba16Sint,
632    Rgba16Float,
633
634    // 128-bit formats
635    Rgba32Uint,
636    Rgba32Sint,
637    Rgba32Float,
638
639    // Normalized 16-bit per channel formats
640    R16Unorm,
641    R16Snorm,
642    Rg16Unorm,
643    Rg16Snorm,
644    Rgba16Unorm,
645    Rgba16Snorm,
646}
647
648/// Sub-class of the image type.
649#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
650#[cfg_attr(feature = "serialize", derive(Serialize))]
651#[cfg_attr(feature = "deserialize", derive(Deserialize))]
652#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
653pub enum ImageClass {
654    /// Regular sampled image.
655    Sampled {
656        /// Kind of values to sample.
657        kind: ScalarKind,
658        /// Multi-sampled image.
659        ///
660        /// A multi-sampled image holds several samples per texel. Multi-sampled
661        /// images cannot have mipmaps.
662        multi: bool,
663    },
664    /// Depth comparison image.
665    Depth {
666        /// Multi-sampled depth image.
667        multi: bool,
668    },
669    /// External texture.
670    External,
671    /// Storage image.
672    Storage {
673        format: StorageFormat,
674        access: StorageAccess,
675    },
676}
677
678/// A data type declared in the module.
679#[derive(Clone, Debug, Eq, Hash, PartialEq)]
680#[cfg_attr(feature = "serialize", derive(Serialize))]
681#[cfg_attr(feature = "deserialize", derive(Deserialize))]
682#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
683pub struct Type {
684    /// The name of the type, if any.
685    pub name: Option<String>,
686    /// Inner structure that depends on the kind of the type.
687    pub inner: TypeInner,
688}
689
690/// Enum with additional information, depending on the kind of type.
691///
692/// Comparison using `==` is not reliable in the case of [`Pointer`],
693/// [`ValuePointer`], or [`Struct`] variants. For these variants,
694/// use [`TypeInner::non_struct_equivalent`] or [`compare_types`].
695///
696/// [`compare_types`]: crate::proc::compare_types
697/// [`ValuePointer`]: TypeInner::ValuePointer
698/// [`Pointer`]: TypeInner::Pointer
699/// [`Struct`]: TypeInner::Struct
700#[derive(Clone, Debug, Eq, Hash, PartialEq)]
701#[cfg_attr(feature = "serialize", derive(Serialize))]
702#[cfg_attr(feature = "deserialize", derive(Deserialize))]
703#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
704pub enum TypeInner {
705    /// Number of integral or floating-point kind.
706    Scalar(Scalar),
707    /// Vector of numbers.
708    Vector { size: VectorSize, scalar: Scalar },
709    /// Matrix of numbers.
710    Matrix {
711        columns: VectorSize,
712        rows: VectorSize,
713        scalar: Scalar,
714    },
715    /// Atomic scalar.
716    Atomic(Scalar),
717    /// Pointer to another type.
718    ///
719    /// Pointers to scalars and vectors should be treated as equivalent to
720    /// [`ValuePointer`] types. Use either [`TypeInner::non_struct_equivalent`]
721    /// or [`compare_types`] to compare types in a way that treats pointers
722    /// correctly.
723    ///
724    /// ## Pointers to non-`SIZED` types
725    ///
726    /// The `base` type of a pointer may be a non-[`SIZED`] type like a
727    /// dynamically-sized [`Array`], or a [`Struct`] whose last member is a
728    /// dynamically sized array. Such pointers occur as the types of
729    /// [`GlobalVariable`] or [`AccessIndex`] expressions referring to
730    /// dynamically-sized arrays.
731    ///
732    /// However, among pointers to non-`SIZED` types, only pointers to `Struct`s
733    /// are [`DATA`]. Pointers to dynamically sized `Array`s cannot be passed as
734    /// arguments, stored in variables, or held in arrays or structures. Their
735    /// only use is as the types of `AccessIndex` expressions.
736    ///
737    /// [`SIZED`]: crate::valid::TypeFlags::SIZED
738    /// [`DATA`]: crate::valid::TypeFlags::DATA
739    /// [`Array`]: TypeInner::Array
740    /// [`Struct`]: TypeInner::Struct
741    /// [`ValuePointer`]: TypeInner::ValuePointer
742    /// [`GlobalVariable`]: Expression::GlobalVariable
743    /// [`AccessIndex`]: Expression::AccessIndex
744    /// [`compare_types`]: crate::proc::compare_types
745    Pointer {
746        base: Handle<Type>,
747        space: AddressSpace,
748    },
749
750    /// Pointer to a scalar or vector.
751    ///
752    /// A `ValuePointer` type is equivalent to a `Pointer` whose `base` is a
753    /// `Scalar` or `Vector` type. This is for use in [`TypeResolution::Value`]
754    /// variants; see the documentation for [`TypeResolution`] for details.
755    ///
756    /// Use [`TypeInner::non_struct_equivalent`] or [`compare_types`] to compare
757    /// types that could be pointers, to ensure that `Pointer` and
758    /// `ValuePointer` types are recognized as equivalent.
759    ///
760    /// [`TypeResolution`]: crate::proc::TypeResolution
761    /// [`TypeResolution::Value`]: crate::proc::TypeResolution::Value
762    /// [`compare_types`]: crate::proc::compare_types
763    ValuePointer {
764        size: Option<VectorSize>,
765        scalar: Scalar,
766        space: AddressSpace,
767    },
768
769    /// Homogeneous list of elements.
770    ///
771    /// The `base` type must be a [`SIZED`], [`DATA`] type.
772    ///
773    /// ## Dynamically sized arrays
774    ///
775    /// An `Array` is [`SIZED`] unless its `size` is [`Dynamic`].
776    /// Dynamically-sized arrays may only appear in a few situations:
777    ///
778    /// -   They may appear as the type of a [`GlobalVariable`], or as the last
779    ///     member of a [`Struct`].
780    ///
781    /// -   They may appear as the base type of a [`Pointer`]. An
782    ///     [`AccessIndex`] expression referring to a struct's final
783    ///     unsized array member would have such a pointer type. However, such
784    ///     pointer types may only appear as the types of such intermediate
785    ///     expressions. They are not [`DATA`], and cannot be stored in
786    ///     variables, held in arrays or structs, or passed as parameters.
787    ///
788    /// [`SIZED`]: crate::valid::TypeFlags::SIZED
789    /// [`DATA`]: crate::valid::TypeFlags::DATA
790    /// [`Dynamic`]: ArraySize::Dynamic
791    /// [`Struct`]: TypeInner::Struct
792    /// [`Pointer`]: TypeInner::Pointer
793    /// [`AccessIndex`]: Expression::AccessIndex
794    Array {
795        base: Handle<Type>,
796        size: ArraySize,
797        stride: u32,
798    },
799
800    /// User-defined structure.
801    ///
802    /// There must always be at least one member.
803    ///
804    /// A `Struct` type is [`DATA`], and the types of its members must be
805    /// `DATA` as well.
806    ///
807    /// Member types must be [`SIZED`], except for the final member of a
808    /// struct, which may be a dynamically sized [`Array`]. The
809    /// `Struct` type itself is `SIZED` when all its members are `SIZED`.
810    ///
811    /// Two structure types with different names are not equivalent. Because
812    /// this variant does not contain the name, it is not possible to use it
813    /// to compare struct types. Use [`compare_types`] to compare two types
814    /// that may be structs.
815    ///
816    /// [`DATA`]: crate::valid::TypeFlags::DATA
817    /// [`SIZED`]: crate::∅TypeFlags::SIZED
818    /// [`Array`]: TypeInner::Array
819    /// [`compare_types`]: crate::proc::compare_types
820    Struct {
821        members: Vec<StructMember>,
822        //TODO: should this be unaligned?
823        span: u32,
824    },
825    /// Possibly multidimensional array of texels.
826    Image {
827        dim: ImageDimension,
828        arrayed: bool,
829        //TODO: consider moving `multisampled: bool` out
830        class: ImageClass,
831    },
832    /// Can be used to sample values from images.
833    Sampler { comparison: bool },
834
835    /// Opaque object representing an acceleration structure of geometry.
836    AccelerationStructure { vertex_return: bool },
837
838    /// Locally used handle for ray queries.
839    RayQuery { vertex_return: bool },
840
841    /// Array of bindings.
842    ///
843    /// A `BindingArray` represents an array where each element draws its value
844    /// from a separate bound resource. The array's element type `base` may be
845    /// [`Image`], [`Sampler`], or any type that would be permitted for a global
846    /// in the [`Uniform`] or [`Storage`] address spaces. Only global variables
847    /// may be binding arrays; on the host side, their values are provided by
848    /// [`TextureViewArray`], [`SamplerArray`], or [`BufferArray`]
849    /// bindings.
850    ///
851    /// Since each element comes from a distinct resource, a binding array of
852    /// images could have images of varying sizes (but not varying dimensions;
853    /// they must all have the same `Image` type). Or, a binding array of
854    /// buffers could have elements that are dynamically sized arrays, each with
855    /// a different length.
856    ///
857    /// Binding arrays are in the same address spaces as their underlying type.
858    /// As such, referring to an array of images produces an [`Image`] value
859    /// directly (as opposed to a pointer). The only operation permitted on
860    /// `BindingArray` values is indexing, which works transparently: indexing
861    /// a binding array of samplers yields a [`Sampler`], indexing a pointer to the
862    /// binding array of storage buffers produces a pointer to the storage struct.
863    ///
864    /// Unlike textures and samplers, binding arrays are not [`ARGUMENT`], so
865    /// they cannot be passed as arguments to functions.
866    ///
867    /// Naga's WGSL front end supports binding arrays with the type syntax
868    /// `binding_array<T, N>`.
869    ///
870    /// [`Image`]: TypeInner::Image
871    /// [`Sampler`]: TypeInner::Sampler
872    /// [`Uniform`]: AddressSpace::Uniform
873    /// [`Storage`]: AddressSpace::Storage
874    /// [`TextureViewArray`]: https://docs.rs/wgpu/latest/wgpu/enum.BindingResource.html#variant.TextureViewArray
875    /// [`SamplerArray`]: https://docs.rs/wgpu/latest/wgpu/enum.BindingResource.html#variant.SamplerArray
876    /// [`BufferArray`]: https://docs.rs/wgpu/latest/wgpu/enum.BindingResource.html#variant.BufferArray
877    /// [`DATA`]: crate::valid::TypeFlags::DATA
878    /// [`ARGUMENT`]: crate::valid::TypeFlags::ARGUMENT
879    /// [naga#1864]: https://github.com/gfx-rs/naga/issues/1864
880    BindingArray { base: Handle<Type>, size: ArraySize },
881}
882
883#[derive(Debug, Clone, Copy, PartialEq, PartialOrd)]
884#[cfg_attr(feature = "serialize", derive(Serialize))]
885#[cfg_attr(feature = "deserialize", derive(Deserialize))]
886#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
887pub enum Literal {
888    /// May not be NaN or infinity.
889    F64(f64),
890    /// May not be NaN or infinity.
891    F32(f32),
892    /// May not be NaN or infinity.
893    F16(f16),
894    U32(u32),
895    I32(i32),
896    U64(u64),
897    I64(i64),
898    Bool(bool),
899    AbstractInt(i64),
900    AbstractFloat(f64),
901}
902
903/// Pipeline-overridable constant.
904#[derive(Clone, Debug, PartialEq)]
905#[cfg_attr(feature = "serialize", derive(Serialize))]
906#[cfg_attr(feature = "deserialize", derive(Deserialize))]
907#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
908pub struct Override {
909    pub name: Option<String>,
910    /// Pipeline Constant ID.
911    pub id: Option<u16>,
912    pub ty: Handle<Type>,
913
914    /// The default value of the pipeline-overridable constant.
915    ///
916    /// This [`Handle`] refers to [`Module::global_expressions`], not
917    /// any [`Function::expressions`] arena.
918    pub init: Option<Handle<Expression>>,
919}
920
921/// Constant value.
922#[derive(Clone, Debug, PartialEq)]
923#[cfg_attr(feature = "serialize", derive(Serialize))]
924#[cfg_attr(feature = "deserialize", derive(Deserialize))]
925#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
926pub struct Constant {
927    pub name: Option<String>,
928    pub ty: Handle<Type>,
929
930    /// The value of the constant.
931    ///
932    /// This [`Handle`] refers to [`Module::global_expressions`], not
933    /// any [`Function::expressions`] arena.
934    pub init: Handle<Expression>,
935}
936
937/// Describes how an input/output variable is to be bound.
938#[derive(Clone, Debug, Eq, PartialEq, Hash)]
939#[cfg_attr(feature = "serialize", derive(Serialize))]
940#[cfg_attr(feature = "deserialize", derive(Deserialize))]
941#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
942pub enum Binding {
943    /// Built-in shader variable.
944    BuiltIn(BuiltIn),
945
946    /// Indexed location.
947    ///
948    /// Values passed from the [`Vertex`] stage to the [`Fragment`] stage must
949    /// have their `interpolation` defaulted (i.e. not `None`) by the front end
950    /// as appropriate for that language.
951    ///
952    /// For other stages, we permit interpolations even though they're ignored.
953    /// When a front end is parsing a struct type, it usually doesn't know what
954    /// stages will be using it for IO, so it's easiest if it can apply the
955    /// defaults to anything with a `Location` binding, just in case.
956    ///
957    /// For anything other than floating-point scalars and vectors, the
958    /// interpolation must be `Flat`.
959    ///
960    /// [`Vertex`]: crate::ShaderStage::Vertex
961    /// [`Fragment`]: crate::ShaderStage::Fragment
962    Location {
963        location: u32,
964        interpolation: Option<Interpolation>,
965        sampling: Option<Sampling>,
966        /// Optional `blend_src` index used for dual source blending.
967        /// See <https://www.w3.org/TR/WGSL/#attribute-blend_src>
968        blend_src: Option<u32>,
969    },
970}
971
972/// Pipeline binding information for global resources.
973#[derive(Copy, Clone, Debug, Eq, Hash, Ord, PartialEq, PartialOrd)]
974#[cfg_attr(feature = "serialize", derive(Serialize))]
975#[cfg_attr(feature = "deserialize", derive(Deserialize))]
976#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
977pub struct ResourceBinding {
978    /// The bind group index.
979    pub group: u32,
980    /// Binding number within the group.
981    pub binding: u32,
982}
983
984/// Variable defined at module level.
985#[derive(Clone, Debug, PartialEq)]
986#[cfg_attr(feature = "serialize", derive(Serialize))]
987#[cfg_attr(feature = "deserialize", derive(Deserialize))]
988#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
989pub struct GlobalVariable {
990    /// Name of the variable, if any.
991    pub name: Option<String>,
992    /// How this variable is to be stored.
993    pub space: AddressSpace,
994    /// For resources, defines the binding point.
995    pub binding: Option<ResourceBinding>,
996    /// The type of this variable.
997    pub ty: Handle<Type>,
998    /// Initial value for this variable.
999    ///
1000    /// This refers to an [`Expression`] in [`Module::global_expressions`].
1001    pub init: Option<Handle<Expression>>,
1002}
1003
1004/// Variable defined at function level.
1005#[derive(Clone, Debug)]
1006#[cfg_attr(feature = "serialize", derive(Serialize))]
1007#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1008#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1009pub struct LocalVariable {
1010    /// Name of the variable, if any.
1011    pub name: Option<String>,
1012    /// The type of this variable.
1013    pub ty: Handle<Type>,
1014    /// Initial value for this variable.
1015    ///
1016    /// This handle refers to an expression in this `LocalVariable`'s function's
1017    /// [`expressions`] arena, but it is required to be an evaluated override
1018    /// expression.
1019    ///
1020    /// [`expressions`]: Function::expressions
1021    pub init: Option<Handle<Expression>>,
1022}
1023
1024/// Operation that can be applied on a single value.
1025#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1026#[cfg_attr(feature = "serialize", derive(Serialize))]
1027#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1028#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1029pub enum UnaryOperator {
1030    Negate,
1031    LogicalNot,
1032    BitwiseNot,
1033}
1034
1035/// Operation that can be applied on two values.
1036///
1037/// ## Arithmetic type rules
1038///
1039/// The arithmetic operations `Add`, `Subtract`, `Multiply`, `Divide`, and
1040/// `Modulo` can all be applied to [`Scalar`] types other than [`Bool`], or
1041/// [`Vector`]s thereof. Both operands must have the same type.
1042///
1043/// `Add` and `Subtract` can also be applied to [`Matrix`] values. Both operands
1044/// must have the same type.
1045///
1046/// `Multiply` supports additional cases:
1047///
1048/// -   A [`Matrix`] or [`Vector`] can be multiplied by a scalar [`Float`],
1049///     either on the left or the right.
1050///
1051/// -   A [`Matrix`] on the left can be multiplied by a [`Vector`] on the right
1052///     if the matrix has as many columns as the vector has components
1053///     (`matCxR * VecC`).
1054///
1055/// -   A [`Vector`] on the left can be multiplied by a [`Matrix`] on the right
1056///     if the matrix has as many rows as the vector has components
1057///     (`VecR * matCxR`).
1058///
1059/// -   Two matrices can be multiplied if the left operand has as many columns
1060///     as the right operand has rows (`matNxR * matCxN`).
1061///
1062/// In all the above `Multiply` cases, the byte widths of the underlying scalar
1063/// types of both operands must be the same.
1064///
1065/// Note that `Multiply` supports mixed vector and scalar operations directly,
1066/// whereas the other arithmetic operations require an explicit [`Splat`] for
1067/// mixed-type use.
1068///
1069/// [`Scalar`]: TypeInner::Scalar
1070/// [`Vector`]: TypeInner::Vector
1071/// [`Matrix`]: TypeInner::Matrix
1072/// [`Float`]: ScalarKind::Float
1073/// [`Bool`]: ScalarKind::Bool
1074/// [`Splat`]: Expression::Splat
1075#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1076#[cfg_attr(feature = "serialize", derive(Serialize))]
1077#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1078#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1079pub enum BinaryOperator {
1080    Add,
1081    Subtract,
1082    Multiply,
1083    Divide,
1084    /// Equivalent of the WGSL's `%` operator or SPIR-V's `OpFRem`
1085    Modulo,
1086    Equal,
1087    NotEqual,
1088    Less,
1089    LessEqual,
1090    Greater,
1091    GreaterEqual,
1092    And,
1093    ExclusiveOr,
1094    InclusiveOr,
1095    LogicalAnd,
1096    LogicalOr,
1097    ShiftLeft,
1098    /// Right shift carries the sign of signed integers only.
1099    ShiftRight,
1100}
1101
1102/// Function on an atomic value.
1103///
1104/// Note: these do not include load/store, which use the existing
1105/// [`Expression::Load`] and [`Statement::Store`].
1106///
1107/// All `Handle<Expression>` values here refer to an expression in
1108/// [`Function::expressions`].
1109#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1110#[cfg_attr(feature = "serialize", derive(Serialize))]
1111#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1112#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1113pub enum AtomicFunction {
1114    Add,
1115    Subtract,
1116    And,
1117    ExclusiveOr,
1118    InclusiveOr,
1119    Min,
1120    Max,
1121    Exchange { compare: Option<Handle<Expression>> },
1122}
1123
1124/// Hint at which precision to compute a derivative.
1125#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1126#[cfg_attr(feature = "serialize", derive(Serialize))]
1127#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1128#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1129pub enum DerivativeControl {
1130    Coarse,
1131    Fine,
1132    None,
1133}
1134
1135/// Axis on which to compute a derivative.
1136#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1137#[cfg_attr(feature = "serialize", derive(Serialize))]
1138#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1139#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1140pub enum DerivativeAxis {
1141    X,
1142    Y,
1143    Width,
1144}
1145
1146/// Built-in shader function for testing relation between values.
1147#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1148#[cfg_attr(feature = "serialize", derive(Serialize))]
1149#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1150#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1151pub enum RelationalFunction {
1152    All,
1153    Any,
1154    IsNan,
1155    IsInf,
1156}
1157
1158/// Built-in shader function for math.
1159#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1160#[cfg_attr(feature = "serialize", derive(Serialize))]
1161#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1162#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1163pub enum MathFunction {
1164    // comparison
1165    Abs,
1166    Min,
1167    Max,
1168    Clamp,
1169    Saturate,
1170    // trigonometry
1171    Cos,
1172    Cosh,
1173    Sin,
1174    Sinh,
1175    Tan,
1176    Tanh,
1177    Acos,
1178    Asin,
1179    Atan,
1180    Atan2,
1181    Asinh,
1182    Acosh,
1183    Atanh,
1184    Radians,
1185    Degrees,
1186    // decomposition
1187    Ceil,
1188    Floor,
1189    Round,
1190    Fract,
1191    Trunc,
1192    Modf,
1193    Frexp,
1194    Ldexp,
1195    // exponent
1196    Exp,
1197    Exp2,
1198    Log,
1199    Log2,
1200    Pow,
1201    // geometry
1202    Dot,
1203    Dot4I8Packed,
1204    Dot4U8Packed,
1205    Outer,
1206    Cross,
1207    Distance,
1208    Length,
1209    Normalize,
1210    FaceForward,
1211    Reflect,
1212    Refract,
1213    // computational
1214    Sign,
1215    Fma,
1216    Mix,
1217    Step,
1218    SmoothStep,
1219    Sqrt,
1220    InverseSqrt,
1221    Inverse,
1222    Transpose,
1223    Determinant,
1224    QuantizeToF16,
1225    // bits
1226    CountTrailingZeros,
1227    CountLeadingZeros,
1228    CountOneBits,
1229    ReverseBits,
1230    ExtractBits,
1231    InsertBits,
1232    FirstTrailingBit,
1233    FirstLeadingBit,
1234    // data packing
1235    Pack4x8snorm,
1236    Pack4x8unorm,
1237    Pack2x16snorm,
1238    Pack2x16unorm,
1239    Pack2x16float,
1240    Pack4xI8,
1241    Pack4xU8,
1242    Pack4xI8Clamp,
1243    Pack4xU8Clamp,
1244    // data unpacking
1245    Unpack4x8snorm,
1246    Unpack4x8unorm,
1247    Unpack2x16snorm,
1248    Unpack2x16unorm,
1249    Unpack2x16float,
1250    Unpack4xI8,
1251    Unpack4xU8,
1252}
1253
1254/// Sampling modifier to control the level of detail.
1255///
1256/// All `Handle<Expression>` values here refer to an expression in
1257/// [`Function::expressions`].
1258#[derive(Clone, Copy, Debug, PartialEq)]
1259#[cfg_attr(feature = "serialize", derive(Serialize))]
1260#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1261#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1262pub enum SampleLevel {
1263    Auto,
1264    Zero,
1265    Exact(Handle<Expression>),
1266    Bias(Handle<Expression>),
1267    Gradient {
1268        x: Handle<Expression>,
1269        y: Handle<Expression>,
1270    },
1271}
1272
1273/// Type of an image query.
1274///
1275/// All `Handle<Expression>` values here refer to an expression in
1276/// [`Function::expressions`].
1277#[derive(Clone, Copy, Debug, PartialEq)]
1278#[cfg_attr(feature = "serialize", derive(Serialize))]
1279#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1280#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1281pub enum ImageQuery {
1282    /// Get the size at the specified level.
1283    ///
1284    /// The return value is a `u32` for 1D images, and a `vecN<u32>`
1285    /// for an image with dimensions N > 2.
1286    Size {
1287        /// If `None`, the base level is considered.
1288        level: Option<Handle<Expression>>,
1289    },
1290    /// Get the number of mipmap levels, a `u32`.
1291    NumLevels,
1292    /// Get the number of array layers, a `u32`.
1293    NumLayers,
1294    /// Get the number of samples, a `u32`.
1295    NumSamples,
1296}
1297
1298/// Component selection for a vector swizzle.
1299#[repr(u8)]
1300#[derive(Clone, Copy, Debug, PartialEq, PartialOrd)]
1301#[cfg_attr(feature = "serialize", derive(Serialize))]
1302#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1303#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1304pub enum SwizzleComponent {
1305    X = 0,
1306    Y = 1,
1307    Z = 2,
1308    W = 3,
1309}
1310
1311/// The specific behavior of a [`SubgroupGather`] statement.
1312///
1313/// All `Handle<Expression>` values here refer to an expression in
1314/// [`Function::expressions`].
1315///
1316/// [`SubgroupGather`]: Statement::SubgroupGather
1317#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1318#[cfg_attr(feature = "serialize", derive(Serialize))]
1319#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1320#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1321pub enum GatherMode {
1322    /// All gather from the active lane with the smallest index
1323    BroadcastFirst,
1324    /// All gather from the same lane at the index given by the expression
1325    Broadcast(Handle<Expression>),
1326    /// Each gathers from a different lane at the index given by the expression
1327    Shuffle(Handle<Expression>),
1328    /// Each gathers from their lane plus the shift given by the expression
1329    ShuffleDown(Handle<Expression>),
1330    /// Each gathers from their lane minus the shift given by the expression
1331    ShuffleUp(Handle<Expression>),
1332    /// Each gathers from their lane xored with the given by the expression
1333    ShuffleXor(Handle<Expression>),
1334    /// All gather from the same quad lane at the index given by the expression
1335    QuadBroadcast(Handle<Expression>),
1336    /// Each gathers from the opposite quad lane along the given direction
1337    QuadSwap(Direction),
1338}
1339
1340#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1341#[cfg_attr(feature = "serialize", derive(Serialize))]
1342#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1343#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1344pub enum Direction {
1345    X = 0,
1346    Y = 1,
1347    Diagonal = 2,
1348}
1349
1350#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1351#[cfg_attr(feature = "serialize", derive(Serialize))]
1352#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1353#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1354pub enum SubgroupOperation {
1355    All = 0,
1356    Any = 1,
1357    Add = 2,
1358    Mul = 3,
1359    Min = 4,
1360    Max = 5,
1361    And = 6,
1362    Or = 7,
1363    Xor = 8,
1364}
1365
1366#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
1367#[cfg_attr(feature = "serialize", derive(Serialize))]
1368#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1369#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1370pub enum CollectiveOperation {
1371    Reduce = 0,
1372    InclusiveScan = 1,
1373    ExclusiveScan = 2,
1374}
1375
1376bitflags::bitflags! {
1377    /// Memory barrier flags.
1378    #[cfg_attr(feature = "serialize", derive(Serialize))]
1379    #[cfg_attr(feature = "deserialize", derive(Deserialize))]
1380    #[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1381    #[derive(Clone, Copy, Debug, Default, Eq, PartialEq)]
1382    pub struct Barrier: u32 {
1383        /// Barrier affects all [`AddressSpace::Storage`] accesses.
1384        const STORAGE = 1 << 0;
1385        /// Barrier affects all [`AddressSpace::WorkGroup`] accesses.
1386        const WORK_GROUP = 1 << 1;
1387        /// Barrier synchronizes execution across all invocations within a subgroup that execute this instruction.
1388        const SUB_GROUP = 1 << 2;
1389        /// Barrier synchronizes texture memory accesses in a workgroup.
1390        const TEXTURE = 1 << 3;
1391    }
1392}
1393
1394/// An expression that can be evaluated to obtain a value.
1395///
1396/// This is a Single Static Assignment (SSA) scheme similar to SPIR-V.
1397///
1398/// When an `Expression` variant holds `Handle<Expression>` fields, they refer
1399/// to another expression in the same arena, unless explicitly noted otherwise.
1400/// One `Arena<Expression>` may only refer to a different arena indirectly, via
1401/// [`Constant`] or [`Override`] expressions, which hold handles for their
1402/// respective types.
1403///
1404/// [`Constant`]: Expression::Constant
1405/// [`Override`]: Expression::Override
1406#[derive(Clone, Debug, PartialEq)]
1407#[cfg_attr(feature = "serialize", derive(Serialize))]
1408#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1409#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1410pub enum Expression {
1411    /// Literal.
1412    Literal(Literal),
1413    /// Constant value.
1414    Constant(Handle<Constant>),
1415    /// Pipeline-overridable constant.
1416    Override(Handle<Override>),
1417    /// Zero value of a type.
1418    ZeroValue(Handle<Type>),
1419    /// Composite expression.
1420    Compose {
1421        ty: Handle<Type>,
1422        components: Vec<Handle<Expression>>,
1423    },
1424
1425    /// Array access with a computed index.
1426    ///
1427    /// ## Typing rules
1428    ///
1429    /// The `base` operand must be some composite type: [`Vector`], [`Matrix`],
1430    /// [`Array`], a [`Pointer`] to one of those, or a [`ValuePointer`] with a
1431    /// `size`.
1432    ///
1433    /// The `index` operand must be an integer, signed or unsigned.
1434    ///
1435    /// Indexing a [`Vector`] or [`Array`] produces a value of its element type.
1436    /// Indexing a [`Matrix`] produces a [`Vector`].
1437    ///
1438    /// Indexing a [`Pointer`] to any of the above produces a pointer to the
1439    /// element/component type, in the same [`space`]. In the case of [`Array`],
1440    /// the result is an actual [`Pointer`], but for vectors and matrices, there
1441    /// may not be any type in the arena representing the component's type, so
1442    /// those produce [`ValuePointer`] types equivalent to the appropriate
1443    /// [`Pointer`].
1444    ///
1445    /// ## Dynamic indexing restrictions
1446    ///
1447    /// To accommodate restrictions in some of the shader languages that Naga
1448    /// targets, it is not permitted to subscript a matrix with a dynamically
1449    /// computed index unless that matrix appears behind a pointer. In other
1450    /// words, if the inner type of `base` is [`Matrix`], then `index` must be a
1451    /// constant. But if the type of `base` is a [`Pointer`] to an matrix, then
1452    /// the index may be any expression of integer type.
1453    ///
1454    /// You can use the [`Expression::is_dynamic_index`] method to determine
1455    /// whether a given index expression requires matrix base operands to be
1456    /// behind a pointer.
1457    ///
1458    /// (It would be simpler to always require the use of `AccessIndex` when
1459    /// subscripting matrices that are not behind pointers, but to accommodate
1460    /// existing front ends, Naga also permits `Access`, with a restricted
1461    /// `index`.)
1462    ///
1463    /// [`Vector`]: TypeInner::Vector
1464    /// [`Matrix`]: TypeInner::Matrix
1465    /// [`Array`]: TypeInner::Array
1466    /// [`Pointer`]: TypeInner::Pointer
1467    /// [`space`]: TypeInner::Pointer::space
1468    /// [`ValuePointer`]: TypeInner::ValuePointer
1469    /// [`Float`]: ScalarKind::Float
1470    Access {
1471        base: Handle<Expression>,
1472        index: Handle<Expression>,
1473    },
1474    /// Access the same types as [`Access`], plus [`Struct`] with a known index.
1475    ///
1476    /// [`Access`]: Expression::Access
1477    /// [`Struct`]: TypeInner::Struct
1478    AccessIndex {
1479        base: Handle<Expression>,
1480        index: u32,
1481    },
1482    /// Splat scalar into a vector.
1483    Splat {
1484        size: VectorSize,
1485        value: Handle<Expression>,
1486    },
1487    /// Vector swizzle.
1488    Swizzle {
1489        size: VectorSize,
1490        vector: Handle<Expression>,
1491        pattern: [SwizzleComponent; 4],
1492    },
1493
1494    /// Reference a function parameter, by its index.
1495    ///
1496    /// A `FunctionArgument` expression evaluates to the argument's value.
1497    FunctionArgument(u32),
1498
1499    /// Reference a global variable.
1500    ///
1501    /// If the given `GlobalVariable`'s [`space`] is [`AddressSpace::Handle`],
1502    /// then the variable stores some opaque type like a sampler or an image,
1503    /// and a `GlobalVariable` expression referring to it produces the
1504    /// variable's value directly.
1505    ///
1506    /// For any other address space, a `GlobalVariable` expression produces a
1507    /// pointer to the variable's value. You must use a [`Load`] expression to
1508    /// retrieve its value, or a [`Store`] statement to assign it a new value.
1509    ///
1510    /// [`space`]: GlobalVariable::space
1511    /// [`Load`]: Expression::Load
1512    /// [`Store`]: Statement::Store
1513    GlobalVariable(Handle<GlobalVariable>),
1514
1515    /// Reference a local variable.
1516    ///
1517    /// A `LocalVariable` expression evaluates to a pointer to the variable's value.
1518    /// You must use a [`Load`](Expression::Load) expression to retrieve its value,
1519    /// or a [`Store`](Statement::Store) statement to assign it a new value.
1520    LocalVariable(Handle<LocalVariable>),
1521
1522    /// Load a value indirectly.
1523    ///
1524    /// For [`TypeInner::Atomic`] the result is a corresponding scalar.
1525    /// For other types behind the `pointer<T>`, the result is `T`.
1526    Load { pointer: Handle<Expression> },
1527    /// Sample a point from a sampled or a depth image.
1528    ImageSample {
1529        image: Handle<Expression>,
1530        sampler: Handle<Expression>,
1531        /// If Some(), this operation is a gather operation
1532        /// on the selected component.
1533        gather: Option<SwizzleComponent>,
1534        coordinate: Handle<Expression>,
1535        array_index: Option<Handle<Expression>>,
1536        /// This must be a const-expression.
1537        offset: Option<Handle<Expression>>,
1538        level: SampleLevel,
1539        depth_ref: Option<Handle<Expression>>,
1540        /// Whether the sampling operation should clamp each component of
1541        /// `coordinate` to the range `[half_texel, 1 - half_texel]`, regardless
1542        /// of `sampler`.
1543        clamp_to_edge: bool,
1544    },
1545
1546    /// Load a texel from an image.
1547    ///
1548    /// For most images, this returns a four-element vector of the same
1549    /// [`ScalarKind`] as the image. If the format of the image does not have
1550    /// four components, default values are provided: the first three components
1551    /// (typically R, G, and B) default to zero, and the final component
1552    /// (typically alpha) defaults to one.
1553    ///
1554    /// However, if the image's [`class`] is [`Depth`], then this returns a
1555    /// [`Float`] scalar value.
1556    ///
1557    /// [`ScalarKind`]: ScalarKind
1558    /// [`class`]: TypeInner::Image::class
1559    /// [`Depth`]: ImageClass::Depth
1560    /// [`Float`]: ScalarKind::Float
1561    ImageLoad {
1562        /// The image to load a texel from. This must have type [`Image`]. (This
1563        /// will necessarily be a [`GlobalVariable`] or [`FunctionArgument`]
1564        /// expression, since no other expressions are allowed to have that
1565        /// type.)
1566        ///
1567        /// [`Image`]: TypeInner::Image
1568        /// [`GlobalVariable`]: Expression::GlobalVariable
1569        /// [`FunctionArgument`]: Expression::FunctionArgument
1570        image: Handle<Expression>,
1571
1572        /// The coordinate of the texel we wish to load. This must be a scalar
1573        /// for [`D1`] images, a [`Bi`] vector for [`D2`] images, and a [`Tri`]
1574        /// vector for [`D3`] images. (Array indices, sample indices, and
1575        /// explicit level-of-detail values are supplied separately.) Its
1576        /// component type must be [`Sint`].
1577        ///
1578        /// [`D1`]: ImageDimension::D1
1579        /// [`D2`]: ImageDimension::D2
1580        /// [`D3`]: ImageDimension::D3
1581        /// [`Bi`]: VectorSize::Bi
1582        /// [`Tri`]: VectorSize::Tri
1583        /// [`Sint`]: ScalarKind::Sint
1584        coordinate: Handle<Expression>,
1585
1586        /// The index into an arrayed image. If the [`arrayed`] flag in
1587        /// `image`'s type is `true`, then this must be `Some(expr)`, where
1588        /// `expr` is a [`Sint`] scalar. Otherwise, it must be `None`.
1589        ///
1590        /// [`arrayed`]: TypeInner::Image::arrayed
1591        /// [`Sint`]: ScalarKind::Sint
1592        array_index: Option<Handle<Expression>>,
1593
1594        /// A sample index, for multisampled [`Sampled`] and [`Depth`] images.
1595        ///
1596        /// [`Sampled`]: ImageClass::Sampled
1597        /// [`Depth`]: ImageClass::Depth
1598        sample: Option<Handle<Expression>>,
1599
1600        /// A level of detail, for mipmapped images.
1601        ///
1602        /// This must be present when accessing non-multisampled
1603        /// [`Sampled`] and [`Depth`] images, even if only the
1604        /// full-resolution level is present (in which case the only
1605        /// valid level is zero).
1606        ///
1607        /// [`Sampled`]: ImageClass::Sampled
1608        /// [`Depth`]: ImageClass::Depth
1609        level: Option<Handle<Expression>>,
1610    },
1611
1612    /// Query information from an image.
1613    ImageQuery {
1614        image: Handle<Expression>,
1615        query: ImageQuery,
1616    },
1617    /// Apply an unary operator.
1618    Unary {
1619        op: UnaryOperator,
1620        expr: Handle<Expression>,
1621    },
1622    /// Apply a binary operator.
1623    Binary {
1624        op: BinaryOperator,
1625        left: Handle<Expression>,
1626        right: Handle<Expression>,
1627    },
1628    /// Select between two values based on a condition.
1629    ///
1630    /// Note that, because expressions have no side effects, it is unobservable
1631    /// whether the non-selected branch is evaluated.
1632    Select {
1633        /// Boolean expression
1634        condition: Handle<Expression>,
1635        accept: Handle<Expression>,
1636        reject: Handle<Expression>,
1637    },
1638    /// Compute the derivative on an axis.
1639    Derivative {
1640        axis: DerivativeAxis,
1641        ctrl: DerivativeControl,
1642        expr: Handle<Expression>,
1643    },
1644    /// Call a relational function.
1645    Relational {
1646        fun: RelationalFunction,
1647        argument: Handle<Expression>,
1648    },
1649    /// Call a math function
1650    Math {
1651        fun: MathFunction,
1652        arg: Handle<Expression>,
1653        arg1: Option<Handle<Expression>>,
1654        arg2: Option<Handle<Expression>>,
1655        arg3: Option<Handle<Expression>>,
1656    },
1657    /// Cast a simple type to another kind.
1658    As {
1659        /// Source expression, which can only be a scalar or a vector.
1660        expr: Handle<Expression>,
1661        /// Target scalar kind.
1662        kind: ScalarKind,
1663        /// If provided, converts to the specified byte width.
1664        /// Otherwise, bitcast.
1665        convert: Option<Bytes>,
1666    },
1667    /// Result of calling another function.
1668    CallResult(Handle<Function>),
1669
1670    /// Result of an atomic operation.
1671    ///
1672    /// This expression must be referred to by the [`result`] field of exactly one
1673    /// [`Atomic`][stmt] statement somewhere in the same function. Let `T` be the
1674    /// scalar type contained by the [`Atomic`][type] value that the statement
1675    /// operates on.
1676    ///
1677    /// If `comparison` is `false`, then `ty` must be the scalar type `T`.
1678    ///
1679    /// If `comparison` is `true`, then `ty` must be a [`Struct`] with two members:
1680    ///
1681    /// - A member named `old_value`, whose type is `T`, and
1682    ///
1683    /// - A member named `exchanged`, of type [`BOOL`].
1684    ///
1685    /// [`result`]: Statement::Atomic::result
1686    /// [stmt]: Statement::Atomic
1687    /// [type]: TypeInner::Atomic
1688    /// [`Struct`]: TypeInner::Struct
1689    /// [`BOOL`]: Scalar::BOOL
1690    AtomicResult { ty: Handle<Type>, comparison: bool },
1691
1692    /// Result of a [`WorkGroupUniformLoad`] statement.
1693    ///
1694    /// [`WorkGroupUniformLoad`]: Statement::WorkGroupUniformLoad
1695    WorkGroupUniformLoadResult {
1696        /// The type of the result
1697        ty: Handle<Type>,
1698    },
1699    /// Get the length of an array.
1700    /// The expression must resolve to a pointer to an array with a dynamic size.
1701    ///
1702    /// This doesn't match the semantics of spirv's `OpArrayLength`, which must be passed
1703    /// a pointer to a structure containing a runtime array in its' last field.
1704    ArrayLength(Handle<Expression>),
1705
1706    /// Get the Positions of the triangle hit by the [`RayQuery`]
1707    ///
1708    /// [`RayQuery`]: Statement::RayQuery
1709    RayQueryVertexPositions {
1710        query: Handle<Expression>,
1711        committed: bool,
1712    },
1713
1714    /// Result of a [`Proceed`] [`RayQuery`] statement.
1715    ///
1716    /// [`Proceed`]: RayQueryFunction::Proceed
1717    /// [`RayQuery`]: Statement::RayQuery
1718    RayQueryProceedResult,
1719
1720    /// Return an intersection found by `query`.
1721    ///
1722    /// If `committed` is true, return the committed result available when
1723    RayQueryGetIntersection {
1724        query: Handle<Expression>,
1725        committed: bool,
1726    },
1727    /// Result of a [`SubgroupBallot`] statement.
1728    ///
1729    /// [`SubgroupBallot`]: Statement::SubgroupBallot
1730    SubgroupBallotResult,
1731    /// Result of a [`SubgroupCollectiveOperation`] or [`SubgroupGather`] statement.
1732    ///
1733    /// [`SubgroupCollectiveOperation`]: Statement::SubgroupCollectiveOperation
1734    /// [`SubgroupGather`]: Statement::SubgroupGather
1735    SubgroupOperationResult { ty: Handle<Type> },
1736}
1737
1738/// The value of the switch case.
1739#[derive(Clone, Copy, Debug, Eq, Hash, PartialEq)]
1740#[cfg_attr(feature = "serialize", derive(Serialize))]
1741#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1742#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1743pub enum SwitchValue {
1744    I32(i32),
1745    U32(u32),
1746    Default,
1747}
1748
1749/// A case for a switch statement.
1750// Clone is used only for error reporting and is not intended for end users
1751#[derive(Clone, Debug)]
1752#[cfg_attr(feature = "serialize", derive(Serialize))]
1753#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1754#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1755pub struct SwitchCase {
1756    /// Value, upon which the case is considered true.
1757    pub value: SwitchValue,
1758    /// Body of the case.
1759    pub body: Block,
1760    /// If true, the control flow continues to the next case in the list,
1761    /// or default.
1762    pub fall_through: bool,
1763}
1764
1765/// An operation that a [`RayQuery` statement] applies to its [`query`] operand.
1766///
1767/// [`RayQuery` statement]: Statement::RayQuery
1768/// [`query`]: Statement::RayQuery::query
1769#[derive(Clone, Debug)]
1770#[cfg_attr(feature = "serialize", derive(Serialize))]
1771#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1772#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1773pub enum RayQueryFunction {
1774    /// Initialize the `RayQuery` object.
1775    Initialize {
1776        /// The acceleration structure within which this query should search for hits.
1777        ///
1778        /// The expression must be an [`AccelerationStructure`].
1779        ///
1780        /// [`AccelerationStructure`]: TypeInner::AccelerationStructure
1781        acceleration_structure: Handle<Expression>,
1782
1783        #[allow(rustdoc::private_intra_doc_links)]
1784        /// A struct of detailed parameters for the ray query.
1785        ///
1786        /// This expression should have the struct type given in
1787        /// [`SpecialTypes::ray_desc`]. This is available in the WGSL
1788        /// front end as the `RayDesc` type.
1789        descriptor: Handle<Expression>,
1790    },
1791
1792    /// Start or continue the query given by the statement's [`query`] operand.
1793    ///
1794    /// After executing this statement, the `result` expression is a
1795    /// [`Bool`] scalar indicating whether there are more intersection
1796    /// candidates to consider.
1797    ///
1798    /// [`query`]: Statement::RayQuery::query
1799    /// [`Bool`]: ScalarKind::Bool
1800    Proceed {
1801        result: Handle<Expression>,
1802    },
1803
1804    /// Add a candidate generated intersection to be included
1805    /// in the determination of the closest hit for a ray query.
1806    GenerateIntersection {
1807        hit_t: Handle<Expression>,
1808    },
1809
1810    /// Confirm a triangle intersection to be included in the determination of
1811    /// the closest hit for a ray query.
1812    ConfirmIntersection,
1813
1814    Terminate,
1815}
1816
1817//TODO: consider removing `Clone`. It's not valid to clone `Statement::Emit` anyway.
1818/// Instructions which make up an executable block.
1819///
1820/// `Handle<Expression>` and `Range<Expression>` values in `Statement` variants
1821/// refer to expressions in [`Function::expressions`], unless otherwise noted.
1822// Clone is used only for error reporting and is not intended for end users
1823#[derive(Clone, Debug)]
1824#[cfg_attr(feature = "serialize", derive(Serialize))]
1825#[cfg_attr(feature = "deserialize", derive(Deserialize))]
1826#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
1827pub enum Statement {
1828    /// Emit a range of expressions, visible to all statements that follow in this block.
1829    ///
1830    /// See the [module-level documentation][emit] for details.
1831    ///
1832    /// [emit]: index.html#expression-evaluation-time
1833    Emit(Range<Expression>),
1834    /// A block containing more statements, to be executed sequentially.
1835    Block(Block),
1836    /// Conditionally executes one of two blocks, based on the value of the condition.
1837    ///
1838    /// Naga IR does not have "phi" instructions. If you need to use
1839    /// values computed in an `accept` or `reject` block after the `If`,
1840    /// store them in a [`LocalVariable`].
1841    If {
1842        condition: Handle<Expression>, //bool
1843        accept: Block,
1844        reject: Block,
1845    },
1846    /// Conditionally executes one of multiple blocks, based on the value of the selector.
1847    ///
1848    /// Each case must have a distinct [`value`], exactly one of which must be
1849    /// [`Default`]. The `Default` may appear at any position, and covers all
1850    /// values not explicitly appearing in other cases. A `Default` appearing in
1851    /// the midst of the list of cases does not shadow the cases that follow.
1852    ///
1853    /// Some backend languages don't support fallthrough (HLSL due to FXC,
1854    /// WGSL), and may translate fallthrough cases in the IR by duplicating
1855    /// code. However, all backend languages do support cases selected by
1856    /// multiple values, like `case 1: case 2: case 3: { ... }`. This is
1857    /// represented in the IR as a series of fallthrough cases with empty
1858    /// bodies, except for the last.
1859    ///
1860    /// Naga IR does not have "phi" instructions. If you need to use
1861    /// values computed in a [`SwitchCase::body`] block after the `Switch`,
1862    /// store them in a [`LocalVariable`].
1863    ///
1864    /// [`value`]: SwitchCase::value
1865    /// [`body`]: SwitchCase::body
1866    /// [`Default`]: SwitchValue::Default
1867    Switch {
1868        selector: Handle<Expression>,
1869        cases: Vec<SwitchCase>,
1870    },
1871
1872    /// Executes a block repeatedly.
1873    ///
1874    /// Each iteration of the loop executes the `body` block, followed by the
1875    /// `continuing` block.
1876    ///
1877    /// Executing a [`Break`], [`Return`] or [`Kill`] statement exits the loop.
1878    ///
1879    /// A [`Continue`] statement in `body` jumps to the `continuing` block. The
1880    /// `continuing` block is meant to be used to represent structures like the
1881    /// third expression of a C-style `for` loop head, to which `continue`
1882    /// statements in the loop's body jump.
1883    ///
1884    /// The `continuing` block and its substatements must not contain `Return`
1885    /// or `Kill` statements, or any `Break` or `Continue` statements targeting
1886    /// this loop. (It may have `Break` and `Continue` statements targeting
1887    /// loops or switches nested within the `continuing` block.) Expressions
1888    /// emitted in `body` are in scope in `continuing`.
1889    ///
1890    /// If present, `break_if` is an expression which is evaluated after the
1891    /// continuing block. Expressions emitted in `body` or `continuing` are
1892    /// considered to be in scope. If the expression's value is true, control
1893    /// continues after the `Loop` statement, rather than branching back to the
1894    /// top of body as usual. The `break_if` expression corresponds to a "break
1895    /// if" statement in WGSL, or a loop whose back edge is an
1896    /// `OpBranchConditional` instruction in SPIR-V.
1897    ///
1898    /// Naga IR does not have "phi" instructions. If you need to use
1899    /// values computed in a `body` or `continuing` block after the
1900    /// `Loop`, store them in a [`LocalVariable`].
1901    ///
1902    /// [`Break`]: Statement::Break
1903    /// [`Continue`]: Statement::Continue
1904    /// [`Kill`]: Statement::Kill
1905    /// [`Return`]: Statement::Return
1906    /// [`break if`]: Self::Loop::break_if
1907    Loop {
1908        body: Block,
1909        continuing: Block,
1910        break_if: Option<Handle<Expression>>,
1911    },
1912
1913    /// Exits the innermost enclosing [`Loop`] or [`Switch`].
1914    ///
1915    /// A `Break` statement may only appear within a [`Loop`] or [`Switch`]
1916    /// statement. It may not break out of a [`Loop`] from within the loop's
1917    /// `continuing` block.
1918    ///
1919    /// [`Loop`]: Statement::Loop
1920    /// [`Switch`]: Statement::Switch
1921    Break,
1922
1923    /// Skips to the `continuing` block of the innermost enclosing [`Loop`].
1924    ///
1925    /// A `Continue` statement may only appear within the `body` block of the
1926    /// innermost enclosing [`Loop`] statement. It must not appear within that
1927    /// loop's `continuing` block.
1928    ///
1929    /// [`Loop`]: Statement::Loop
1930    Continue,
1931
1932    /// Returns from the function (possibly with a value).
1933    ///
1934    /// `Return` statements are forbidden within the `continuing` block of a
1935    /// [`Loop`] statement.
1936    ///
1937    /// [`Loop`]: Statement::Loop
1938    Return { value: Option<Handle<Expression>> },
1939
1940    /// Aborts the current shader execution.
1941    ///
1942    /// `Kill` statements are forbidden within the `continuing` block of a
1943    /// [`Loop`] statement.
1944    ///
1945    /// [`Loop`]: Statement::Loop
1946    Kill,
1947
1948    /// Synchronize invocations within the work group.
1949    /// The `Barrier` flags control which memory accesses should be synchronized.
1950    /// If empty, this becomes purely an execution barrier.
1951    ControlBarrier(Barrier),
1952
1953    /// Synchronize invocations within the work group.
1954    /// The `Barrier` flags control which memory accesses should be synchronized.
1955    MemoryBarrier(Barrier),
1956
1957    /// Stores a value at an address.
1958    ///
1959    /// For [`TypeInner::Atomic`] type behind the pointer, the value
1960    /// has to be a corresponding scalar.
1961    /// For other types behind the `pointer<T>`, the value is `T`.
1962    ///
1963    /// This statement is a barrier for any operations on the
1964    /// `Expression::LocalVariable` or `Expression::GlobalVariable`
1965    /// that is the destination of an access chain, started
1966    /// from the `pointer`.
1967    Store {
1968        pointer: Handle<Expression>,
1969        value: Handle<Expression>,
1970    },
1971    /// Stores a texel value to an image.
1972    ///
1973    /// The `image`, `coordinate`, and `array_index` fields have the same
1974    /// meanings as the corresponding operands of an [`ImageLoad`] expression;
1975    /// see that documentation for details. Storing into multisampled images or
1976    /// images with mipmaps is not supported, so there are no `level` or
1977    /// `sample` operands.
1978    ///
1979    /// This statement is a barrier for any operations on the corresponding
1980    /// [`Expression::GlobalVariable`] for this image.
1981    ///
1982    /// [`ImageLoad`]: Expression::ImageLoad
1983    ImageStore {
1984        image: Handle<Expression>,
1985        coordinate: Handle<Expression>,
1986        array_index: Option<Handle<Expression>>,
1987        value: Handle<Expression>,
1988    },
1989    /// Atomic function.
1990    Atomic {
1991        /// Pointer to an atomic value.
1992        ///
1993        /// This must be a [`Pointer`] to an [`Atomic`] value. The atomic's
1994        /// scalar type may be [`I32`] or [`U32`].
1995        ///
1996        /// If [`SHADER_INT64_ATOMIC_MIN_MAX`] or [`SHADER_INT64_ATOMIC_ALL_OPS`] are
1997        /// enabled, this may also be [`I64`] or [`U64`].
1998        ///
1999        /// If [`SHADER_FLOAT32_ATOMIC`] is enabled, this may be [`F32`].
2000        ///
2001        /// [`Pointer`]: TypeInner::Pointer
2002        /// [`Atomic`]: TypeInner::Atomic
2003        /// [`I32`]: Scalar::I32
2004        /// [`U32`]: Scalar::U32
2005        /// [`SHADER_INT64_ATOMIC_MIN_MAX`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX
2006        /// [`SHADER_INT64_ATOMIC_ALL_OPS`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
2007        /// [`SHADER_FLOAT32_ATOMIC`]: crate::valid::Capabilities::SHADER_FLOAT32_ATOMIC
2008        /// [`I64`]: Scalar::I64
2009        /// [`U64`]: Scalar::U64
2010        /// [`F32`]: Scalar::F32
2011        pointer: Handle<Expression>,
2012
2013        /// Function to run on the atomic value.
2014        ///
2015        /// If [`pointer`] refers to a 64-bit atomic value, then:
2016        ///
2017        /// - The [`SHADER_INT64_ATOMIC_ALL_OPS`] capability allows any [`AtomicFunction`]
2018        ///   value here.
2019        ///
2020        /// - The [`SHADER_INT64_ATOMIC_MIN_MAX`] capability allows
2021        ///   [`AtomicFunction::Min`] and [`AtomicFunction::Max`]
2022        ///   in the [`Storage`] address space here.
2023        ///
2024        /// - If neither of those capabilities are present, then 64-bit scalar
2025        ///   atomics are not allowed.
2026        ///
2027        /// If [`pointer`] refers to a 32-bit floating-point atomic value, then:
2028        ///
2029        /// - The [`SHADER_FLOAT32_ATOMIC`] capability allows [`AtomicFunction::Add`],
2030        ///   [`AtomicFunction::Subtract`], and [`AtomicFunction::Exchange { compare: None }`]
2031        ///   in the [`Storage`] address space here.
2032        ///
2033        /// [`AtomicFunction::Exchange { compare: None }`]: AtomicFunction::Exchange
2034        /// [`pointer`]: Statement::Atomic::pointer
2035        /// [`Storage`]: AddressSpace::Storage
2036        /// [`SHADER_INT64_ATOMIC_MIN_MAX`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX
2037        /// [`SHADER_INT64_ATOMIC_ALL_OPS`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
2038        /// [`SHADER_FLOAT32_ATOMIC`]: crate::valid::Capabilities::SHADER_FLOAT32_ATOMIC
2039        fun: AtomicFunction,
2040
2041        /// Value to use in the function.
2042        ///
2043        /// This must be a scalar of the same type as [`pointer`]'s atomic's scalar type.
2044        ///
2045        /// [`pointer`]: Statement::Atomic::pointer
2046        value: Handle<Expression>,
2047
2048        /// [`AtomicResult`] expression representing this function's result.
2049        ///
2050        /// If [`fun`] is [`Exchange { compare: None }`], this must be `Some`,
2051        /// as otherwise that operation would be equivalent to a simple [`Store`]
2052        /// to the atomic.
2053        ///
2054        /// Otherwise, this may be `None` if the return value of the operation is not needed.
2055        ///
2056        /// If `pointer` refers to a 64-bit atomic value, [`SHADER_INT64_ATOMIC_MIN_MAX`]
2057        /// is enabled, and [`SHADER_INT64_ATOMIC_ALL_OPS`] is not, this must be `None`.
2058        ///
2059        /// [`AtomicResult`]: crate::Expression::AtomicResult
2060        /// [`fun`]: Statement::Atomic::fun
2061        /// [`Store`]: Statement::Store
2062        /// [`Exchange { compare: None }`]: AtomicFunction::Exchange
2063        /// [`SHADER_INT64_ATOMIC_MIN_MAX`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX
2064        /// [`SHADER_INT64_ATOMIC_ALL_OPS`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
2065        result: Option<Handle<Expression>>,
2066    },
2067    /// Performs an atomic operation on a texel value of an image.
2068    ///
2069    /// Doing atomics on images with mipmaps is not supported, so there is no
2070    /// `level` operand.
2071    ImageAtomic {
2072        /// The image to perform an atomic operation on. This must have type
2073        /// [`Image`]. (This will necessarily be a [`GlobalVariable`] or
2074        /// [`FunctionArgument`] expression, since no other expressions are
2075        /// allowed to have that type.)
2076        ///
2077        /// [`Image`]: TypeInner::Image
2078        /// [`GlobalVariable`]: Expression::GlobalVariable
2079        /// [`FunctionArgument`]: Expression::FunctionArgument
2080        image: Handle<Expression>,
2081
2082        /// The coordinate of the texel we wish to load. This must be a scalar
2083        /// for [`D1`] images, a [`Bi`] vector for [`D2`] images, and a [`Tri`]
2084        /// vector for [`D3`] images. (Array indices, sample indices, and
2085        /// explicit level-of-detail values are supplied separately.) Its
2086        /// component type must be [`Sint`].
2087        ///
2088        /// [`D1`]: ImageDimension::D1
2089        /// [`D2`]: ImageDimension::D2
2090        /// [`D3`]: ImageDimension::D3
2091        /// [`Bi`]: VectorSize::Bi
2092        /// [`Tri`]: VectorSize::Tri
2093        /// [`Sint`]: ScalarKind::Sint
2094        coordinate: Handle<Expression>,
2095
2096        /// The index into an arrayed image. If the [`arrayed`] flag in
2097        /// `image`'s type is `true`, then this must be `Some(expr)`, where
2098        /// `expr` is a [`Sint`] scalar. Otherwise, it must be `None`.
2099        ///
2100        /// [`arrayed`]: TypeInner::Image::arrayed
2101        /// [`Sint`]: ScalarKind::Sint
2102        array_index: Option<Handle<Expression>>,
2103
2104        /// The kind of atomic operation to perform on the texel.
2105        fun: AtomicFunction,
2106
2107        /// The value with which to perform the atomic operation.
2108        value: Handle<Expression>,
2109    },
2110    /// Load uniformly from a uniform pointer in the workgroup address space.
2111    ///
2112    /// Corresponds to the [`workgroupUniformLoad`](https://www.w3.org/TR/WGSL/#workgroupUniformLoad-builtin)
2113    /// built-in function of wgsl, and has the same barrier semantics
2114    WorkGroupUniformLoad {
2115        /// This must be of type [`Pointer`] in the [`WorkGroup`] address space
2116        ///
2117        /// [`Pointer`]: TypeInner::Pointer
2118        /// [`WorkGroup`]: AddressSpace::WorkGroup
2119        pointer: Handle<Expression>,
2120        /// The [`WorkGroupUniformLoadResult`] expression representing this load's result.
2121        ///
2122        /// [`WorkGroupUniformLoadResult`]: Expression::WorkGroupUniformLoadResult
2123        result: Handle<Expression>,
2124    },
2125    /// Calls a function.
2126    ///
2127    /// If the `result` is `Some`, the corresponding expression has to be
2128    /// `Expression::CallResult`, and this statement serves as a barrier for any
2129    /// operations on that expression.
2130    Call {
2131        function: Handle<Function>,
2132        arguments: Vec<Handle<Expression>>,
2133        result: Option<Handle<Expression>>,
2134    },
2135    RayQuery {
2136        /// The [`RayQuery`] object this statement operates on.
2137        ///
2138        /// [`RayQuery`]: TypeInner::RayQuery
2139        query: Handle<Expression>,
2140
2141        /// The specific operation we're performing on `query`.
2142        fun: RayQueryFunction,
2143    },
2144    /// Calculate a bitmask using a boolean from each active thread in the subgroup
2145    SubgroupBallot {
2146        /// The [`SubgroupBallotResult`] expression representing this load's result.
2147        ///
2148        /// [`SubgroupBallotResult`]: Expression::SubgroupBallotResult
2149        result: Handle<Expression>,
2150        /// The value from this thread to store in the ballot
2151        predicate: Option<Handle<Expression>>,
2152    },
2153    /// Gather a value from another active thread in the subgroup
2154    SubgroupGather {
2155        /// Specifies which thread to gather from
2156        mode: GatherMode,
2157        /// The value to broadcast over
2158        argument: Handle<Expression>,
2159        /// The [`SubgroupOperationResult`] expression representing this load's result.
2160        ///
2161        /// [`SubgroupOperationResult`]: Expression::SubgroupOperationResult
2162        result: Handle<Expression>,
2163    },
2164    /// Compute a collective operation across all active threads in the subgroup
2165    SubgroupCollectiveOperation {
2166        /// What operation to compute
2167        op: SubgroupOperation,
2168        /// How to combine the results
2169        collective_op: CollectiveOperation,
2170        /// The value to compute over
2171        argument: Handle<Expression>,
2172        /// The [`SubgroupOperationResult`] expression representing this load's result.
2173        ///
2174        /// [`SubgroupOperationResult`]: Expression::SubgroupOperationResult
2175        result: Handle<Expression>,
2176    },
2177}
2178
2179/// A function argument.
2180#[derive(Clone, Debug)]
2181#[cfg_attr(feature = "serialize", derive(Serialize))]
2182#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2183#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2184pub struct FunctionArgument {
2185    /// Name of the argument, if any.
2186    pub name: Option<String>,
2187    /// Type of the argument.
2188    pub ty: Handle<Type>,
2189    /// For entry points, an argument has to have a binding
2190    /// unless it's a structure.
2191    pub binding: Option<Binding>,
2192}
2193
2194/// A function result.
2195#[derive(Clone, Debug)]
2196#[cfg_attr(feature = "serialize", derive(Serialize))]
2197#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2198#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2199pub struct FunctionResult {
2200    /// Type of the result.
2201    pub ty: Handle<Type>,
2202    /// For entry points, the result has to have a binding
2203    /// unless it's a structure.
2204    pub binding: Option<Binding>,
2205}
2206
2207/// A function defined in the module.
2208#[derive(Debug, Default, Clone)]
2209#[cfg_attr(feature = "serialize", derive(Serialize))]
2210#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2211#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2212pub struct Function {
2213    /// Name of the function, if any.
2214    pub name: Option<String>,
2215    /// Information about function argument.
2216    pub arguments: Vec<FunctionArgument>,
2217    /// The result of this function, if any.
2218    pub result: Option<FunctionResult>,
2219    /// Local variables defined and used in the function.
2220    pub local_variables: Arena<LocalVariable>,
2221    /// Expressions used inside this function.
2222    ///
2223    /// Unless explicitly stated otherwise, if an [`Expression`] is in this
2224    /// arena, then its subexpressions are in this arena too. In other words,
2225    /// every `Handle<Expression>` in this arena refers to an [`Expression`] in
2226    /// this arena too.
2227    ///
2228    /// The main ways this arena refers to [`Module::global_expressions`] are:
2229    ///
2230    /// - [`Constant`], [`Override`], and [`GlobalVariable`] expressions hold
2231    ///   handles for their respective types, whose initializer expressions are
2232    ///   in [`Module::global_expressions`].
2233    ///
2234    /// - Various expressions hold [`Type`] handles, and [`Type`]s may refer to
2235    ///   global expressions, for things like array lengths.
2236    ///
2237    /// An [`Expression`] must occur before all other [`Expression`]s that use
2238    /// its value.
2239    ///
2240    /// [`Constant`]: Expression::Constant
2241    /// [`Override`]: Expression::Override
2242    /// [`GlobalVariable`]: Expression::GlobalVariable
2243    pub expressions: Arena<Expression>,
2244    /// Map of expressions that have associated variable names
2245    pub named_expressions: NamedExpressions,
2246    /// Block of instructions comprising the body of the function.
2247    pub body: Block,
2248    /// The leaf of all diagnostic filter rules tree (stored in [`Module::diagnostic_filters`])
2249    /// parsed on this function.
2250    ///
2251    /// In WGSL, this corresponds to `@diagnostic(…)` attributes.
2252    ///
2253    /// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
2254    /// validation.
2255    pub diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
2256}
2257
2258/// The main function for a pipeline stage.
2259///
2260/// An [`EntryPoint`] is a [`Function`] that serves as the main function for a
2261/// graphics or compute pipeline stage. For example, an `EntryPoint` whose
2262/// [`stage`] is [`ShaderStage::Vertex`] can serve as a graphics pipeline's
2263/// vertex shader.
2264///
2265/// Since an entry point is called directly by the graphics or compute pipeline,
2266/// not by other WGSL functions, you must specify what the pipeline should pass
2267/// as the entry point's arguments, and what values it will return. For example,
2268/// a vertex shader needs a vertex's attributes as its arguments, but if it's
2269/// used for instanced draw calls, it will also want to know the instance id.
2270/// The vertex shader's return value will usually include an output vertex
2271/// position, and possibly other attributes to be interpolated and passed along
2272/// to a fragment shader.
2273///
2274/// To specify this, the arguments and result of an `EntryPoint`'s [`function`]
2275/// must each have a [`Binding`], or be structs whose members all have
2276/// `Binding`s. This associates every value passed to or returned from the entry
2277/// point with either a [`BuiltIn`] or a [`Location`]:
2278///
2279/// -   A [`BuiltIn`] has special semantics, usually specific to its pipeline
2280///     stage. For example, the result of a vertex shader can include a
2281///     [`BuiltIn::Position`] value, which determines the position of a vertex
2282///     of a rendered primitive. Or, a compute shader might take an argument
2283///     whose binding is [`BuiltIn::WorkGroupSize`], through which the compute
2284///     pipeline would pass the number of invocations in your workgroup.
2285///
2286/// -   A [`Location`] indicates user-defined IO to be passed from one pipeline
2287///     stage to the next. For example, a vertex shader might also produce a
2288///     `uv` texture location as a user-defined IO value.
2289///
2290/// In other words, the pipeline stage's input and output interface are
2291/// determined by the bindings of the arguments and result of the `EntryPoint`'s
2292/// [`function`].
2293///
2294/// [`Function`]: crate::Function
2295/// [`Location`]: Binding::Location
2296/// [`function`]: EntryPoint::function
2297/// [`stage`]: EntryPoint::stage
2298#[derive(Debug, Clone)]
2299#[cfg_attr(feature = "serialize", derive(Serialize))]
2300#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2301#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2302pub struct EntryPoint {
2303    /// Name of this entry point, visible externally.
2304    ///
2305    /// Entry point names for a given `stage` must be distinct within a module.
2306    pub name: String,
2307    /// Shader stage.
2308    pub stage: ShaderStage,
2309    /// Early depth test for fragment stages.
2310    pub early_depth_test: Option<EarlyDepthTest>,
2311    /// Workgroup size for compute stages
2312    pub workgroup_size: [u32; 3],
2313    /// Override expressions for workgroup size in the global_expressions arena
2314    pub workgroup_size_overrides: Option<[Option<Handle<Expression>>; 3]>,
2315    /// The entrance function.
2316    pub function: Function,
2317}
2318
2319/// Return types predeclared for the frexp, modf, and atomicCompareExchangeWeak built-in functions.
2320///
2321/// These cannot be spelled in WGSL source.
2322///
2323/// Stored in [`SpecialTypes::predeclared_types`] and created by [`Module::generate_predeclared_type`].
2324#[derive(Debug, PartialEq, Eq, Hash, Clone)]
2325#[cfg_attr(feature = "serialize", derive(Serialize))]
2326#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2327#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2328pub enum PredeclaredType {
2329    AtomicCompareExchangeWeakResult(Scalar),
2330    ModfResult {
2331        size: Option<VectorSize>,
2332        scalar: Scalar,
2333    },
2334    FrexpResult {
2335        size: Option<VectorSize>,
2336        scalar: Scalar,
2337    },
2338}
2339
2340/// Set of special types that can be optionally generated by the frontends.
2341#[derive(Debug, Default, Clone)]
2342#[cfg_attr(feature = "serialize", derive(Serialize))]
2343#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2344#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2345pub struct SpecialTypes {
2346    /// Type for `RayDesc`.
2347    ///
2348    /// Call [`Module::generate_ray_desc_type`] to populate this if
2349    /// needed and return the handle.
2350    pub ray_desc: Option<Handle<Type>>,
2351
2352    /// Type for `RayIntersection`.
2353    ///
2354    /// Call [`Module::generate_ray_intersection_type`] to populate
2355    /// this if needed and return the handle.
2356    pub ray_intersection: Option<Handle<Type>>,
2357
2358    /// Type for `RayVertexReturn`.
2359    ///
2360    /// Call [`Module::generate_vertex_return_type`]
2361    pub ray_vertex_return: Option<Handle<Type>>,
2362
2363    /// Struct containing parameters required by some backends to emit code for
2364    /// [`ImageClass::External`] textures.
2365    ///
2366    /// See `wgpu_core::device::resource::ExternalTextureParams` for the
2367    /// documentation of each field.
2368    ///
2369    /// In WGSL, this type would be:
2370    ///
2371    /// ```ignore
2372    /// struct NagaExternalTextureParams {         // align size offset
2373    ///     yuv_conversion_matrix: mat4x4<f32>,    //    16   64      0
2374    ///     gamut_conversion_matrix: mat3x3<f32>,  //    16   48     64
2375    ///     src_tf: NagaExternalTextureTransferFn, //     4   16    112
2376    ///     dst_tf: NagaExternalTextureTransferFn, //     4   16    128
2377    ///     sample_transform: mat3x2<f32>,         //     8   24    144
2378    ///     load_transform: mat3x2<f32>,           //     8   24    168
2379    ///     size: vec2<u32>,                       //     8    8    192
2380    ///     num_planes: u32,                       //     4    4    200
2381    /// }                            // whole struct:    16  208
2382    /// ```
2383    ///
2384    /// Call [`Module::generate_external_texture_types`] to populate this if
2385    /// needed.
2386    pub external_texture_params: Option<Handle<Type>>,
2387
2388    /// Struct describing a gamma encoding transfer function. Member of
2389    /// `NagaExternalTextureParams`, describing how the backend should perform
2390    /// color space conversion when sampling from [`ImageClass::External`]
2391    /// textures.
2392    ///
2393    /// In WGSL, this type would be:
2394    ///
2395    /// ```ignore
2396    /// struct NagaExternalTextureTransferFn { // align size offset
2397    ///     a: f32,                            //     4    4      0
2398    ///     b: f32,                            //     4    4      4
2399    ///     g: f32,                            //     4    4      8
2400    ///     k: f32,                            //     4    4     12
2401    /// }                         // whole struct:    4   16
2402    /// ```
2403    ///
2404    /// Call [`Module::generate_external_texture_types`] to populate this if
2405    /// needed.
2406    pub external_texture_transfer_function: Option<Handle<Type>>,
2407
2408    /// Types for predeclared wgsl types instantiated on demand.
2409    ///
2410    /// Call [`Module::generate_predeclared_type`] to populate this if
2411    /// needed and return the handle.
2412    pub predeclared_types: FastIndexMap<PredeclaredType, Handle<Type>>,
2413}
2414
2415bitflags::bitflags! {
2416    /// Ray flags used when casting rays.
2417    /// Matching vulkan constants can be found in
2418    /// https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/KHR/ray_common/ray_flags_section.txt
2419    #[cfg_attr(feature = "serialize", derive(Serialize))]
2420    #[cfg_attr(feature = "deserialize", derive(Deserialize))]
2421    #[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2422    #[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)]
2423    pub struct RayFlag: u32 {
2424        /// Force all intersections to be treated as opaque.
2425        const FORCE_OPAQUE = 0x1;
2426        /// Force all intersections to be treated as non-opaque.
2427        const FORCE_NO_OPAQUE = 0x2;
2428        /// Stop traversal after the first hit.
2429        const TERMINATE_ON_FIRST_HIT = 0x4;
2430        /// Don't execute the closest hit shader.
2431        const SKIP_CLOSEST_HIT_SHADER = 0x8;
2432        /// Cull back facing geometry.
2433        const CULL_BACK_FACING = 0x10;
2434        /// Cull front facing geometry.
2435        const CULL_FRONT_FACING = 0x20;
2436        /// Cull opaque geometry.
2437        const CULL_OPAQUE = 0x40;
2438        /// Cull non-opaque geometry.
2439        const CULL_NO_OPAQUE = 0x80;
2440        /// Skip triangular geometry.
2441        const SKIP_TRIANGLES = 0x100;
2442        /// Skip axis-aligned bounding boxes.
2443        const SKIP_AABBS = 0x200;
2444    }
2445}
2446
2447/// Type of a ray query intersection.
2448/// Matching vulkan constants can be found in
2449/// <https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/KHR/SPV_KHR_ray_query.asciidoc>
2450/// but the actual values are different for candidate intersections.
2451#[cfg_attr(feature = "serialize", derive(Serialize))]
2452#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2453#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2454#[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)]
2455pub enum RayQueryIntersection {
2456    /// No intersection found.
2457    /// Matches `RayQueryCommittedIntersectionNoneKHR`.
2458    #[default]
2459    None = 0,
2460    /// Intersecting with triangles.
2461    /// Matches `RayQueryCommittedIntersectionTriangleKHR` and `RayQueryCandidateIntersectionTriangleKHR`.
2462    Triangle = 1,
2463    /// Intersecting with generated primitives.
2464    /// Matches `RayQueryCommittedIntersectionGeneratedKHR`.
2465    Generated = 2,
2466    /// Intersecting with Axis Aligned Bounding Boxes.
2467    /// Matches `RayQueryCandidateIntersectionAABBKHR`.
2468    Aabb = 3,
2469}
2470
2471/// Doc comments preceding items.
2472///
2473/// These can be used to generate automated documentation,
2474/// IDE hover information or translate shaders with their context comments.
2475#[derive(Debug, Default, Clone)]
2476#[cfg_attr(feature = "serialize", derive(Serialize))]
2477#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2478#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2479pub struct DocComments {
2480    pub types: FastIndexMap<Handle<Type>, Vec<String>>,
2481    // The key is:
2482    // - key.0: the handle to the Struct
2483    // - key.1: the index of the `StructMember`.
2484    pub struct_members: FastIndexMap<(Handle<Type>, usize), Vec<String>>,
2485    pub entry_points: FastIndexMap<usize, Vec<String>>,
2486    pub functions: FastIndexMap<Handle<Function>, Vec<String>>,
2487    pub constants: FastIndexMap<Handle<Constant>, Vec<String>>,
2488    pub global_variables: FastIndexMap<Handle<GlobalVariable>, Vec<String>>,
2489    // Top level comments, appearing before any space.
2490    pub module: Vec<String>,
2491}
2492
2493/// Shader module.
2494///
2495/// A module is a set of constants, global variables and functions, as well as
2496/// the types required to define them.
2497///
2498/// Some functions are marked as entry points, to be used in a certain shader stage.
2499///
2500/// To create a new module, use the `Default` implementation.
2501/// Alternatively, you can load an existing shader using one of the [available front ends].
2502///
2503/// When finished, you can export modules using one of the [available backends].
2504///
2505/// ## Module arenas
2506///
2507/// Most module contents are stored in [`Arena`]s. In a valid module, arena
2508/// elements only refer to prior arena elements. That is, whenever an element in
2509/// some `Arena<T>` contains a `Handle<T>` referring to another element the same
2510/// arena, the handle's referent always precedes the element containing the
2511/// handle.
2512///
2513/// The elements of [`Module::types`] may refer to [`Expression`]s in
2514/// [`Module::global_expressions`], and those expressions may in turn refer back
2515/// to [`Type`]s in [`Module::types`]. In a valid module, there exists an order
2516/// in which all types and global expressions can be visited such that:
2517///
2518/// - types and expressions are visited in the order in which they appear in
2519///   their arenas, and
2520///
2521/// - every element refers only to previously visited elements.
2522///
2523/// This implies that the graph of types and global expressions is acyclic.
2524/// (However, it is a stronger condition: there are cycle-free arrangements of
2525/// types and expressions for which an order like the one described above does
2526/// not exist. Modules arranged in such a way are not valid.)
2527///
2528/// [available front ends]: crate::front
2529/// [available backends]: crate::back
2530#[derive(Debug, Default, Clone)]
2531#[cfg_attr(feature = "serialize", derive(Serialize))]
2532#[cfg_attr(feature = "deserialize", derive(Deserialize))]
2533#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
2534pub struct Module {
2535    /// Arena for the types defined in this module.
2536    ///
2537    /// See the [`Module`] docs for more details about this field.
2538    pub types: UniqueArena<Type>,
2539    /// Dictionary of special type handles.
2540    pub special_types: SpecialTypes,
2541    /// Arena for the constants defined in this module.
2542    pub constants: Arena<Constant>,
2543    /// Arena for the pipeline-overridable constants defined in this module.
2544    pub overrides: Arena<Override>,
2545    /// Arena for the global variables defined in this module.
2546    pub global_variables: Arena<GlobalVariable>,
2547    /// [Constant expressions] and [override expressions] used by this module.
2548    ///
2549    /// If an expression is in this arena, then its subexpressions are in this
2550    /// arena too. In other words, every `Handle<Expression>` in this arena
2551    /// refers to an [`Expression`] in this arena too.
2552    ///
2553    /// See the [`Module`] docs for more details about this field.
2554    ///
2555    /// [Constant expressions]: index.html#constant-expressions
2556    /// [override expressions]: index.html#override-expressions
2557    pub global_expressions: Arena<Expression>,
2558    /// Arena for the functions defined in this module.
2559    ///
2560    /// Each function must appear in this arena strictly before all its callers.
2561    /// Recursion is not supported.
2562    pub functions: Arena<Function>,
2563    /// Entry points.
2564    pub entry_points: Vec<EntryPoint>,
2565    /// Arena for all diagnostic filter rules parsed in this module, including those in functions
2566    /// and statements.
2567    ///
2568    /// This arena contains elements of a _tree_ of diagnostic filter rules. When nodes are built
2569    /// by a front-end, they refer to a parent scope
2570    pub diagnostic_filters: Arena<DiagnosticFilterNode>,
2571    /// The leaf of all diagnostic filter rules tree parsed from directives in this module.
2572    ///
2573    /// In WGSL, this corresponds to `diagnostic(…);` directives.
2574    ///
2575    /// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
2576    /// validation.
2577    pub diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
2578    /// Doc comments.
2579    pub doc_comments: Option<Box<DocComments>>,
2580}