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