naga/valid/
interface.rs

1use alloc::vec::Vec;
2
3use bit_set::BitSet;
4
5use super::{
6    analyzer::{FunctionInfo, GlobalUse},
7    Capabilities, Disalignment, FunctionError, ImmediateError, ModuleInfo,
8};
9use crate::arena::{Handle, UniqueArena};
10use crate::span::{AddSpan as _, MapErrWithSpan as _, SpanProvider as _, WithSpan};
11
12const MAX_WORKGROUP_SIZE: u32 = 0x4000;
13
14#[derive(Clone, Debug, thiserror::Error)]
15#[cfg_attr(test, derive(PartialEq))]
16pub enum GlobalVariableError {
17    #[error("Usage isn't compatible with address space {0:?}")]
18    InvalidUsage(crate::AddressSpace),
19    #[error("Type isn't compatible with address space {0:?}")]
20    InvalidType(crate::AddressSpace),
21    #[error("Type {0:?} isn't compatible with binding arrays")]
22    InvalidBindingArray(Handle<crate::Type>),
23    #[error("Type flags {seen:?} do not meet the required {required:?}")]
24    MissingTypeFlags {
25        required: super::TypeFlags,
26        seen: super::TypeFlags,
27    },
28    #[error("Capability {0:?} is not supported")]
29    UnsupportedCapability(Capabilities),
30    #[error("Binding decoration is missing or not applicable")]
31    InvalidBinding,
32    #[error("Alignment requirements for address space {0:?} are not met by {1:?}")]
33    Alignment(
34        crate::AddressSpace,
35        Handle<crate::Type>,
36        #[source] Disalignment,
37    ),
38    #[error("Initializer must be an override-expression")]
39    InitializerExprType,
40    #[error("Initializer doesn't match the variable type")]
41    InitializerType,
42    #[error("Initializer can't be used with address space {0:?}")]
43    InitializerNotAllowed(crate::AddressSpace),
44    #[error("Storage address space doesn't support write-only access")]
45    StorageAddressSpaceWriteOnlyNotSupported,
46    #[error("Type is not valid for use as a immediate data")]
47    InvalidImmediateType(#[source] ImmediateError),
48    #[error("Task payload must not be zero-sized")]
49    ZeroSizedTaskPayload,
50}
51
52#[derive(Clone, Debug, thiserror::Error)]
53#[cfg_attr(test, derive(PartialEq))]
54pub enum VaryingError {
55    #[error("The type {0:?} does not match the varying")]
56    InvalidType(Handle<crate::Type>),
57    #[error("The type {0:?} cannot be used for user-defined entry point inputs or outputs")]
58    NotIOShareableType(Handle<crate::Type>),
59    #[error("Interpolation is not valid")]
60    InvalidInterpolation,
61    #[error("Interpolation {0:?} is only valid for stage {1:?}")]
62    InvalidInterpolationInStage(crate::Interpolation, crate::ShaderStage),
63    #[error("Cannot combine {interpolation:?} interpolation with the {sampling:?} sample type")]
64    InvalidInterpolationSamplingCombination {
65        interpolation: crate::Interpolation,
66        sampling: crate::Sampling,
67    },
68    #[error("Interpolation must be specified on vertex shader outputs and fragment shader inputs")]
69    MissingInterpolation,
70    #[error("Built-in {0:?} is not available at this stage")]
71    InvalidBuiltInStage(crate::BuiltIn),
72    #[error("Built-in type for {0:?} is invalid. Found {1:?}")]
73    InvalidBuiltInType(crate::BuiltIn, crate::TypeInner),
74    #[error("Entry point arguments and return values must all have bindings")]
75    MissingBinding,
76    #[error("Struct member {0} is missing a binding")]
77    MemberMissingBinding(u32),
78    #[error("Multiple bindings at location {location} are present")]
79    BindingCollision { location: u32 },
80    #[error("Multiple bindings use the same `blend_src` {blend_src}")]
81    BindingCollisionBlendSrc { blend_src: u32 },
82    #[error("Built-in {0:?} is present more than once")]
83    DuplicateBuiltIn(crate::BuiltIn),
84    #[error("Capability {0:?} is not supported")]
85    UnsupportedCapability(Capabilities),
86    #[error("The attribute {0:?} is only valid as an output for stage {1:?}")]
87    InvalidInputAttributeInStage(&'static str, crate::ShaderStage),
88    #[error("The attribute {0:?} is not valid for stage {1:?}")]
89    InvalidAttributeInStage(&'static str, crate::ShaderStage),
90    #[error("The `blend_src` attribute can only be used on location 0, only indices 0 and 1 are valid. Location was {location}, index was {blend_src}.")]
91    InvalidBlendSrcIndex { location: u32, blend_src: u32 },
92    #[error("If `blend_src` is used, there must be exactly two outputs both with location 0, one with `blend_src(0)` and the other with `blend_src(1)`.")]
93    IncompleteBlendSrcUsage,
94    #[error("If `blend_src` is used, both outputs must have the same type. `blend_src(0)` has type {blend_src_0_type:?} and `blend_src(1)` has type {blend_src_1_type:?}.")]
95    BlendSrcOutputTypeMismatch {
96        blend_src_0_type: Handle<crate::Type>,
97        blend_src_1_type: Handle<crate::Type>,
98    },
99    #[error("Workgroup size is multi dimensional, `@builtin(subgroup_id)` and `@builtin(subgroup_invocation_id)` are not supported.")]
100    InvalidMultiDimensionalSubgroupBuiltIn,
101    #[error("The `@per_primitive` attribute can only be used in fragment shader inputs or mesh shader primitive outputs")]
102    InvalidPerPrimitive,
103    #[error("Non-builtin members of a mesh primitive output struct must be decorated with `@per_primitive`")]
104    MissingPerPrimitive,
105    #[error("Per vertex fragment inputs must be an array of length 3.")]
106    PerVertexNotArrayOfThree,
107}
108
109#[derive(Clone, Debug, thiserror::Error)]
110#[cfg_attr(test, derive(PartialEq))]
111pub enum EntryPointError {
112    #[error("Multiple conflicting entry points")]
113    Conflict,
114    #[error("Vertex shaders must return a `@builtin(position)` output value")]
115    MissingVertexOutputPosition,
116    #[error("Early depth test is not applicable")]
117    UnexpectedEarlyDepthTest,
118    #[error("Workgroup size is not applicable")]
119    UnexpectedWorkgroupSize,
120    #[error("Workgroup size is out of range")]
121    OutOfRangeWorkgroupSize,
122    #[error("Uses operations forbidden at this stage")]
123    ForbiddenStageOperations,
124    #[error("Global variable {0:?} is used incorrectly as {1:?}")]
125    InvalidGlobalUsage(Handle<crate::GlobalVariable>, GlobalUse),
126    #[error("More than 1 immediate data variable is used")]
127    MoreThanOneImmediateUsed,
128    #[error("Bindings for {0:?} conflict with other resource")]
129    BindingCollision(Handle<crate::GlobalVariable>),
130    #[error("Argument {0} varying error")]
131    Argument(u32, #[source] VaryingError),
132    #[error(transparent)]
133    Result(#[from] VaryingError),
134    #[error("Location {location} interpolation of an integer has to be flat")]
135    InvalidIntegerInterpolation { location: u32 },
136    #[error(transparent)]
137    Function(#[from] FunctionError),
138    #[error("Capability {0:?} is not supported")]
139    UnsupportedCapability(Capabilities),
140
141    #[error("mesh shader entry point missing mesh shader attributes")]
142    ExpectedMeshShaderAttributes,
143    #[error("Non mesh shader entry point cannot have mesh shader attributes")]
144    UnexpectedMeshShaderAttributes,
145    #[error("Non mesh/task shader entry point cannot have task payload attribute")]
146    UnexpectedTaskPayload,
147    #[error("Task payload must be declared with `var<task_payload>`")]
148    TaskPayloadWrongAddressSpace,
149    #[error("For a task payload to be used, it must be declared with @payload")]
150    WrongTaskPayloadUsed,
151    #[error("Task shader entry point must return @builtin(mesh_task_size) vec3<u32>")]
152    WrongTaskShaderEntryResult,
153    #[error("Task shaders must declare a task payload output")]
154    ExpectedTaskPayload,
155    #[error(
156        "Mesh shader output variable must be a struct with fields that are all allowed builtins"
157    )]
158    BadMeshOutputVariableType,
159    #[error("Mesh shader output variable fields must have types that are in accordance with the mesh shader spec")]
160    BadMeshOutputVariableField,
161    #[error("Mesh shader entry point cannot have a return type")]
162    UnexpectedMeshShaderEntryResult,
163    #[error(
164        "Mesh output type must be a user-defined struct with fields in alignment with the mesh shader spec"
165    )]
166    InvalidMeshOutputType,
167    #[error("Mesh primitive outputs must have exactly one of `@builtin(triangle_indices)`, `@builtin(line_indices)`, or `@builtin(point_index)`")]
168    InvalidMeshPrimitiveOutputType,
169    #[error("Mesh output global variable must live in the workgroup address space")]
170    WrongMeshOutputAddressSpace,
171    #[error("Task payload must be at least 4 bytes, but is {0} bytes")]
172    TaskPayloadTooSmall(u32),
173    #[error("Only the `ray_generation`, `closest_hit`, and `any_hit` shader stages can access a global variable in the `ray_payload` address space")]
174    RayPayloadInInvalidStage(crate::ShaderStage),
175    #[error("Only the `closest_hit`, `any_hit`, and `miss` shader stages can access a global variable in the `incoming_ray_payload` address space")]
176    IncomingRayPayloadInInvalidStage(crate::ShaderStage),
177}
178
179fn storage_usage(access: crate::StorageAccess) -> GlobalUse {
180    let mut storage_usage = GlobalUse::QUERY;
181    if access.contains(crate::StorageAccess::LOAD) {
182        storage_usage |= GlobalUse::READ;
183    }
184    if access.contains(crate::StorageAccess::STORE) {
185        storage_usage |= GlobalUse::WRITE;
186    }
187    if access.contains(crate::StorageAccess::ATOMIC) {
188        storage_usage |= GlobalUse::ATOMIC;
189    }
190    storage_usage
191}
192
193#[derive(Clone, Copy, Debug, PartialEq, Eq)]
194enum MeshOutputType {
195    None,
196    VertexOutput,
197    PrimitiveOutput,
198}
199
200struct VaryingContext<'a> {
201    stage: crate::ShaderStage,
202    output: bool,
203    types: &'a UniqueArena<crate::Type>,
204    type_info: &'a Vec<super::r#type::TypeInfo>,
205    location_mask: &'a mut BitSet,
206    blend_src_mask: &'a mut BitSet,
207    built_ins: &'a mut crate::FastHashSet<crate::BuiltIn>,
208    capabilities: Capabilities,
209    flags: super::ValidationFlags,
210    mesh_output_type: MeshOutputType,
211    has_task_payload: bool,
212}
213
214impl VaryingContext<'_> {
215    fn validate_impl(
216        &mut self,
217        ep: &crate::EntryPoint,
218        ty: Handle<crate::Type>,
219        binding: &crate::Binding,
220    ) -> Result<(), VaryingError> {
221        use crate::{BuiltIn as Bi, ShaderStage as St, TypeInner as Ti, VectorSize as Vs};
222
223        let ty_inner = &self.types[ty].inner;
224        match *binding {
225            crate::Binding::BuiltIn(built_in) => {
226                // Ignore the `invariant` field for the sake of duplicate checks,
227                // but use the original in error messages.
228                let canonical = match built_in {
229                    crate::BuiltIn::Position { .. } => {
230                        crate::BuiltIn::Position { invariant: false }
231                    }
232                    crate::BuiltIn::Barycentric { .. } => {
233                        crate::BuiltIn::Barycentric { perspective: false }
234                    }
235                    x => x,
236                };
237
238                if self.built_ins.contains(&canonical) {
239                    return Err(VaryingError::DuplicateBuiltIn(built_in));
240                }
241                self.built_ins.insert(canonical);
242
243                let required = match built_in {
244                    Bi::ClipDistance => Capabilities::CLIP_DISTANCE,
245                    Bi::CullDistance => Capabilities::CULL_DISTANCE,
246                    // Primitive index is allowed w/o any other extensions in any- and closest-hit shaders
247                    Bi::PrimitiveIndex if !matches!(ep.stage, St::AnyHit | St::ClosestHit) => {
248                        Capabilities::PRIMITIVE_INDEX
249                    }
250                    Bi::Barycentric { .. } => Capabilities::SHADER_BARYCENTRICS,
251                    Bi::ViewIndex => Capabilities::MULTIVIEW,
252                    Bi::SampleIndex => Capabilities::MULTISAMPLED_SHADING,
253                    Bi::NumSubgroups
254                    | Bi::SubgroupId
255                    | Bi::SubgroupSize
256                    | Bi::SubgroupInvocationId => Capabilities::SUBGROUP,
257                    Bi::DrawIndex => Capabilities::DRAW_INDEX,
258                    _ => Capabilities::empty(),
259                };
260                if !self.capabilities.contains(required) {
261                    return Err(VaryingError::UnsupportedCapability(required));
262                }
263
264                if matches!(
265                    built_in,
266                    crate::BuiltIn::SubgroupId | crate::BuiltIn::SubgroupInvocationId
267                ) && ep.workgroup_size[1..].iter().any(|&s| s > 1)
268                {
269                    return Err(VaryingError::InvalidMultiDimensionalSubgroupBuiltIn);
270                }
271
272                let (visible, type_good) = match built_in {
273                    Bi::BaseInstance | Bi::BaseVertex | Bi::VertexIndex => (
274                        self.stage == St::Vertex && !self.output,
275                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
276                    ),
277                    Bi::InstanceIndex => (
278                        matches!(self.stage, St::Vertex | St::AnyHit | St::ClosestHit)
279                            && !self.output,
280                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
281                    ),
282                    Bi::DrawIndex => (
283                        // Always allowed in task/vertex stage. Allowed in mesh stage if there is no task stage in the pipeline.
284                        (self.stage == St::Vertex
285                            || self.stage == St::Task
286                            || (self.stage == St::Mesh && !self.has_task_payload))
287                            && !self.output,
288                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
289                    ),
290                    Bi::ClipDistance | Bi::CullDistance => (
291                        (self.stage == St::Vertex || self.stage == St::Mesh) && self.output,
292                        match *ty_inner {
293                            Ti::Array { base, size, .. } => {
294                                self.types[base].inner == Ti::Scalar(crate::Scalar::F32)
295                                    && match size {
296                                        crate::ArraySize::Constant(non_zero) => non_zero.get() <= 8,
297                                        _ => false,
298                                    }
299                            }
300                            _ => false,
301                        },
302                    ),
303                    Bi::PointSize => (
304                        (self.stage == St::Vertex || self.stage == St::Mesh) && self.output,
305                        *ty_inner == Ti::Scalar(crate::Scalar::F32),
306                    ),
307                    Bi::PointCoord => (
308                        self.stage == St::Fragment && !self.output,
309                        *ty_inner
310                            == Ti::Vector {
311                                size: Vs::Bi,
312                                scalar: crate::Scalar::F32,
313                            },
314                    ),
315                    Bi::Position { .. } => (
316                        match self.stage {
317                            St::Vertex | St::Mesh => self.output,
318                            St::Fragment => !self.output,
319                            St::Compute | St::Task => false,
320                            St::RayGeneration | St::AnyHit | St::ClosestHit | St::Miss => false,
321                        },
322                        *ty_inner
323                            == Ti::Vector {
324                                size: Vs::Quad,
325                                scalar: crate::Scalar::F32,
326                            },
327                    ),
328                    Bi::ViewIndex => (
329                        match self.stage {
330                            St::Vertex | St::Fragment | St::Task | St::Mesh => !self.output,
331                            St::Compute
332                            | St::RayGeneration
333                            | St::AnyHit
334                            | St::ClosestHit
335                            | St::Miss => false,
336                        },
337                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
338                    ),
339                    Bi::FragDepth => (
340                        self.stage == St::Fragment && self.output,
341                        *ty_inner == Ti::Scalar(crate::Scalar::F32),
342                    ),
343                    Bi::FrontFacing => (
344                        self.stage == St::Fragment && !self.output,
345                        *ty_inner == Ti::Scalar(crate::Scalar::BOOL),
346                    ),
347                    Bi::PrimitiveIndex => (
348                        (matches!(self.stage, St::Fragment | St::AnyHit | St::ClosestHit)
349                            && !self.output)
350                            || (self.stage == St::Mesh
351                                && self.output
352                                && self.mesh_output_type == MeshOutputType::PrimitiveOutput),
353                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
354                    ),
355                    Bi::Barycentric { .. } => (
356                        self.stage == St::Fragment && !self.output,
357                        *ty_inner
358                            == Ti::Vector {
359                                size: Vs::Tri,
360                                scalar: crate::Scalar::F32,
361                            },
362                    ),
363                    Bi::SampleIndex => (
364                        self.stage == St::Fragment && !self.output,
365                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
366                    ),
367                    Bi::SampleMask => (
368                        self.stage == St::Fragment,
369                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
370                    ),
371                    Bi::LocalInvocationIndex => (
372                        self.stage.compute_like() && !self.output,
373                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
374                    ),
375                    Bi::GlobalInvocationId
376                    | Bi::LocalInvocationId
377                    | Bi::WorkGroupId
378                    | Bi::WorkGroupSize
379                    | Bi::NumWorkGroups => (
380                        self.stage.compute_like() && !self.output,
381                        *ty_inner
382                            == Ti::Vector {
383                                size: Vs::Tri,
384                                scalar: crate::Scalar::U32,
385                            },
386                    ),
387                    Bi::NumSubgroups | Bi::SubgroupId => (
388                        self.stage.compute_like() && !self.output,
389                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
390                    ),
391                    Bi::SubgroupSize | Bi::SubgroupInvocationId => (
392                        match self.stage {
393                            St::Compute
394                            | St::Fragment
395                            | St::Task
396                            | St::Mesh
397                            | St::RayGeneration
398                            | St::AnyHit
399                            | St::ClosestHit
400                            | St::Miss => !self.output,
401                            St::Vertex => false,
402                        },
403                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
404                    ),
405                    Bi::CullPrimitive => (
406                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
407                        *ty_inner == Ti::Scalar(crate::Scalar::BOOL),
408                    ),
409                    Bi::PointIndex => (
410                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
411                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
412                    ),
413                    Bi::LineIndices => (
414                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
415                        *ty_inner
416                            == Ti::Vector {
417                                size: Vs::Bi,
418                                scalar: crate::Scalar::U32,
419                            },
420                    ),
421                    Bi::TriangleIndices => (
422                        self.mesh_output_type == MeshOutputType::PrimitiveOutput,
423                        *ty_inner
424                            == Ti::Vector {
425                                size: Vs::Tri,
426                                scalar: crate::Scalar::U32,
427                            },
428                    ),
429                    Bi::MeshTaskSize => (
430                        self.stage == St::Task && self.output,
431                        *ty_inner
432                            == Ti::Vector {
433                                size: Vs::Tri,
434                                scalar: crate::Scalar::U32,
435                            },
436                    ),
437                    Bi::RayInvocationId => (
438                        match self.stage {
439                            St::Vertex | St::Fragment | St::Compute | St::Mesh | St::Task => false,
440                            St::RayGeneration | St::AnyHit | St::ClosestHit | St::Miss => true,
441                        },
442                        *ty_inner
443                            == Ti::Vector {
444                                size: Vs::Tri,
445                                scalar: crate::Scalar::U32,
446                            },
447                    ),
448                    Bi::NumRayInvocations => (
449                        match self.stage {
450                            St::Vertex | St::Fragment | St::Compute | St::Mesh | St::Task => false,
451                            St::RayGeneration | St::AnyHit | St::ClosestHit | St::Miss => true,
452                        },
453                        *ty_inner
454                            == Ti::Vector {
455                                size: Vs::Tri,
456                                scalar: crate::Scalar::U32,
457                            },
458                    ),
459                    Bi::InstanceCustomData => (
460                        match self.stage {
461                            St::RayGeneration
462                            | St::Miss
463                            | St::Vertex
464                            | St::Fragment
465                            | St::Compute
466                            | St::Mesh
467                            | St::Task => false,
468                            St::AnyHit | St::ClosestHit => true,
469                        },
470                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
471                    ),
472                    Bi::GeometryIndex => (
473                        match self.stage {
474                            St::RayGeneration
475                            | St::Miss
476                            | St::Vertex
477                            | St::Fragment
478                            | St::Compute
479                            | St::Mesh
480                            | St::Task => false,
481                            St::AnyHit | St::ClosestHit => true,
482                        },
483                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
484                    ),
485                    Bi::WorldRayOrigin => (
486                        match self.stage {
487                            St::RayGeneration
488                            | St::Vertex
489                            | St::Fragment
490                            | St::Compute
491                            | St::Mesh
492                            | St::Task => false,
493                            St::AnyHit | St::ClosestHit | St::Miss => true,
494                        },
495                        *ty_inner
496                            == Ti::Vector {
497                                size: Vs::Tri,
498                                scalar: crate::Scalar::F32,
499                            },
500                    ),
501                    Bi::WorldRayDirection => (
502                        match self.stage {
503                            St::RayGeneration
504                            | St::Vertex
505                            | St::Fragment
506                            | St::Compute
507                            | St::Mesh
508                            | St::Task => false,
509                            St::AnyHit | St::ClosestHit | St::Miss => true,
510                        },
511                        *ty_inner
512                            == Ti::Vector {
513                                size: Vs::Tri,
514                                scalar: crate::Scalar::F32,
515                            },
516                    ),
517                    Bi::ObjectRayOrigin => (
518                        match self.stage {
519                            St::RayGeneration
520                            | St::Miss
521                            | St::Vertex
522                            | St::Fragment
523                            | St::Compute
524                            | St::Mesh
525                            | St::Task => false,
526                            St::AnyHit | St::ClosestHit => true,
527                        },
528                        *ty_inner
529                            == Ti::Vector {
530                                size: Vs::Tri,
531                                scalar: crate::Scalar::F32,
532                            },
533                    ),
534                    Bi::ObjectRayDirection => (
535                        match self.stage {
536                            St::RayGeneration
537                            | St::Miss
538                            | St::Vertex
539                            | St::Fragment
540                            | St::Compute
541                            | St::Mesh
542                            | St::Task => false,
543                            St::AnyHit | St::ClosestHit => true,
544                        },
545                        *ty_inner
546                            == Ti::Vector {
547                                size: Vs::Tri,
548                                scalar: crate::Scalar::F32,
549                            },
550                    ),
551                    Bi::RayTmin => (
552                        match self.stage {
553                            St::RayGeneration
554                            | St::Vertex
555                            | St::Fragment
556                            | St::Compute
557                            | St::Mesh
558                            | St::Task => false,
559                            St::AnyHit | St::ClosestHit | St::Miss => true,
560                        },
561                        *ty_inner == Ti::Scalar(crate::Scalar::F32),
562                    ),
563                    Bi::RayTCurrentMax => (
564                        match self.stage {
565                            St::RayGeneration
566                            | St::Vertex
567                            | St::Fragment
568                            | St::Compute
569                            | St::Mesh
570                            | St::Task => false,
571                            St::AnyHit | St::ClosestHit | St::Miss => true,
572                        },
573                        *ty_inner == Ti::Scalar(crate::Scalar::F32),
574                    ),
575                    Bi::ObjectToWorld => (
576                        match self.stage {
577                            St::RayGeneration
578                            | St::Miss
579                            | St::Vertex
580                            | St::Fragment
581                            | St::Compute
582                            | St::Mesh
583                            | St::Task => false,
584                            St::AnyHit | St::ClosestHit => true,
585                        },
586                        *ty_inner
587                            == Ti::Matrix {
588                                columns: crate::VectorSize::Quad,
589                                rows: crate::VectorSize::Tri,
590                                scalar: crate::Scalar::F32,
591                            },
592                    ),
593                    Bi::WorldToObject => (
594                        match self.stage {
595                            St::RayGeneration
596                            | St::Miss
597                            | St::Vertex
598                            | St::Fragment
599                            | St::Compute
600                            | St::Mesh
601                            | St::Task => false,
602                            St::AnyHit | St::ClosestHit => true,
603                        },
604                        *ty_inner
605                            == Ti::Matrix {
606                                columns: crate::VectorSize::Quad,
607                                rows: crate::VectorSize::Tri,
608                                scalar: crate::Scalar::F32,
609                            },
610                    ),
611                    Bi::HitKind => (
612                        match self.stage {
613                            St::RayGeneration
614                            | St::Miss
615                            | St::Vertex
616                            | St::Fragment
617                            | St::Compute
618                            | St::Mesh
619                            | St::Task => false,
620                            St::AnyHit | St::ClosestHit => true,
621                        },
622                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
623                    ),
624                    // Validated elsewhere, shouldn't be here
625                    Bi::VertexCount | Bi::PrimitiveCount | Bi::Vertices | Bi::Primitives => {
626                        (false, true)
627                    }
628                };
629                match built_in {
630                    Bi::CullPrimitive
631                    | Bi::PointIndex
632                    | Bi::LineIndices
633                    | Bi::TriangleIndices
634                    | Bi::MeshTaskSize
635                    | Bi::VertexCount
636                    | Bi::PrimitiveCount
637                    | Bi::Vertices
638                    | Bi::Primitives => {
639                        if !self.capabilities.contains(Capabilities::MESH_SHADER) {
640                            return Err(VaryingError::UnsupportedCapability(
641                                Capabilities::MESH_SHADER,
642                            ));
643                        }
644                    }
645                    _ => (),
646                }
647
648                if !visible {
649                    return Err(VaryingError::InvalidBuiltInStage(built_in));
650                }
651                if !type_good {
652                    return Err(VaryingError::InvalidBuiltInType(built_in, ty_inner.clone()));
653                }
654            }
655            crate::Binding::Location {
656                location,
657                interpolation,
658                sampling,
659                blend_src,
660                per_primitive,
661            } => {
662                if per_primitive && !self.capabilities.contains(Capabilities::MESH_SHADER) {
663                    return Err(VaryingError::UnsupportedCapability(
664                        Capabilities::MESH_SHADER,
665                    ));
666                }
667                if interpolation == Some(crate::Interpolation::PerVertex) {
668                    if self.stage != crate::ShaderStage::Fragment {
669                        return Err(VaryingError::InvalidInterpolationInStage(
670                            crate::Interpolation::PerVertex,
671                            crate::ShaderStage::Fragment,
672                        ));
673                    }
674                    if !self.capabilities.contains(Capabilities::PER_VERTEX) {
675                        return Err(VaryingError::UnsupportedCapability(
676                            Capabilities::PER_VERTEX,
677                        ));
678                    }
679                }
680                // If this is per-vertex, we change the type we validate to the inner type, otherwise we leave it be.
681                // This lets all validation be done on the inner type once we've ensured the per-vertex is array<T, 3>
682                let (ty, ty_inner) = if interpolation == Some(crate::Interpolation::PerVertex) {
683                    let three = crate::ArraySize::Constant(core::num::NonZeroU32::new(3).unwrap());
684                    match ty_inner {
685                        &Ti::Array { base, size, .. } if size == three => {
686                            (base, &self.types[base].inner)
687                        }
688                        _ => return Err(VaryingError::PerVertexNotArrayOfThree),
689                    }
690                } else {
691                    (ty, ty_inner)
692                };
693
694                // Only IO-shareable types may be stored in locations.
695                if !self.type_info[ty.index()]
696                    .flags
697                    .contains(super::TypeFlags::IO_SHAREABLE)
698                {
699                    return Err(VaryingError::NotIOShareableType(ty));
700                }
701
702                // Check whether `per_primitive` is appropriate for this stage and direction.
703                if self.mesh_output_type == MeshOutputType::PrimitiveOutput {
704                    // All mesh shader `Location` outputs must be `per_primitive`.
705                    if !per_primitive {
706                        return Err(VaryingError::MissingPerPrimitive);
707                    }
708                } else if self.stage == crate::ShaderStage::Fragment && !self.output {
709                    // Fragment stage inputs may be `per_primitive`. We'll only
710                    // know if these are correct when the whole mesh pipeline is
711                    // created and we're paired with a specific mesh or vertex
712                    // shader.
713                } else if per_primitive {
714                    // All other `Location` bindings must not be `per_primitive`.
715                    return Err(VaryingError::InvalidPerPrimitive);
716                }
717
718                if let Some(blend_src) = blend_src {
719                    // `blend_src` is only valid if dual source blending was explicitly enabled,
720                    // see https://www.w3.org/TR/WGSL/#extension-dual_source_blending
721                    if !self
722                        .capabilities
723                        .contains(Capabilities::DUAL_SOURCE_BLENDING)
724                    {
725                        return Err(VaryingError::UnsupportedCapability(
726                            Capabilities::DUAL_SOURCE_BLENDING,
727                        ));
728                    }
729                    if self.stage != crate::ShaderStage::Fragment {
730                        return Err(VaryingError::InvalidAttributeInStage(
731                            "blend_src",
732                            self.stage,
733                        ));
734                    }
735                    if !self.output {
736                        return Err(VaryingError::InvalidInputAttributeInStage(
737                            "blend_src",
738                            self.stage,
739                        ));
740                    }
741                    if (blend_src != 0 && blend_src != 1) || location != 0 {
742                        return Err(VaryingError::InvalidBlendSrcIndex {
743                            location,
744                            blend_src,
745                        });
746                    }
747                    if !self.blend_src_mask.insert(blend_src as usize) {
748                        return Err(VaryingError::BindingCollisionBlendSrc { blend_src });
749                    }
750                } else if !self.location_mask.insert(location as usize)
751                    && self.flags.contains(super::ValidationFlags::BINDINGS)
752                {
753                    return Err(VaryingError::BindingCollision { location });
754                }
755
756                if let Some(interpolation) = interpolation {
757                    let invalid_sampling = match (interpolation, sampling) {
758                        (_, None)
759                        | (
760                            crate::Interpolation::Perspective | crate::Interpolation::Linear,
761                            Some(
762                                crate::Sampling::Center
763                                | crate::Sampling::Centroid
764                                | crate::Sampling::Sample,
765                            ),
766                        )
767                        | (
768                            crate::Interpolation::Flat,
769                            Some(crate::Sampling::First | crate::Sampling::Either),
770                        ) => None,
771                        (_, Some(invalid_sampling)) => Some(invalid_sampling),
772                    };
773                    if let Some(sampling) = invalid_sampling {
774                        return Err(VaryingError::InvalidInterpolationSamplingCombination {
775                            interpolation,
776                            sampling,
777                        });
778                    }
779                }
780
781                let needs_interpolation = match self.stage {
782                    crate::ShaderStage::Vertex => self.output,
783                    crate::ShaderStage::Fragment => !self.output && !per_primitive,
784                    crate::ShaderStage::Compute
785                    | crate::ShaderStage::Task
786                    | crate::ShaderStage::RayGeneration
787                    | crate::ShaderStage::AnyHit
788                    | crate::ShaderStage::ClosestHit
789                    | crate::ShaderStage::Miss => false,
790                    crate::ShaderStage::Mesh => self.output,
791                };
792
793                // It doesn't make sense to specify a sampling when `interpolation` is `Flat`, but
794                // SPIR-V and GLSL both explicitly tolerate such combinations of decorators /
795                // qualifiers, so we won't complain about that here.
796                let _ = sampling;
797
798                let required = match sampling {
799                    Some(crate::Sampling::Sample) => Capabilities::MULTISAMPLED_SHADING,
800                    _ => Capabilities::empty(),
801                };
802                if !self.capabilities.contains(required) {
803                    return Err(VaryingError::UnsupportedCapability(required));
804                }
805
806                if interpolation != Some(crate::Interpolation::PerVertex) {
807                    match ty_inner.scalar_kind() {
808                        Some(crate::ScalarKind::Float) => {
809                            if needs_interpolation && interpolation.is_none() {
810                                return Err(VaryingError::MissingInterpolation);
811                            }
812                        }
813                        Some(_) => {
814                            if needs_interpolation
815                                && interpolation != Some(crate::Interpolation::Flat)
816                            {
817                                return Err(VaryingError::InvalidInterpolation);
818                            }
819                        }
820                        None => return Err(VaryingError::InvalidType(ty)),
821                    }
822                }
823            }
824        }
825
826        Ok(())
827    }
828
829    fn validate(
830        &mut self,
831        ep: &crate::EntryPoint,
832        ty: Handle<crate::Type>,
833        binding: Option<&crate::Binding>,
834    ) -> Result<(), WithSpan<VaryingError>> {
835        let span_context = self.types.get_span_context(ty);
836        match binding {
837            Some(binding) => self
838                .validate_impl(ep, ty, binding)
839                .map_err(|e| e.with_span_context(span_context)),
840            None => {
841                let crate::TypeInner::Struct { ref members, .. } = self.types[ty].inner else {
842                    if self.flags.contains(super::ValidationFlags::BINDINGS) {
843                        return Err(VaryingError::MissingBinding.with_span());
844                    } else {
845                        return Ok(());
846                    }
847                };
848
849                for (index, member) in members.iter().enumerate() {
850                    let span_context = self.types.get_span_context(ty);
851                    match member.binding {
852                        None => {
853                            if self.flags.contains(super::ValidationFlags::BINDINGS) {
854                                return Err(VaryingError::MemberMissingBinding(index as u32)
855                                    .with_span_context(span_context));
856                            }
857                        }
858                        Some(ref binding) => self
859                            .validate_impl(ep, member.ty, binding)
860                            .map_err(|e| e.with_span_context(span_context))?,
861                    }
862                }
863
864                if !self.blend_src_mask.is_empty() {
865                    let span_context = self.types.get_span_context(ty);
866
867                    // If there's any blend_src usage, it must apply to all members of which there must be exactly 2.
868                    if members.len() != 2 || self.blend_src_mask.len() != 2 {
869                        return Err(
870                            VaryingError::IncompleteBlendSrcUsage.with_span_context(span_context)
871                        );
872                    }
873                    // Also, all members must have the same type.
874                    if members[0].ty != members[1].ty {
875                        return Err(VaryingError::BlendSrcOutputTypeMismatch {
876                            blend_src_0_type: members[0].ty,
877                            blend_src_1_type: members[1].ty,
878                        }
879                        .with_span_context(span_context));
880                    }
881                }
882                Ok(())
883            }
884        }
885    }
886}
887
888impl super::Validator {
889    pub(super) fn validate_global_var(
890        &self,
891        var: &crate::GlobalVariable,
892        gctx: crate::proc::GlobalCtx,
893        mod_info: &ModuleInfo,
894        global_expr_kind: &crate::proc::ExpressionKindTracker,
895    ) -> Result<(), GlobalVariableError> {
896        use super::TypeFlags;
897
898        log::debug!("var {var:?}");
899        let inner_ty = match gctx.types[var.ty].inner {
900            // A binding array is (mostly) supposed to behave the same as a
901            // series of individually bound resources, so we can (mostly)
902            // validate a `binding_array<T>` as if it were just a plain `T`.
903            crate::TypeInner::BindingArray { base, .. } => match var.space {
904                crate::AddressSpace::Storage { .. } => {
905                    if !self
906                        .capabilities
907                        .contains(Capabilities::STORAGE_BUFFER_BINDING_ARRAY)
908                    {
909                        return Err(GlobalVariableError::UnsupportedCapability(
910                            Capabilities::STORAGE_BUFFER_BINDING_ARRAY,
911                        ));
912                    }
913                    base
914                }
915                crate::AddressSpace::Uniform => {
916                    if !self
917                        .capabilities
918                        .contains(Capabilities::BUFFER_BINDING_ARRAY)
919                    {
920                        return Err(GlobalVariableError::UnsupportedCapability(
921                            Capabilities::BUFFER_BINDING_ARRAY,
922                        ));
923                    }
924                    base
925                }
926                crate::AddressSpace::Handle => {
927                    match gctx.types[base].inner {
928                        crate::TypeInner::Image { class, .. } => match class {
929                            crate::ImageClass::Storage { .. } => {
930                                if !self
931                                    .capabilities
932                                    .contains(Capabilities::STORAGE_TEXTURE_BINDING_ARRAY)
933                                {
934                                    return Err(GlobalVariableError::UnsupportedCapability(
935                                        Capabilities::STORAGE_TEXTURE_BINDING_ARRAY,
936                                    ));
937                                }
938                            }
939                            crate::ImageClass::Sampled { .. } | crate::ImageClass::Depth { .. } => {
940                                if !self
941                                    .capabilities
942                                    .contains(Capabilities::TEXTURE_AND_SAMPLER_BINDING_ARRAY)
943                                {
944                                    return Err(GlobalVariableError::UnsupportedCapability(
945                                        Capabilities::TEXTURE_AND_SAMPLER_BINDING_ARRAY,
946                                    ));
947                                }
948                            }
949                            crate::ImageClass::External => {
950                                // This should have been rejected in `validate_type`.
951                                unreachable!("binding arrays of external images are not supported");
952                            }
953                        },
954                        crate::TypeInner::Sampler { .. } => {
955                            if !self
956                                .capabilities
957                                .contains(Capabilities::TEXTURE_AND_SAMPLER_BINDING_ARRAY)
958                            {
959                                return Err(GlobalVariableError::UnsupportedCapability(
960                                    Capabilities::TEXTURE_AND_SAMPLER_BINDING_ARRAY,
961                                ));
962                            }
963                        }
964                        crate::TypeInner::AccelerationStructure { .. } => {
965                            return Err(GlobalVariableError::InvalidBindingArray(base));
966                        }
967                        crate::TypeInner::RayQuery { .. } => {
968                            // This should have been rejected in `validate_type`.
969                            unreachable!("binding arrays of ray queries are not supported");
970                        }
971                        _ => {
972                            // Fall through to the regular validation, which will reject `base`
973                            // as invalid in `AddressSpace::Handle`.
974                        }
975                    }
976                    base
977                }
978                _ => return Err(GlobalVariableError::InvalidUsage(var.space)),
979            },
980            _ => var.ty,
981        };
982        let type_info = &self.types[inner_ty.index()];
983
984        let (required_type_flags, is_resource) = match var.space {
985            crate::AddressSpace::Function => {
986                return Err(GlobalVariableError::InvalidUsage(var.space))
987            }
988            crate::AddressSpace::Storage { access } => {
989                if let Err((ty_handle, disalignment)) = type_info.storage_layout {
990                    if self.flags.contains(super::ValidationFlags::STRUCT_LAYOUTS) {
991                        return Err(GlobalVariableError::Alignment(
992                            var.space,
993                            ty_handle,
994                            disalignment,
995                        ));
996                    }
997                }
998                if access == crate::StorageAccess::STORE {
999                    return Err(GlobalVariableError::StorageAddressSpaceWriteOnlyNotSupported);
1000                }
1001                (
1002                    TypeFlags::DATA | TypeFlags::HOST_SHAREABLE | TypeFlags::CREATION_RESOLVED,
1003                    true,
1004                )
1005            }
1006            crate::AddressSpace::Uniform => {
1007                if let Err((ty_handle, disalignment)) = type_info.uniform_layout {
1008                    if self.flags.contains(super::ValidationFlags::STRUCT_LAYOUTS) {
1009                        return Err(GlobalVariableError::Alignment(
1010                            var.space,
1011                            ty_handle,
1012                            disalignment,
1013                        ));
1014                    }
1015                }
1016                (
1017                    TypeFlags::DATA
1018                        | TypeFlags::COPY
1019                        | TypeFlags::SIZED
1020                        | TypeFlags::HOST_SHAREABLE
1021                        | TypeFlags::CREATION_RESOLVED,
1022                    true,
1023                )
1024            }
1025            crate::AddressSpace::Handle => {
1026                match gctx.types[inner_ty].inner {
1027                    crate::TypeInner::Image { class, .. } => match class {
1028                        crate::ImageClass::Storage {
1029                            format:
1030                                crate::StorageFormat::R16Unorm
1031                                | crate::StorageFormat::R16Snorm
1032                                | crate::StorageFormat::Rg16Unorm
1033                                | crate::StorageFormat::Rg16Snorm
1034                                | crate::StorageFormat::Rgba16Unorm
1035                                | crate::StorageFormat::Rgba16Snorm,
1036                            ..
1037                        } => {
1038                            if !self
1039                                .capabilities
1040                                .contains(Capabilities::STORAGE_TEXTURE_16BIT_NORM_FORMATS)
1041                            {
1042                                return Err(GlobalVariableError::UnsupportedCapability(
1043                                    Capabilities::STORAGE_TEXTURE_16BIT_NORM_FORMATS,
1044                                ));
1045                            }
1046                        }
1047                        _ => {}
1048                    },
1049                    crate::TypeInner::Sampler { .. }
1050                    | crate::TypeInner::AccelerationStructure { .. }
1051                    | crate::TypeInner::RayQuery { .. } => {}
1052                    _ => {
1053                        return Err(GlobalVariableError::InvalidType(var.space));
1054                    }
1055                }
1056
1057                (TypeFlags::empty(), true)
1058            }
1059            crate::AddressSpace::Private => (
1060                TypeFlags::CONSTRUCTIBLE | TypeFlags::CREATION_RESOLVED,
1061                false,
1062            ),
1063            crate::AddressSpace::WorkGroup => (TypeFlags::DATA | TypeFlags::SIZED, false),
1064            crate::AddressSpace::TaskPayload => {
1065                if !self.capabilities.contains(Capabilities::MESH_SHADER) {
1066                    return Err(GlobalVariableError::UnsupportedCapability(
1067                        Capabilities::MESH_SHADER,
1068                    ));
1069                }
1070                (TypeFlags::DATA | TypeFlags::SIZED, false)
1071            }
1072            crate::AddressSpace::Immediate => {
1073                if !self.capabilities.contains(Capabilities::IMMEDIATES) {
1074                    return Err(GlobalVariableError::UnsupportedCapability(
1075                        Capabilities::IMMEDIATES,
1076                    ));
1077                }
1078                if let Err(ref err) = type_info.immediates_compatibility {
1079                    return Err(GlobalVariableError::InvalidImmediateType(err.clone()));
1080                }
1081                (
1082                    TypeFlags::DATA
1083                        | TypeFlags::COPY
1084                        | TypeFlags::HOST_SHAREABLE
1085                        | TypeFlags::SIZED,
1086                    false,
1087                )
1088            }
1089            crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => {
1090                if !self
1091                    .capabilities
1092                    .contains(Capabilities::RAY_TRACING_PIPELINE)
1093                {
1094                    return Err(GlobalVariableError::UnsupportedCapability(
1095                        Capabilities::RAY_TRACING_PIPELINE,
1096                    ));
1097                }
1098                (TypeFlags::DATA | TypeFlags::SIZED, false)
1099            }
1100        };
1101
1102        if !type_info.flags.contains(required_type_flags) {
1103            return Err(GlobalVariableError::MissingTypeFlags {
1104                seen: type_info.flags,
1105                required: required_type_flags,
1106            });
1107        }
1108
1109        if is_resource != var.binding.is_some() {
1110            if self.flags.contains(super::ValidationFlags::BINDINGS) {
1111                return Err(GlobalVariableError::InvalidBinding);
1112            }
1113        }
1114
1115        if var.space == crate::AddressSpace::TaskPayload {
1116            let ty = &gctx.types[var.ty].inner;
1117            // HLSL doesn't allow zero sized payloads.
1118            if ty.try_size(gctx) == Some(0) {
1119                return Err(GlobalVariableError::ZeroSizedTaskPayload);
1120            }
1121        }
1122
1123        if let Some(init) = var.init {
1124            match var.space {
1125                crate::AddressSpace::Private | crate::AddressSpace::Function => {}
1126                _ => {
1127                    return Err(GlobalVariableError::InitializerNotAllowed(var.space));
1128                }
1129            }
1130
1131            if !global_expr_kind.is_const_or_override(init) {
1132                return Err(GlobalVariableError::InitializerExprType);
1133            }
1134
1135            if !gctx.compare_types(
1136                &crate::proc::TypeResolution::Handle(var.ty),
1137                &mod_info[init],
1138            ) {
1139                return Err(GlobalVariableError::InitializerType);
1140            }
1141        }
1142
1143        Ok(())
1144    }
1145
1146    /// Validate the mesh shader output type `ty`, used as `mesh_output_type`.
1147    fn validate_mesh_output_type(
1148        &mut self,
1149        ep: &crate::EntryPoint,
1150        module: &crate::Module,
1151        ty: Handle<crate::Type>,
1152        mesh_output_type: MeshOutputType,
1153    ) -> Result<(), WithSpan<EntryPointError>> {
1154        if !matches!(module.types[ty].inner, crate::TypeInner::Struct { .. }) {
1155            return Err(EntryPointError::InvalidMeshOutputType.with_span_handle(ty, &module.types));
1156        }
1157        let mut result_built_ins = crate::FastHashSet::default();
1158        let mut ctx = VaryingContext {
1159            stage: ep.stage,
1160            output: true,
1161            types: &module.types,
1162            type_info: &self.types,
1163            location_mask: &mut self.location_mask,
1164            blend_src_mask: &mut self.blend_src_mask,
1165            built_ins: &mut result_built_ins,
1166            capabilities: self.capabilities,
1167            flags: self.flags,
1168            mesh_output_type,
1169            has_task_payload: ep.task_payload.is_some(),
1170        };
1171        ctx.validate(ep, ty, None)
1172            .map_err_inner(|e| EntryPointError::Result(e).with_span())?;
1173        if mesh_output_type == MeshOutputType::PrimitiveOutput {
1174            let mut num_indices_builtins = 0;
1175            if result_built_ins.contains(&crate::BuiltIn::PointIndex) {
1176                num_indices_builtins += 1;
1177            }
1178            if result_built_ins.contains(&crate::BuiltIn::LineIndices) {
1179                num_indices_builtins += 1;
1180            }
1181            if result_built_ins.contains(&crate::BuiltIn::TriangleIndices) {
1182                num_indices_builtins += 1;
1183            }
1184            if num_indices_builtins != 1 {
1185                return Err(EntryPointError::InvalidMeshPrimitiveOutputType
1186                    .with_span_handle(ty, &module.types));
1187            }
1188        } else if mesh_output_type == MeshOutputType::VertexOutput
1189            && !result_built_ins.contains(&crate::BuiltIn::Position { invariant: false })
1190        {
1191            return Err(
1192                EntryPointError::MissingVertexOutputPosition.with_span_handle(ty, &module.types)
1193            );
1194        }
1195
1196        Ok(())
1197    }
1198
1199    pub(super) fn validate_entry_point(
1200        &mut self,
1201        ep: &crate::EntryPoint,
1202        module: &crate::Module,
1203        mod_info: &ModuleInfo,
1204    ) -> Result<FunctionInfo, WithSpan<EntryPointError>> {
1205        match ep.stage {
1206            crate::ShaderStage::Task | crate::ShaderStage::Mesh
1207                if !self.capabilities.contains(Capabilities::MESH_SHADER) =>
1208            {
1209                return Err(
1210                    EntryPointError::UnsupportedCapability(Capabilities::MESH_SHADER).with_span(),
1211                );
1212            }
1213            crate::ShaderStage::RayGeneration
1214            | crate::ShaderStage::AnyHit
1215            | crate::ShaderStage::ClosestHit
1216            | crate::ShaderStage::Miss
1217                if !self
1218                    .capabilities
1219                    .contains(Capabilities::RAY_TRACING_PIPELINE) =>
1220            {
1221                return Err(EntryPointError::UnsupportedCapability(
1222                    Capabilities::RAY_TRACING_PIPELINE,
1223                )
1224                .with_span());
1225            }
1226            _ => {}
1227        }
1228        if ep.early_depth_test.is_some() {
1229            let required = Capabilities::EARLY_DEPTH_TEST;
1230            if !self.capabilities.contains(required) {
1231                return Err(
1232                    EntryPointError::Result(VaryingError::UnsupportedCapability(required))
1233                        .with_span(),
1234                );
1235            }
1236
1237            if ep.stage != crate::ShaderStage::Fragment {
1238                return Err(EntryPointError::UnexpectedEarlyDepthTest.with_span());
1239            }
1240        }
1241
1242        if ep.stage.compute_like() {
1243            if ep
1244                .workgroup_size
1245                .iter()
1246                .any(|&s| s == 0 || s > MAX_WORKGROUP_SIZE)
1247            {
1248                return Err(EntryPointError::OutOfRangeWorkgroupSize.with_span());
1249            }
1250        } else if ep.workgroup_size != [0; 3] {
1251            return Err(EntryPointError::UnexpectedWorkgroupSize.with_span());
1252        }
1253
1254        match (ep.stage, &ep.mesh_info) {
1255            (crate::ShaderStage::Mesh, &None) => {
1256                return Err(EntryPointError::ExpectedMeshShaderAttributes.with_span());
1257            }
1258            (crate::ShaderStage::Mesh, &Some(..)) => {}
1259            (_, &Some(_)) => {
1260                return Err(EntryPointError::UnexpectedMeshShaderAttributes.with_span());
1261            }
1262            (_, _) => {}
1263        }
1264
1265        let mut info = self
1266            .validate_function(&ep.function, module, mod_info, true)
1267            .map_err(WithSpan::into_other)?;
1268
1269        // Validate the task shader payload.
1270        match ep.stage {
1271            // Task shaders must produce a payload.
1272            crate::ShaderStage::Task => {
1273                let Some(handle) = ep.task_payload else {
1274                    return Err(EntryPointError::ExpectedTaskPayload.with_span());
1275                };
1276                if module.global_variables[handle].space != crate::AddressSpace::TaskPayload {
1277                    return Err(EntryPointError::TaskPayloadWrongAddressSpace
1278                        .with_span_handle(handle, &module.global_variables));
1279                }
1280                info.insert_global_use(GlobalUse::READ | GlobalUse::WRITE, handle);
1281            }
1282
1283            // Mesh shaders may accept a payload.
1284            crate::ShaderStage::Mesh => {
1285                if let Some(handle) = ep.task_payload {
1286                    if module.global_variables[handle].space != crate::AddressSpace::TaskPayload {
1287                        return Err(EntryPointError::TaskPayloadWrongAddressSpace
1288                            .with_span_handle(handle, &module.global_variables));
1289                    }
1290                    info.insert_global_use(GlobalUse::READ, handle);
1291                }
1292                if let Some(ref mesh_info) = ep.mesh_info {
1293                    info.insert_global_use(GlobalUse::READ, mesh_info.output_variable);
1294                }
1295            }
1296
1297            // Other stages must not have a payload.
1298            _ => {
1299                if let Some(handle) = ep.task_payload {
1300                    return Err(EntryPointError::UnexpectedTaskPayload
1301                        .with_span_handle(handle, &module.global_variables));
1302                }
1303            }
1304        }
1305
1306        {
1307            use super::ShaderStages;
1308
1309            let stage_bit = match ep.stage {
1310                crate::ShaderStage::Vertex => ShaderStages::VERTEX,
1311                crate::ShaderStage::Fragment => ShaderStages::FRAGMENT,
1312                crate::ShaderStage::Compute => ShaderStages::COMPUTE,
1313                crate::ShaderStage::Mesh => ShaderStages::MESH,
1314                crate::ShaderStage::Task => ShaderStages::TASK,
1315                crate::ShaderStage::RayGeneration => ShaderStages::RAY_GENERATION,
1316                crate::ShaderStage::AnyHit => ShaderStages::ANY_HIT,
1317                crate::ShaderStage::ClosestHit => ShaderStages::CLOSEST_HIT,
1318                crate::ShaderStage::Miss => ShaderStages::MISS,
1319            };
1320
1321            if !info.available_stages.contains(stage_bit) {
1322                return Err(EntryPointError::ForbiddenStageOperations.with_span());
1323            }
1324        }
1325
1326        self.location_mask.clear();
1327        let mut argument_built_ins = crate::FastHashSet::default();
1328        // TODO: add span info to function arguments
1329        for (index, fa) in ep.function.arguments.iter().enumerate() {
1330            let mut ctx = VaryingContext {
1331                stage: ep.stage,
1332                output: false,
1333                types: &module.types,
1334                type_info: &self.types,
1335                location_mask: &mut self.location_mask,
1336                blend_src_mask: &mut self.blend_src_mask,
1337                built_ins: &mut argument_built_ins,
1338                capabilities: self.capabilities,
1339                flags: self.flags,
1340                mesh_output_type: MeshOutputType::None,
1341                has_task_payload: ep.task_payload.is_some(),
1342            };
1343            ctx.validate(ep, fa.ty, fa.binding.as_ref())
1344                .map_err_inner(|e| EntryPointError::Argument(index as u32, e).with_span())?;
1345        }
1346
1347        self.location_mask.clear();
1348        if let Some(ref fr) = ep.function.result {
1349            let mut result_built_ins = crate::FastHashSet::default();
1350            let mut ctx = VaryingContext {
1351                stage: ep.stage,
1352                output: true,
1353                types: &module.types,
1354                type_info: &self.types,
1355                location_mask: &mut self.location_mask,
1356                blend_src_mask: &mut self.blend_src_mask,
1357                built_ins: &mut result_built_ins,
1358                capabilities: self.capabilities,
1359                flags: self.flags,
1360                mesh_output_type: MeshOutputType::None,
1361                has_task_payload: ep.task_payload.is_some(),
1362            };
1363            ctx.validate(ep, fr.ty, fr.binding.as_ref())
1364                .map_err_inner(|e| EntryPointError::Result(e).with_span())?;
1365            if ep.stage == crate::ShaderStage::Vertex
1366                && !result_built_ins.contains(&crate::BuiltIn::Position { invariant: false })
1367            {
1368                return Err(EntryPointError::MissingVertexOutputPosition.with_span());
1369            }
1370            if ep.stage == crate::ShaderStage::Mesh {
1371                return Err(EntryPointError::UnexpectedMeshShaderEntryResult.with_span());
1372            }
1373            // Task shaders must have a single `MeshTaskSize` output, and nothing else.
1374            if ep.stage == crate::ShaderStage::Task {
1375                let ok = module.types[fr.ty].inner
1376                    == crate::TypeInner::Vector {
1377                        size: crate::VectorSize::Tri,
1378                        scalar: crate::Scalar::U32,
1379                    };
1380                if !ok {
1381                    return Err(EntryPointError::WrongTaskShaderEntryResult.with_span());
1382                }
1383            }
1384            if !self.blend_src_mask.is_empty() {
1385                info.dual_source_blending = true;
1386            }
1387        } else if ep.stage == crate::ShaderStage::Vertex {
1388            return Err(EntryPointError::MissingVertexOutputPosition.with_span());
1389        } else if ep.stage == crate::ShaderStage::Task {
1390            return Err(EntryPointError::WrongTaskShaderEntryResult.with_span());
1391        }
1392
1393        {
1394            let mut used_immediates = module
1395                .global_variables
1396                .iter()
1397                .filter(|&(_, var)| var.space == crate::AddressSpace::Immediate)
1398                .map(|(handle, _)| handle)
1399                .filter(|&handle| !info[handle].is_empty());
1400            // Check if there is more than one immediate data, and error if so.
1401            // Use a loop for when returning multiple errors is supported.
1402            if let Some(handle) = used_immediates.nth(1) {
1403                return Err(EntryPointError::MoreThanOneImmediateUsed
1404                    .with_span_handle(handle, &module.global_variables));
1405            }
1406        }
1407
1408        self.ep_resource_bindings.clear();
1409        for (var_handle, var) in module.global_variables.iter() {
1410            let usage = info[var_handle];
1411            if usage.is_empty() {
1412                continue;
1413            }
1414
1415            if var.space == crate::AddressSpace::TaskPayload {
1416                if ep.task_payload != Some(var_handle) {
1417                    return Err(EntryPointError::WrongTaskPayloadUsed
1418                        .with_span_handle(var_handle, &module.global_variables));
1419                }
1420                let size = module.types[var.ty].inner.size(module.to_ctx());
1421                if size < 4 {
1422                    return Err(EntryPointError::TaskPayloadTooSmall(size)
1423                        .with_span_handle(var_handle, &module.global_variables));
1424                }
1425            }
1426
1427            let allowed_usage = match var.space {
1428                crate::AddressSpace::Function => unreachable!(),
1429                crate::AddressSpace::Uniform => GlobalUse::READ | GlobalUse::QUERY,
1430                crate::AddressSpace::Storage { access } => storage_usage(access),
1431                crate::AddressSpace::Handle => match module.types[var.ty].inner {
1432                    crate::TypeInner::BindingArray { base, .. } => match module.types[base].inner {
1433                        crate::TypeInner::Image {
1434                            class: crate::ImageClass::Storage { access, .. },
1435                            ..
1436                        } => storage_usage(access),
1437                        _ => GlobalUse::READ | GlobalUse::QUERY,
1438                    },
1439                    crate::TypeInner::Image {
1440                        class: crate::ImageClass::Storage { access, .. },
1441                        ..
1442                    } => storage_usage(access),
1443                    _ => GlobalUse::READ | GlobalUse::QUERY,
1444                },
1445                crate::AddressSpace::Private | crate::AddressSpace::WorkGroup => {
1446                    GlobalUse::READ | GlobalUse::WRITE | GlobalUse::QUERY
1447                }
1448                crate::AddressSpace::TaskPayload => {
1449                    GlobalUse::READ
1450                        | GlobalUse::QUERY
1451                        | if ep.stage == crate::ShaderStage::Task {
1452                            GlobalUse::WRITE
1453                        } else {
1454                            GlobalUse::empty()
1455                        }
1456                }
1457                crate::AddressSpace::Immediate => GlobalUse::READ,
1458                crate::AddressSpace::RayPayload => {
1459                    if !matches!(
1460                        ep.stage,
1461                        crate::ShaderStage::RayGeneration
1462                            | crate::ShaderStage::ClosestHit
1463                            | crate::ShaderStage::Miss
1464                    ) {
1465                        return Err(EntryPointError::RayPayloadInInvalidStage(ep.stage)
1466                            .with_span_handle(var_handle, &module.global_variables));
1467                    }
1468                    GlobalUse::READ | GlobalUse::QUERY | GlobalUse::WRITE
1469                }
1470                crate::AddressSpace::IncomingRayPayload => {
1471                    if !matches!(
1472                        ep.stage,
1473                        crate::ShaderStage::AnyHit
1474                            | crate::ShaderStage::ClosestHit
1475                            | crate::ShaderStage::Miss
1476                    ) {
1477                        return Err(EntryPointError::IncomingRayPayloadInInvalidStage(ep.stage)
1478                            .with_span_handle(var_handle, &module.global_variables));
1479                    }
1480                    GlobalUse::READ | GlobalUse::QUERY | GlobalUse::WRITE
1481                }
1482            };
1483            if !allowed_usage.contains(usage) {
1484                log::warn!("\tUsage error for: {var:?}");
1485                log::warn!("\tAllowed usage: {allowed_usage:?}, requested: {usage:?}");
1486                return Err(EntryPointError::InvalidGlobalUsage(var_handle, usage)
1487                    .with_span_handle(var_handle, &module.global_variables));
1488            }
1489
1490            if let Some(ref bind) = var.binding {
1491                if !self.ep_resource_bindings.insert(*bind) {
1492                    if self.flags.contains(super::ValidationFlags::BINDINGS) {
1493                        return Err(EntryPointError::BindingCollision(var_handle)
1494                            .with_span_handle(var_handle, &module.global_variables));
1495                    }
1496                }
1497            }
1498        }
1499
1500        // If this is a `Mesh` entry point, check its vertex and primitive output types.
1501        // We verified previously that only mesh shaders can have `mesh_info`.
1502        if let &Some(ref mesh_info) = &ep.mesh_info {
1503            if module.global_variables[mesh_info.output_variable].space
1504                != crate::AddressSpace::WorkGroup
1505            {
1506                return Err(EntryPointError::WrongMeshOutputAddressSpace.with_span());
1507            }
1508
1509            let mut implied = module.analyze_mesh_shader_info(mesh_info.output_variable);
1510            if let Some(e) = implied.2 {
1511                return Err(e);
1512            }
1513
1514            if let Some(e) = mesh_info.max_vertices_override {
1515                if let crate::Expression::Override(o) = module.global_expressions[e] {
1516                    if implied.1[0] != Some(o) {
1517                        return Err(EntryPointError::BadMeshOutputVariableType.with_span());
1518                    }
1519                }
1520            }
1521            if let Some(e) = mesh_info.max_primitives_override {
1522                if let crate::Expression::Override(o) = module.global_expressions[e] {
1523                    if implied.1[1] != Some(o) {
1524                        return Err(EntryPointError::BadMeshOutputVariableType.with_span());
1525                    }
1526                }
1527            }
1528
1529            implied.0.max_vertices_override = mesh_info.max_vertices_override;
1530            implied.0.max_primitives_override = mesh_info.max_primitives_override;
1531            if implied.0 != *mesh_info {
1532                return Err(EntryPointError::BadMeshOutputVariableType.with_span());
1533            }
1534            if mesh_info.topology == crate::MeshOutputTopology::Points
1535                && !self
1536                    .capabilities
1537                    .contains(Capabilities::MESH_SHADER_POINT_TOPOLOGY)
1538            {
1539                return Err(EntryPointError::UnsupportedCapability(
1540                    Capabilities::MESH_SHADER_POINT_TOPOLOGY,
1541                )
1542                .with_span());
1543            }
1544
1545            self.validate_mesh_output_type(
1546                ep,
1547                module,
1548                mesh_info.vertex_output_type,
1549                MeshOutputType::VertexOutput,
1550            )?;
1551            self.validate_mesh_output_type(
1552                ep,
1553                module,
1554                mesh_info.primitive_output_type,
1555                MeshOutputType::PrimitiveOutput,
1556            )?;
1557        }
1558
1559        Ok(info)
1560    }
1561}
naga/valid/interface.rs

naga/valid/
interface.rs
