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