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 if access.contains(crate::StorageAccess::ATOMIC) {
133 storage_usage |= GlobalUse::ATOMIC;
134 }
135 storage_usage
136}
137
138struct VaryingContext<'a> {
139 stage: crate::ShaderStage,
140 output: bool,
141 second_blend_source: bool,
142 types: &'a UniqueArena<crate::Type>,
143 type_info: &'a Vec<super::r#type::TypeInfo>,
144 location_mask: &'a mut BitSet,
145 built_ins: &'a mut crate::FastHashSet<crate::BuiltIn>,
146 capabilities: Capabilities,
147 flags: super::ValidationFlags,
148}
149
150impl VaryingContext<'_> {
151 fn validate_impl(
152 &mut self,
153 ep: &crate::EntryPoint,
154 ty: Handle<crate::Type>,
155 binding: &crate::Binding,
156 ) -> Result<(), VaryingError> {
157 use crate::{BuiltIn as Bi, ShaderStage as St, TypeInner as Ti, VectorSize as Vs};
158
159 let ty_inner = &self.types[ty].inner;
160 match *binding {
161 crate::Binding::BuiltIn(built_in) => {
162 let canonical = if let crate::BuiltIn::Position { .. } = built_in {
165 crate::BuiltIn::Position { invariant: false }
166 } else {
167 built_in
168 };
169
170 if self.built_ins.contains(&canonical) {
171 return Err(VaryingError::DuplicateBuiltIn(built_in));
172 }
173 self.built_ins.insert(canonical);
174
175 let required = match built_in {
176 Bi::ClipDistance => Capabilities::CLIP_DISTANCE,
177 Bi::CullDistance => Capabilities::CULL_DISTANCE,
178 Bi::PrimitiveIndex => Capabilities::PRIMITIVE_INDEX,
179 Bi::ViewIndex => Capabilities::MULTIVIEW,
180 Bi::SampleIndex => Capabilities::MULTISAMPLED_SHADING,
181 Bi::NumSubgroups
182 | Bi::SubgroupId
183 | Bi::SubgroupSize
184 | Bi::SubgroupInvocationId => Capabilities::SUBGROUP,
185 _ => Capabilities::empty(),
186 };
187 if !self.capabilities.contains(required) {
188 return Err(VaryingError::UnsupportedCapability(required));
189 }
190
191 if matches!(
192 built_in,
193 crate::BuiltIn::SubgroupId | crate::BuiltIn::SubgroupInvocationId
194 ) && ep.workgroup_size[1..].iter().any(|&s| s > 1)
195 {
196 return Err(VaryingError::InvalidMultiDimensionalSubgroupBuiltIn);
197 }
198
199 let (visible, type_good) = match built_in {
200 Bi::BaseInstance
201 | Bi::BaseVertex
202 | Bi::InstanceIndex
203 | Bi::VertexIndex
204 | Bi::DrawID => (
205 self.stage == St::Vertex && !self.output,
206 *ty_inner == Ti::Scalar(crate::Scalar::U32),
207 ),
208 Bi::ClipDistance | Bi::CullDistance => (
209 self.stage == St::Vertex && self.output,
210 match *ty_inner {
211 Ti::Array { base, .. } => {
212 self.types[base].inner == Ti::Scalar(crate::Scalar::F32)
213 }
214 _ => false,
215 },
216 ),
217 Bi::PointSize => (
218 self.stage == St::Vertex && self.output,
219 *ty_inner == Ti::Scalar(crate::Scalar::F32),
220 ),
221 Bi::PointCoord => (
222 self.stage == St::Fragment && !self.output,
223 *ty_inner
224 == Ti::Vector {
225 size: Vs::Bi,
226 scalar: crate::Scalar::F32,
227 },
228 ),
229 Bi::Position { .. } => (
230 match self.stage {
231 St::Vertex => self.output,
232 St::Fragment => !self.output,
233 St::Compute => false,
234 },
235 *ty_inner
236 == Ti::Vector {
237 size: Vs::Quad,
238 scalar: crate::Scalar::F32,
239 },
240 ),
241 Bi::ViewIndex => (
242 match self.stage {
243 St::Vertex | St::Fragment => !self.output,
244 St::Compute => false,
245 },
246 *ty_inner == Ti::Scalar(crate::Scalar::I32),
247 ),
248 Bi::FragDepth => (
249 self.stage == St::Fragment && self.output,
250 *ty_inner == Ti::Scalar(crate::Scalar::F32),
251 ),
252 Bi::FrontFacing => (
253 self.stage == St::Fragment && !self.output,
254 *ty_inner == Ti::Scalar(crate::Scalar::BOOL),
255 ),
256 Bi::PrimitiveIndex => (
257 self.stage == St::Fragment && !self.output,
258 *ty_inner == Ti::Scalar(crate::Scalar::U32),
259 ),
260 Bi::SampleIndex => (
261 self.stage == St::Fragment && !self.output,
262 *ty_inner == Ti::Scalar(crate::Scalar::U32),
263 ),
264 Bi::SampleMask => (
265 self.stage == St::Fragment,
266 *ty_inner == Ti::Scalar(crate::Scalar::U32),
267 ),
268 Bi::LocalInvocationIndex => (
269 self.stage == St::Compute && !self.output,
270 *ty_inner == Ti::Scalar(crate::Scalar::U32),
271 ),
272 Bi::GlobalInvocationId
273 | Bi::LocalInvocationId
274 | Bi::WorkGroupId
275 | Bi::WorkGroupSize
276 | Bi::NumWorkGroups => (
277 self.stage == St::Compute && !self.output,
278 *ty_inner
279 == Ti::Vector {
280 size: Vs::Tri,
281 scalar: crate::Scalar::U32,
282 },
283 ),
284 Bi::NumSubgroups | Bi::SubgroupId => (
285 self.stage == St::Compute && !self.output,
286 *ty_inner == Ti::Scalar(crate::Scalar::U32),
287 ),
288 Bi::SubgroupSize | Bi::SubgroupInvocationId => (
289 match self.stage {
290 St::Compute | St::Fragment => !self.output,
291 St::Vertex => false,
292 },
293 *ty_inner == Ti::Scalar(crate::Scalar::U32),
294 ),
295 };
296
297 if !visible {
298 return Err(VaryingError::InvalidBuiltInStage(built_in));
299 }
300 if !type_good {
301 log::warn!("Wrong builtin type: {:?}", ty_inner);
302 return Err(VaryingError::InvalidBuiltInType(built_in));
303 }
304 }
305 crate::Binding::Location {
306 location,
307 interpolation,
308 sampling,
309 second_blend_source,
310 } => {
311 if !self.type_info[ty.index()]
313 .flags
314 .contains(super::TypeFlags::IO_SHAREABLE)
315 {
316 return Err(VaryingError::NotIOShareableType(ty));
317 }
318
319 if second_blend_source {
320 if !self
321 .capabilities
322 .contains(Capabilities::DUAL_SOURCE_BLENDING)
323 {
324 return Err(VaryingError::UnsupportedCapability(
325 Capabilities::DUAL_SOURCE_BLENDING,
326 ));
327 }
328 if self.stage != crate::ShaderStage::Fragment {
329 return Err(VaryingError::InvalidAttributeInStage(
330 "second_blend_source",
331 self.stage,
332 ));
333 }
334 if !self.output {
335 return Err(VaryingError::InvalidInputAttributeInStage(
336 "second_blend_source",
337 self.stage,
338 ));
339 }
340 if location != 0 {
341 return Err(VaryingError::InvalidLocationAttributeCombination {
342 location,
343 attribute: "second_blend_source",
344 });
345 }
346
347 self.second_blend_source = true;
348 } else if !self.location_mask.insert(location as usize) {
349 if self.flags.contains(super::ValidationFlags::BINDINGS) {
350 return Err(VaryingError::BindingCollision { location });
351 }
352 }
353
354 if let Some(interpolation) = interpolation {
355 let invalid_sampling = match (interpolation, sampling) {
356 (_, None)
357 | (
358 crate::Interpolation::Perspective | crate::Interpolation::Linear,
359 Some(
360 crate::Sampling::Center
361 | crate::Sampling::Centroid
362 | crate::Sampling::Sample,
363 ),
364 )
365 | (
366 crate::Interpolation::Flat,
367 Some(crate::Sampling::First | crate::Sampling::Either),
368 ) => None,
369 (_, Some(invalid_sampling)) => Some(invalid_sampling),
370 };
371 if let Some(sampling) = invalid_sampling {
372 return Err(VaryingError::InvalidInterpolationSamplingCombination {
373 interpolation,
374 sampling,
375 });
376 }
377 }
378
379 let needs_interpolation = match self.stage {
380 crate::ShaderStage::Vertex => self.output,
381 crate::ShaderStage::Fragment => !self.output,
382 crate::ShaderStage::Compute => false,
383 };
384
385 let _ = sampling;
389
390 let required = match sampling {
391 Some(crate::Sampling::Sample) => Capabilities::MULTISAMPLED_SHADING,
392 _ => Capabilities::empty(),
393 };
394 if !self.capabilities.contains(required) {
395 return Err(VaryingError::UnsupportedCapability(required));
396 }
397
398 match ty_inner.scalar_kind() {
399 Some(crate::ScalarKind::Float) => {
400 if needs_interpolation && interpolation.is_none() {
401 return Err(VaryingError::MissingInterpolation);
402 }
403 }
404 Some(_) => {
405 if needs_interpolation && interpolation != Some(crate::Interpolation::Flat)
406 {
407 return Err(VaryingError::InvalidInterpolation);
408 }
409 }
410 None => return Err(VaryingError::InvalidType(ty)),
411 }
412 }
413 }
414
415 Ok(())
416 }
417
418 fn validate(
419 &mut self,
420 ep: &crate::EntryPoint,
421 ty: Handle<crate::Type>,
422 binding: Option<&crate::Binding>,
423 ) -> Result<(), WithSpan<VaryingError>> {
424 let span_context = self.types.get_span_context(ty);
425 match binding {
426 Some(binding) => self
427 .validate_impl(ep, ty, binding)
428 .map_err(|e| e.with_span_context(span_context)),
429 None => {
430 match self.types[ty].inner {
431 crate::TypeInner::Struct { ref members, .. } => {
432 for (index, member) in members.iter().enumerate() {
433 let span_context = self.types.get_span_context(ty);
434 match member.binding {
435 None => {
436 if self.flags.contains(super::ValidationFlags::BINDINGS) {
437 return Err(VaryingError::MemberMissingBinding(
438 index as u32,
439 )
440 .with_span_context(span_context));
441 }
442 }
443 Some(ref binding) => self
444 .validate_impl(ep, member.ty, binding)
445 .map_err(|e| e.with_span_context(span_context))?,
446 }
447 }
448 }
449 _ => {
450 if self.flags.contains(super::ValidationFlags::BINDINGS) {
451 return Err(VaryingError::MissingBinding.with_span());
452 }
453 }
454 }
455 Ok(())
456 }
457 }
458 }
459}
460
461impl super::Validator {
462 pub(super) fn validate_global_var(
463 &self,
464 var: &crate::GlobalVariable,
465 gctx: crate::proc::GlobalCtx,
466 mod_info: &ModuleInfo,
467 global_expr_kind: &crate::proc::ExpressionKindTracker,
468 ) -> Result<(), GlobalVariableError> {
469 use super::TypeFlags;
470
471 log::debug!("var {:?}", var);
472 let inner_ty = match gctx.types[var.ty].inner {
473 crate::TypeInner::BindingArray { base, .. } => match var.space {
477 crate::AddressSpace::Storage { .. }
478 | crate::AddressSpace::Uniform
479 | crate::AddressSpace::Handle => base,
480 _ => return Err(GlobalVariableError::InvalidUsage(var.space)),
481 },
482 _ => var.ty,
483 };
484 let type_info = &self.types[inner_ty.index()];
485
486 let (required_type_flags, is_resource) = match var.space {
487 crate::AddressSpace::Function => {
488 return Err(GlobalVariableError::InvalidUsage(var.space))
489 }
490 crate::AddressSpace::Storage { access } => {
491 if let Err((ty_handle, disalignment)) = type_info.storage_layout {
492 if self.flags.contains(super::ValidationFlags::STRUCT_LAYOUTS) {
493 return Err(GlobalVariableError::Alignment(
494 var.space,
495 ty_handle,
496 disalignment,
497 ));
498 }
499 }
500 if access == crate::StorageAccess::STORE {
501 return Err(GlobalVariableError::StorageAddressSpaceWriteOnlyNotSupported);
502 }
503 (
504 TypeFlags::DATA | TypeFlags::HOST_SHAREABLE | TypeFlags::CREATION_RESOLVED,
505 true,
506 )
507 }
508 crate::AddressSpace::Uniform => {
509 if let Err((ty_handle, disalignment)) = type_info.uniform_layout {
510 if self.flags.contains(super::ValidationFlags::STRUCT_LAYOUTS) {
511 return Err(GlobalVariableError::Alignment(
512 var.space,
513 ty_handle,
514 disalignment,
515 ));
516 }
517 }
518 (
519 TypeFlags::DATA
520 | TypeFlags::COPY
521 | TypeFlags::SIZED
522 | TypeFlags::HOST_SHAREABLE
523 | TypeFlags::CREATION_RESOLVED,
524 true,
525 )
526 }
527 crate::AddressSpace::Handle => {
528 match gctx.types[inner_ty].inner {
529 crate::TypeInner::Image { class, .. } => match class {
530 crate::ImageClass::Storage {
531 format:
532 crate::StorageFormat::R16Unorm
533 | crate::StorageFormat::R16Snorm
534 | crate::StorageFormat::Rg16Unorm
535 | crate::StorageFormat::Rg16Snorm
536 | crate::StorageFormat::Rgba16Unorm
537 | crate::StorageFormat::Rgba16Snorm,
538 ..
539 } => {
540 if !self
541 .capabilities
542 .contains(Capabilities::STORAGE_TEXTURE_16BIT_NORM_FORMATS)
543 {
544 return Err(GlobalVariableError::UnsupportedCapability(
545 Capabilities::STORAGE_TEXTURE_16BIT_NORM_FORMATS,
546 ));
547 }
548 }
549 _ => {}
550 },
551 crate::TypeInner::Sampler { .. }
552 | crate::TypeInner::AccelerationStructure
553 | crate::TypeInner::RayQuery => {}
554 _ => {
555 return Err(GlobalVariableError::InvalidType(var.space));
556 }
557 }
558
559 (TypeFlags::empty(), true)
560 }
561 crate::AddressSpace::Private => (
562 TypeFlags::CONSTRUCTIBLE | TypeFlags::CREATION_RESOLVED,
563 false,
564 ),
565 crate::AddressSpace::WorkGroup => (TypeFlags::DATA | TypeFlags::SIZED, false),
566 crate::AddressSpace::PushConstant => {
567 if !self.capabilities.contains(Capabilities::PUSH_CONSTANT) {
568 return Err(GlobalVariableError::UnsupportedCapability(
569 Capabilities::PUSH_CONSTANT,
570 ));
571 }
572 (
573 TypeFlags::DATA
574 | TypeFlags::COPY
575 | TypeFlags::HOST_SHAREABLE
576 | TypeFlags::SIZED,
577 false,
578 )
579 }
580 };
581
582 if !type_info.flags.contains(required_type_flags) {
583 return Err(GlobalVariableError::MissingTypeFlags {
584 seen: type_info.flags,
585 required: required_type_flags,
586 });
587 }
588
589 if is_resource != var.binding.is_some() {
590 if self.flags.contains(super::ValidationFlags::BINDINGS) {
591 return Err(GlobalVariableError::InvalidBinding);
592 }
593 }
594
595 if let Some(init) = var.init {
596 match var.space {
597 crate::AddressSpace::Private | crate::AddressSpace::Function => {}
598 _ => {
599 return Err(GlobalVariableError::InitializerNotAllowed(var.space));
600 }
601 }
602
603 if !global_expr_kind.is_const_or_override(init) {
604 return Err(GlobalVariableError::InitializerExprType);
605 }
606
607 let decl_ty = &gctx.types[var.ty].inner;
608 let init_ty = mod_info[init].inner_with(gctx.types);
609 if !decl_ty.equivalent(init_ty, gctx.types) {
610 return Err(GlobalVariableError::InitializerType);
611 }
612 }
613
614 Ok(())
615 }
616
617 pub(super) fn validate_entry_point(
618 &mut self,
619 ep: &crate::EntryPoint,
620 module: &crate::Module,
621 mod_info: &ModuleInfo,
622 global_expr_kind: &crate::proc::ExpressionKindTracker,
623 ) -> Result<FunctionInfo, WithSpan<EntryPointError>> {
624 if ep.early_depth_test.is_some() {
625 let required = Capabilities::EARLY_DEPTH_TEST;
626 if !self.capabilities.contains(required) {
627 return Err(
628 EntryPointError::Result(VaryingError::UnsupportedCapability(required))
629 .with_span(),
630 );
631 }
632
633 if ep.stage != crate::ShaderStage::Fragment {
634 return Err(EntryPointError::UnexpectedEarlyDepthTest.with_span());
635 }
636 }
637
638 if ep.stage == crate::ShaderStage::Compute {
639 if ep
640 .workgroup_size
641 .iter()
642 .any(|&s| s == 0 || s > MAX_WORKGROUP_SIZE)
643 {
644 return Err(EntryPointError::OutOfRangeWorkgroupSize.with_span());
645 }
646 } else if ep.workgroup_size != [0; 3] {
647 return Err(EntryPointError::UnexpectedWorkgroupSize.with_span());
648 }
649
650 let mut info = self
651 .validate_function(&ep.function, module, mod_info, true, global_expr_kind)
652 .map_err(WithSpan::into_other)?;
653
654 {
655 use super::ShaderStages;
656
657 let stage_bit = match ep.stage {
658 crate::ShaderStage::Vertex => ShaderStages::VERTEX,
659 crate::ShaderStage::Fragment => ShaderStages::FRAGMENT,
660 crate::ShaderStage::Compute => ShaderStages::COMPUTE,
661 };
662
663 if !info.available_stages.contains(stage_bit) {
664 return Err(EntryPointError::ForbiddenStageOperations.with_span());
665 }
666 }
667
668 self.location_mask.clear();
669 let mut argument_built_ins = crate::FastHashSet::default();
670 for (index, fa) in ep.function.arguments.iter().enumerate() {
672 let mut ctx = VaryingContext {
673 stage: ep.stage,
674 output: false,
675 second_blend_source: false,
676 types: &module.types,
677 type_info: &self.types,
678 location_mask: &mut self.location_mask,
679 built_ins: &mut argument_built_ins,
680 capabilities: self.capabilities,
681 flags: self.flags,
682 };
683 ctx.validate(ep, fa.ty, fa.binding.as_ref())
684 .map_err_inner(|e| EntryPointError::Argument(index as u32, e).with_span())?;
685 }
686
687 self.location_mask.clear();
688 if let Some(ref fr) = ep.function.result {
689 let mut result_built_ins = crate::FastHashSet::default();
690 let mut ctx = VaryingContext {
691 stage: ep.stage,
692 output: true,
693 second_blend_source: false,
694 types: &module.types,
695 type_info: &self.types,
696 location_mask: &mut self.location_mask,
697 built_ins: &mut result_built_ins,
698 capabilities: self.capabilities,
699 flags: self.flags,
700 };
701 ctx.validate(ep, fr.ty, fr.binding.as_ref())
702 .map_err_inner(|e| EntryPointError::Result(e).with_span())?;
703 if ctx.second_blend_source {
704 if ctx.location_mask.len() == 1 && ctx.location_mask.contains(0) {
706 info.dual_source_blending = true;
707 } else {
708 return Err(EntryPointError::InvalidLocationsWhileDualSourceBlending {
709 location_mask: self.location_mask.clone(),
710 }
711 .with_span());
712 }
713 }
714
715 if ep.stage == crate::ShaderStage::Vertex
716 && !result_built_ins.contains(&crate::BuiltIn::Position { invariant: false })
717 {
718 return Err(EntryPointError::MissingVertexOutputPosition.with_span());
719 }
720 } else if ep.stage == crate::ShaderStage::Vertex {
721 return Err(EntryPointError::MissingVertexOutputPosition.with_span());
722 }
723
724 {
725 let mut used_push_constants = module
726 .global_variables
727 .iter()
728 .filter(|&(_, var)| var.space == crate::AddressSpace::PushConstant)
729 .map(|(handle, _)| handle)
730 .filter(|&handle| !info[handle].is_empty());
731 if let Some(handle) = used_push_constants.nth(1) {
734 return Err(EntryPointError::MoreThanOnePushConstantUsed
735 .with_span_handle(handle, &module.global_variables));
736 }
737 }
738
739 self.ep_resource_bindings.clear();
740 for (var_handle, var) in module.global_variables.iter() {
741 let usage = info[var_handle];
742 if usage.is_empty() {
743 continue;
744 }
745
746 let allowed_usage = match var.space {
747 crate::AddressSpace::Function => unreachable!(),
748 crate::AddressSpace::Uniform => GlobalUse::READ | GlobalUse::QUERY,
749 crate::AddressSpace::Storage { access } => storage_usage(access),
750 crate::AddressSpace::Handle => match module.types[var.ty].inner {
751 crate::TypeInner::BindingArray { base, .. } => match module.types[base].inner {
752 crate::TypeInner::Image {
753 class: crate::ImageClass::Storage { access, .. },
754 ..
755 } => storage_usage(access),
756 _ => GlobalUse::READ | GlobalUse::QUERY,
757 },
758 crate::TypeInner::Image {
759 class: crate::ImageClass::Storage { access, .. },
760 ..
761 } => storage_usage(access),
762 _ => GlobalUse::READ | GlobalUse::QUERY,
763 },
764 crate::AddressSpace::Private | crate::AddressSpace::WorkGroup => {
765 GlobalUse::READ | GlobalUse::WRITE | GlobalUse::QUERY
766 }
767 crate::AddressSpace::PushConstant => GlobalUse::READ,
768 };
769 if !allowed_usage.contains(usage) {
770 log::warn!("\tUsage error for: {:?}", var);
771 log::warn!(
772 "\tAllowed usage: {:?}, requested: {:?}",
773 allowed_usage,
774 usage
775 );
776 return Err(EntryPointError::InvalidGlobalUsage(var_handle, usage)
777 .with_span_handle(var_handle, &module.global_variables));
778 }
779
780 if let Some(ref bind) = var.binding {
781 if !self.ep_resource_bindings.insert(bind.clone()) {
782 if self.flags.contains(super::ValidationFlags::BINDINGS) {
783 return Err(EntryPointError::BindingCollision(var_handle)
784 .with_span_handle(var_handle, &module.global_variables));
785 }
786 }
787 }
788 }
789
790 Ok(info)
791 }
792}