From 4666a9c722b8eb5d41674493d8cee16979bbdc58 Mon Sep 17 00:00:00 2001 From: Rua Date: Sun, 3 Mar 2024 23:29:54 +0100 Subject: [PATCH] Add runtime SPIR-V validation (#2460) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * Add runtime SPIR-V validation * Remove copyright message * fmt * Up vk version in shader compilation (#2467) * Up vk version in shader compilation * Update test in codegen.rs * cargo +nightly fmt * Update lib.rs * clippy + fmt fixes * simplify changes * post merge * #2467 changelog * Remove windows specific dnd disable (#2474) * Fix unnecessarily strict validation for DRM format modifiers (#2469) * #2469 changelog * Add support for querying memory requirements directly from the device (#2470) * #2470 changelog * Make image_index and final_views accessible, and add new example. (#2473) * Make image_index and final_views accessible, and new example. The first 2 changes should make creating frame buffers easier. The new example should make it easier to learn vulkano-util. * Remove unnecessary imports, and run clippy. * Run fmt. * .acquire() no longer returns image_index * rename final_views() to swapchain_image_views() The name change makes it more consistent with swapchain_image_view(). Personally I don't understand why the field name is final_views, yet we externally in function names refer to it as swapchain image views and such like. * Fractal example no longer creates framebuffer every frame. * Game of life example no longer creates framebuffer every frame. (Also removed a piece of code I had commented out, but had forgotten to remove from the fractal example.) * Rename if_recreate_swapchain to on_recreate_swapchain and update acquire() documentation. to on_recreate_swapchain * on_recreate_swapchain is now impl FnOnce instead of generics based FnMut Thanks marc0246! Co-authored-by: marc0246 <40955683+marc0246@users.noreply.github.com> * Replace empty comment with an actual comment. --------- Co-authored-by: marc0246 <40955683+marc0246@users.noreply.github.com> * Fix `VulkanoWindowRenderer::swapchain_image_views` return type I have only noticed this as I was writing the changelog. @coolcatcoder for future reference, `&Vec` is an anti-pattern. There's nothing more you can do with it than with `&[T]` (because the reference is immutable) and it means that we can't use a different underlying buffer without a breaking change. * #2473 changelog * Fix `VulkanoWindowRenderer::acquire` taking `&Vec` as well * Replace cgmath with glam in the examples (#2475) * Replace cgmath with glam in the examples * Implement type_for_format! for glam * Remove comment where I'm freaking out because of OpenGL flashbacks * Update Cargo.toml Co-authored-by: marc0246 <40955683+marc0246@users.noreply.github.com> * Update vulkano/autogen/formats.rs Co-authored-by: marc0246 <40955683+marc0246@users.noreply.github.com> * Fix glam type_for_format * Format the code --------- Co-authored-by: marc0246 <40955683+marc0246@users.noreply.github.com> * #2475 changelog * Fix alignment checks when allocating buffers (#2476) * #2476 changelog * Add `DepthState::reverse` helper method (#2483) * #2483 changelog * Add runtime SPIR-V validation * Remove copyright message * fmt --------- Co-authored-by: maratik123 Co-authored-by: Okko Hakola Co-authored-by: marc0246 <40955683+marc0246@users.noreply.github.com> Co-authored-by: Katt <51190960+coolcatcoder@users.noreply.github.com> Co-authored-by: stefnotch Co-authored-by: José Miguel Sánchez García --- vulkano/autogen/spirv_parse.rs | 326 +- .../src/pipeline/shader/inout_interface.rs | 324 +- vulkano/src/pipeline/shader/mod.rs | 4 + .../src/pipeline/shader/validate_runtime.rs | 2925 +++++++++++++++++ vulkano/src/shader/mod.rs | 203 +- vulkano/src/shader/reflect.rs | 91 + 6 files changed, 3677 insertions(+), 196 deletions(-) create mode 100644 vulkano/src/pipeline/shader/validate_runtime.rs diff --git a/vulkano/autogen/spirv_parse.rs b/vulkano/autogen/spirv_parse.rs index cecbda24..6e8afe6c 100644 --- a/vulkano/autogen/spirv_parse.rs +++ b/vulkano/autogen/spirv_parse.rs @@ -4,6 +4,7 @@ use heck::ToSnakeCase; use once_cell::sync::Lazy; use proc_macro2::{Ident, TokenStream}; use quote::{format_ident, quote}; +use std::borrow::Cow; // From the documentation of the OpSpecConstantOp instruction. // The instructions requiring the Kernel capability are not listed, @@ -88,8 +89,19 @@ pub fn write(grammar: &SpirvGrammar) { #[derive(Clone, Debug)] struct InstructionMember { name: Ident, + is_atomic_operation: bool, + is_cooperative_matrix: bool, + is_cooperative_matrix_nv: bool, + is_group_operation: bool, + is_quad_group_operation: bool, + is_image_gather: bool, + is_image_fetch: bool, + is_image_sample: bool, has_result_id: bool, has_result_type_id: bool, + has_execution_scope_id: bool, + has_memory_scope_id: bool, + has_image_operands: Option, opcode: u16, operands: Vec, } @@ -187,6 +199,153 @@ fn instruction_output(members: &[InstructionMember], spec_constant: bool) -> Tok } }, ); + let is_cooperative_matrix_items = members.iter().filter_map( + |InstructionMember { + name, + is_cooperative_matrix, + .. + }| { + if *is_cooperative_matrix { + Some(quote! { Self::#name { .. } }) + } else { + None + } + }, + ); + let is_cooperative_matrix_nv_items = members.iter().filter_map( + |InstructionMember { + name, + is_cooperative_matrix_nv, + .. + }| { + if *is_cooperative_matrix_nv { + Some(quote! { Self::#name { .. } }) + } else { + None + } + }, + ); + let is_group_operation_items = members.iter().filter_map( + |InstructionMember { + name, + is_group_operation, + .. + }| { + if *is_group_operation { + Some(quote! { Self::#name { .. } }) + } else { + None + } + }, + ); + let is_quad_group_operation_items = members.iter().filter_map( + |InstructionMember { + name, + is_quad_group_operation, + .. + }| { + if *is_quad_group_operation { + Some(quote! { Self::#name { .. } }) + } else { + None + } + }, + ); + let is_image_fetch_items = members.iter().filter_map( + |InstructionMember { + name, + is_image_fetch, + .. + }| { + if *is_image_fetch { + Some(quote! { Self::#name { .. } }) + } else { + None + } + }, + ); + let is_image_gather_items = members.iter().filter_map( + |InstructionMember { + name, + is_image_gather, + .. + }| { + if *is_image_gather { + Some(quote! { Self::#name { .. } }) + } else { + None + } + }, + ); + let is_image_sample_items = members.iter().filter_map( + |InstructionMember { + name, + is_image_sample, + .. + }| { + if *is_image_sample { + Some(quote! { Self::#name { .. } }) + } else { + None + } + }, + ); + let atomic_pointer_id_items = members.iter().filter_map( + |InstructionMember { + name, + is_atomic_operation, + .. + }| { + if *is_atomic_operation { + Some(quote! { Self::#name { pointer, .. } }) + } else { + None + } + }, + ); + let execution_scope_id_items = members.iter().filter_map( + |InstructionMember { + name, + has_execution_scope_id, + .. + }| { + if *has_execution_scope_id { + Some(quote! { Self::#name { execution, .. } }) + } else { + None + } + }, + ); + let memory_scope_id_items = members.iter().filter_map( + |InstructionMember { + name, + has_memory_scope_id, + .. + }| { + if *has_memory_scope_id { + Some(quote! { Self::#name { memory, .. } }) + } else { + None + } + }, + ); + let image_operands_items = members.iter().filter_map( + |InstructionMember { + name, + has_image_operands, + .. + }| { + if let Some(has_image_operands) = *has_image_operands { + if has_image_operands { + Some(quote! { Self::#name { image_operands: Some(image_operands), .. } }) + } else { + Some(quote! { Self::#name { image_operands, .. } }) + } + } else { + None + } + }, + ); quote! { /// Returns the `Id` that is assigned by this instruction, if any. @@ -204,6 +363,94 @@ fn instruction_output(members: &[InstructionMember], spec_constant: bool) -> Tok _ => None } } + + /// Returns the `Id` of the pointer in an atomic operation, if any. + pub fn atomic_pointer_id(&self) -> Option { + match self { + #(#atomic_pointer_id_items)|* => Some(*pointer), + _ => None + } + } + + /// Returns whether the instruction is a cooperative matrix instruction. + pub fn is_cooperative_matrix(&self) -> bool { + matches!( + self, + #(#is_cooperative_matrix_items)|* + ) + } + + /// Returns whether the instruction is an NV cooperative matrix instruction. + pub fn is_cooperative_matrix_nv(&self) -> bool { + matches!( + self, + #(#is_cooperative_matrix_nv_items)|* + ) + } + + /// Returns whether the instruction is a group operation instruction. + pub fn is_group_operation(&self) -> bool { + matches!( + self, + #(#is_group_operation_items)|* + ) + } + + /// Returns whether the instruction is a quad group operation instruction. + pub fn is_quad_group_operation(&self) -> bool { + matches!( + self, + #(#is_quad_group_operation_items)|* + ) + } + + /// Returns whether the instruction is an `ImageFetch*` instruction. + pub fn is_image_fetch(&self) -> bool { + matches!( + self, + #(#is_image_fetch_items)|* + ) + } + + /// Returns whether the instruction is an `Image*Gather` instruction. + pub fn is_image_gather(&self) -> bool { + matches!( + self, + #(#is_image_gather_items)|* + ) + } + + /// Returns whether the instruction is an `ImageSample*` instruction. + pub fn is_image_sample(&self) -> bool { + matches!( + self, + #(#is_image_sample_items)|* + ) + } + + /// Returns the `Id` of the execution scope ID operand, if any. + pub fn execution_scope_id(&self) -> Option { + match self { + #(#execution_scope_id_items)|* => Some(*execution), + _ => None + } + } + + /// Returns the `Id` of the memory scope ID operand, if any. + pub fn memory_scope_id(&self) -> Option { + match self { + #(#memory_scope_id_items)|* => Some(*memory), + _ => None + } + } + + /// Returns the image operands, if any. + pub fn image_operands(&self) -> Option<&ImageOperands> { + match self { + #(#image_operands_items)|* => Some(image_operands), + _ => None + } + } } }; @@ -242,9 +489,23 @@ fn instruction_members(grammar: &SpirvGrammar) -> Vec { .instructions .iter() .map(|instruction| { - let name = format_ident!("{}", instruction.opname.strip_prefix("Op").unwrap()); + let name = instruction.opname.strip_prefix("Op").unwrap(); + let is_atomic_operation = instruction.class == "Atomic"; + let is_cooperative_matrix = + name.starts_with("CooperativeMatrix") && !name.ends_with("NV"); + let is_cooperative_matrix_nv = + name.starts_with("CooperativeMatrix") && name.ends_with("NV"); + let is_group_operation = + instruction.class == "Group" || instruction.class == "Non-Uniform"; + let is_quad_group_operation = is_group_operation && instruction.opname.contains("Quad"); + let is_image_fetch = name.starts_with("ImageFetch"); + let is_image_gather = name.starts_with("Image") && name.ends_with("Gather"); + let is_image_sample = name.starts_with("ImageSample"); let mut has_result_id = false; let mut has_result_type_id = false; + let mut has_execution_scope_id = false; + let mut has_memory_scope_id = false; + let mut has_image_operands = None; let mut operand_names = HashMap::default(); let mut operands = instruction @@ -258,7 +519,23 @@ fn instruction_members(grammar: &SpirvGrammar) -> Vec { has_result_type_id = true; format_ident!("result_type_id") } else { - to_member_name(&operand.kind, operand.name.as_deref()) + let member_name = to_member_name(&operand.kind, operand.name.as_deref()); + + if operand.kind == "IdScope" { + if member_name == "execution" { + has_execution_scope_id = true; + } else if member_name == "memory" { + has_memory_scope_id = true; + } + } else if operand.kind == "ImageOperands" { + if operand.quantifier == Some('?') { + has_image_operands = Some(true); + } else { + has_image_operands = Some(false); + } + } + + format_ident!("{}", member_name) }; *operand_names.entry(name.clone()).or_insert(0) += 1; @@ -305,9 +582,20 @@ fn instruction_members(grammar: &SpirvGrammar) -> Vec { } InstructionMember { - name, + name: format_ident!("{}", name), + is_atomic_operation, + is_cooperative_matrix, + is_cooperative_matrix_nv, + is_group_operation, + is_quad_group_operation, + is_image_fetch, + is_image_gather, + is_image_sample, has_result_id, has_result_type_id, + has_execution_scope_id, + has_memory_scope_id, + has_image_operands, opcode: instruction.opcode, operands, } @@ -465,7 +753,10 @@ fn bit_enum_members(grammar: &SpirvGrammar) -> Vec<(Ident, Vec)> .parameters .iter() .map(|param| { - let name = to_member_name(¶m.kind, param.name.as_deref()); + let name = format_ident!( + "{}", + to_member_name(¶m.kind, param.name.as_deref()) + ); let (ty, parse) = parameter_kinds[param.kind.as_str()].clone(); OperandMember { name, ty, parse } @@ -614,7 +905,10 @@ fn value_enum_members(grammar: &SpirvGrammar) -> Vec<(Ident, Vec .parameters .iter() .map(|param| { - let name = to_member_name(¶m.kind, param.name.as_deref()); + let name = format_ident!( + "{}", + to_member_name(¶m.kind, param.name.as_deref()) + ); let (ty, parse) = parameter_kinds[param.kind.as_str()].clone(); OperandMember { name, ty, parse } @@ -634,24 +928,24 @@ fn value_enum_members(grammar: &SpirvGrammar) -> Vec<(Ident, Vec .collect() } -fn to_member_name(kind: &str, name: Option<&str>) -> Ident { +fn to_member_name(kind: &str, name: Option<&str>) -> Cow<'static, str> { if let Some(name) = name { let name = name.to_snake_case(); // Fix some weird names match name.as_str() { - "argument_0_argument_1" => format_ident!("arguments"), - "member_0_type_member_1_type" => format_ident!("member_types"), - "operand_1_operand_2" => format_ident!("operands"), - "parameter_0_type_parameter_1_type" => format_ident!("parameter_types"), - "the_name_of_the_opaque_type" => format_ident!("name"), - "d_ref" => format_ident!("dref"), - "type" => format_ident!("ty"), // type is a keyword - "use" => format_ident!("usage"), // use is a keyword - _ => format_ident!("{}", name.replace("operand_", "operand")), + "argument_0_argument_1" => "arguments".into(), + "member_0_type_member_1_type" => "member_types".into(), + "operand_1_operand_2" => "operands".into(), + "parameter_0_type_parameter_1_type" => "parameter_types".into(), + "the_name_of_the_opaque_type" => "name".into(), + "d_ref" => "dref".into(), + "type" => "ty".into(), // type is a keyword + "use" => "usage".into(), // use is a keyword + _ => name.replace("operand_", "operand").into(), } } else { - format_ident!("{}", kind.to_snake_case()) + kind.to_snake_case().into() } } diff --git a/vulkano/src/pipeline/shader/inout_interface.rs b/vulkano/src/pipeline/shader/inout_interface.rs index 10ab4ade..a0075e7b 100644 --- a/vulkano/src/pipeline/shader/inout_interface.rs +++ b/vulkano/src/pipeline/shader/inout_interface.rs @@ -209,11 +209,11 @@ fn get_variables_by_key<'a>( variable_id, filter_storage_class, |key, data| { - if let InputOutputKey::Location { + if let InputOutputKey::User(InputOutputUserKey { location, component, .. - } = key + }) = key { let InputOutputData { variable_id, @@ -646,6 +646,15 @@ pub(crate) enum ShaderInterfaceLocationWidth { Bits64, } +impl ShaderInterfaceLocationWidth { + pub(crate) fn component_count(self) -> u32 { + match self { + ShaderInterfaceLocationWidth::Bits32 => 1, + ShaderInterfaceLocationWidth::Bits64 => 2, + } + } +} + impl From for ShaderInterfaceLocationWidth { #[inline] fn from(value: u32) -> Self { @@ -662,138 +671,6 @@ pub(crate) fn shader_interface_location_info( entry_point_id: Id, filter_storage_class: StorageClass, ) -> HashMap { - fn add_type( - locations: &mut HashMap, - spirv: &Spirv, - mut location: u32, - mut component: u32, - index: u32, - type_id: Id, - ) -> (u32, u32) { - debug_assert!(component < 4); - - let mut add_scalar = |numeric_type: NumericType, width: u32| -> (u32, u32) { - let width = ShaderInterfaceLocationWidth::from(width); - let components_to_add = match width { - ShaderInterfaceLocationWidth::Bits32 => { - ColorComponents::from_index(component as usize) - } - ShaderInterfaceLocationWidth::Bits64 => { - debug_assert!(component & 1 == 0); - ColorComponents::from_index(component as usize) - | ColorComponents::from_index(component as usize + 1) - } - }; - - let location_info = match locations.entry(location) { - Entry::Occupied(entry) => { - let location_info = entry.into_mut(); - debug_assert_eq!(location_info.numeric_type, numeric_type); - debug_assert_eq!(location_info.width, width); - location_info - } - Entry::Vacant(entry) => entry.insert(ShaderInterfaceLocationInfo { - numeric_type, - width, - components: [ColorComponents::empty(); 2], - }), - }; - - let components = &mut location_info.components[index as usize]; - debug_assert!(!components.intersects(components_to_add)); - *components |= components_to_add; - - (components_to_add.count(), 1) - }; - - match *spirv.id(type_id).instruction() { - Instruction::TypeInt { - width, signedness, .. - } => { - let numeric_type = if signedness == 1 { - NumericType::Int - } else { - NumericType::Uint - }; - - add_scalar(numeric_type, width) - } - Instruction::TypeFloat { width, .. } => add_scalar(NumericType::Float, width), - Instruction::TypeVector { - component_type, - component_count, - .. - } => { - let mut total_locations_added = 1; - - for _ in 0..component_count { - // Overflow into next location - if component == 4 { - component = 0; - location += 1; - total_locations_added += 1; - } else { - debug_assert!(component < 4); - } - - let (_, components_added) = - add_type(locations, spirv, location, component, index, component_type); - component += components_added; - } - - (total_locations_added, 0) - } - Instruction::TypeMatrix { - column_type, - column_count, - .. - } => { - let mut total_locations_added = 0; - - for _ in 0..column_count { - let (locations_added, _) = - add_type(locations, spirv, location, component, index, column_type); - location += locations_added; - total_locations_added += locations_added; - } - - (total_locations_added, 0) - } - Instruction::TypeArray { - element_type, - length, - .. - } => { - let length = get_constant(spirv, length).unwrap(); - let mut total_locations_added = 0; - - for _ in 0..length { - let (locations_added, _) = - add_type(locations, spirv, location, component, index, element_type); - location += locations_added; - total_locations_added += locations_added; - } - - (total_locations_added, 0) - } - Instruction::TypeStruct { - ref member_types, .. - } => { - let mut total_locations_added = 0; - - for &member_type in member_types { - let (locations_added, _) = - add_type(locations, spirv, location, component, index, member_type); - location += locations_added; - total_locations_added += locations_added; - } - - (total_locations_added, 0) - } - _ => unimplemented!(), - } - } - let (execution_model, interface) = match spirv.function(entry_point_id).entry_point() { Some(&Instruction::EntryPoint { execution_model, @@ -803,7 +680,42 @@ pub(crate) fn shader_interface_location_info( _ => unreachable!(), }; - let mut locations = HashMap::default(); + let mut locations: HashMap = HashMap::default(); + let mut scalar_func = |key: InputOutputUserKey, + width: ShaderInterfaceLocationWidth, + numeric_type: NumericType| { + let InputOutputUserKey { + location, + component, + index, + } = key; + + let location_info = match locations.entry(location) { + Entry::Occupied(entry) => { + let location_info = entry.into_mut(); + debug_assert_eq!(location_info.numeric_type, numeric_type); + debug_assert_eq!(location_info.width, width); + location_info + } + Entry::Vacant(entry) => entry.insert(ShaderInterfaceLocationInfo { + numeric_type, + width, + components: [ColorComponents::empty(); 2], + }), + }; + let components = &mut location_info.components[index as usize]; + + let components_to_add = match width { + ShaderInterfaceLocationWidth::Bits32 => ColorComponents::from_index(component as usize), + ShaderInterfaceLocationWidth::Bits64 => { + debug_assert!(component & 1 == 0); + ColorComponents::from_index(component as usize) + | ColorComponents::from_index(component as usize + 1) + } + }; + debug_assert!(!components.intersects(components_to_add)); + *components |= components_to_add; + }; for &variable_id in interface { input_output_map( @@ -812,14 +724,9 @@ pub(crate) fn shader_interface_location_info( variable_id, filter_storage_class, |key, data| { - if let InputOutputKey::Location { - location, - component, - index, - } = key - { + if let InputOutputKey::User(key) = key { let InputOutputData { type_id, .. } = data; - add_type(&mut locations, spirv, location, component, index, type_id); + shader_interface_analyze_type(spirv, type_id, key, &mut scalar_func); } }, ); @@ -828,16 +735,121 @@ pub(crate) fn shader_interface_location_info( locations } +/// Recursively analyzes the type `type_id` with the given `key`. Calls `scalar_func` on every +/// scalar type that is encountered, and returns the number of locations and components to advance. +pub(crate) fn shader_interface_analyze_type( + spirv: &Spirv, + type_id: Id, + mut key: InputOutputUserKey, + scalar_func: &mut impl FnMut(InputOutputUserKey, ShaderInterfaceLocationWidth, NumericType), +) -> (u32, u32) { + debug_assert!(key.component < 4); + + match *spirv.id(type_id).instruction() { + Instruction::TypeInt { + width, signedness, .. + } => { + let numeric_type = if signedness == 1 { + NumericType::Int + } else { + NumericType::Uint + }; + + let width = ShaderInterfaceLocationWidth::from(width); + scalar_func(key, width, numeric_type); + (1, width.component_count()) + } + Instruction::TypeFloat { width, .. } => { + let width = ShaderInterfaceLocationWidth::from(width); + scalar_func(key, width, NumericType::Float); + (1, width.component_count()) + } + Instruction::TypeVector { + component_type, + component_count, + .. + } => { + let mut total_locations_added = 1; + + for _ in 0..component_count { + // Overflow into next location + if key.component == 4 { + key.component = 0; + key.location += 1; + total_locations_added += 1; + } else { + debug_assert!(key.component < 4); + } + + let (_, components_added) = + shader_interface_analyze_type(spirv, component_type, key, scalar_func); + key.component += components_added; + } + + (total_locations_added, 0) + } + Instruction::TypeMatrix { + column_type, + column_count, + .. + } => { + let mut total_locations_added = 0; + + for _ in 0..column_count { + let (locations_added, _) = + shader_interface_analyze_type(spirv, column_type, key, scalar_func); + key.location += locations_added; + total_locations_added += locations_added; + } + + (total_locations_added, 0) + } + Instruction::TypeArray { + element_type, + length, + .. + } => { + let length = get_constant(spirv, length).unwrap(); + let mut total_locations_added = 0; + + for _ in 0..length { + let (locations_added, _) = + shader_interface_analyze_type(spirv, element_type, key, scalar_func); + key.location += locations_added; + total_locations_added += locations_added; + } + + (total_locations_added, 0) + } + Instruction::TypeStruct { + ref member_types, .. + } => { + let mut total_locations_added = 0; + + for &member_type in member_types { + let (locations_added, _) = + shader_interface_analyze_type(spirv, member_type, key, scalar_func); + key.location += locations_added; + total_locations_added += locations_added; + } + + (total_locations_added, 0) + } + _ => unimplemented!(), + } +} + #[derive(Clone, Copy, Debug, PartialEq, Eq, Hash)] pub(crate) enum InputOutputKey { - Location { - location: u32, - component: u32, - index: u32, - }, - BuiltIn { - built_in: BuiltIn, - }, + User(InputOutputUserKey), + BuiltIn(BuiltIn), +} + +#[derive(Clone, Copy, Debug, Default, PartialEq, Eq, Hash)] +pub(crate) struct InputOutputUserKey { + pub(crate) location: u32, + pub(crate) component: u32, + pub(crate) index: u32, } pub(crate) struct InputOutputData { @@ -895,11 +907,11 @@ pub(crate) fn input_output_map( if let Some(location) = location { func( - InputOutputKey::Location { + InputOutputKey::User(InputOutputUserKey { location, component, index, - }, + }), InputOutputData { variable_id, pointer_type_id, @@ -909,7 +921,7 @@ pub(crate) fn input_output_map( ); } else if let Some(built_in) = built_in { func( - InputOutputKey::BuiltIn { built_in }, + InputOutputKey::BuiltIn(built_in), InputOutputData { variable_id, pointer_type_id, @@ -949,11 +961,11 @@ pub(crate) fn input_output_map( if let Some(location) = location { func( - InputOutputKey::Location { + InputOutputKey::User(InputOutputUserKey { location, component, index, - }, + }), InputOutputData { variable_id, pointer_type_id, @@ -966,7 +978,7 @@ pub(crate) fn input_output_map( ); } else if let Some(built_in) = built_in { func( - InputOutputKey::BuiltIn { built_in }, + InputOutputKey::BuiltIn(built_in), InputOutputData { variable_id, pointer_type_id, diff --git a/vulkano/src/pipeline/shader/mod.rs b/vulkano/src/pipeline/shader/mod.rs index 5a55f4be..dc85cd1d 100644 --- a/vulkano/src/pipeline/shader/mod.rs +++ b/vulkano/src/pipeline/shader/mod.rs @@ -9,6 +9,7 @@ use crate::{ }; pub(crate) mod inout_interface; +pub(crate) mod validate_runtime; /// Specifies a single shader stage when creating a pipeline. #[derive(Clone, Debug)] @@ -62,6 +63,9 @@ impl PipelineShaderStageCreateInfo { } = self; let spirv = entry_point.module().spirv(); + validate_runtime::validate_runtime(device, spirv, entry_point.id()) + .map_err(|err| err.add_context("entry_point"))?; + let properties = device.physical_device().properties(); flags.validate_device(device).map_err(|err| { diff --git a/vulkano/src/pipeline/shader/validate_runtime.rs b/vulkano/src/pipeline/shader/validate_runtime.rs new file mode 100644 index 00000000..7d915d34 --- /dev/null +++ b/vulkano/src/pipeline/shader/validate_runtime.rs @@ -0,0 +1,2925 @@ +use crate::{ + descriptor_set::layout::DescriptorType, + device::{physical::ShaderFloatControlsIndependence, Device, Features}, + pipeline::inout_interface::{ + input_output_map, shader_interface_analyze_type, InputOutputData, InputOutputKey, + }, + shader::{ + reflect::{ + get_constant, get_constant_composite, get_constant_composite_composite, + get_constant_float_composite, get_constant_maybe_composite, size_of_type, + }, + spirv::{ + Capability, Decoration, Dim, ExecutionMode, ExecutionModel, FunctionInfo, Id, + ImageFormat, Instruction, Scope, Spirv, StorageClass, + }, + ShaderStage, + }, + DeviceSize, Requires, RequiresAllOf, RequiresOneOf, ValidationError, Version, +}; +use ahash::HashMap; +use std::cmp::max; + +pub(crate) fn validate_runtime( + device: &Device, + spirv: &Spirv, + entry_point: Id, +) -> Result<(), Box> { + let entry_point_info = spirv.function(entry_point); + let Instruction::EntryPoint { + execution_model, + ref interface, + .. + } = *entry_point_info.entry_point().unwrap() + else { + unreachable!() + }; + + let mut validator = RuntimeValidator { + device, + spirv, + entry_point_info, + execution_model, + interface, + + first_emitted_stream: None, + output_primitives: None, + }; + + // Ordering is important + validator.validate_capabilities()?; + validator.validate_decorations()?; + validator.validate_execution_modes()?; + validator.validate_types()?; + validator.validate_global_variables()?; + validator.validate_functions()?; + + Ok(()) +} + +struct RuntimeValidator<'a> { + device: &'a Device, + spirv: &'a Spirv, + entry_point_info: &'a FunctionInfo, + execution_model: ExecutionModel, + interface: &'a [Id], + + first_emitted_stream: Option, + output_primitives: Option<&'a ExecutionMode>, +} + +impl<'a> RuntimeValidator<'a> { + fn validate_capabilities(&self) -> Result<(), Box> { + for instruction in self.spirv.capabilities() { + let capability = match *instruction { + Instruction::Capability { capability } => capability, + _ => continue, + }; + + #[allow(clippy::single_match)] + match capability { + Capability::InterpolationFunction => { + if self.device.enabled_extensions().khr_portability_subset + && !self + .device + .enabled_features() + .shader_sample_rate_interpolation_functions + { + return Err(Box::new(ValidationError { + problem: "this device is a portability subset device, and the shader \ + uses the `InterpolationFunction` capability" + .into(), + vuids: &[ + "VUID-RuntimeSpirv-shaderSampleRateInterpolationFunctions-06325", + ], + ..Default::default() + })); + } + } + _ => (), + } + } + + Ok(()) + } + + fn validate_decorations(&self) -> Result<(), Box> { + let properties = self.device.physical_device().properties(); + + for instruction in self.spirv.decorations() { + let decoration = match instruction { + Instruction::Decorate { decoration, .. } + | Instruction::DecorateId { decoration, .. } + | Instruction::DecorateString { decoration, .. } + | Instruction::MemberDecorate { decoration, .. } + | Instruction::MemberDecorateString { decoration, .. } => decoration, + _ => continue, + }; + + match *decoration { + Decoration::XfbStride { xfb_stride } => { + if xfb_stride + > properties + .max_transform_feedback_buffer_data_stride + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the `XfbStride` decoration is used, but its `xfb_stride` \ + value is greater than the \ + `max_transform_feedback_buffer_data_stride` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-XfbStride-06313"], + ..Default::default() + })); + } + } + Decoration::Stream { stream_number } => { + if stream_number + >= properties + .max_transform_feedback_streams + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the `Stream` decoration is used, but its `stream_number` \ + value is not less than the `max_transform_feedback_streams` \ + device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Stream-06312"], + ..Default::default() + })); + } + } + _ => (), + } + } + + Ok(()) + } + + fn validate_execution_modes(&mut self) -> Result<(), Box> { + let properties = self.device.physical_device().properties(); + + #[derive(PartialEq, Eq)] + enum DenormMode { + Preserve, + FlushToZero, + } + + #[allow(clippy::upper_case_acronyms)] + #[derive(PartialEq, Eq)] + enum RoundingMode { + RTE, + RTZ, + } + + let mut denorm_mode_16 = None; + let mut denorm_mode_32 = None; + let mut denorm_mode_64 = None; + let mut rounding_mode_16 = None; + let mut rounding_mode_32 = None; + let mut rounding_mode_64 = None; + + for instruction in self.entry_point_info.execution_modes() { + let execution_mode = match instruction { + Instruction::ExecutionMode { mode, .. } + | Instruction::ExecutionModeId { mode, .. } => mode, + _ => continue, + }; + + match *execution_mode { + ExecutionMode::SignedZeroInfNanPreserve { target_width } => match target_width { + 16 => { + if !properties + .shader_signed_zero_inf_nan_preserve_float16 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `SignedZeroInfNanPreserve` \ + execution mode with a `target_width` of 16, but the \ + `shader_signed_zero_inf_nan_preserve_float16` \ + device property is `false`" + .into(), + vuids: &[ + "VUID-RuntimeSpirv-shaderSignedZeroInfNanPreserveFloat16-06293", + ], + ..Default::default() + })); + } + } + 32 => { + if !properties + .shader_signed_zero_inf_nan_preserve_float32 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `SignedZeroInfNanPreserve` \ + execution mode with a `target_width` of 32, but the \ + `shader_signed_zero_inf_nan_preserve_float32` \ + device property is `false`" + .into(), + vuids: &[ + "VUID-RuntimeSpirv-shaderSignedZeroInfNanPreserveFloat32-06294", + ], + ..Default::default() + })); + } + } + 64 => { + if !properties + .shader_signed_zero_inf_nan_preserve_float64 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `SignedZeroInfNanPreserve` \ + execution mode with a `target_width` of 64, but the \ + `shader_signed_zero_inf_nan_preserve_float64` \ + device property is `false`" + .into(), + vuids: &[ + "VUID-RuntimeSpirv-shaderSignedZeroInfNanPreserveFloat64-06295", + ], + ..Default::default() + })); + } + } + _ => (), + }, + ExecutionMode::DenormPreserve { target_width } => match target_width { + 16 => { + denorm_mode_16 = Some(DenormMode::Preserve); + + if !properties + .shader_denorm_preserve_float16 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `DenormPreserve` \ + execution mode with a `target_width` of 16, but the \ + `shader_denorm_preserve_float16` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderDenormPreserveFloat16-06296"], + ..Default::default() + })); + } + } + 32 => { + denorm_mode_32 = Some(DenormMode::Preserve); + + if !properties + .shader_denorm_preserve_float32 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `DenormPreserve` \ + execution mode with a `target_width` of 32, but the \ + `shader_denorm_preserve_float32` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderDenormPreserveFloat32-06297"], + ..Default::default() + })); + } + } + 64 => { + denorm_mode_64 = Some(DenormMode::Preserve); + + if !properties + .shader_denorm_preserve_float64 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `DenormPreserve` \ + execution mode with a `target_width` of 64, but the \ + `shader_denorm_preserve_float64` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderDenormPreserveFloat64-06298"], + ..Default::default() + })); + } + } + _ => (), + }, + ExecutionMode::DenormFlushToZero { target_width } => match target_width { + 16 => { + denorm_mode_16 = Some(DenormMode::FlushToZero); + + if !properties + .shader_denorm_flush_to_zero_float16 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `DenormFlushToZero` \ + execution mode with a `target_width` of 16, but the \ + `shader_denorm_flush_to_zero_float16` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderDenormFlushToZeroFloat16-06299"], + ..Default::default() + })); + } + } + 32 => { + denorm_mode_32 = Some(DenormMode::FlushToZero); + + if !properties + .shader_denorm_flush_to_zero_float32 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `DenormFlushToZero` \ + execution mode with a `target_width` of 32, but the \ + `shader_denorm_flush_to_zero_float32` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderDenormFlushToZeroFloat32-06300"], + ..Default::default() + })); + } + } + 64 => { + denorm_mode_64 = Some(DenormMode::FlushToZero); + + if !properties + .shader_denorm_flush_to_zero_float64 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `DenormFlushToZero` \ + execution mode with a `target_width` of 64, but the \ + `shader_denorm_flush_to_zero_float64` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderDenormFlushToZeroFloat64-06301"], + ..Default::default() + })); + } + } + _ => (), + }, + ExecutionMode::RoundingModeRTE { target_width } => match target_width { + 16 => { + rounding_mode_16 = Some(RoundingMode::RTE); + + if !properties + .shader_rounding_mode_rte_float16 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `RoundingModeRTE` \ + execution mode with a `target_width` of 16, but the \ + `shader_rounding_mode_rte_float16` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderRoundingModeRTEFloat16-06302"], + ..Default::default() + })); + } + } + 32 => { + rounding_mode_32 = Some(RoundingMode::RTE); + + if !properties + .shader_rounding_mode_rte_float32 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `RoundingModeRTE` \ + execution mode with a `target_width` of 32, but the \ + `shader_rounding_mode_rte_float32` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderRoundingModeRTEFloat32-06303"], + ..Default::default() + })); + } + } + 64 => { + rounding_mode_64 = Some(RoundingMode::RTE); + + if !properties + .shader_rounding_mode_rte_float64 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `RoundingModeRTE` \ + execution mode with a `target_width` of 64, but the \ + `shader_rounding_mode_rte_float64` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderRoundingModeRTEFloat64-06304"], + ..Default::default() + })); + } + } + _ => (), + }, + ExecutionMode::RoundingModeRTZ { target_width } => match target_width { + 16 => { + rounding_mode_16 = Some(RoundingMode::RTZ); + + if !properties + .shader_rounding_mode_rtz_float16 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `RoundingModeRTZ` \ + execution mode with a `target_width` of 16, but the \ + `shader_rounding_mode_rtz_float16` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderRoundingModeRTZFloat16-06305"], + ..Default::default() + })); + } + } + 32 => { + rounding_mode_32 = Some(RoundingMode::RTZ); + + if !properties + .shader_rounding_mode_rtz_float32 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `RoundingModeRTZ` \ + execution mode with a `target_width` of 32, but the \ + `shader_rounding_mode_rtz_float32` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderRoundingModeRTZFloat32-06306"], + ..Default::default() + })); + } + } + 64 => { + rounding_mode_64 = Some(RoundingMode::RTZ); + + if !properties + .shader_rounding_mode_rtz_float64 + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `RoundingModeRTZ` \ + execution mode with a `target_width` of 64, but the \ + `shader_rounding_mode_rtz_float64` \ + device property is `false`" + .into(), + vuids: &["VUID-RuntimeSpirv-shaderRoundingModeRTZFloat64-06307"], + ..Default::default() + })); + } + } + _ => (), + }, + ExecutionMode::Isolines => { + if self.device.enabled_extensions().khr_portability_subset + && self.device.enabled_features().tessellation_shader + && !self.device.enabled_features().tessellation_isolines + { + return Err(Box::new(ValidationError { + problem: "this device is a portability subset device, and \ + the entry point has an `IsoLines` execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "tessellation_isolines", + )])]), + vuids: &["VUID-RuntimeSpirv-tessellationShader-06326"], + ..Default::default() + })); + } + } + ExecutionMode::PointMode => { + if self.device.enabled_extensions().khr_portability_subset + && self.device.enabled_features().tessellation_shader + && !self.device.enabled_features().tessellation_point_mode + { + return Err(Box::new(ValidationError { + problem: "this device is a portability subset device, and \ + the entry point has an `PointMode` execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "tessellation_point_mode", + )])]), + vuids: &["VUID-RuntimeSpirv-tessellationShader-06327"], + ..Default::default() + })); + } + } + ExecutionMode::LocalSize { .. } | ExecutionMode::LocalSizeId { .. } => { + let local_size = match *execution_mode { + ExecutionMode::LocalSize { + x_size, + y_size, + z_size, + } => [x_size as u64, y_size as u64, z_size as u64], + ExecutionMode::LocalSizeId { + x_size, + y_size, + z_size, + } => { + if !self.device.enabled_features().maintenance4 { + return Err(Box::new(ValidationError { + problem: "the entry point has a `LocalSizeId` execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("maintenance4"), + ])]), + vuids: &["VUID-RuntimeSpirv-LocalSizeId-06434"], + ..Default::default() + })); + } + + let x_size = get_constant(self.spirv, x_size).unwrap(); + let y_size = get_constant(self.spirv, y_size).unwrap(); + let z_size = get_constant(self.spirv, z_size).unwrap(); + + [x_size, y_size, z_size] + } + _ => unreachable!(), + }; + let workgroup_size = local_size.into_iter().try_fold(1, |t, x| { + u32::try_from(x).ok().and_then(|x| x.checked_mul(t)) + }); + + match self.execution_model { + ExecutionModel::GLCompute => { + if u32::try_from(local_size[0]).map_or(true, |size| { + size > properties.max_compute_work_group_size[0] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `GLCompute`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_x` is greater than the \ + `max_compute_work_group_size[0]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-x-06429"], + ..Default::default() + })); + } + + if u32::try_from(local_size[1]).map_or(true, |size| { + size > properties.max_compute_work_group_size[1] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `GLCompute`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_y` is greater than the \ + `max_compute_work_group_size[1]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-x-06430"], + ..Default::default() + })); + } + + if u32::try_from(local_size[2]).map_or(true, |size| { + size > properties.max_compute_work_group_size[2] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `GLCompute`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_z` is greater than the \ + `max_compute_work_group_size[2]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-x-06431"], + ..Default::default() + })); + } + + if workgroup_size.map_or(true, |size| { + size > properties.max_compute_work_group_invocations + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `GLCompute`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but the product of its `size_x`, `size_y` and `size_z` is \ + greater than the `max_compute_work_group_invocations` \ + device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-x-06432"], + ..Default::default() + })); + } + } + ExecutionModel::TaskEXT => { + if u32::try_from(local_size[0]).map_or(true, |size| { + size > properties.max_task_work_group_size.unwrap_or_default()[0] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_x` is greater than the \ + `max_task_work_group_size[0]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07291"], + ..Default::default() + })); + } + + if u32::try_from(local_size[1]).map_or(true, |size| { + size > properties.max_task_work_group_size.unwrap_or_default()[1] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_y` is greater than the \ + `max_task_work_group_size[1]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07292"], + ..Default::default() + })); + } + + if u32::try_from(local_size[2]).map_or(true, |size| { + size > properties.max_task_work_group_size.unwrap_or_default()[2] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_z` is greater than the \ + `max_task_work_group_size[2]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07293"], + ..Default::default() + })); + } + + if workgroup_size.map_or(true, |size| { + size > properties + .max_task_work_group_invocations + .unwrap_or_default() + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but the product of its `size_x`, `size_y` and `size_z` is \ + greater than the `max_task_work_group_invocations` \ + device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07294"], + ..Default::default() + })); + } + } + ExecutionModel::MeshEXT => { + if u32::try_from(local_size[0]).map_or(true, |size| { + size > properties.max_mesh_work_group_size.unwrap_or_default()[0] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_x` is greater than the \ + `max_mesh_work_group_size[0]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-MeshEXT-07295"], + ..Default::default() + })); + } + + if u32::try_from(local_size[1]).map_or(true, |size| { + size > properties.max_mesh_work_group_size.unwrap_or_default()[1] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_y` is greater than the \ + `max_mesh_work_group_size[1]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-MeshEXT-07296"], + ..Default::default() + })); + } + + if u32::try_from(local_size[2]).map_or(true, |size| { + size > properties.max_mesh_work_group_size.unwrap_or_default()[2] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but its `size_z` is greater than the \ + `max_mesh_work_group_size[2]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-MeshEXT-07297"], + ..Default::default() + })); + } + + if workgroup_size.map_or(true, |size| { + size > properties + .max_mesh_work_group_invocations + .unwrap_or_default() + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, \ + and it has a `LocalSize` or `LocalSizeId` execution mode, \ + but the product of its `size_x`, `size_y` and `size_z` is \ + greater than the `max_mesh_work_group_invocations` \ + device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-MeshEXT-07298"], + ..Default::default() + })); + } + } + _ => (), + } + } + ExecutionMode::SubgroupUniformControlFlowKHR => { + if !self + .device + .enabled_features() + .shader_subgroup_uniform_control_flow + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `SubgroupUniformControlFlowKHR` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_subgroup_uniform_control_flow", + )])]), + vuids: &["VUID-RuntimeSpirv-SubgroupUniformControlFlowKHR-06379"], + ..Default::default() + })); + } + + if !properties + .supported_stages + .unwrap_or_default() + .contains_enum(ShaderStage::from(self.execution_model)) + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `SubgroupUniformControlFlowKHR` \ + execution mode, but the `supported_stages` device property does \ + not contain the shader stage of the entry point's execution model" + .into(), + vuids: &["VUID-RuntimeSpirv-SubgroupUniformControlFlowKHR-06379"], + ..Default::default() + })); + } + } + ExecutionMode::EarlyAndLateFragmentTestsAMD => { + if !self + .device + .enabled_features() + .shader_early_and_late_fragment_tests + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `EarlyAndLateFragmentTestsAMD` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_early_and_late_fragment_tests", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderEarlyAndLateFragmentTests-06767"], + ..Default::default() + })); + } + } + ExecutionMode::StencilRefUnchangedFrontAMD => { + if !self + .device + .enabled_features() + .shader_early_and_late_fragment_tests + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `StencilRefUnchangedFrontAMD` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_early_and_late_fragment_tests", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderEarlyAndLateFragmentTests-06768"], + ..Default::default() + })); + } + } + ExecutionMode::StencilRefUnchangedBackAMD => { + if !self + .device + .enabled_features() + .shader_early_and_late_fragment_tests + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `StencilRefUnchangedBackAMD` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_early_and_late_fragment_tests", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderEarlyAndLateFragmentTests-06769"], + ..Default::default() + })); + } + } + ExecutionMode::StencilRefGreaterFrontAMD => { + if !self + .device + .enabled_features() + .shader_early_and_late_fragment_tests + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `StencilRefGreaterFrontAMD` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_early_and_late_fragment_tests", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderEarlyAndLateFragmentTests-06770"], + ..Default::default() + })); + } + } + ExecutionMode::StencilRefGreaterBackAMD => { + if !self + .device + .enabled_features() + .shader_early_and_late_fragment_tests + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `StencilRefGreaterBackAMD` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_early_and_late_fragment_tests", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderEarlyAndLateFragmentTests-06771"], + ..Default::default() + })); + } + } + ExecutionMode::StencilRefLessFrontAMD => { + if !self + .device + .enabled_features() + .shader_early_and_late_fragment_tests + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `StencilRefLessFrontAMD` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_early_and_late_fragment_tests", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderEarlyAndLateFragmentTests-06772"], + ..Default::default() + })); + } + } + ExecutionMode::StencilRefLessBackAMD => { + if !self + .device + .enabled_features() + .shader_early_and_late_fragment_tests + { + return Err(Box::new(ValidationError { + problem: "the entry point has a `StencilRefLessBackAMD` \ + execution mode" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_early_and_late_fragment_tests", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderEarlyAndLateFragmentTests-06773"], + ..Default::default() + })); + } + } + ExecutionMode::OutputPoints + | ExecutionMode::OutputLineStrip + | ExecutionMode::OutputTriangleStrip => { + self.output_primitives = Some(execution_mode); + } + ExecutionMode::OutputVertices { vertex_count } => { + match self.execution_model { + ExecutionModel::MeshNV => { + // TODO: needs VK_NV_mesh_shader support + // VUID-RuntimeSpirv-MeshNV-07113 + } + ExecutionModel::MeshEXT => { + if vertex_count + > properties.max_mesh_output_vertices.unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, and \ + it has an `OutputVertices` execution mode, but its \ + `vertex_count` is greater than the \ + `max_mesh_output_vertices` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-MeshEXT-07115"], + ..Default::default() + })); + } + } + _ => (), + } + } + ExecutionMode::OutputPrimitivesEXT { primitive_count } => { + match self.execution_model { + ExecutionModel::MeshNV => { + // TODO: needs VK_NV_mesh_shader support + // VUID-RuntimeSpirv-MeshNV-07114 + } + ExecutionModel::MeshEXT => { + if primitive_count + > properties.max_mesh_output_primitives.unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, and \ + it has an `OutputPrimitivesEXT` execution mode, but its \ + `primitive_count` is greater than the \ + `max_mesh_output_primitives` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-MeshEXT-07116"], + ..Default::default() + })); + } + } + _ => (), + } + } + _ => (), + } + } + + match properties.denorm_behavior_independence { + Some(ShaderFloatControlsIndependence::Float32Only) => { + if denorm_mode_16 != denorm_mode_64 { + return Err(Box::new(ValidationError { + problem: "the `denorm_behavior_independence` device property is \ + `ShaderFloatControlsIndependence::Float32Only`, but the entry point \ + does not have the same denormals execution mode for \ + both 16-bit and 64-bit values" + .into(), + vuids: &["VUID-RuntimeSpirv-denormBehaviorIndependence-06289"], + ..Default::default() + })); + } + } + Some(ShaderFloatControlsIndependence::None) => { + if denorm_mode_16 != denorm_mode_32 || denorm_mode_16 != denorm_mode_64 { + return Err(Box::new(ValidationError { + problem: "the `denorm_behavior_independence` device property is \ + `ShaderFloatControlsIndependence::None`, but the entry point \ + does not have the same denormals execution mode for \ + 16-bit, 32-bit and 64-bit values" + .into(), + vuids: &["VUID-RuntimeSpirv-denormBehaviorIndependence-06290"], + ..Default::default() + })); + } + } + _ => (), + } + + match properties.rounding_mode_independence { + Some(ShaderFloatControlsIndependence::Float32Only) => { + if rounding_mode_16 != rounding_mode_64 { + return Err(Box::new(ValidationError { + problem: "the `rounding_mode_independence` device property is \ + `ShaderFloatControlsIndependence::Float32Only`, but the entry point \ + does not have the same rounding execution mode for \ + both 16-bit and 64-bit values" + .into(), + vuids: &["VUID-RuntimeSpirv-roundingModeIndependence-06291"], + ..Default::default() + })); + } + } + Some(ShaderFloatControlsIndependence::None) => { + if rounding_mode_16 != rounding_mode_32 || rounding_mode_16 != rounding_mode_64 { + return Err(Box::new(ValidationError { + problem: "the `rounding_mode_independence` device property is \ + `ShaderFloatControlsIndependence::None`, but the entry point \ + does not have the same rounding execution mode for \ + 16-bit, 32-bit and 64-bit values" + .into(), + vuids: &["VUID-RuntimeSpirv-roundingModeIndependence-06292"], + ..Default::default() + })); + } + } + _ => (), + } + + Ok(()) + } + + fn validate_types(&self) -> Result<(), Box> { + let properties = self.device.physical_device().properties(); + + for instruction in self.spirv.types() { + match *instruction { + Instruction::TypeCooperativeMatrixKHR { .. } => { + // TODO: needs VK_KHR_cooperative_matrix support + // VUID-RuntimeSpirv-OpTypeCooperativeMatrixKHR-08974 + + if !properties + .cooperative_matrix_supported_stages + .unwrap() + .contains_enum(ShaderStage::from(self.execution_model)) + { + return Err(Box::new(ValidationError { + problem: "an `OpTypeCooperativeMatrixKHR` type is declared, but \ + the `cooperative_matrix_supported_stages` device property does \ + not contain the shader stage of the entry point's execution model" + .into(), + vuids: &["VUID-RuntimeSpirv-cooperativeMatrixSupportedStages-08985"], + ..Default::default() + })); + } + } + Instruction::TypeCooperativeMatrixNV { .. } => { + // TODO: needs VK_NV_cooperative_matrix support + // VUID-RuntimeSpirv-OpTypeCooperativeMatrixNV-06316 + // VUID-RuntimeSpirv-OpTypeCooperativeMatrixNV-06322 + } + _ => (), + } + } + + Ok(()) + } + + fn validate_global_variables(&self) -> Result<(), Box> { + let properties = self.device.physical_device().properties(); + + // Graphics stages except task. + let mut input_locations_required = 0; + let mut output_locations_required = 0; + let mut per_patch_output_locations_required = 0; + + // Only TaskEXT and MeshEXT. + let mut task_payload_workgroup_memory_size = 0; + + // Only GLCompute, TaskEXT and MeshEXT. + let mut workgroup_memory_size = 0; + + let mut stream_sizes: HashMap> = HashMap::default(); + + for instruction in self.spirv.global_variables() { + let Instruction::Variable { + result_type_id, + result_id, + storage_class, + initializer, + } = *instruction + else { + continue; + }; + + let is_in_interface = self.interface.contains(&result_id); + let Instruction::TypePointer { ty: type_id, .. } = + *self.spirv.id(result_type_id).instruction() + else { + unreachable!() + }; + + fn get_bits(has_8bit: &mut bool, has_16bit: &mut bool, spirv: &Spirv, ty: Id) { + match *spirv.id(ty).instruction() { + Instruction::TypeInt { width, .. } | Instruction::TypeFloat { width, .. } => { + match width { + 8 => *has_8bit = true, + 16 => *has_16bit = true, + _ => (), + } + } + Instruction::TypePointer { ty, .. } + | Instruction::TypeArray { + element_type: ty, .. + } + | Instruction::TypeRuntimeArray { + element_type: ty, .. + } + | Instruction::TypeVector { + component_type: ty, .. + } + | Instruction::TypeMatrix { + column_type: ty, .. + } => get_bits(has_8bit, has_16bit, spirv, ty), + Instruction::TypeStruct { + ref member_types, .. + } => { + for &ty in member_types { + get_bits(has_8bit, has_16bit, spirv, ty) + } + } + _ => (), + } + } + + let mut has_8bit = false; + let mut has_16bit = false; + get_bits(&mut has_8bit, &mut has_16bit, self.spirv, type_id); + + let mut has_aliased = false; + let mut has_block = false; + let mut has_buffer_block = false; + let mut has_non_readable = false; + let mut has_non_writable = false; + let mut has_patch = true; + + let mut offset = None; + let mut stream = 0; + let mut xfb_buffer = None; + + for instruction in self.spirv.id(result_id).decorations() { + if let Instruction::Decorate { decoration, .. } = instruction { + match *decoration { + Decoration::Aliased => has_aliased = true, + Decoration::Block => has_block = true, + Decoration::BufferBlock => has_buffer_block = true, + Decoration::NonReadable => has_non_readable = true, + Decoration::NonWritable => has_non_writable = true, + Decoration::Offset { byte_offset } => offset = Some(byte_offset), + Decoration::Patch => has_patch = true, + Decoration::Stream { stream_number } => stream = stream_number, + Decoration::XfbBuffer { xfb_buffer_number } => { + xfb_buffer = Some(xfb_buffer_number) + } + _ => (), + } + } + } + + match storage_class { + StorageClass::Workgroup => { + if matches!( + self.execution_model, + ExecutionModel::GLCompute + | ExecutionModel::TaskEXT + | ExecutionModel::MeshEXT + ) { + if let Some(size) = size_of_type(self.spirv, type_id) { + if has_aliased { + workgroup_memory_size = workgroup_memory_size.max(size); + } else { + workgroup_memory_size += size; + } + } + } + + if initializer.is_some() + && !self + .device + .enabled_features() + .shader_zero_initialize_workgroup_memory + { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} has a storage class of `StorageClass::Workgroup`, \ + and has an initializer operand", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_zero_initialize_workgroup_memory", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderZeroInitializeWorkgroupMemory-06372"], + ..Default::default() + })); + } + } + + StorageClass::TaskPayloadWorkgroupEXT => { + if matches!( + self.execution_model, + ExecutionModel::TaskEXT | ExecutionModel::MeshEXT + ) { + if let Some(size) = size_of_type(self.spirv, type_id) { + task_payload_workgroup_memory_size += size; + } + } + } + + StorageClass::StorageBuffer + | StorageClass::ShaderRecordBufferKHR + | StorageClass::PhysicalStorageBuffer => { + if has_8bit && !self.device.enabled_features().storage_buffer8_bit_access { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} contains an 8-bit integer value, \ + and has a storage class of `StorageClass::StorageBuffer`, \ + `StorageClass::ShaderRecordBufferKHR` or \ + `StorageClass::PhysicalStorageBuffer`", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "storage_buffer8_bit_access", + )])]), + vuids: &["VUID-RuntimeSpirv-storageBuffer8BitAccess-06328"], + ..Default::default() + })); + } + + if has_16bit && !self.device.enabled_features().storage_buffer16_bit_access { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} contains an 16-bit integer or floating point value, \ + and has a storage class of `StorageClass::StorageBuffer`, \ + `StorageClass::ShaderRecordBufferKHR` or \ + `StorageClass::PhysicalStorageBuffer`", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "storage_buffer16_bit_access", + )])]), + vuids: &["VUID-RuntimeSpirv-storageBuffer16BitAccess-06331"], + ..Default::default() + })); + } + } + + StorageClass::Uniform => { + if has_block { + if has_8bit + && !self + .device + .enabled_features() + .uniform_and_storage_buffer8_bit_access + { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} contains an 8-bit integer value, and has a \ + storage class of `StorageClass::Uniform` and \ + is decorated with `Decoration::Block`", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("uniform_and_storage_buffer8_bit_access"), + ])]), + vuids: &[ + "VUID-RuntimeSpirv-uniformAndStorageBuffer8BitAccess-06329", + ], + ..Default::default() + })); + } + + if has_16bit + && !self + .device + .enabled_features() + .uniform_and_storage_buffer16_bit_access + { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} contains an 16-bit integer or floating point \ + value, and has a storage class of `StorageClass::Uniform` and \ + is decorated with `Decoration::Block`", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("uniform_and_storage_buffer16_bit_access"), + ])]), + vuids: &[ + "VUID-RuntimeSpirv-uniformAndStorageBuffer16BitAccess-06332", + ], + ..Default::default() + })); + } + } + } + + StorageClass::PushConstant => { + if has_8bit && !self.device.enabled_features().storage_push_constant8 { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} contains an 8-bit integer value, and \ + has a storage class of `StorageClass::PushConstant`", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "storage_push_constant8", + )])]), + vuids: &["VUID-RuntimeSpirv-storagePushConstant8-06330"], + ..Default::default() + })); + } + + if has_16bit && !self.device.enabled_features().storage_push_constant16 { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} contains an 16-bit integer or floating point value, \ + and has a storage class of `StorageClass::PushConstant`", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "storage_push_constant16", + )])]), + vuids: &["VUID-RuntimeSpirv-storagePushConstant16-06333"], + ..Default::default() + })); + } + } + + StorageClass::Input | StorageClass::Output => { + if has_16bit && !self.device.enabled_features().storage_input_output16 { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} contains an 16-bit integer or floating point value, \ + and has a storage class of `StorageClass::Input` or \ + `StorageClass::Output`", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "storage_input_output16", + )])]), + vuids: &["VUID-RuntimeSpirv-storageInputOutput16-06334"], + ..Default::default() + })); + } + + if matches!( + self.execution_model, + ExecutionModel::Vertex + | ExecutionModel::TessellationControl + | ExecutionModel::TessellationEvaluation + | ExecutionModel::Geometry + | ExecutionModel::Fragment + ) || self.execution_model == ExecutionModel::MeshEXT + && storage_class == StorageClass::Output + { + let locations_required = match storage_class { + StorageClass::Input => &mut input_locations_required, + StorageClass::Output => { + if has_patch { + &mut per_patch_output_locations_required + } else { + &mut output_locations_required + } + } + _ => unreachable!(), + }; + let mut dummy_scalar_func = |_, _, _| {}; + + input_output_map( + self.spirv, + self.execution_model, + result_id, + storage_class, + |key, data| { + let InputOutputData { type_id, .. } = data; + + match key { + InputOutputKey::User(key) => { + let (locations, _) = shader_interface_analyze_type( + self.spirv, + type_id, + key, + &mut dummy_scalar_func, + ); + + *locations_required = + max(*locations_required, key.location + locations); + } + InputOutputKey::BuiltIn(_) => { + // TODO: The spec doesn't currently say how to count this. + // https://github.com/KhronosGroup/Vulkan-Docs/issues/2293 + } + } + }, + ); + } + + if is_in_interface && storage_class == StorageClass::Output { + let mut check = |offset: u32, xfb_buffer: u32, stream: u32| { + if let Some(size) = size_of_type(self.spirv, result_id) { + let required_size = offset as DeviceSize + size; + + if required_size + > properties + .max_transform_feedback_buffer_data_size + .unwrap_or_default() + as DeviceSize + { + return Err(Box::new(ValidationError { + problem: format!( + "for the value written to transform feedback buffer \ + {} at offset {}, the offset plus the size of the \ + value is greater than the \ + `max_transform_feedback_buffer_data_size` device limit", + xfb_buffer, offset, + ) + .into(), + vuids: &["VUID-RuntimeSpirv-Offset-06308"], + ..Default::default() + })); + } + + let buffer_data_size = stream_sizes + .entry(stream) + .or_default() + .entry(xfb_buffer) + .or_insert(0); + *buffer_data_size = max(*buffer_data_size, required_size); + } + + Ok(()) + }; + + if let (Some(offset), Some(xfb_buffer)) = (offset, xfb_buffer) { + check(offset, xfb_buffer, stream)?; + } else if let Instruction::TypeStruct { .. } = + self.spirv.id(type_id).instruction() + { + for member_info in self.spirv.id(type_id).members() { + let mut member_offset = None; + let mut member_stream = None; + let mut member_xfb_buffer = None; + + for instruction in member_info.decorations() { + if let Instruction::Decorate { decoration, .. } = instruction { + match *decoration { + Decoration::Offset { byte_offset } => { + member_offset = Some(byte_offset) + } + Decoration::Stream { stream_number } => { + member_stream = Some(stream_number) + } + Decoration::XfbBuffer { xfb_buffer_number } => { + member_xfb_buffer = Some(xfb_buffer_number) + } + _ => (), + } + } + } + + // Inherit the XfbBuffer and Stream of the parent variable if there + // is one. + if let (Some(offset), Some(xfb_buffer)) = + (member_offset, member_xfb_buffer.or(xfb_buffer)) + { + check(offset, xfb_buffer, member_stream.unwrap_or(stream))?; + } + } + } + } + } + + _ => (), + } + + let descriptor_type = match storage_class { + StorageClass::StorageBuffer | StorageClass::PhysicalStorageBuffer => { + match has_block { + true => Some(DescriptorType::StorageBuffer), + false => Some(DescriptorType::UniformBuffer), + } + } + StorageClass::Uniform => match has_buffer_block { + true => Some(DescriptorType::StorageBuffer), + false => Some(DescriptorType::UniformBuffer), + }, + StorageClass::UniformConstant => { + let base_type = match *self.spirv.id(type_id).instruction() { + Instruction::TypeArray { element_type, .. } => element_type, + _ => type_id, + }; + + match *self.spirv.id(base_type).instruction() { + Instruction::TypeImage { dim, sampled, .. } => match dim { + Dim::Dim1D | Dim::Dim2D | Dim::Dim3D | Dim::Cube | Dim::Rect => { + if sampled == 1 { + Some(DescriptorType::SampledImage) + } else { + Some(DescriptorType::StorageImage) + } + } + Dim::Buffer => { + if sampled == 1 { + Some(DescriptorType::UniformTexelBuffer) + } else { + Some(DescriptorType::StorageTexelBuffer) + } + } + Dim::SubpassData => Some(DescriptorType::InputAttachment), + Dim::TileImageDataEXT => None, + }, + _ => None, + } + } + _ => None, + }; + + if !has_non_writable + && matches!( + descriptor_type, + Some( + DescriptorType::StorageImage + | DescriptorType::StorageTexelBuffer + | DescriptorType::StorageBuffer + ) + ) + { + match self.execution_model { + ExecutionModel::Fragment => { + if !self.device.enabled_features().fragment_stores_and_atomics { + return Err(Box::new(ValidationError { + problem: format!( + "the entry point's execution model is `Fragment`, and \ + variable {} is a storage image, storage texel buffer or \ + storage buffer variable, and does not have a `NonWritable` \ + decoration", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("fragment_stores_and_atomics"), + ])]), + vuids: &["VUID-RuntimeSpirv-NonWritable-06340"], + ..Default::default() + })); + } + } + ExecutionModel::Vertex + | ExecutionModel::TessellationControl + | ExecutionModel::TessellationEvaluation + | ExecutionModel::Geometry => { + if !self + .device + .enabled_features() + .vertex_pipeline_stores_and_atomics + { + return Err(Box::new(ValidationError { + problem: format!( + "the entry point's execution model is `Vertex`, \ + `TessellationControl`, `TessellationEvaluation` or \ + `Geometry`, and variable {} is a storage image, storage texel \ + buffer or storage buffer variable, and does not have a \ + `NonWritable` decoration", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("vertex_pipeline_stores_and_atomics"), + ])]), + vuids: &["VUID-RuntimeSpirv-NonWritable-06341"], + ..Default::default() + })); + } + } + _ => (), + } + } + + if self.device.physical_device().api_version() < Version::V1_3 + && !self + .device + .physical_device() + .supported_extensions() + .khr_format_feature_flags2 + { + let base_type = match *self.spirv.id(type_id).instruction() { + Instruction::TypeArray { element_type, .. } => element_type, + _ => type_id, + }; + + if matches!( + *self.spirv.id(base_type).instruction(), + Instruction::TypeImage { + sampled: 2, + image_format: ImageFormat::Unknown, + .. + } + ) { + if !has_non_writable + && self + .device + .enabled_features() + .shader_storage_image_write_without_format + { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} is an image with a `sampled` operand of 2, and an \ + `image_format` operand of `ImageFormat::Unknown`, and does not \ + have a `NonWritable` decoration", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::APIVersion(Version::V1_3)]), + RequiresAllOf(&[Requires::DeviceExtension( + "khr_format_feature_flags2", + )]), + RequiresAllOf(&[Requires::Feature("storage_input_output16")]), + ]), + vuids: &["VUID-RuntimeSpirv-apiVersion-07954"], + ..Default::default() + })); + } + + if !has_non_readable + && self + .device + .enabled_features() + .shader_storage_image_read_without_format + { + return Err(Box::new(ValidationError { + problem: format!( + "variable {} is an image with a `sampled` operand of 2, and an \ + `image_format` operand of `ImageFormat::Unknown`, and does not \ + have a `NonReadable` decoration", + result_id, + ) + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::APIVersion(Version::V1_3)]), + RequiresAllOf(&[Requires::DeviceExtension( + "khr_format_feature_flags2", + )]), + RequiresAllOf(&[Requires::Feature("storage_input_output16")]), + ]), + vuids: &["VUID-RuntimeSpirv-apiVersion-07955"], + ..Default::default() + })); + } + } + } + } + + for (stream, buffer_data_sizes) in stream_sizes { + let required_size: DeviceSize = buffer_data_sizes.values().sum(); + + if required_size + > properties + .max_transform_feedback_buffer_data_size + .unwrap_or_default() as DeviceSize + { + return Err(Box::new(ValidationError { + problem: format!( + "the sum of the data sizes of all transform feedback buffers \ + associated with stream {} is greater than the \ + `max_transform_feedback_buffer_data_size` device limit", + stream, + ) + .into(), + vuids: &["VUID-RuntimeSpirv-XfbBuffer-06309"], + ..Default::default() + })); + } + } + + match self.execution_model { + ExecutionModel::Vertex => { + if input_locations_required > properties.max_vertex_input_attributes { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `Vertex`, but \ + the number of input locations required is greater than the \ + `max_vertex_input_attributes` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if output_locations_required > properties.max_vertex_output_components / 4 { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `Vertex`, but \ + the number of output components required is greater than the \ + `max_vertex_output_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + } + ExecutionModel::TessellationControl => { + if input_locations_required + > properties.max_tessellation_control_per_vertex_input_components / 4 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TessellationControl`, but \ + the number of input components required is greater than the \ + `max_tessellation_control_per_vertex_input_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if output_locations_required + > properties.max_tessellation_control_per_vertex_output_components / 4 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TessellationControl`, but \ + the number of per-vertex output components required is greater than \ + the `max_tessellation_control_per_vertex_output_components` device \ + limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if per_patch_output_locations_required + > properties.max_tessellation_control_per_patch_output_components / 4 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TessellationControl`, but \ + the number of per-patch output components required is greater than \ + the `max_tessellation_control_per_patch_output_components` device \ + limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if max( + output_locations_required, + per_patch_output_locations_required, + ) > properties.max_tessellation_control_total_output_components / 4 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TessellationControl`, but \ + the number of output components required is greater than the \ + `max_tessellation_control_total_output_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + } + ExecutionModel::TessellationEvaluation => { + if input_locations_required + > properties.max_tessellation_evaluation_input_components / 4 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TessellationEvaluation`, \ + but the number of input components required is greater than the \ + `max_tessellation_evaluation_input_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if output_locations_required + > properties.max_tessellation_evaluation_output_components / 4 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TessellationEvaluation`, \ + but the number of output components required is greater than the \ + `max_tessellation_evaluation_output_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + } + ExecutionModel::Geometry => { + if input_locations_required > properties.max_geometry_input_components / 4 { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `Geometry`, but \ + the number of input components required is greater than the \ + `max_geometry_input_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if output_locations_required > properties.max_geometry_output_components / 4 { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `Geometry`, but \ + the number of output components required is greater than the \ + `max_geometry_output_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + // TODO: max_geometry_total_output_components + } + ExecutionModel::Fragment => { + if input_locations_required > properties.max_fragment_input_components / 4 { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `Fragment`, but \ + the number of input components required is greater than the \ + `max_fragment_input_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if output_locations_required > properties.max_fragment_output_attachments { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `Fragment`, but \ + the number of output locations required is greater than the \ + `max_fragment_output_attachments` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + } + ExecutionModel::GLCompute => { + if workgroup_memory_size > properties.max_compute_shared_memory_size as u64 { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `GLCompute`, but \ + the total size of all variables in the `Workgroup` storage class is \ + greater than the `max_compute_shared_memory_size` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Workgroup-06530"], + ..Default::default() + })); + } + } + ExecutionModel::MeshEXT => { + if output_locations_required + > properties.max_mesh_output_components.unwrap_or_default() / 4 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, but \ + the number of output components required is greater than the \ + `max_mesh_output_components` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-Location-06272"], + ..Default::default() + })); + } + + if workgroup_memory_size + > properties.max_mesh_shared_memory_size.unwrap_or_default() as u64 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, but \ + the total size of all variables in the `Workgroup` storage class is \ + greater than the `max_mesh_shared_memory_size` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-maxMeshSharedMemorySize-08754"], + ..Default::default() + })); + } + + if task_payload_workgroup_memory_size + workgroup_memory_size + > properties + .max_mesh_payload_and_shared_memory_size + .unwrap_or_default() as u64 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `MeshEXT`, but \ + the total size of all variables in the `TaskPayloadWorkGroupEXT` or \ + `Workgroup` storage classes is greater than the \ + `max_mesh_payload_and_shared_memory_size` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-maxMeshPayloadAndSharedMemorySize-08755"], + ..Default::default() + })); + } + } + ExecutionModel::TaskEXT => { + if task_payload_workgroup_memory_size + > properties.max_task_payload_size.unwrap_or_default() as u64 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, but \ + the total size of all variables in the `TaskPayloadWorkgroupEXT` \ + storage class is greater than the `max_task_payload_size` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-maxTaskPayloadSize-08758"], + ..Default::default() + })); + } + + if workgroup_memory_size + > properties.max_task_shared_memory_size.unwrap_or_default() as u64 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, but \ + the total size of all variables in the `Workgroup` storage class is \ + greater than the `max_task_shared_memory_size` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-maxTaskSharedMemorySize-08759"], + ..Default::default() + })); + } + + if task_payload_workgroup_memory_size + workgroup_memory_size + > properties + .max_task_payload_and_shared_memory_size + .unwrap_or_default() as u64 + { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, but \ + the total size of all variables in the `TaskPayloadWorkgroupEXT` or \ + `Workgroup` storage class is greater than the \ + `max_task_payload_and_shared_memory_size` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-maxTaskPayloadAndSharedMemorySize-08760"], + ..Default::default() + })); + } + } + _ => (), + } + + Ok(()) + } + + fn validate_functions(&mut self) -> Result<(), Box> { + for &function in self.spirv.functions().keys() { + self.validate_function(function)?; + } + + Ok(()) + } + + fn validate_function(&mut self, function: Id) -> Result<(), Box> { + let properties = self.device.physical_device().properties(); + + for instruction in self.spirv.function(function).instructions() { + if let Some(pointer) = instruction.atomic_pointer_id() { + let (storage_class, ty) = + match (self.spirv.id(pointer).instruction().result_type_id()) + .map(|id| self.spirv.id(id).instruction()) + { + Some(&Instruction::TypePointer { + storage_class, ty, .. + }) => (storage_class, ty), + _ => unreachable!(), + }; + + match *self.spirv.id(ty).instruction() { + Instruction::TypeInt { width: 64, .. } => match storage_class { + StorageClass::StorageBuffer | StorageClass::Uniform => { + if !self.device.enabled_features().shader_buffer_int64_atomics { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + 64-bit integer value with a storage class of \ + `StorageClass::StorageBuffer` or `StorageClass::Uniform`" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_buffer_int64_atomics"), + ])]), + vuids: &["VUID-RuntimeSpirv-None-06278"], + ..Default::default() + })); + } + } + StorageClass::Workgroup => { + if !self.device.enabled_features().shader_shared_int64_atomics { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + 64-bit integer value with a storage class of \ + `StorageClass::Workgroup`" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_shared_int64_atomics"), + ])]), + vuids: &["VUID-RuntimeSpirv-None-06279"], + ..Default::default() + })); + } + } + StorageClass::Image => { + if !self.device.enabled_features().shader_image_int64_atomics { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + 64-bit integer value with a storage class of \ + `StorageClass::Image`" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_image_int64_atomics"), + ])]), + vuids: &["VUID-RuntimeSpirv-None-06288"], + ..Default::default() + })); + } + } + _ => (), + }, + Instruction::TypeFloat { width, .. } => { + match width { + 16 => { + if !self.device.enabled_features().intersects(&Features { + shader_buffer_float16_atomics: true, + shader_buffer_float16_atomic_add: true, + shader_buffer_float16_atomic_min_max: true, + shader_shared_float16_atomics: true, + shader_shared_float16_atomic_add: true, + shader_shared_float16_atomic_min_max: true, + ..Features::empty() + }) { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + 16-bit floating point value" + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float16_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float16_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float16_atomic_min_max", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float16_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float16_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float16_atomic_min_max", + )]), + ]), + vuids: &["VUID-RuntimeSpirv-None-06337"], + ..Default::default() + })); + } + } + 32 => { + if !self.device.enabled_features().intersects(&Features { + shader_buffer_float32_atomics: true, + shader_buffer_float32_atomic_add: true, + shader_buffer_float32_atomic_min_max: true, + shader_shared_float32_atomics: true, + shader_shared_float32_atomic_add: true, + shader_shared_float32_atomic_min_max: true, + ..Features::empty() + }) { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + 32-bit floating point value" + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float32_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float32_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float32_atomic_min_max", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float32_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float32_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float32_atomic_min_max", + )]), + ]), + vuids: &["VUID-RuntimeSpirv-None-06338"], + ..Default::default() + })); + } + } + 64 => { + if !self.device.enabled_features().intersects(&Features { + shader_buffer_float64_atomics: true, + shader_buffer_float64_atomic_add: true, + shader_buffer_float64_atomic_min_max: true, + shader_shared_float64_atomics: true, + shader_shared_float64_atomic_add: true, + shader_shared_float64_atomic_min_max: true, + ..Features::empty() + }) { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + 64-bit floating point value" + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float64_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float64_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float64_atomic_min_max", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float64_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float64_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float64_atomic_min_max", + )]), + ]), + vuids: &["VUID-RuntimeSpirv-None-06339"], + ..Default::default() + })); + } + } + _ => (), + } + + match storage_class { + StorageClass::StorageBuffer => { + if !self.device.enabled_features().intersects(&Features { + shader_buffer_float16_atomics: true, + shader_buffer_float16_atomic_add: true, + shader_buffer_float16_atomic_min_max: true, + shader_buffer_float32_atomics: true, + shader_buffer_float32_atomic_add: true, + shader_buffer_float32_atomic_min_max: true, + shader_buffer_float64_atomics: true, + shader_buffer_float64_atomic_add: true, + shader_buffer_float64_atomic_min_max: true, + ..Features::empty() + }) { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + floating point value with a storage class of \ + `StorageClass::StorageBuffer`" + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float16_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float16_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float16_atomic_min_max", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float32_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float32_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float32_atomic_min_max", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float64_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float64_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_buffer_float64_atomic_min_max", + )]), + ]), + vuids: &["VUID-RuntimeSpirv-None-06284"], + ..Default::default() + })); + } + } + StorageClass::Workgroup => { + if !self.device.enabled_features().intersects(&Features { + shader_shared_float16_atomics: true, + shader_shared_float16_atomic_add: true, + shader_shared_float16_atomic_min_max: true, + shader_shared_float32_atomics: true, + shader_shared_float32_atomic_add: true, + shader_shared_float32_atomic_min_max: true, + shader_shared_float64_atomics: true, + shader_shared_float64_atomic_add: true, + shader_shared_float64_atomic_min_max: true, + ..Features::empty() + }) { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + floating point value with a storage class of \ + `StorageClass::Workgroup`" + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::Feature( + "shader_shared_float16_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float16_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float16_atomic_min_max", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float32_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float32_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float32_atomic_min_max", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float64_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float64_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_shared_float64_atomic_min_max", + )]), + ]), + vuids: &["VUID-RuntimeSpirv-None-06285"], + ..Default::default() + })); + } + } + StorageClass::Image => { + if width == 32 + && !self.device.enabled_features().intersects(&Features { + shader_image_float32_atomics: true, + shader_image_float32_atomic_add: true, + shader_image_float32_atomic_min_max: true, + ..Features::empty() + }) + { + return Err(Box::new(ValidationError { + problem: "an atomic operation is performed on a \ + 32-bit floating point value with a storage \ + class of `StorageClass::Image`" + .into(), + requires_one_of: RequiresOneOf(&[ + RequiresAllOf(&[Requires::Feature( + "shader_image_float32_atomics", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_image_float32_atomic_add", + )]), + RequiresAllOf(&[Requires::Feature( + "shader_image_float32_atomic_min_max", + )]), + ]), + vuids: &["VUID-RuntimeSpirv-None-06286"], + ..Default::default() + })); + } + } + _ => (), + } + } + _ => (), + } + } + + if instruction.is_cooperative_matrix() { + if !properties + .cooperative_matrix_supported_stages + .unwrap() + .contains_enum(ShaderStage::from(self.execution_model)) + { + return Err(Box::new(ValidationError { + problem: "a cooperative matrix operation is performed, but \ + the `cooperative_matrix_supported_stages` device property does not \ + contain the shader stage of the entry point's execution model" + .into(), + vuids: &["VUID-RuntimeSpirv-cooperativeMatrixSupportedStages-08985"], + ..Default::default() + })); + } + } else if instruction.is_cooperative_matrix_nv() { + // TODO: needs VK_NV_cooperative_matrix support + // VUID-RuntimeSpirv-OpTypeCooperativeMatrixNV-06322 + // OpTypeCooperativeMatrixNV and OpCooperativeMatrix* instructions + // must not be used in shader stages not included in + // VkPhysicalDeviceCooperativeMatrixPropertiesNV::cooperativeMatrixSupportedStages + } + + if let Some(scope) = instruction + .memory_scope_id() + .and_then(|scope| get_constant(self.spirv, scope)) + .and_then(|scope| Scope::try_from(scope as u32).ok()) + { + match scope { + Scope::Device => { + if self.device.enabled_features().vulkan_memory_model + && !self + .device + .enabled_features() + .vulkan_memory_model_device_scope + { + return Err(Box::new(ValidationError { + problem: "an instruction uses `Device` as the memory scope, and \ + the `vulkan_memory_model` feature is enabled" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("vulkan_memory_model_device_scope"), + ])]), + vuids: &["VUID-RuntimeSpirv-vulkanMemoryModel-06265"], + ..Default::default() + })); + } + } + Scope::QueueFamily => { + if !self.device.enabled_features().vulkan_memory_model { + return Err(Box::new(ValidationError { + problem: "an instruction uses `QueueFamily` as the memory scope" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("vulkan_memory_model"), + ])]), + vuids: &["VUID-RuntimeSpirv-vulkanMemoryModel-06266"], + ..Default::default() + })); + } + } + _ => (), + } + } + + if instruction.is_group_operation() { + if !properties + .subgroup_supported_stages + .unwrap_or_default() + .contains_enum(ShaderStage::from(self.execution_model)) + { + let execution_scope = if let Some(scope) = instruction.execution_scope_id() { + get_constant(self.spirv, scope) + .and_then(|scope| Scope::try_from(scope as u32).ok()) + } else if matches!(instruction, Instruction::GroupNonUniformPartitionNV { .. }) + { + Some(Scope::Subgroup) + } else { + todo!( + "Encountered an unknown group instruction without an `execution` \ + operand. This is a Vulkano bug and should be reported.\n\ + Instruction::{:?}", + instruction + ) + }; + + if let Some(scope) = execution_scope { + if scope == Scope::Subgroup { + return Err(Box::new(ValidationError { + problem: "a group operation instruction is performed \ + with an execution scope of `Scope::Subgroup`, but \ + the `subgroup_supported_stages` device property does not contain \ + the shader stage of the entry point's execution model" + .into(), + vuids: &["VUID-RuntimeSpirv-None-06343"], + ..Default::default() + })); + } + } + } + + if !self + .device + .enabled_features() + .shader_subgroup_extended_types + { + if let Some(mut result_type_id) = instruction.result_type_id() { + if let Instruction::TypeVector { component_type, .. } = + *self.spirv.id(result_type_id).instruction() + { + result_type_id = component_type; + } + + match *self.spirv.id(result_type_id).instruction() { + Instruction::TypeInt { width, .. } => match width { + 8 => { + return Err(Box::new(ValidationError { + problem: "a group operation instruction is performed \ + on an 8-bit integer or vector" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_subgroup_extended_types"), + ])]), + vuids: &["VUID-RuntimeSpirv-None-06275"], + ..Default::default() + })); + } + 16 => { + return Err(Box::new(ValidationError { + problem: "a group operation instruction is performed \ + on a 16-bit integer or vector" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_subgroup_extended_types"), + ])]), + vuids: &["VUID-RuntimeSpirv-None-06275"], + ..Default::default() + })); + } + 64 => { + return Err(Box::new(ValidationError { + problem: "a group operation instruction is performed \ + on an 64-bit integer or vector" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_subgroup_extended_types"), + ])]), + vuids: &["VUID-RuntimeSpirv-None-06275"], + ..Default::default() + })); + } + _ => (), + }, + Instruction::TypeFloat { width: 16, .. } => { + return Err(Box::new(ValidationError { + problem: "a group operation instruction is performed \ + on an 16-bit floating point scalar or vector" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_subgroup_extended_types"), + ])]), + vuids: &["VUID-RuntimeSpirv-None-06275"], + ..Default::default() + })); + } + _ => (), + } + } + } + + if instruction.is_quad_group_operation() + && !properties + .subgroup_quad_operations_in_all_stages + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "a quad group operation instruction is performed, and \ + the `subgroup_quad_operations_in_all_stages` device property is \ + `false`, but entry point's execution model is not `Fragment` or \ + `GLCompute`" + .into(), + vuids: &["VUID-RuntimeSpirv-None-06342"], + ..Default::default() + })); + } + } + + if instruction.is_image_gather() { + if let Some(image_operands) = instruction.image_operands() { + if let Some(components) = + (image_operands.const_offset.or(image_operands.offset)) + .and_then(|offset| get_constant_maybe_composite(self.spirv, offset)) + { + for offset in components { + if offset < properties.min_texel_gather_offset as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImage*Gather` instruction is performed, but \ + its `Offset`, `ConstOffset` or `ConstOffsets` \ + image operand contains a value that is less than the \ + `min_texel_gather_offset` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpImage-06376"], + ..Default::default() + })); + } + + if offset > properties.max_texel_gather_offset as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImage*Gather` instruction is performed, but \ + its `Offset`, `ConstOffset` or `ConstOffsets` \ + image operand contains a value that is greater than the \ + `max_texel_gather_offset` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpImage-06377"], + ..Default::default() + })); + } + } + } else if let Some(elements) = image_operands + .const_offsets + .and_then(|id| get_constant_composite_composite(self.spirv, id)) + { + for components in elements { + for offset in components { + if offset < properties.min_texel_gather_offset as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImage*Gather` instruction is performed, \ + but its `Offset`, `ConstOffset` or `ConstOffsets` \ + image operand contains a value that is less than \ + the `min_texel_gather_offset` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpImage-06376"], + ..Default::default() + })); + } + + if offset > properties.max_texel_gather_offset as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImage*Gather` instruction is performed, \ + but its `Offset`, `ConstOffset` or `ConstOffsets` \ + image operand contains a value that is greater than \ + the `max_texel_gather_offset` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpImage-06377"], + ..Default::default() + })); + } + } + } + } + } + } + + if instruction.is_image_sample() || instruction.is_image_fetch() { + if let Some(image_operands) = instruction.image_operands() { + if let Some(components) = (image_operands.const_offset) + .and_then(|offset| get_constant_maybe_composite(self.spirv, offset)) + { + for offset in components { + if offset < properties.min_texel_offset as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImageSample*` or `OpImageFetch*` instruction \ + is performed, but its `ConstOffset` image operand \ + contains a value that is less than the \ + `min_texel_offset` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpImageSample-06435"], + ..Default::default() + })); + } + + if offset > properties.max_texel_offset as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImageSample*` or `OpImageFetch*` instruction \ + is performed, but its `ConstOffset` image operand \ + contains a value that is greater than the \ + `max_texel_offset` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpImageSample-06436"], + ..Default::default() + })); + } + } + } + } + } + + match *instruction { + Instruction::ReadClockKHR { scope, .. } => { + let scope = get_constant(self.spirv, scope) + .and_then(|scope| Scope::try_from(scope as u32).ok()); + + if let Some(scope) = scope { + match scope { + Scope::Subgroup => { + if self.device.enabled_features().shader_subgroup_clock { + return Err(Box::new(ValidationError { + problem: "an `OpReadClockKHR` instruction is performed \ + with a scope of `Scope::Subgroup`" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_subgroup_clock"), + ])]), + vuids: &["VUID-RuntimeSpirv-shaderSubgroupClock-06267"], + ..Default::default() + })); + } + } + Scope::Device => { + if self.device.enabled_features().shader_device_clock { + return Err(Box::new(ValidationError { + problem: "an `OpReadClockKHR` instruction is performed \ + with a scope of `Scope::Device`" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[ + Requires::Feature("shader_device_clock"), + ])]), + vuids: &["VUID-RuntimeSpirv-shaderDeviceClock-06268"], + ..Default::default() + })); + } + } + _ => (), + } + } + } + Instruction::GroupNonUniformQuadBroadcast { index, .. } => { + if !self.device.enabled_features().subgroup_broadcast_dynamic_id + && !matches!( + self.spirv.id(index).instruction(), + Instruction::Constant { .. } + ) + { + return Err(Box::new(ValidationError { + problem: "an `OpGroupNonUniformQuadBroadcast` instruction is \ + performed, and its `index` operand is not a constant" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "subgroup_broadcast_dynamic_id", + )])]), + vuids: &["VUID-RuntimeSpirv-subgroupBroadcastDynamicId-06276"], + ..Default::default() + })); + } + } + Instruction::GroupNonUniformBroadcast { id, .. } => { + if !self.device.enabled_features().subgroup_broadcast_dynamic_id + && !matches!( + self.spirv.id(id).instruction(), + Instruction::Constant { .. } + ) + { + return Err(Box::new(ValidationError { + problem: "an `OpGroupNonUniformBroadcast` instruction is \ + performed, and its `id` operand is not a constant" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "subgroup_broadcast_dynamic_id", + )])]), + vuids: &["VUID-RuntimeSpirv-subgroupBroadcastDynamicId-06277"], + ..Default::default() + })); + } + } + Instruction::CooperativeMatrixMulAddNV { .. } => { + // TODO: needs VK_NV_cooperative_matrix support + // VUID-RuntimeSpirv-OpCooperativeMatrixMulAddNV-06317 + // VUID-RuntimeSpirv-OpCooperativeMatrixMulAddNV-06318 + // VUID-RuntimeSpirv-OpCooperativeMatrixMulAddNV-06319 + // VUID-RuntimeSpirv-OpCooperativeMatrixMulAddNV-06320 + // VUID-RuntimeSpirv-OpCooperativeMatrixMulAddNV-06321 + } + Instruction::CooperativeMatrixMulAddKHR { .. } => { + // TODO: needs VK_KHR_cooperative_matrix support + // VUID-RuntimeSpirv-MSize-08975 + // VUID-RuntimeSpirv-KSize-08977 + // VUID-RuntimeSpirv-MSize-08979 + // VUID-RuntimeSpirv-MSize-08981 + // VUID-RuntimeSpirv-saturatingAccumulation-08983 + } + Instruction::EmitStreamVertex { stream } + | Instruction::EndStreamPrimitive { stream } => { + let stream = get_constant(self.spirv, stream).unwrap(); + + if u32::try_from(stream).map_or(true, |stream| { + stream + >= properties + .max_transform_feedback_streams + .unwrap_or_default() + }) { + return Err(Box::new(ValidationError { + problem: "an `OpEmitStreamVertex` or `OpEndStreamPrimitive` \ + instruction is performed, but the value of the `stream` operand \ + is not less than the `max_transform_feedback_streams` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpEmitStreamVertex-06310"], + ..Default::default() + })); + } + + if let (&Instruction::EmitStreamVertex { stream }, false) = ( + instruction, + matches!(self.output_primitives, Some(ExecutionMode::OutputPoints)), + ) { + let stream = get_constant(self.spirv, stream).unwrap() as u32; + + match self.first_emitted_stream { + Some(first_emitted_stream) => { + if stream != first_emitted_stream + && !properties + .transform_feedback_streams_lines_triangles + .unwrap_or_default() + { + return Err(Box::new(ValidationError { + problem: "the shader emits to more than one vertex \ + stream, and the entry point does not have an \ + `OutputPoints` execution mode, but the \ + `transform_feedback_streams_lines_triangles` device \ + property is `false`".into(), + vuids: &["VUID-RuntimeSpirv-transformFeedbackStreamsLinesTriangles-06311"], + ..Default::default() + })); + } + } + None => self.first_emitted_stream = Some(stream), + } + } + } + Instruction::ImageBoxFilterQCOM { box_size, .. } => { + if let Some(box_size) = get_constant_float_composite(self.spirv, box_size) { + if box_size[1] + > properties.max_box_filter_block_size.unwrap_or_default()[1] as f64 + { + return Err(Box::new(ValidationError { + problem: "an `OpImageBoxFilterQCOM` instruction is performed, but \ + the `y` component of the `box_size` operand is greater than \ + the `max_box_filter_block_size[1]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-OpImageBoxFilterQCOM-06989"], + ..Default::default() + })); + } + } + } + Instruction::EmitMeshTasksEXT { + group_count_x, + group_count_y, + group_count_z, + .. + } => { + // TODO: If the shader has multiple entry points with different execution + // models, then we really need to use the entry point's call + // tree, instead of a flat iteration over all functions. + if self.execution_model == ExecutionModel::MeshEXT { + let group_count_x = get_constant(self.spirv, group_count_x); + let group_count_y = get_constant(self.spirv, group_count_y); + let group_count_z = get_constant(self.spirv, group_count_z); + let mut product: Option = Some(1); + + if let Some(count) = group_count_x { + product = product + .zip(count.try_into().ok()) + .and_then(|(product, count)| product.checked_mul(count)); + + if u32::try_from(count).map_or(true, |count| { + count > properties.max_mesh_work_group_count.unwrap_or_default()[0] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, and \ + an `OpEmitMeshTasksEXT` instruction is performed, but \ + the value of the `group_count_x` operand is greater than \ + the `max_mesh_work_group_count[0]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07299"], + ..Default::default() + })); + } + } + + if let Some(count) = group_count_y { + product = product + .zip(count.try_into().ok()) + .and_then(|(product, count)| product.checked_mul(count)); + + if u32::try_from(count).map_or(true, |count| { + count > properties.max_mesh_work_group_count.unwrap_or_default()[1] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, and \ + an `OpEmitMeshTasksEXT` instruction is performed, but \ + the value of the `group_count_y` operand is greater than \ + the `max_mesh_work_group_count[1]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07300"], + ..Default::default() + })); + } + } + + if let Some(count) = group_count_z { + product = product + .zip(count.try_into().ok()) + .and_then(|(product, count)| product.checked_mul(count)); + + if u32::try_from(count).map_or(true, |count| { + count > properties.max_mesh_work_group_count.unwrap_or_default()[2] + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, and \ + an `OpEmitMeshTasksEXT` instruction is performed, but \ + the value of the `group_count_z` operand is greater than \ + the `max_mesh_work_group_count[2]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07301"], + ..Default::default() + })); + } + } + + if product.map_or(true, |product| { + product + > properties + .max_mesh_work_group_total_count + .unwrap_or_default() + }) { + return Err(Box::new(ValidationError { + problem: "the entry point's execution model is `TaskEXT`, and \ + an `OpEmitMeshTasksEXT` instruction is performed, but \ + the product of its `group_count_x`, `group_count_y` and \ + `group_count_z` operands is greater than the \ + `max_mesh_work_group_total_count` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-TaskEXT-07302"], + ..Default::default() + })); + } + } + } + Instruction::ColorAttachmentReadEXT { .. } => { + if !self + .device + .enabled_features() + .shader_tile_image_color_read_access + { + return Err(Box::new(ValidationError { + problem: "an `OpColorAttachmentReadEXT` instruction is performed" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_tile_image_color_read_access", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderTileImageColorReadAccess-08728"], + ..Default::default() + })); + } + } + Instruction::DepthAttachmentReadEXT { .. } => { + if !self + .device + .enabled_features() + .shader_tile_image_depth_read_access + { + return Err(Box::new(ValidationError { + problem: "an `OpDepthAttachmentReadEXT` instruction is performed" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_tile_image_depth_read_access", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderTileImageDepthReadAccess-08729"], + ..Default::default() + })); + } + } + Instruction::StencilAttachmentReadEXT { .. } => { + if !self + .device + .enabled_features() + .shader_tile_image_stencil_read_access + { + return Err(Box::new(ValidationError { + problem: "an `OpStencilAttachmentReadEXT` instruction is performed" + .into(), + requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature( + "shader_tile_image_stencil_read_access", + )])]), + vuids: &["VUID-RuntimeSpirv-shaderTileImageStencilReadAccess-08730"], + ..Default::default() + })); + } + } + Instruction::ImageBlockMatchSADQCOM { block_size, .. } + | Instruction::ImageBlockMatchSSDQCOM { block_size, .. } => { + let block_size = get_constant_composite(self.spirv, block_size); + + if let Some(block_size) = block_size { + let max_block_match_region = + properties.max_block_match_region.unwrap_or_default(); + + if block_size[0] > max_block_match_region[0] as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImageBlockMatchSADQCOM` or \ + `OpImageBlockMatchSSDQCOM` instruction is performed, but \ + the `x` component of the `block_size` operand is greater than \ + the `max_block_match_region[0]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-maxBlockMatchRegion-09225"], + ..Default::default() + })); + } + + if block_size[1] > max_block_match_region[1] as u64 { + return Err(Box::new(ValidationError { + problem: "an `OpImageBlockMatchSADQCOM` or \ + `OpImageBlockMatchSSDQCOM` instruction is performed, but \ + the `y` component of the `block_size` operand is greater than \ + the `max_block_match_region[1]` device limit" + .into(), + vuids: &["VUID-RuntimeSpirv-maxBlockMatchRegion-09225"], + ..Default::default() + })); + } + } + } + _ => (), + } + } + + Ok(()) + } +} + +// TODO: spec clarification + +// VUID-RuntimeSpirv-maxMeshOutputMemorySize-08756 +// VUID-RuntimeSpirv-maxMeshPayloadAndOutputMemorySize-08757 +// VUID-RuntimeSpirv-Location-06428 +// VUID-RuntimeSpirv-maxExecutionGraphShaderPayloadSize-09193 +// VUID-RuntimeSpirv-maxExecutionGraphShaderPayloadSize-09194 +// VUID-RuntimeSpirv-maxExecutionGraphShaderPayloadSize-09195 +// VUID-RuntimeSpirv-maxExecutionGraphShaderPayloadCount-09196 +// VUID-RuntimeSpirv-maxExecutionGraphShaderOutputNodes-09197 + +// TODO: depends on descriptor resources + +// VUID-RuntimeSpirv-None-06287 +// VUID-RuntimeSpirv-OpEntryPoint-08727 + +// TODO: requires items that are not implemented + +// VUID-RuntimeSpirv-OpTraceRayMotionNV-06367 +// VUID-RuntimeSpirv-OpHitObjectTraceRayMotionNV-07711 +// VUID-RuntimeSpirv-OpHitObjectTraceRayMotionNV-07704 +// VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07715 +// VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07716 +// VUID-RuntimeSpirv-flags-08761 +// VUID-RuntimeSpirv-OpImageBlockMatchSSDQCOM-06985 +// VUID-RuntimeSpirv-OpImageBlockMatchSSDQCOM-06986 +// VUID-RuntimeSpirv-OpImageBlockMatchSSDQCOM-06987 +// VUID-RuntimeSpirv-OpImageBlockMatchSSDQCOM-06988 +// VUID-RuntimeSpirv-OpImageBlockMatchWindow-09223 +// VUID-RuntimeSpirv-OpImageBlockMatchWindow-09224 +// VUID-RuntimeSpirv-pNext-09226 +// VUID-RuntimeSpirv-minSampleShading-08731 +// VUID-RuntimeSpirv-minSampleShading-08732 diff --git a/vulkano/src/shader/mod.rs b/vulkano/src/shader/mod.rs index 54b5253d..d4c60fe0 100644 --- a/vulkano/src/shader/mod.rs +++ b/vulkano/src/shader/mod.rs @@ -138,13 +138,23 @@ //! then if the shader accesses a descriptor in that binding, the descriptor must be initialized //! and contain a valid resource. //! -//! ## Buffers +//! ## Buffers and memory accesses //! //! - If the [`robust_buffer_access`](Features::robust_buffer_access) feature is not enabled on the //! device, then the shader must not access any values outside the range of the buffer, as //! specified when writing the descriptor set. [\[06935\]] [\[06936\]] -//! - If any `PhysicalStorageBuffer` pointers to device memory are dereferenced in the shader, then -//! they must point to valid buffer memory of the correct type. +//! - If any `PhysicalStorageBuffer` pointers to device memory are dereferenced in the shader, +//! then: +//! - The pointer must point to valid memory of the correct type. +//! - The pointer must be aligned to a multiple of the largest scalar type within the type that +//! it points to. [\[06314\]] +//! - If the instruction has `Aligned` as one of its memory operands, the pointer must be aligned +//! to the specified alignment. [\[06315\]] +//! - For `OpCooperativeMatrixLoadKHR`, `OpCooperativeMatrixStoreKHR`, `OpCooperativeMatrixLoadNV` +//! and `OpCooperativeMatrixStoreNV` instructions, the `Pointer` and `Stride` operands must both +//! be aligned to the minimum of either 16 bytes or the number of bytes per row/column of the +//! matrix (depending on the `ColumnMajor` and `RowMajor` decorations). [\[06324\]] +//! [\[08986\]] //! //! ## Image views and buffer views //! @@ -160,6 +170,8 @@ //! only if the format of the bound image view or buffer view also has a 64-bit component. //! Otherwise, it must have a `Width` of 32. [\[04470\]] [\[04471\]] [\[04472\]] //! [\[04473\]] +//! - The [`samples`](Image::samples) of the underlying image of the bound image view must match +//! the `MS` operand of the `OpImageType`. [\[08725\]] [\[08726\]] //! - For a storage image/texel buffer declared with `OpTypeImage` with an `Unknown` format: //! - If it is written to in the shader, the format of the bound image view or buffer view must //! have the [`FormatFeatures::STORAGE_WRITE_WITHOUT_FORMAT`] format feature. [\[07027\]] @@ -209,54 +221,197 @@ //! - The sampler must not be used with the `ConstOffset` or `Offset` image operands. //! [\[06551\]] //! -//! ## Acceleration structures +//! ## Mesh shading //! +//! - If the shader declares the `OutputPoints` execution mode with a value greater than 0, and the +//! [`maintenance5`](Features::maintenance5) feature is not enabled on the device, then the +//! shader must write to a variable decorated with `PointSize` for each output point. +//! [\[09218\]] +//! +//! For `OpSetMeshOutputsEXT` instructions: +//! +//! - The `Vertex Count` operand must be less than or equal to the value declared with the shader's +//! `OutputVertices` execution mode. [\[07332\]] +//! - The `Primitive Count` operand must be less than or equal to the value declared with the +//! shader's `OutputPrimitivesEXT` execution mode. [\[07333\]] +//! +//! ## Acceleration structures, ray queries and ray tracing +//! +//! - Acceleration structures that are used as operands to an instruction must have been built as a +//! top-level acceleration structure. [\[06352\]] [\[06359\]] [\[06365\]] [\[07709\]] //! - In any top-level acceleration structure, the pointers that refer to the contained -//! bottom-level acceleration structure instances must point to valid acceleration structures. +//! bottom-level acceleration structure instances must point to valid bottom-level acceleration +//! structures. +//! +//! For `OpRayQueryInitializeKHR` and `OpTraceRayKHR` instructions: +//! +//! - The `Rayflags` operand must not contain more than one of: +//! - `SkipTrianglesKHR`, `CullBackFacingTrianglesKHR` and `CullFrontFacingTrianglesKHR` +//! [\[06889\]] [\[06892\]] +//! - `SkipTrianglesKHR` and `SkipAABBsKHR` [\[06890\]] [\[06552\]] [\[07712\]] +//! - `OpaqueKHR`, `NoOpaqueKHR`, `CullOpaqueKHR`, and `CullNoOpaqueKHR` [\[06891\]] +//! [\[06893\]] +//! - The `RayOrigin` and `RayDirection` operands must not contain infinite or NaN values. +//! [\[06348\]] [\[06351\]] [\[06355\]] [\[06358\]] +//! - The `RayTmin` and `RayTmax` operands must not contain negative or NaN values, and `RayTmin` +//! must be less than or equal to `RayTmax`. [\[06349\]] [\[06350\]] [\[06351\]] +//! [\[06356\]] [\[06357\]] [\[06358\]] +//! +//! For `OpRayQueryGenerateIntersectionKHR` instructions: +//! +//! - The `Hit T` operand must be greater than or equal to the value that would be returned by +//! `OpRayQueryGetRayTMinKHR`. [\[06353\]] +//! - The `Hit T` operand must be less than or equal to the value that would be returned by +//! `OpRayQueryGetIntersectionTKHR` for the current committed intersection. +//! [\[06353\]] +//! +//! For `OpReportIntersectionKHR` instructions: +//! +//! - The `Hit Kind` operand must be between 0 and 127 inclusive. [\[06998\]] +//! +//! ## Dynamically uniform values and control flow +//! +//! In a shader, a value (expression, variable) is *[dynamically uniform]* if its value is the same +//! for all shader invocations within an *invocation group*. What counts as an invocation group +//! depends on the type of shader being executed: +//! +//! - For compute, task and mesh shaders, an invocation group is the same as the (local) workgroup. +//! A single `dispatch` command value spawns one distinct invocation group for every element in +//! the product of the given `group_counts` argument. +//! - For all other graphics shaders, an invocation group is all shaders invoked by a single draw +//! command. For indirect draws, each element of the indirect buffer creates one draw call. +//! - For ray tracing shaders, an invocation group is an implementation-dependent subset of the +//! shaders invoked by a single ray tracing command. +//! +//! Vulkan and SPIR-V assume that certain values within a shader are dynamically uniform, and will +//! optimize the generated shader code accordingly. If such a value is not actually dynamically +//! uniform, this results in undefined behavior. This concerns the following values: +//! +//! - The index into an arrayed descriptor binding. If the index is not dynamically uniform, you +//! must explicitly mark it with the `NonUniform` decoration in SPIR-V, or the `nonuniformEXT` +//! function in GLSL. [\[06274\]] +//! - The `Index` argument of the `OpGroupNonUniformQuadBroadcast` instruction. +//! [\[06276\]] +//! - The `Id` argument of the `OpGroupNonUniformBroadcast` instruction. [\[06277\]] +//! - The arguments of the `OpEmitMeshTasksEXT` and `OpSetMeshOutputsEXT` instructions. +//! [\[07117\]] [\[07118\]] +//! - The `Texture Sampled Image` and `Weight Image` arguments of the `OpImageWeightedSampleQCOM` +//! instruction. [\[06979\]] +//! - The `Texture Sampled Image`, `Reference Sampled Image` and `Block Size` arguments of the +//! `OpImageBlockMatchSADQCOM` and `OpImageBlockMatchSSDQCOM` instructions. +//! [\[06982\]] +//! - The `Sampled Texture Image` and `Box Size` arguments of the `OpImageBoxFilterQCOM` +//! instruction. [\[06990\]] +//! - The `Target Sampled Image`, `Reference Sampled Image` and `Block Size` arguments of any +//! `OpImageBlockMatchWindow*QCOM` or `OpImageBlockMatchGather*QCOM` instructions. +//! [\[09219\]] +//! +//! Some operations have specific requirements for control flow within the shader: +//! +//! - The `OpEmitMeshTasksEXT` and `OpSetMeshOutputsEXT` instructions must be executed uniformly +//! within the invocation group. That means that, either all shader invocations within the +//! invocation group must execute the instruction, or none of them must execute it. +//! [\[07117\]] [\[07118\]] +//! - If the `PointSize` built-in is written to, then all execution paths must write to it. +//! [\[09190\]] //! //! [alignment rules]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/chap15.html#interfaces-resources-layout //! [`GL_EXT_scalar_block_layout`]: https://github.com/KhronosGroup/GLSL/blob/master/extensions/ext/GL_EXT_scalar_block_layout.txt -//! [`scalar_block_layout`]: crate::device::Features::scalar_block_layout -//! [`uniform_buffer_standard_layout`]: crate::device::Features::uniform_buffer_standard_layout -//! [\[06935\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-uniformBuffers-06935 -//! [\[06936\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-storageBuffers-06936 -//! [\[07752\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-viewType-07752 -//! [\[07753\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-format-07753 +//! [`scalar_block_layout`]: Features::scalar_block_layout +//! [`uniform_buffer_standard_layout`]: Features::uniform_buffer_standard_layout +//! [dynamically uniform]: https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_uniformity +//! [\[02691\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-02691 +//! [\[02692\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-02692 +//! [\[02694\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-filterCubic-02694 +//! [\[02695\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-filterCubicMinmax-02695 //! [\[04469\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpImageWrite-04469 -//! [\[08795\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpImageWrite-08795 -//! [\[08796\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpImageWrite-08796 //! [\[04470\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-SampledType-04470 //! [\[04471\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-SampledType-04471 //! [\[04472\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-SampledType-04472 //! [\[04473\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-SampledType-04473 +//! [\[04553\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-magFilter-04553 +//! [\[04770\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-mipmapMode-04770 +//! [\[06274\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-NonUniform-06274 +//! [\[06276\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-subgroupBroadcastDynamicId-06276 +//! [\[06277\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-subgroupBroadcastDynamicId-06277 +//! [\[06314\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-PhysicalStorageBuffer64-06314 +//! [\[06315\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-PhysicalStorageBuffer64-06315 +//! [\[06324\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpCooperativeMatrixLoadNV-06324 +//! [\[06348\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06348 +//! [\[06349\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06349 +//! [\[06350\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06350 +//! [\[06351\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06351 +//! [\[06352\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06352 +//! [\[06353\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryGenerateIntersectionKHR-06353 +//! [\[06355\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06355 +//! [\[06356\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06356 +//! [\[06357\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06357 +//! [\[06358\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06358 +//! [\[06359\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06359 +//! [\[06361\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayMotionNV-06361 +//! [\[06362\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayMotionNV-06362 +//! [\[06363\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayMotionNV-06363 +//! [\[06364\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayMotionNV-06364 +//! [\[06365\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayMotionNV-06365 +//! [\[06366\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayMotionNV-06366 +//! [\[06479\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-06479 +//! [\[06550\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-06550 +//! [\[06551\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-ConstOffset-06551 +//! [\[06552\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06552 +//! [\[06889\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06889 +//! [\[06890\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06890 +//! [\[06891\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06891 +//! [\[06892\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06892 +//! [\[06893\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpTraceRayKHR-06893 +//! [\[06935\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-uniformBuffers-06935 +//! [\[06936\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-storageBuffers-06936 +//! [\[06979\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpImageWeightedSampleQCOM-06979 +//! [\[06982\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpImageBlockMatchSADQCOM-06982 +//! [\[06990\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpImageBoxFilterQCOM-06990 +//! [\[06998\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpReportIntersectionKHR-06998 //! [\[07027\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpTypeImage-07027 //! [\[07029\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpTypeImage-07029 //! [\[07028\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpTypeImage-07028 //! [\[07030\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpTypeImage-07030 -//! [\[02691\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-02691 +//! [\[07117\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-TaskEXT-07117 +//! [\[07118\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-MeshEXT-07118 +//! [\[07332\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-MeshEXT-07332 +//! [\[07333\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-MeshEXT-07333 +//! [\[07705\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07705 +//! [\[07706\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07706 +//! [\[07707\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07707 +//! [\[07708\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07708 +//! [\[07709\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayMotionNV-07709 +//! [\[07710\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07710 +//! [\[07712\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07712 +//! [\[07713\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07713 +//! [\[07714\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpHitObjectTraceRayNV-07714 +//! [\[07752\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-viewType-07752 +//! [\[07753\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-format-07753 //! [\[07888\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-07888 -//! [\[04553\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-magFilter-04553 -//! [\[04770\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-mipmapMode-04770 -//! [\[02692\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-02692 -//! [\[02694\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-filterCubic-02694 -//! [\[02695\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-filterCubicMinmax-02695 -//! [\[06479\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-06479 //! [\[08609\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-08609 //! [\[08610\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-08610 //! [\[08611\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-08611 -//! [\[06550\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-None-06550 -//! [\[06551\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-ConstOffset-06551 +//! [\[08725\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-samples-08725 +//! [\[08726\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-samples-08726 +//! [\[08795\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpImageWrite-08795 +//! [\[08796\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-vkCmdDispatch-OpImageWrite-08796 +//! [\[08986\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpCooperativeMatrixLoadKHR-08986 +//! [\[09190\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-maintenance5-09190 +//! [\[09218\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-MeshEXT-09218 +//! [\[09219\]]: https://registry.khronos.org/vulkan/specs/1.3-extensions/html/vkspec.html#VUID-RuntimeSpirv-OpImageBlockMatchWindow-09219 use self::spirv::{Id, Instruction}; #[cfg(doc)] use crate::{ + acceleration_structure::BuildAccelerationStructureFlags, descriptor_set::layout::DescriptorBindingFlags, - device::{physical::PhysicalDevice, Features}, + device::{physical::PhysicalDevice, Features, Properties}, format::FormatFeatures, image::{ sampler::{Filter, Sampler, SamplerCreateInfo, SamplerMipmapMode, SamplerReductionMode}, view::ImageView, - ImageFormatProperties, + Image, ImageFormatProperties, }, }; use crate::{ diff --git a/vulkano/src/shader/reflect.rs b/vulkano/src/shader/reflect.rs index 6b130ed9..607778f6 100644 --- a/vulkano/src/shader/reflect.rs +++ b/vulkano/src/shader/reflect.rs @@ -15,6 +15,7 @@ use crate::{ }; use ahash::{HashMap, HashSet}; use half::f16; +use smallvec::{smallvec, SmallVec}; use std::borrow::Cow; /// Returns an iterator over all entry points in `spirv`, with information about the entry point. @@ -1450,6 +1451,96 @@ pub(crate) fn get_constant(spirv: &Spirv, id: Id) -> Option { } } +pub(crate) fn get_constant_composite(spirv: &Spirv, id: Id) -> Option> { + match spirv.id(id).instruction() { + Instruction::ConstantComposite { constituents, .. } => Some( + constituents + .iter() + .map(|&id| match spirv.id(id).instruction() { + Instruction::Constant { value, .. } => match value.len() { + 1 => value[0] as u64, + 2 => value[0] as u64 | (value[1] as u64) << 32, + _ => panic!("constant {} is larger than 64 bits", id), + }, + _ => unreachable!(), + }) + .collect(), + ), + _ => None, + } +} + +pub(crate) fn get_constant_float_composite(spirv: &Spirv, id: Id) -> Option> { + match spirv.id(id).instruction() { + Instruction::ConstantComposite { constituents, .. } => Some( + constituents + .iter() + .map(|&id| match spirv.id(id).instruction() { + Instruction::Constant { value, .. } => match value.len() { + 1 => f32::from_bits(value[0]) as f64, + 2 => f64::from_bits(value[0] as u64 | (value[1] as u64) << 32), + _ => panic!("constant {} is larger than 64 bits", id), + }, + _ => unreachable!(), + }) + .collect(), + ), + _ => None, + } +} + +pub(crate) fn get_constant_maybe_composite(spirv: &Spirv, id: Id) -> Option> { + match spirv.id(id).instruction() { + Instruction::Constant { value, .. } => match value.len() { + 1 => Some(smallvec![value[0] as u64]), + 2 => Some(smallvec![value[0] as u64 | (value[1] as u64) << 32]), + _ => panic!("constant {} is larger than 64 bits", id), + }, + Instruction::ConstantComposite { constituents, .. } => Some( + constituents + .iter() + .map(|&id| match spirv.id(id).instruction() { + Instruction::Constant { value, .. } => match value.len() { + 1 => value[0] as u64, + 2 => value[0] as u64 | (value[1] as u64) << 32, + _ => panic!("constant {} is larger than 64 bits", id), + }, + _ => unreachable!(), + }) + .collect(), + ), + _ => None, + } +} + +pub(crate) fn get_constant_composite_composite( + spirv: &Spirv, + id: Id, +) -> Option; 4]>> { + match spirv.id(id).instruction() { + Instruction::ConstantComposite { constituents, .. } => Some( + constituents + .iter() + .map(|&id| match spirv.id(id).instruction() { + Instruction::ConstantComposite { constituents, .. } => constituents + .iter() + .map(|&id| match spirv.id(id).instruction() { + Instruction::Constant { value, .. } => match value.len() { + 1 => value[0] as u64, + 2 => value[0] as u64 | (value[1] as u64) << 32, + _ => panic!("constant {} is larger than 64 bits", id), + }, + _ => unreachable!(), + }) + .collect(), + _ => unreachable!(), + }) + .collect(), + ), + _ => None, + } +} + #[cfg(test)] mod tests { use super::{HashMap, PushConstantRange, ShaderStages, Version};