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