Skip to main content

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