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