naga/valid/
interface.rs

1use super::{
2    analyzer::{FunctionInfo, GlobalUse},
3    Capabilities, Disalignment, FunctionError, ModuleInfo,
4};
5use crate::arena::{Handle, UniqueArena};
6
7use crate::span::{AddSpan as _, MapErrWithSpan as _, SpanProvider as _, WithSpan};
8use bit_set::BitSet;
9
10const MAX_WORKGROUP_SIZE: u32 = 0x4000;
11
12#[derive(Clone, Debug, thiserror::Error)]
13#[cfg_attr(test, derive(PartialEq))]
14pub enum GlobalVariableError {
15    #[error("Usage isn't compatible with address space {0:?}")]
16    InvalidUsage(crate::AddressSpace),
17    #[error("Type isn't compatible with address space {0:?}")]
18    InvalidType(crate::AddressSpace),
19    #[error("Type flags {seen:?} do not meet the required {required:?}")]
20    MissingTypeFlags {
21        required: super::TypeFlags,
22        seen: super::TypeFlags,
23    },
24    #[error("Capability {0:?} is not supported")]
25    UnsupportedCapability(Capabilities),
26    #[error("Binding decoration is missing or not applicable")]
27    InvalidBinding,
28    #[error("Alignment requirements for address space {0:?} are not met by {1:?}")]
29    Alignment(
30        crate::AddressSpace,
31        Handle<crate::Type>,
32        #[source] Disalignment,
33    ),
34    #[error("Initializer must be an override-expression")]
35    InitializerExprType,
36    #[error("Initializer doesn't match the variable type")]
37    InitializerType,
38    #[error("Initializer can't be used with address space {0:?}")]
39    InitializerNotAllowed(crate::AddressSpace),
40    #[error("Storage address space doesn't support write-only access")]
41    StorageAddressSpaceWriteOnlyNotSupported,
42}
43
44#[derive(Clone, Debug, thiserror::Error)]
45#[cfg_attr(test, derive(PartialEq))]
46pub enum VaryingError {
47    #[error("The type {0:?} does not match the varying")]
48    InvalidType(Handle<crate::Type>),
49    #[error("The type {0:?} cannot be used for user-defined entry point inputs or outputs")]
50    NotIOShareableType(Handle<crate::Type>),
51    #[error("Interpolation is not valid")]
52    InvalidInterpolation,
53    #[error("Cannot combine {interpolation:?} interpolation with the {sampling:?} sample type")]
54    InvalidInterpolationSamplingCombination {
55        interpolation: crate::Interpolation,
56        sampling: crate::Sampling,
57    },
58    #[error("Interpolation must be specified on vertex shader outputs and fragment shader inputs")]
59    MissingInterpolation,
60    #[error("Built-in {0:?} is not available at this stage")]
61    InvalidBuiltInStage(crate::BuiltIn),
62    #[error("Built-in type for {0:?} is invalid")]
63    InvalidBuiltInType(crate::BuiltIn),
64    #[error("Entry point arguments and return values must all have bindings")]
65    MissingBinding,
66    #[error("Struct member {0} is missing a binding")]
67    MemberMissingBinding(u32),
68    #[error("Multiple bindings at location {location} are present")]
69    BindingCollision { location: u32 },
70    #[error("Built-in {0:?} is present more than once")]
71    DuplicateBuiltIn(crate::BuiltIn),
72    #[error("Capability {0:?} is not supported")]
73    UnsupportedCapability(Capabilities),
74    #[error("The attribute {0:?} is only valid as an output for stage {1:?}")]
75    InvalidInputAttributeInStage(&'static str, crate::ShaderStage),
76    #[error("The attribute {0:?} is not valid for stage {1:?}")]
77    InvalidAttributeInStage(&'static str, crate::ShaderStage),
78    #[error(
79        "The location index {location} cannot be used together with the attribute {attribute:?}"
80    )]
81    InvalidLocationAttributeCombination {
82        location: u32,
83        attribute: &'static str,
84    },
85    #[error("Workgroup size is multi dimensional, @builtin(subgroup_id) and @builtin(subgroup_invocation_id) are not supported.")]
86    InvalidMultiDimensionalSubgroupBuiltIn,
87}
88
89#[derive(Clone, Debug, thiserror::Error)]
90#[cfg_attr(test, derive(PartialEq))]
91pub enum EntryPointError {
92    #[error("Multiple conflicting entry points")]
93    Conflict,
94    #[error("Vertex shaders must return a `@builtin(position)` output value")]
95    MissingVertexOutputPosition,
96    #[error("Early depth test is not applicable")]
97    UnexpectedEarlyDepthTest,
98    #[error("Workgroup size is not applicable")]
99    UnexpectedWorkgroupSize,
100    #[error("Workgroup size is out of range")]
101    OutOfRangeWorkgroupSize,
102    #[error("Uses operations forbidden at this stage")]
103    ForbiddenStageOperations,
104    #[error("Global variable {0:?} is used incorrectly as {1:?}")]
105    InvalidGlobalUsage(Handle<crate::GlobalVariable>, GlobalUse),
106    #[error("More than 1 push constant variable is used")]
107    MoreThanOnePushConstantUsed,
108    #[error("Bindings for {0:?} conflict with other resource")]
109    BindingCollision(Handle<crate::GlobalVariable>),
110    #[error("Argument {0} varying error")]
111    Argument(u32, #[source] VaryingError),
112    #[error(transparent)]
113    Result(#[from] VaryingError),
114    #[error("Location {location} interpolation of an integer has to be flat")]
115    InvalidIntegerInterpolation { location: u32 },
116    #[error(transparent)]
117    Function(#[from] FunctionError),
118    #[error(
119        "Invalid locations {location_mask:?} are set while dual source blending. Only location 0 may be set."
120    )]
121    InvalidLocationsWhileDualSourceBlending { location_mask: BitSet },
122}
123
124fn storage_usage(access: crate::StorageAccess) -> GlobalUse {
125    let mut storage_usage = GlobalUse::QUERY;
126    if access.contains(crate::StorageAccess::LOAD) {
127        storage_usage |= GlobalUse::READ;
128    }
129    if access.contains(crate::StorageAccess::STORE) {
130        storage_usage |= GlobalUse::WRITE;
131    }
132    storage_usage
133}
134
135struct VaryingContext<'a> {
136    stage: crate::ShaderStage,
137    output: bool,
138    second_blend_source: bool,
139    types: &'a UniqueArena<crate::Type>,
140    type_info: &'a Vec<super::r#type::TypeInfo>,
141    location_mask: &'a mut BitSet,
142    built_ins: &'a mut crate::FastHashSet<crate::BuiltIn>,
143    capabilities: Capabilities,
144    flags: super::ValidationFlags,
145}
146
147impl VaryingContext<'_> {
148    fn validate_impl(
149        &mut self,
150        ep: &crate::EntryPoint,
151        ty: Handle<crate::Type>,
152        binding: &crate::Binding,
153    ) -> Result<(), VaryingError> {
154        use crate::{BuiltIn as Bi, ShaderStage as St, TypeInner as Ti, VectorSize as Vs};
155
156        let ty_inner = &self.types[ty].inner;
157        match *binding {
158            crate::Binding::BuiltIn(built_in) => {
159                // Ignore the `invariant` field for the sake of duplicate checks,
160                // but use the original in error messages.
161                let canonical = if let crate::BuiltIn::Position { .. } = built_in {
162                    crate::BuiltIn::Position { invariant: false }
163                } else {
164                    built_in
165                };
166
167                if self.built_ins.contains(&canonical) {
168                    return Err(VaryingError::DuplicateBuiltIn(built_in));
169                }
170                self.built_ins.insert(canonical);
171
172                let required = match built_in {
173                    Bi::ClipDistance => Capabilities::CLIP_DISTANCE,
174                    Bi::CullDistance => Capabilities::CULL_DISTANCE,
175                    Bi::PrimitiveIndex => Capabilities::PRIMITIVE_INDEX,
176                    Bi::ViewIndex => Capabilities::MULTIVIEW,
177                    Bi::SampleIndex => Capabilities::MULTISAMPLED_SHADING,
178                    Bi::NumSubgroups
179                    | Bi::SubgroupId
180                    | Bi::SubgroupSize
181                    | Bi::SubgroupInvocationId => Capabilities::SUBGROUP,
182                    _ => Capabilities::empty(),
183                };
184                if !self.capabilities.contains(required) {
185                    return Err(VaryingError::UnsupportedCapability(required));
186                }
187
188                if matches!(
189                    built_in,
190                    crate::BuiltIn::SubgroupId | crate::BuiltIn::SubgroupInvocationId
191                ) && ep.workgroup_size[1..].iter().any(|&s| s > 1)
192                {
193                    return Err(VaryingError::InvalidMultiDimensionalSubgroupBuiltIn);
194                }
195
196                let (visible, type_good) = match built_in {
197                    Bi::BaseInstance
198                    | Bi::BaseVertex
199                    | Bi::InstanceIndex
200                    | Bi::VertexIndex
201                    | Bi::DrawID => (
202                        self.stage == St::Vertex && !self.output,
203                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
204                    ),
205                    Bi::ClipDistance | Bi::CullDistance => (
206                        self.stage == St::Vertex && self.output,
207                        match *ty_inner {
208                            Ti::Array { base, .. } => {
209                                self.types[base].inner == Ti::Scalar(crate::Scalar::F32)
210                            }
211                            _ => false,
212                        },
213                    ),
214                    Bi::PointSize => (
215                        self.stage == St::Vertex && self.output,
216                        *ty_inner == Ti::Scalar(crate::Scalar::F32),
217                    ),
218                    Bi::PointCoord => (
219                        self.stage == St::Fragment && !self.output,
220                        *ty_inner
221                            == Ti::Vector {
222                                size: Vs::Bi,
223                                scalar: crate::Scalar::F32,
224                            },
225                    ),
226                    Bi::Position { .. } => (
227                        match self.stage {
228                            St::Vertex => self.output,
229                            St::Fragment => !self.output,
230                            St::Compute => false,
231                        },
232                        *ty_inner
233                            == Ti::Vector {
234                                size: Vs::Quad,
235                                scalar: crate::Scalar::F32,
236                            },
237                    ),
238                    Bi::ViewIndex => (
239                        match self.stage {
240                            St::Vertex | St::Fragment => !self.output,
241                            St::Compute => false,
242                        },
243                        *ty_inner == Ti::Scalar(crate::Scalar::I32),
244                    ),
245                    Bi::FragDepth => (
246                        self.stage == St::Fragment && self.output,
247                        *ty_inner == Ti::Scalar(crate::Scalar::F32),
248                    ),
249                    Bi::FrontFacing => (
250                        self.stage == St::Fragment && !self.output,
251                        *ty_inner == Ti::Scalar(crate::Scalar::BOOL),
252                    ),
253                    Bi::PrimitiveIndex => (
254                        self.stage == St::Fragment && !self.output,
255                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
256                    ),
257                    Bi::SampleIndex => (
258                        self.stage == St::Fragment && !self.output,
259                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
260                    ),
261                    Bi::SampleMask => (
262                        self.stage == St::Fragment,
263                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
264                    ),
265                    Bi::LocalInvocationIndex => (
266                        self.stage == St::Compute && !self.output,
267                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
268                    ),
269                    Bi::GlobalInvocationId
270                    | Bi::LocalInvocationId
271                    | Bi::WorkGroupId
272                    | Bi::WorkGroupSize
273                    | Bi::NumWorkGroups => (
274                        self.stage == St::Compute && !self.output,
275                        *ty_inner
276                            == Ti::Vector {
277                                size: Vs::Tri,
278                                scalar: crate::Scalar::U32,
279                            },
280                    ),
281                    Bi::NumSubgroups | Bi::SubgroupId => (
282                        self.stage == St::Compute && !self.output,
283                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
284                    ),
285                    Bi::SubgroupSize | Bi::SubgroupInvocationId => (
286                        match self.stage {
287                            St::Compute | St::Fragment => !self.output,
288                            St::Vertex => false,
289                        },
290                        *ty_inner == Ti::Scalar(crate::Scalar::U32),
291                    ),
292                };
293
294                if !visible {
295                    return Err(VaryingError::InvalidBuiltInStage(built_in));
296                }
297                if !type_good {
298                    log::warn!("Wrong builtin type: {:?}", ty_inner);
299                    return Err(VaryingError::InvalidBuiltInType(built_in));
300                }
301            }
302            crate::Binding::Location {
303                location,
304                interpolation,
305                sampling,
306                second_blend_source,
307            } => {
308                // Only IO-shareable types may be stored in locations.
309                if !self.type_info[ty.index()]
310                    .flags
311                    .contains(super::TypeFlags::IO_SHAREABLE)
312                {
313                    return Err(VaryingError::NotIOShareableType(ty));
314                }
315
316                if second_blend_source {
317                    if !self
318                        .capabilities
319                        .contains(Capabilities::DUAL_SOURCE_BLENDING)
320                    {
321                        return Err(VaryingError::UnsupportedCapability(
322                            Capabilities::DUAL_SOURCE_BLENDING,
323                        ));
324                    }
325                    if self.stage != crate::ShaderStage::Fragment {
326                        return Err(VaryingError::InvalidAttributeInStage(
327                            "second_blend_source",
328                            self.stage,
329                        ));
330                    }
331                    if !self.output {
332                        return Err(VaryingError::InvalidInputAttributeInStage(
333                            "second_blend_source",
334                            self.stage,
335                        ));
336                    }
337                    if location != 0 {
338                        return Err(VaryingError::InvalidLocationAttributeCombination {
339                            location,
340                            attribute: "second_blend_source",
341                        });
342                    }
343
344                    self.second_blend_source = true;
345                } else if !self.location_mask.insert(location as usize) {
346                    if self.flags.contains(super::ValidationFlags::BINDINGS) {
347                        return Err(VaryingError::BindingCollision { location });
348                    }
349                }
350
351                if let Some(interpolation) = interpolation {
352                    let invalid_sampling = match (interpolation, sampling) {
353                        (_, None)
354                        | (
355                            crate::Interpolation::Perspective | crate::Interpolation::Linear,
356                            Some(
357                                crate::Sampling::Center
358                                | crate::Sampling::Centroid
359                                | crate::Sampling::Sample,
360                            ),
361                        )
362                        | (
363                            crate::Interpolation::Flat,
364                            Some(crate::Sampling::First | crate::Sampling::Either),
365                        ) => None,
366                        (_, Some(invalid_sampling)) => Some(invalid_sampling),
367                    };
368                    if let Some(sampling) = invalid_sampling {
369                        return Err(VaryingError::InvalidInterpolationSamplingCombination {
370                            interpolation,
371                            sampling,
372                        });
373                    }
374                }
375
376                let needs_interpolation = match self.stage {
377                    crate::ShaderStage::Vertex => self.output,
378                    crate::ShaderStage::Fragment => !self.output,
379                    crate::ShaderStage::Compute => false,
380                };
381
382                // It doesn't make sense to specify a sampling when `interpolation` is `Flat`, but
383                // SPIR-V and GLSL both explicitly tolerate such combinations of decorators /
384                // qualifiers, so we won't complain about that here.
385                let _ = sampling;
386
387                let required = match sampling {
388                    Some(crate::Sampling::Sample) => Capabilities::MULTISAMPLED_SHADING,
389                    _ => Capabilities::empty(),
390                };
391                if !self.capabilities.contains(required) {
392                    return Err(VaryingError::UnsupportedCapability(required));
393                }
394
395                match ty_inner.scalar_kind() {
396                    Some(crate::ScalarKind::Float) => {
397                        if needs_interpolation && interpolation.is_none() {
398                            return Err(VaryingError::MissingInterpolation);
399                        }
400                    }
401                    Some(_) => {
402                        if needs_interpolation && interpolation != Some(crate::Interpolation::Flat)
403                        {
404                            return Err(VaryingError::InvalidInterpolation);
405                        }
406                    }
407                    None => return Err(VaryingError::InvalidType(ty)),
408                }
409            }
410        }
411
412        Ok(())
413    }
414
415    fn validate(
416        &mut self,
417        ep: &crate::EntryPoint,
418        ty: Handle<crate::Type>,
419        binding: Option<&crate::Binding>,
420    ) -> Result<(), WithSpan<VaryingError>> {
421        let span_context = self.types.get_span_context(ty);
422        match binding {
423            Some(binding) => self
424                .validate_impl(ep, ty, binding)
425                .map_err(|e| e.with_span_context(span_context)),
426            None => {
427                match self.types[ty].inner {
428                    crate::TypeInner::Struct { ref members, .. } => {
429                        for (index, member) in members.iter().enumerate() {
430                            let span_context = self.types.get_span_context(ty);
431                            match member.binding {
432                                None => {
433                                    if self.flags.contains(super::ValidationFlags::BINDINGS) {
434                                        return Err(VaryingError::MemberMissingBinding(
435                                            index as u32,
436                                        )
437                                        .with_span_context(span_context));
438                                    }
439                                }
440                                Some(ref binding) => self
441                                    .validate_impl(ep, member.ty, binding)
442                                    .map_err(|e| e.with_span_context(span_context))?,
443                            }
444                        }
445                    }
446                    _ => {
447                        if self.flags.contains(super::ValidationFlags::BINDINGS) {
448                            return Err(VaryingError::MissingBinding.with_span());
449                        }
450                    }
451                }
452                Ok(())
453            }
454        }
455    }
456}
457
458impl super::Validator {
459    pub(super) fn validate_global_var(
460        &self,
461        var: &crate::GlobalVariable,
462        gctx: crate::proc::GlobalCtx,
463        mod_info: &ModuleInfo,
464        global_expr_kind: &crate::proc::ExpressionKindTracker,
465    ) -> Result<(), GlobalVariableError> {
466        use super::TypeFlags;
467
468        log::debug!("var {:?}", var);
469        let inner_ty = match gctx.types[var.ty].inner {
470            // A binding array is (mostly) supposed to behave the same as a
471            // series of individually bound resources, so we can (mostly)
472            // validate a `binding_array<T>` as if it were just a plain `T`.
473            crate::TypeInner::BindingArray { base, .. } => match var.space {
474                crate::AddressSpace::Storage { .. }
475                | crate::AddressSpace::Uniform
476                | crate::AddressSpace::Handle => base,
477                _ => return Err(GlobalVariableError::InvalidUsage(var.space)),
478            },
479            _ => var.ty,
480        };
481        let type_info = &self.types[inner_ty.index()];
482
483        let (required_type_flags, is_resource) = match var.space {
484            crate::AddressSpace::Function => {
485                return Err(GlobalVariableError::InvalidUsage(var.space))
486            }
487            crate::AddressSpace::Storage { access } => {
488                if let Err((ty_handle, disalignment)) = type_info.storage_layout {
489                    if self.flags.contains(super::ValidationFlags::STRUCT_LAYOUTS) {
490                        return Err(GlobalVariableError::Alignment(
491                            var.space,
492                            ty_handle,
493                            disalignment,
494                        ));
495                    }
496                }
497                if access == crate::StorageAccess::STORE {
498                    return Err(GlobalVariableError::StorageAddressSpaceWriteOnlyNotSupported);
499                }
500                (TypeFlags::DATA | TypeFlags::HOST_SHAREABLE, true)
501            }
502            crate::AddressSpace::Uniform => {
503                if let Err((ty_handle, disalignment)) = type_info.uniform_layout {
504                    if self.flags.contains(super::ValidationFlags::STRUCT_LAYOUTS) {
505                        return Err(GlobalVariableError::Alignment(
506                            var.space,
507                            ty_handle,
508                            disalignment,
509                        ));
510                    }
511                }
512                (
513                    TypeFlags::DATA
514                        | TypeFlags::COPY
515                        | TypeFlags::SIZED
516                        | TypeFlags::HOST_SHAREABLE,
517                    true,
518                )
519            }
520            crate::AddressSpace::Handle => {
521                match gctx.types[inner_ty].inner {
522                    crate::TypeInner::Image { class, .. } => match class {
523                        crate::ImageClass::Storage {
524                            format:
525                                crate::StorageFormat::R16Unorm
526                                | crate::StorageFormat::R16Snorm
527                                | crate::StorageFormat::Rg16Unorm
528                                | crate::StorageFormat::Rg16Snorm
529                                | crate::StorageFormat::Rgba16Unorm
530                                | crate::StorageFormat::Rgba16Snorm,
531                            ..
532                        } => {
533                            if !self
534                                .capabilities
535                                .contains(Capabilities::STORAGE_TEXTURE_16BIT_NORM_FORMATS)
536                            {
537                                return Err(GlobalVariableError::UnsupportedCapability(
538                                    Capabilities::STORAGE_TEXTURE_16BIT_NORM_FORMATS,
539                                ));
540                            }
541                        }
542                        _ => {}
543                    },
544                    crate::TypeInner::Sampler { .. }
545                    | crate::TypeInner::AccelerationStructure
546                    | crate::TypeInner::RayQuery => {}
547                    _ => {
548                        return Err(GlobalVariableError::InvalidType(var.space));
549                    }
550                }
551
552                (TypeFlags::empty(), true)
553            }
554            crate::AddressSpace::Private => (TypeFlags::CONSTRUCTIBLE, false),
555            crate::AddressSpace::WorkGroup => (TypeFlags::DATA | TypeFlags::SIZED, false),
556            crate::AddressSpace::PushConstant => {
557                if !self.capabilities.contains(Capabilities::PUSH_CONSTANT) {
558                    return Err(GlobalVariableError::UnsupportedCapability(
559                        Capabilities::PUSH_CONSTANT,
560                    ));
561                }
562                (
563                    TypeFlags::DATA
564                        | TypeFlags::COPY
565                        | TypeFlags::HOST_SHAREABLE
566                        | TypeFlags::SIZED,
567                    false,
568                )
569            }
570        };
571
572        if !type_info.flags.contains(required_type_flags) {
573            return Err(GlobalVariableError::MissingTypeFlags {
574                seen: type_info.flags,
575                required: required_type_flags,
576            });
577        }
578
579        if is_resource != var.binding.is_some() {
580            if self.flags.contains(super::ValidationFlags::BINDINGS) {
581                return Err(GlobalVariableError::InvalidBinding);
582            }
583        }
584
585        if let Some(init) = var.init {
586            match var.space {
587                crate::AddressSpace::Private | crate::AddressSpace::Function => {}
588                _ => {
589                    return Err(GlobalVariableError::InitializerNotAllowed(var.space));
590                }
591            }
592
593            if !global_expr_kind.is_const_or_override(init) {
594                return Err(GlobalVariableError::InitializerExprType);
595            }
596
597            let decl_ty = &gctx.types[var.ty].inner;
598            let init_ty = mod_info[init].inner_with(gctx.types);
599            if !decl_ty.equivalent(init_ty, gctx.types) {
600                return Err(GlobalVariableError::InitializerType);
601            }
602        }
603
604        Ok(())
605    }
606
607    pub(super) fn validate_entry_point(
608        &mut self,
609        ep: &crate::EntryPoint,
610        module: &crate::Module,
611        mod_info: &ModuleInfo,
612        global_expr_kind: &crate::proc::ExpressionKindTracker,
613    ) -> Result<FunctionInfo, WithSpan<EntryPointError>> {
614        if ep.early_depth_test.is_some() {
615            let required = Capabilities::EARLY_DEPTH_TEST;
616            if !self.capabilities.contains(required) {
617                return Err(
618                    EntryPointError::Result(VaryingError::UnsupportedCapability(required))
619                        .with_span(),
620                );
621            }
622
623            if ep.stage != crate::ShaderStage::Fragment {
624                return Err(EntryPointError::UnexpectedEarlyDepthTest.with_span());
625            }
626        }
627
628        if ep.stage == crate::ShaderStage::Compute {
629            if ep
630                .workgroup_size
631                .iter()
632                .any(|&s| s == 0 || s > MAX_WORKGROUP_SIZE)
633            {
634                return Err(EntryPointError::OutOfRangeWorkgroupSize.with_span());
635            }
636        } else if ep.workgroup_size != [0; 3] {
637            return Err(EntryPointError::UnexpectedWorkgroupSize.with_span());
638        }
639
640        let mut info = self
641            .validate_function(&ep.function, module, mod_info, true, global_expr_kind)
642            .map_err(WithSpan::into_other)?;
643
644        {
645            use super::ShaderStages;
646
647            let stage_bit = match ep.stage {
648                crate::ShaderStage::Vertex => ShaderStages::VERTEX,
649                crate::ShaderStage::Fragment => ShaderStages::FRAGMENT,
650                crate::ShaderStage::Compute => ShaderStages::COMPUTE,
651            };
652
653            if !info.available_stages.contains(stage_bit) {
654                return Err(EntryPointError::ForbiddenStageOperations.with_span());
655            }
656        }
657
658        self.location_mask.clear();
659        let mut argument_built_ins = crate::FastHashSet::default();
660        // TODO: add span info to function arguments
661        for (index, fa) in ep.function.arguments.iter().enumerate() {
662            let mut ctx = VaryingContext {
663                stage: ep.stage,
664                output: false,
665                second_blend_source: false,
666                types: &module.types,
667                type_info: &self.types,
668                location_mask: &mut self.location_mask,
669                built_ins: &mut argument_built_ins,
670                capabilities: self.capabilities,
671                flags: self.flags,
672            };
673            ctx.validate(ep, fa.ty, fa.binding.as_ref())
674                .map_err_inner(|e| EntryPointError::Argument(index as u32, e).with_span())?;
675        }
676
677        self.location_mask.clear();
678        if let Some(ref fr) = ep.function.result {
679            let mut result_built_ins = crate::FastHashSet::default();
680            let mut ctx = VaryingContext {
681                stage: ep.stage,
682                output: true,
683                second_blend_source: false,
684                types: &module.types,
685                type_info: &self.types,
686                location_mask: &mut self.location_mask,
687                built_ins: &mut result_built_ins,
688                capabilities: self.capabilities,
689                flags: self.flags,
690            };
691            ctx.validate(ep, fr.ty, fr.binding.as_ref())
692                .map_err_inner(|e| EntryPointError::Result(e).with_span())?;
693            if ctx.second_blend_source {
694                // Only the first location may be used when dual source blending
695                if ctx.location_mask.len() == 1 && ctx.location_mask.contains(0) {
696                    info.dual_source_blending = true;
697                } else {
698                    return Err(EntryPointError::InvalidLocationsWhileDualSourceBlending {
699                        location_mask: self.location_mask.clone(),
700                    }
701                    .with_span());
702                }
703            }
704
705            if ep.stage == crate::ShaderStage::Vertex
706                && !result_built_ins.contains(&crate::BuiltIn::Position { invariant: false })
707            {
708                return Err(EntryPointError::MissingVertexOutputPosition.with_span());
709            }
710        } else if ep.stage == crate::ShaderStage::Vertex {
711            return Err(EntryPointError::MissingVertexOutputPosition.with_span());
712        }
713
714        {
715            let mut used_push_constants = module
716                .global_variables
717                .iter()
718                .filter(|&(_, var)| var.space == crate::AddressSpace::PushConstant)
719                .map(|(handle, _)| handle)
720                .filter(|&handle| !info[handle].is_empty());
721            // Check if there is more than one push constant, and error if so.
722            // Use a loop for when returning multiple errors is supported.
723            if let Some(handle) = used_push_constants.nth(1) {
724                return Err(EntryPointError::MoreThanOnePushConstantUsed
725                    .with_span_handle(handle, &module.global_variables));
726            }
727        }
728
729        self.ep_resource_bindings.clear();
730        for (var_handle, var) in module.global_variables.iter() {
731            let usage = info[var_handle];
732            if usage.is_empty() {
733                continue;
734            }
735
736            let allowed_usage = match var.space {
737                crate::AddressSpace::Function => unreachable!(),
738                crate::AddressSpace::Uniform => GlobalUse::READ | GlobalUse::QUERY,
739                crate::AddressSpace::Storage { access } => storage_usage(access),
740                crate::AddressSpace::Handle => match module.types[var.ty].inner {
741                    crate::TypeInner::BindingArray { base, .. } => match module.types[base].inner {
742                        crate::TypeInner::Image {
743                            class: crate::ImageClass::Storage { access, .. },
744                            ..
745                        } => storage_usage(access),
746                        _ => GlobalUse::READ | GlobalUse::QUERY,
747                    },
748                    crate::TypeInner::Image {
749                        class: crate::ImageClass::Storage { access, .. },
750                        ..
751                    } => storage_usage(access),
752                    _ => GlobalUse::READ | GlobalUse::QUERY,
753                },
754                crate::AddressSpace::Private | crate::AddressSpace::WorkGroup => GlobalUse::all(),
755                crate::AddressSpace::PushConstant => GlobalUse::READ,
756            };
757            if !allowed_usage.contains(usage) {
758                log::warn!("\tUsage error for: {:?}", var);
759                log::warn!(
760                    "\tAllowed usage: {:?}, requested: {:?}",
761                    allowed_usage,
762                    usage
763                );
764                return Err(EntryPointError::InvalidGlobalUsage(var_handle, usage)
765                    .with_span_handle(var_handle, &module.global_variables));
766            }
767
768            if let Some(ref bind) = var.binding {
769                if !self.ep_resource_bindings.insert(bind.clone()) {
770                    if self.flags.contains(super::ValidationFlags::BINDINGS) {
771                        return Err(EntryPointError::BindingCollision(var_handle)
772                            .with_span_handle(var_handle, &module.global_variables));
773                    }
774                }
775            }
776        }
777
778        Ok(info)
779    }
780}
naga/valid/interface.rs

naga/valid/
interface.rs
