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