naga/ir/
mod.rs

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