Minor reorganization of shader validation code (#2445)

* Minor reorganization of shader validation code

* Track call tree of SPIR-V functions

* Restore fixme

vulkano/src/pipeline/graphics/color_blend.rs

@@ -17,7 +17,7 @@ use crate::{
     device::Device,
     format::Format,
     macros::{vulkan_bitflags, vulkan_enum},
-    pipeline::{ShaderInterfaceLocationInfo, ShaderInterfaceLocationWidth},
+    pipeline::inout_interface::{ShaderInterfaceLocationInfo, ShaderInterfaceLocationWidth},
     Requires, RequiresAllOf, RequiresOneOf, ValidationError,
 };
 use ahash::HashMap;

vulkano/src/pipeline/graphics/mod.rs

@@ -90,9 +90,11 @@ use self::{
     viewport::ViewportState,
 };
 use super::{
-    cache::PipelineCache, shader::validate_interfaces_compatible, DynamicState, Pipeline,
+    cache::PipelineCache,
-    PipelineBindPoint, PipelineCreateFlags, PipelineLayout, PipelineShaderStageCreateInfo,
+    inout_interface::{shader_interface_location_info, ShaderInterfaceLocationInfo},
-    ShaderInterfaceLocationInfo,
+    shader::inout_interface::validate_interfaces_compatible,
+    DynamicState, Pipeline, PipelineBindPoint, PipelineCreateFlags, PipelineLayout,
+    PipelineShaderStageCreateInfo,
 };
 use crate::{
     device::{Device, DeviceOwned, DeviceOwnedDebugWrapper},
@@ -971,7 +973,7 @@ impl GraphicsPipeline {
             match entry_point_info.execution_model {
                 ExecutionModel::Vertex => {
                     if vertex_input_state.is_none() {
-                        required_vertex_inputs = Some(super::shader_interface_location_info(
+                        required_vertex_inputs = Some(shader_interface_location_info(
                             entry_point.module().spirv(),
                             entry_point.id(),
                             StorageClass::Input,
@@ -2415,7 +2417,7 @@ impl GraphicsPipelineCreateInfo {
         */
 
         if let (Some(vertex_stage), Some(vertex_input_state)) = (vertex_stage, vertex_input_state) {
-            let required_vertex_inputs = super::shader_interface_location_info(
+            let required_vertex_inputs = shader_interface_location_info(
                 vertex_stage.entry_point.module().spirv(),
                 vertex_stage.entry_point.id(),
                 StorageClass::Input,
@@ -2506,7 +2508,7 @@ impl GraphicsPipelineCreateInfo {
         if let (Some(fragment_stage), Some(color_blend_state), Some(subpass)) =
             (fragment_stage, color_blend_state, subpass)
         {
-            let fragment_shader_outputs = super::shader_interface_location_info(
+            let fragment_shader_outputs = shader_interface_location_info(
                 fragment_stage.entry_point.module().spirv(),
                 fragment_stage.entry_point.id(),
                 StorageClass::Output,

vulkano/src/pipeline/graphics/vertex_input/mod.rs

@@ -101,7 +101,7 @@ use super::color_blend::ColorComponents;
 use crate::{
     device::Device,
     format::{Format, FormatFeatures},
-    pipeline::{ShaderInterfaceLocationInfo, ShaderInterfaceLocationWidth},
+    pipeline::inout_interface::{ShaderInterfaceLocationInfo, ShaderInterfaceLocationWidth},
     DeviceSize, Requires, RequiresAllOf, RequiresOneOf, ValidationError,
 };
 use ahash::HashMap;

vulkano/src/pipeline/shader/mod.rs

@@ -0,0 +1,517 @@
+use crate::{
+    device::Device,
+    macros::vulkan_bitflags,
+    shader::{
+        spirv::{BuiltIn, Decoration, ExecutionMode, Id, Instruction},
+        EntryPoint, ShaderStage,
+    },
+    Requires, RequiresAllOf, RequiresOneOf, ValidationError,
+};
+
+pub(crate) mod inout_interface;
+
+/// Specifies a single shader stage when creating a pipeline.
+#[derive(Clone, Debug)]
+pub struct PipelineShaderStageCreateInfo {
+    /// Additional properties of the shader stage.
+    ///
+    /// The default value is empty.
+    pub flags: PipelineShaderStageCreateFlags,
+
+    /// The shader entry point for the stage, which includes any specialization constants.
+    ///
+    /// There is no default value.
+    pub entry_point: EntryPoint,
+
+    /// The required subgroup size.
+    ///
+    /// Requires [`subgroup_size_control`](crate::device::Features::subgroup_size_control). The
+    /// shader stage must be included in
+    /// [`required_subgroup_size_stages`](crate::device::Properties::required_subgroup_size_stages).
+    /// Subgroup size must be power of 2 and within
+    /// [`min_subgroup_size`](crate::device::Properties::min_subgroup_size)
+    /// and [`max_subgroup_size`](crate::device::Properties::max_subgroup_size).
+    ///
+    /// For compute shaders, `max_compute_workgroup_subgroups * required_subgroup_size` must be
+    /// greater than or equal to `workgroup_size.x * workgroup_size.y * workgroup_size.z`.
+    ///
+    /// The default value is None.
+    pub required_subgroup_size: Option<u32>,
+
+    pub _ne: crate::NonExhaustive,
+}
+
+impl PipelineShaderStageCreateInfo {
+    /// Returns a `PipelineShaderStageCreateInfo` with the specified `entry_point`.
+    #[inline]
+    pub fn new(entry_point: EntryPoint) -> Self {
+        Self {
+            flags: PipelineShaderStageCreateFlags::empty(),
+            entry_point,
+            required_subgroup_size: None,
+            _ne: crate::NonExhaustive(()),
+        }
+    }
+
+    pub(crate) fn validate(&self, device: &Device) -> Result<(), Box<ValidationError>> {
+        let &Self {
+            flags,
+            ref entry_point,
+            required_subgroup_size,
+            _ne: _,
+        } = self;
+
+        let spirv = entry_point.module().spirv();
+        let properties = device.physical_device().properties();
+
+        flags.validate_device(device).map_err(|err| {
+            err.add_context("flags")
+                .set_vuids(&["VUID-VkPipelineShaderStageCreateInfo-flags-parameter"])
+        })?;
+
+        let execution_model = entry_point.info().execution_model;
+        let stage_enum = ShaderStage::from(execution_model);
+
+        stage_enum.validate_device(device).map_err(|err| {
+            err.add_context("entry_point.info().execution")
+                .set_vuids(&["VUID-VkPipelineShaderStageCreateInfo-stage-parameter"])
+        })?;
+
+        // VUID-VkPipelineShaderStageCreateInfo-pName-00707
+        // Guaranteed by definition of `EntryPoint`.
+
+        // TODO:
+        // VUID-VkPipelineShaderStageCreateInfo-maxClipDistances-00708
+        // VUID-VkPipelineShaderStageCreateInfo-maxCullDistances-00709
+        // VUID-VkPipelineShaderStageCreateInfo-maxCombinedClipAndCullDistances-00710
+        // VUID-VkPipelineShaderStageCreateInfo-maxSampleMaskWords-00711
+        // VUID-VkPipelineShaderStageCreateInfo-stage-02596
+        // VUID-VkPipelineShaderStageCreateInfo-stage-02597
+
+        match stage_enum {
+            ShaderStage::Vertex => {
+                // VUID-VkPipelineShaderStageCreateInfo-stage-00712
+                // TODO:
+            }
+            ShaderStage::TessellationControl | ShaderStage::TessellationEvaluation => {
+                if !device.enabled_features().tessellation_shader {
+                    return Err(Box::new(ValidationError {
+                        context: "entry_point".into(),
+                        problem: "specifies a `ShaderStage::TessellationControl` or \
+                            `ShaderStage::TessellationEvaluation` entry point"
+                            .into(),
+                        requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature(
+                            "tessellation_shader",
+                        )])]),
+                        vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-00705"],
+                    }));
+                }
+
+                // VUID-VkPipelineShaderStageCreateInfo-stage-00713
+                // TODO:
+            }
+            ShaderStage::Geometry => {
+                if !device.enabled_features().geometry_shader {
+                    return Err(Box::new(ValidationError {
+                        context: "entry_point".into(),
+                        problem: "specifies a `ShaderStage::Geometry` entry point".into(),
+                        requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature(
+                            "geometry_shader",
+                        )])]),
+                        vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-00704"],
+                    }));
+                }
+
+                // TODO:
+                // VUID-VkPipelineShaderStageCreateInfo-stage-00714
+                // VUID-VkPipelineShaderStageCreateInfo-stage-00715
+            }
+            ShaderStage::Fragment => {
+                // TODO:
+                // VUID-VkPipelineShaderStageCreateInfo-stage-00718
+                // VUID-VkPipelineShaderStageCreateInfo-stage-06685
+                // VUID-VkPipelineShaderStageCreateInfo-stage-06686
+            }
+            ShaderStage::Compute => (),
+            ShaderStage::Raygen => (),
+            ShaderStage::AnyHit => (),
+            ShaderStage::ClosestHit => (),
+            ShaderStage::Miss => (),
+            ShaderStage::Intersection => (),
+            ShaderStage::Callable => (),
+            ShaderStage::Mesh => {
+                if !device.enabled_features().mesh_shader {
+                    return Err(Box::new(ValidationError {
+                        context: "entry_point".into(),
+                        problem: "specifies a `ShaderStage::Mesh` entry point".into(),
+                        requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature(
+                            "mesh_shader",
+                        )])]),
+                        vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-02091"],
+                    }));
+                }
+            }
+            ShaderStage::Task => {
+                if !device.enabled_features().task_shader {
+                    return Err(Box::new(ValidationError {
+                        context: "entry_point".into(),
+                        problem: "specifies a `ShaderStage::Task` entry point".into(),
+                        requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature(
+                            "task_shader",
+                        )])]),
+                        vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-02092"],
+                    }));
+                }
+            }
+            ShaderStage::SubpassShading => (),
+        }
+
+        let mut clip_distance_array_size = 0;
+        let mut cull_distance_array_size = 0;
+
+        for instruction in spirv.decorations() {
+            if let Instruction::Decorate {
+                target,
+                decoration: Decoration::BuiltIn { built_in },
+            } = *instruction
+            {
+                let variable_array_size = |variable| {
+                    let result_type_id = match *spirv.id(variable).instruction() {
+                        Instruction::Variable { result_type_id, .. } => result_type_id,
+                        _ => return None,
+                    };
+
+                    let length = match *spirv.id(result_type_id).instruction() {
+                        Instruction::TypeArray { length, .. } => length,
+                        _ => return None,
+                    };
+
+                    let value = match *spirv.id(length).instruction() {
+                        Instruction::Constant { ref value, .. } => {
+                            if value.len() > 1 {
+                                u32::MAX
+                            } else {
+                                value[0]
+                            }
+                        }
+                        _ => return None,
+                    };
+
+                    Some(value)
+                };
+
+                match built_in {
+                    BuiltIn::ClipDistance => {
+                        clip_distance_array_size = variable_array_size(target).unwrap();
+
+                        if clip_distance_array_size > properties.max_clip_distances {
+                            return Err(Box::new(ValidationError {
+                                context: "entry_point".into(),
+                                problem: "the number of elements in the `ClipDistance` built-in \
+                                    variable is greater than the \
+                                    `max_clip_distances` device limit"
+                                    .into(),
+                                vuids: &[
+                                    "VUID-VkPipelineShaderStageCreateInfo-maxClipDistances-00708",
+                                ],
+                                ..Default::default()
+                            }));
+                        }
+                    }
+                    BuiltIn::CullDistance => {
+                        cull_distance_array_size = variable_array_size(target).unwrap();
+
+                        if cull_distance_array_size > properties.max_cull_distances {
+                            return Err(Box::new(ValidationError {
+                                context: "entry_point".into(),
+                                problem: "the number of elements in the `CullDistance` built-in \
+                                    variable is greater than the \
+                                    `max_cull_distances` device limit"
+                                    .into(),
+                                vuids: &[
+                                    "VUID-VkPipelineShaderStageCreateInfo-maxCullDistances-00709",
+                                ],
+                                ..Default::default()
+                            }));
+                        }
+                    }
+                    BuiltIn::SampleMask => {
+                        if variable_array_size(target).unwrap() > properties.max_sample_mask_words {
+                            return Err(Box::new(ValidationError {
+                                context: "entry_point".into(),
+                                problem: "the number of elements in the `SampleMask` built-in \
+                                    variable is greater than the \
+                                    `max_sample_mask_words` device limit"
+                                    .into(),
+                                vuids: &[
+                                    "VUID-VkPipelineShaderStageCreateInfo-maxSampleMaskWords-00711",
+                                ],
+                                ..Default::default()
+                            }));
+                        }
+                    }
+                    _ => (),
+                }
+            }
+        }
+
+        if clip_distance_array_size
+            .checked_add(cull_distance_array_size)
+            .map_or(true, |sum| {
+                sum > properties.max_combined_clip_and_cull_distances
+            })
+        {
+            return Err(Box::new(ValidationError {
+                context: "entry_point".into(),
+                problem: "the sum of the number of elements in the `ClipDistance` and \
+                    `CullDistance` built-in variables is greater than the \
+                    `max_combined_clip_and_cull_distances` device limit"
+                    .into(),
+                vuids: &[
+                    "VUID-VkPipelineShaderStageCreateInfo-maxCombinedClipAndCullDistances-00710",
+                ],
+                ..Default::default()
+            }));
+        }
+
+        let entry_point_function = spirv.function(entry_point.id());
+
+        for instruction in entry_point_function.execution_modes() {
+            if let Instruction::ExecutionMode {
+                mode: ExecutionMode::OutputVertices { vertex_count },
+                ..
+            } = *instruction
+            {
+                match stage_enum {
+                    ShaderStage::TessellationControl | ShaderStage::TessellationEvaluation => {
+                        if vertex_count == 0 {
+                            return Err(Box::new(ValidationError {
+                                context: "entry_point".into(),
+                                problem: "the `vertex_count` of the \
+                                    `ExecutionMode::OutputVertices` is zero"
+                                    .into(),
+                                vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-00713"],
+                                ..Default::default()
+                            }));
+                        }
+
+                        if vertex_count > properties.max_tessellation_patch_size {
+                            return Err(Box::new(ValidationError {
+                                context: "entry_point".into(),
+                                problem: "the `vertex_count` of the \
+                                    `ExecutionMode::OutputVertices` is greater than the \
+                                    `max_tessellation_patch_size` device limit"
+                                    .into(),
+                                vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-00713"],
+                                ..Default::default()
+                            }));
+                        }
+                    }
+                    ShaderStage::Geometry => {
+                        if vertex_count == 0 {
+                            return Err(Box::new(ValidationError {
+                                context: "entry_point".into(),
+                                problem: "the `vertex_count` of the \
+                                    `ExecutionMode::OutputVertices` is zero"
+                                    .into(),
+                                vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-00714"],
+                                ..Default::default()
+                            }));
+                        }
+
+                        if vertex_count > properties.max_geometry_output_vertices {
+                            return Err(Box::new(ValidationError {
+                                context: "entry_point".into(),
+                                problem: "the `vertex_count` of the \
+                                    `ExecutionMode::OutputVertices` is greater than the \
+                                    `max_geometry_output_vertices` device limit"
+                                    .into(),
+                                vuids: &["VUID-VkPipelineShaderStageCreateInfo-stage-00714"],
+                                ..Default::default()
+                            }));
+                        }
+                    }
+                    _ => (),
+                }
+            }
+        }
+
+        let local_size = (spirv
+            .decorations()
+            .iter()
+            .find_map(|instruction| match *instruction {
+                Instruction::Decorate {
+                    target,
+                    decoration:
+                        Decoration::BuiltIn {
+                            built_in: BuiltIn::WorkgroupSize,
+                        },
+                } => {
+                    let constituents: &[Id; 3] = match *spirv.id(target).instruction() {
+                        Instruction::ConstantComposite {
+                            ref constituents, ..
+                        } => constituents.as_slice().try_into().unwrap(),
+                        _ => unreachable!(),
+                    };
+
+                    let local_size = constituents.map(|id| match *spirv.id(id).instruction() {
+                        Instruction::Constant { ref value, .. } => {
+                            assert!(value.len() == 1);
+                            value[0]
+                        }
+                        _ => unreachable!(),
+                    });
+
+                    Some(local_size)
+                }
+                _ => None,
+            }))
+        .or_else(|| {
+            entry_point_function
+                .execution_modes()
+                .iter()
+                .find_map(|instruction| match *instruction {
+                    Instruction::ExecutionMode {
+                        mode:
+                            ExecutionMode::LocalSize {
+                                x_size,
+                                y_size,
+                                z_size,
+                            },
+                        ..
+                    } => Some([x_size, y_size, z_size]),
+                    Instruction::ExecutionModeId {
+                        mode:
+                            ExecutionMode::LocalSizeId {
+                                x_size,
+                                y_size,
+                                z_size,
+                            },
+                        ..
+                    } => Some([x_size, y_size, z_size].map(
+                        |id| match *spirv.id(id).instruction() {
+                            Instruction::Constant { ref value, .. } => {
+                                assert!(value.len() == 1);
+                                value[0]
+                            }
+                            _ => unreachable!(),
+                        },
+                    )),
+                    _ => None,
+                })
+        })
+        .unwrap_or_default();
+        let workgroup_size = local_size
+            .into_iter()
+            .try_fold(1, u32::checked_mul)
+            .unwrap();
+
+        if let Some(required_subgroup_size) = required_subgroup_size {
+            if !device.enabled_features().subgroup_size_control {
+                return Err(Box::new(ValidationError {
+                    context: "required_subgroup_size".into(),
+                    problem: "is `Some`".into(),
+                    requires_one_of: RequiresOneOf(&[RequiresAllOf(&[Requires::Feature(
+                        "subgroup_size_control",
+                    )])]),
+                    vuids: &["VUID-VkPipelineShaderStageCreateInfo-pNext-02755"],
+                }));
+            }
+
+            if !properties
+                .required_subgroup_size_stages
+                .unwrap_or_default()
+                .contains_enum(stage_enum)
+            {
+                return Err(Box::new(ValidationError {
+                    problem: "`required_subgroup_size` is `Some`, but the \
+                        `required_subgroup_size_stages` device property does not contain the \
+                        shader stage of `entry_point`"
+                        .into(),
+                    vuids: &["VUID-VkPipelineShaderStageCreateInfo-pNext-02755"],
+                    ..Default::default()
+                }));
+            }
+
+            if !required_subgroup_size.is_power_of_two() {
+                return Err(Box::new(ValidationError {
+                    context: "required_subgroup_size".into(),
+                    problem: "is not a power of 2".into(),
+                    vuids: &["VUID-VkPipelineShaderStageRequiredSubgroupSizeCreateInfo-requiredSubgroupSize-02760"],
+                    ..Default::default()
+                }));
+            }
+
+            if required_subgroup_size < properties.min_subgroup_size.unwrap_or(1) {
+                return Err(Box::new(ValidationError {
+                    context: "required_subgroup_size".into(),
+                    problem: "is less than the `min_subgroup_size` device limit".into(),
+                    vuids: &["VUID-VkPipelineShaderStageRequiredSubgroupSizeCreateInfo-requiredSubgroupSize-02761"],
+                    ..Default::default()
+                }));
+            }
+
+            if required_subgroup_size > properties.max_subgroup_size.unwrap_or(128) {
+                return Err(Box::new(ValidationError {
+                    context: "required_subgroup_size".into(),
+                    problem: "is greater than the `max_subgroup_size` device limit".into(),
+                    vuids: &["VUID-VkPipelineShaderStageRequiredSubgroupSizeCreateInfo-requiredSubgroupSize-02762"],
+                    ..Default::default()
+                }));
+            }
+
+            if matches!(
+                stage_enum,
+                ShaderStage::Compute | ShaderStage::Mesh | ShaderStage::Task
+            ) && workgroup_size
+                > required_subgroup_size
+                    .checked_mul(
+                        properties
+                            .max_compute_workgroup_subgroups
+                            .unwrap_or_default(),
+                    )
+                    .unwrap_or(u32::MAX)
+            {
+                return Err(Box::new(ValidationError {
+                    problem: "the product of the `local_size_x`, `local_size_y` and \
+                        `local_size_z` of `entry_point` is greater than the the product \
+                        of `required_subgroup_size` and the \
+                        `max_compute_workgroup_subgroups` device limit"
+                        .into(),
+                    vuids: &["VUID-VkPipelineShaderStageCreateInfo-pNext-02756"],
+                    ..Default::default()
+                }));
+            }
+        }
+
+        // TODO:
+        // VUID-VkPipelineShaderStageCreateInfo-module-08987
+
+        Ok(())
+    }
+}
+
+vulkan_bitflags! {
+    #[non_exhaustive]
+
+    /// Flags specifying additional properties of a pipeline shader stage.
+    PipelineShaderStageCreateFlags = PipelineShaderStageCreateFlags(u32);
+
+    /* TODO: enable
+    // TODO: document
+    ALLOW_VARYING_SUBGROUP_SIZE = ALLOW_VARYING_SUBGROUP_SIZE
+    RequiresOneOf([
+        RequiresAllOf([APIVersion(V1_3)]),
+        RequiresAllOf([DeviceExtension(ext_subgroup_size_control)]),
+    ]),
+    */
+
+    /* TODO: enable
+    // TODO: document
+    REQUIRE_FULL_SUBGROUPS = REQUIRE_FULL_SUBGROUPS
+    RequiresOneOf([
+        RequiresAllOf([APIVersion(V1_3)]),
+        RequiresAllOf([DeviceExtension(ext_subgroup_size_control)]),
+    ]),
+    */
+}

vulkano/src/shader/reflect.rs

@@ -147,7 +147,7 @@ fn inspect_entry_point(
             &mut self,
             chain: [fn(&Spirv, Id) -> Option<Id>; N],
             id: Id,
-        ) -> Option<(&mut DescriptorBindingVariable, Option<u32>)> {
+        ) -> Option<(&mut DescriptorBindingVariable, Option<u64>)> {
             let mut id = chain
                 .into_iter()
                 .try_fold(id, |id, func| func(self.spirv, id))?;
@@ -169,10 +169,7 @@ fn inspect_entry_point(
                     // Variable was accessed with an access chain.
                     // Retrieve index from instruction if it's a constant value.
                     // TODO: handle a `None` index too?
-                    let index = match *self.spirv.id(*indexes.first().unwrap()).instruction() {
+                    let index = get_constant(self.spirv, *indexes.first().unwrap());
-                        Instruction::Constant { ref value, .. } => Some(value[0]),
-                        _ => None,
-                    };
                     let variable = self.result.entry(id).or_insert_with(|| variable.clone());
                     variable.reqs.stages = self.stage.into();
                     return Some((variable, index));
@@ -184,10 +181,15 @@ fn inspect_entry_point(
 
         fn inspect_entry_point_r(&mut self, function: Id) {
             fn desc_reqs(
-                descriptor_variable: Option<(&mut DescriptorBindingVariable, Option<u32>)>,
+                descriptor_variable: Option<(&mut DescriptorBindingVariable, Option<u64>)>,
             ) -> Option<&mut DescriptorRequirements> {
-                descriptor_variable
+                descriptor_variable.map(|(variable, index)| {
-                    .map(|(variable, index)| variable.reqs.descriptors.entry(index).or_default())
+                    variable
+                        .reqs
+                        .descriptors
+                        .entry(index.map(|index| index.try_into().unwrap()))
+                        .or_default()
+                })
             }
 
             fn inst_image_texel_pointer(spirv: &Spirv, id: Id) -> Option<Id> {
@@ -561,7 +563,7 @@ fn inspect_entry_point(
                             Some((variable, Some(index))) => DescriptorIdentifier {
                                 set: variable.set,
                                 binding: variable.binding,
-                                index,
+                                index: index.try_into().unwrap(),
                             },
                             _ => continue,
                         };
@@ -787,12 +789,7 @@ fn descriptor_binding_requirements_of(spirv: &Spirv, variable_id: Id) -> Descrip
                 reqs.descriptor_types
                     .retain(|&d| d != DescriptorType::InlineUniformBlock);
 
-                let len = match spirv.id(length).instruction() {
+                let len = get_constant(spirv, length).expect("failed to find array length");
-                    Instruction::Constant { value, .. } => {
-                        value.iter().rev().fold(0, |a, &b| (a << 32) | b as u64)
-                    }
-                    _ => panic!("failed to find array length"),
-                };
 
                 if let Some(count) = reqs.descriptor_count.as_mut() {
                     *count *= len as u32
@@ -1178,18 +1175,22 @@ fn shader_interface(
 }
 
 /// Returns the size of a type, or `None` if its size cannot be determined.
-fn size_of_type(spirv: &Spirv, id: Id) -> Option<DeviceSize> {
+pub(crate) fn size_of_type(spirv: &Spirv, id: Id) -> Option<DeviceSize> {
     let id_info = spirv.id(id);
 
     match *id_info.instruction() {
-        Instruction::TypeBool { .. } => {
+        Instruction::TypeVoid { .. } => Some(0),
-            panic!("Can't put booleans in structs")
+        Instruction::TypeBool { .. } => Some(4),
-        }
         Instruction::TypeInt { width, .. } | Instruction::TypeFloat { width, .. } => {
             assert!(width % 8 == 0);
             Some(width as DeviceSize / 8)
         }
-        Instruction::TypePointer { .. } => Some(8),
+        Instruction::TypePointer {
+            storage_class, ty, ..
+        } => match storage_class {
+            StorageClass::PhysicalStorageBuffer => Some(8),
+            _ => size_of_type(spirv, ty),
+        },
         Instruction::TypeVector {
             component_type,
             component_count,
@@ -1202,73 +1203,76 @@ fn size_of_type(spirv: &Spirv, id: Id) -> Option<DeviceSize> {
             ..
         } => {
             // FIXME: row-major or column-major
-            size_of_type(spirv, column_type)
+            // FIXME: `MatrixStride` applies to a struct member containing the matrix, not the
-                .map(|column_size| column_size * column_count as DeviceSize)
+            // matrix type itself.
-        }
+            id_info
-        Instruction::TypeArray { length, .. } => {
-            let stride = id_info
                 .decorations()
                 .iter()
                 .find_map(|instruction| match *instruction {
                     Instruction::Decorate {
-                        decoration: Decoration::ArrayStride { array_stride },
+                        decoration: Decoration::MatrixStride { matrix_stride },
                         ..
-                    } => Some(array_stride),
+                    } => Some(matrix_stride as DeviceSize),
                     _ => None,
                 })
-                .unwrap();
+                .or_else(|| size_of_type(spirv, column_type))
-            let length = match spirv.id(length).instruction() {
+                .map(|stride| stride * column_count as DeviceSize)
-                Instruction::Constant { value, .. } => Some(
-                    value
-                        .iter()
-                        .rev()
-                        .fold(0u64, |a, &b| (a << 32) | b as DeviceSize),
-                ),
-                _ => None,
-            }
-            .unwrap();
-
-            Some(stride as DeviceSize * length)
         }
+        Instruction::TypeArray {
+            element_type,
+            length,
+            ..
+        } => id_info
+            .decorations()
+            .iter()
+            .find_map(|instruction| match *instruction {
+                Instruction::Decorate {
+                    decoration: Decoration::ArrayStride { array_stride },
+                    ..
+                } => Some(array_stride as DeviceSize),
+                _ => None,
+            })
+            .or_else(|| size_of_type(spirv, element_type))
+            .map(|stride| {
+                let length = get_constant(spirv, length).unwrap();
+                stride * length
+            }),
         Instruction::TypeRuntimeArray { .. } => None,
         Instruction::TypeStruct {
             ref member_types, ..
         } => {
-            let mut end_of_struct = 0;
+            member_types.iter().zip(id_info.members()).try_fold(
-
+                0,
-            for (&member, member_info) in member_types.iter().zip(id_info.members()) {
+                |end_of_struct, (&member, member_info)| {
-                // Built-ins have an unknown size.
+                    let offset = member_info
-                if member_info.decorations().iter().any(|instruction| {
+                        .decorations()
-                    matches!(
+                        .iter()
-                        instruction,
+                        .find_map(|instruction| {
-                        Instruction::MemberDecorate {
+                            match *instruction {
-                            decoration: Decoration::BuiltIn { .. },
+                                // Built-ins have an unknown size.
-                            ..
+                                Instruction::MemberDecorate {
-                        }
+                                    decoration: Decoration::BuiltIn { .. },
-                    )
+                                    ..
-                }) {
+                                } => Some(None),
-                    return None;
+                                Instruction::MemberDecorate {
-                }
+                                    decoration: Decoration::Offset { byte_offset },
-
+                                    ..
-                // Some structs don't have `Offset` decorations, in the case they are used as local
+                                } => Some(Some(byte_offset as DeviceSize)),
-                // variables only. Ignoring these.
+                                _ => None,
-                let offset =
+                            }
-                    member_info.decorations().iter().find_map(
+                        })
-                        |instruction| match *instruction {
+                        .unwrap_or(Some(end_of_struct))?;
-                            Instruction::MemberDecorate {
+                    let size = size_of_type(spirv, member)?;
-                                decoration: Decoration::Offset { byte_offset },
+                    Some(end_of_struct.max(offset + size))
-                                ..
+                },
-                            } => Some(byte_offset),
+            )
-                            _ => None,
-                        },
-                    )?;
-                let size = size_of_type(spirv, member)?;
-                end_of_struct = end_of_struct.max(offset as DeviceSize + size);
-            }
-
-            Some(end_of_struct)
         }
-        _ => panic!("Type {} not found", id),
+        ref instruction => todo!(
+            "An unknown type was passed to `size_of_type`. \
+            This is a Vulkano bug and should be reported.\n
+            Instruction::{:?}",
+            instruction
+        ),
     }
 }
 
@@ -1366,23 +1370,8 @@ fn shader_interface_type_of(
                 shader_interface_type_of(spirv, element_type, false)
             } else {
                 let mut ty = shader_interface_type_of(spirv, element_type, false);
-                let num_elements = spirv
+                let length = get_constant(spirv, length).expect("failed to find array length");
-                    .constants()
+                ty.num_elements *= length as u32;
-                    .iter()
-                    .find_map(|instruction| match *instruction {
-                        Instruction::Constant {
-                            result_id,
-                            ref value,
-                            ..
-                        } if result_id == length => Some(value.clone()),
-                        _ => None,
-                    })
-                    .expect("failed to find array length")
-                    .iter()
-                    .rev()
-                    .fold(0u64, |a, &b| (a << 32) | b as u64)
-                    as u32;
-                ty.num_elements *= num_elements;
                 ty
             }
         }

vulkano/src/shader/spirv/mod.rs

@@ -7,7 +7,7 @@
 //! [SPIR-V specification](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html).
 
 use crate::{shader::SpecializationConstant, Version};
-use ahash::HashMap;
+use ahash::{HashMap, HashSet};
 use smallvec::{smallvec, SmallVec};
 use std::{
     borrow::Cow,
@@ -77,7 +77,7 @@ impl Spirv {
         let mut global_variables = Vec::new();
 
         let mut functions = HashMap::default();
-        let mut current_function: Option<&mut Vec<Instruction>> = None;
+        let mut current_function: Option<&mut FunctionInfo> = None;
 
         for instruction in iter_instructions(&words[5..]) {
             let instruction = instruction?;
@@ -116,123 +116,133 @@ impl Spirv {
                 continue;
             }
 
-            if current_function.is_some() {
+            match instruction {
-                match instruction {
+                Instruction::Function { result_id, .. } => {
-                    Instruction::FunctionEnd { .. } => {
+                    current_function = None;
-                        current_function.take().unwrap().push(instruction);
+                    let function = functions.entry(result_id).or_insert_with(|| {
-                    }
+                        let entry_point = entry_points
-                    _ => current_function.as_mut().unwrap().push(instruction),
+                            .iter()
+                            .find(|instruction| {
+                                matches!(
+                                    **instruction,
+                                    Instruction::EntryPoint { entry_point, .. }
+                                    if entry_point == result_id
+                                )
+                            })
+                            .cloned();
+                        let execution_modes = execution_modes
+                            .iter()
+                            .filter(|instruction| {
+                                matches!(
+                                    **instruction,
+                                    Instruction::ExecutionMode { entry_point, .. }
+                                    | Instruction::ExecutionModeId { entry_point, .. }
+                                    if entry_point == result_id
+                                )
+                            })
+                            .cloned()
+                            .collect();
+
+                        FunctionInfo {
+                            instructions: Vec::new(),
+                            called_functions: HashSet::default(),
+                            entry_point,
+                            execution_modes,
+                        }
+                    });
+                    let current_function = current_function.insert(function);
+                    current_function.instructions.push(instruction);
                 }
-            } else {
+                Instruction::FunctionEnd { .. } => {
-                let destination = match instruction {
+                    let current_function = current_function.take().unwrap();
-                    Instruction::Function { result_id, .. } => {
+                    current_function.instructions.push(instruction);
-                        current_function = None;
+                }
-                        let function = functions.entry(result_id).or_insert_with(|| {
+                _ => {
-                            let entry_point = entry_points
+                    if let Some(current_function) = current_function.as_mut() {
-                                .iter()
+                        if let Instruction::FunctionCall { function, .. } = instruction {
-                                .find(|instruction| {
+                            current_function.called_functions.insert(function);
-                                    matches!(
+                        }
-                                        **instruction,
-                                        Instruction::EntryPoint { entry_point, .. }
-                                        if entry_point == result_id
-                                    )
-                                })
-                                .cloned();
-                            let execution_modes = execution_modes
-                                .iter()
-                                .filter(|instruction| {
-                                    matches!(
-                                        **instruction,
-                                        Instruction::ExecutionMode { entry_point, .. }
-                                        | Instruction::ExecutionModeId { entry_point, .. }
-                                        if entry_point == result_id
-                                    )
-                                })
-                                .cloned()
-                                .collect();
 
-                            FunctionInfo {
+                        current_function.instructions.push(instruction);
-                                instructions: Vec::new(),
+                    } else {
-                                entry_point,
+                        let destination = match instruction {
-                                execution_modes,
+                            Instruction::Capability { .. } => &mut capabilities,
+                            Instruction::Extension { .. } => &mut extensions,
+                            Instruction::ExtInstImport { .. } => &mut ext_inst_imports,
+                            Instruction::MemoryModel { .. } => &mut memory_models,
+                            Instruction::EntryPoint { .. } => &mut entry_points,
+                            Instruction::ExecutionMode { .. }
+                            | Instruction::ExecutionModeId { .. } => &mut execution_modes,
+                            Instruction::Name { .. } | Instruction::MemberName { .. } => &mut names,
+                            Instruction::Decorate { .. }
+                            | Instruction::MemberDecorate { .. }
+                            | Instruction::DecorationGroup { .. }
+                            | Instruction::GroupDecorate { .. }
+                            | Instruction::GroupMemberDecorate { .. }
+                            | Instruction::DecorateId { .. }
+                            | Instruction::DecorateString { .. }
+                            | Instruction::MemberDecorateString { .. } => &mut decorations,
+                            Instruction::TypeVoid { .. }
+                            | Instruction::TypeBool { .. }
+                            | Instruction::TypeInt { .. }
+                            | Instruction::TypeFloat { .. }
+                            | Instruction::TypeVector { .. }
+                            | Instruction::TypeMatrix { .. }
+                            | Instruction::TypeImage { .. }
+                            | Instruction::TypeSampler { .. }
+                            | Instruction::TypeSampledImage { .. }
+                            | Instruction::TypeArray { .. }
+                            | Instruction::TypeRuntimeArray { .. }
+                            | Instruction::TypeStruct { .. }
+                            | Instruction::TypeOpaque { .. }
+                            | Instruction::TypePointer { .. }
+                            | Instruction::TypeFunction { .. }
+                            | Instruction::TypeEvent { .. }
+                            | Instruction::TypeDeviceEvent { .. }
+                            | Instruction::TypeReserveId { .. }
+                            | Instruction::TypeQueue { .. }
+                            | Instruction::TypePipe { .. }
+                            | Instruction::TypeForwardPointer { .. }
+                            | Instruction::TypePipeStorage { .. }
+                            | Instruction::TypeNamedBarrier { .. }
+                            | Instruction::TypeRayQueryKHR { .. }
+                            | Instruction::TypeAccelerationStructureKHR { .. }
+                            | Instruction::TypeCooperativeMatrixNV { .. }
+                            | Instruction::TypeVmeImageINTEL { .. }
+                            | Instruction::TypeAvcImePayloadINTEL { .. }
+                            | Instruction::TypeAvcRefPayloadINTEL { .. }
+                            | Instruction::TypeAvcSicPayloadINTEL { .. }
+                            | Instruction::TypeAvcMcePayloadINTEL { .. }
+                            | Instruction::TypeAvcMceResultINTEL { .. }
+                            | Instruction::TypeAvcImeResultINTEL { .. }
+                            | Instruction::TypeAvcImeResultSingleReferenceStreamoutINTEL {
+                                ..
                             }
-                        });
+                            | Instruction::TypeAvcImeResultDualReferenceStreamoutINTEL { .. }
-                        current_function.insert(&mut function.instructions)
+                            | Instruction::TypeAvcImeSingleReferenceStreaminINTEL { .. }
-                    }
+                            | Instruction::TypeAvcImeDualReferenceStreaminINTEL { .. }
-                    Instruction::Capability { .. } => &mut capabilities,
+                            | Instruction::TypeAvcRefResultINTEL { .. }
-                    Instruction::Extension { .. } => &mut extensions,
+                            | Instruction::TypeAvcSicResultINTEL { .. } => &mut types,
-                    Instruction::ExtInstImport { .. } => &mut ext_inst_imports,
+                            Instruction::ConstantTrue { .. }
-                    Instruction::MemoryModel { .. } => &mut memory_models,
+                            | Instruction::ConstantFalse { .. }
-                    Instruction::EntryPoint { .. } => &mut entry_points,
+                            | Instruction::Constant { .. }
-                    Instruction::ExecutionMode { .. } | Instruction::ExecutionModeId { .. } => {
+                            | Instruction::ConstantComposite { .. }
-                        &mut execution_modes
+                            | Instruction::ConstantSampler { .. }
-                    }
+                            | Instruction::ConstantNull { .. }
-                    Instruction::Name { .. } | Instruction::MemberName { .. } => &mut names,
+                            | Instruction::ConstantPipeStorage { .. }
-                    Instruction::Decorate { .. }
+                            | Instruction::SpecConstantTrue { .. }
-                    | Instruction::MemberDecorate { .. }
+                            | Instruction::SpecConstantFalse { .. }
-                    | Instruction::DecorationGroup { .. }
+                            | Instruction::SpecConstant { .. }
-                    | Instruction::GroupDecorate { .. }
+                            | Instruction::SpecConstantComposite { .. }
-                    | Instruction::GroupMemberDecorate { .. }
+                            | Instruction::SpecConstantOp { .. }
-                    | Instruction::DecorateId { .. }
+                            | Instruction::Undef { .. } => &mut constants,
-                    | Instruction::DecorateString { .. }
+                            Instruction::Variable { .. } => &mut global_variables,
-                    | Instruction::MemberDecorateString { .. } => &mut decorations,
+                            _ => continue,
-                    Instruction::TypeVoid { .. }
+                        };
-                    | Instruction::TypeBool { .. }
-                    | Instruction::TypeInt { .. }
-                    | Instruction::TypeFloat { .. }
-                    | Instruction::TypeVector { .. }
-                    | Instruction::TypeMatrix { .. }
-                    | Instruction::TypeImage { .. }
-                    | Instruction::TypeSampler { .. }
-                    | Instruction::TypeSampledImage { .. }
-                    | Instruction::TypeArray { .. }
-                    | Instruction::TypeRuntimeArray { .. }
-                    | Instruction::TypeStruct { .. }
-                    | Instruction::TypeOpaque { .. }
-                    | Instruction::TypePointer { .. }
-                    | Instruction::TypeFunction { .. }
-                    | Instruction::TypeEvent { .. }
-                    | Instruction::TypeDeviceEvent { .. }
-                    | Instruction::TypeReserveId { .. }
-                    | Instruction::TypeQueue { .. }
-                    | Instruction::TypePipe { .. }
-                    | Instruction::TypeForwardPointer { .. }
-                    | Instruction::TypePipeStorage { .. }
-                    | Instruction::TypeNamedBarrier { .. }
-                    | Instruction::TypeRayQueryKHR { .. }
-                    | Instruction::TypeAccelerationStructureKHR { .. }
-                    | Instruction::TypeCooperativeMatrixNV { .. }
-                    | Instruction::TypeVmeImageINTEL { .. }
-                    | Instruction::TypeAvcImePayloadINTEL { .. }
-                    | Instruction::TypeAvcRefPayloadINTEL { .. }
-                    | Instruction::TypeAvcSicPayloadINTEL { .. }
-                    | Instruction::TypeAvcMcePayloadINTEL { .. }
-                    | Instruction::TypeAvcMceResultINTEL { .. }
-                    | Instruction::TypeAvcImeResultINTEL { .. }
-                    | Instruction::TypeAvcImeResultSingleReferenceStreamoutINTEL { .. }
-                    | Instruction::TypeAvcImeResultDualReferenceStreamoutINTEL { .. }
-                    | Instruction::TypeAvcImeSingleReferenceStreaminINTEL { .. }
-                    | Instruction::TypeAvcImeDualReferenceStreaminINTEL { .. }
-                    | Instruction::TypeAvcRefResultINTEL { .. }
-                    | Instruction::TypeAvcSicResultINTEL { .. } => &mut types,
-                    Instruction::ConstantTrue { .. }
-                    | Instruction::ConstantFalse { .. }
-                    | Instruction::Constant { .. }
-                    | Instruction::ConstantComposite { .. }
-                    | Instruction::ConstantSampler { .. }
-                    | Instruction::ConstantNull { .. }
-                    | Instruction::ConstantPipeStorage { .. }
-                    | Instruction::SpecConstantTrue { .. }
-                    | Instruction::SpecConstantFalse { .. }
-                    | Instruction::SpecConstant { .. }
-                    | Instruction::SpecConstantComposite { .. }
-                    | Instruction::SpecConstantOp { .. }
-                    | Instruction::Undef { .. } => &mut constants,
-                    Instruction::Variable { .. } => &mut global_variables,
-                    _ => continue,
-                };
 
-                destination.push(instruction);
+                        destination.push(instruction);
+                    }
+                }
             }
         }
 
@@ -647,6 +657,7 @@ impl StructMemberInfo {
 #[derive(Clone, Debug)]
 pub struct FunctionInfo {
     instructions: Vec<Instruction>,
+    called_functions: HashSet<Id>,
     entry_point: Option<Instruction>,
     execution_modes: Vec<Instruction>,
 }
@@ -658,6 +669,13 @@ impl FunctionInfo {
         &self.instructions
     }
 
+    /// Returns `Id`s of all functions that are called by this function.
+    /// This may include recursive function calls.
+    #[inline]
+    pub fn called_functions(&self) -> &HashSet<Id> {
+        &self.called_functions
+    }
+
     /// Returns the `EntryPoint` instruction that targets this function, if there is one.
     #[inline]
     pub fn entry_point(&self) -> Option<&Instruction> {