naga/back/spv/

writer.rs

use alloc::{format, string::String, vec, vec::Vec};

use arrayvec::ArrayVec;
use hashbrown::hash_map::Entry;
use spirv::Word;

use super::{
    block::DebugInfoInner,
    helpers::{contains_builtin, global_needs_wrapper, map_storage_class},
    Block, BlockContext, CachedConstant, CachedExpressions, CooperativeType, DebugInfo,
    EntryPointContext, Error, Function, FunctionArgument, GlobalVariable, IdGenerator, Instruction,
    LocalImageType, LocalType, LocalVariable, LogicalLayout, LookupFunctionType, LookupType,
    NumericType, Options, PhysicalLayout, PipelineOptions, ResultMember, Writer, WriterFlags,
    BITS_PER_BYTE,
};
use crate::{
    arena::{Handle, HandleVec, UniqueArena},
    back::spv::{
        helpers::{is_uniform_matcx2_struct_member_access, BindingDecorations},
        BindingInfo, Std140CompatTypeInfo, WrappedFunction,
    },
    common::ForDebugWithTypes as _,
    proc::{Alignment, TypeResolution},
    valid::{FunctionInfo, ModuleInfo},
};

pub struct FunctionInterface<'a> {
    pub varying_ids: &'a mut Vec<Word>,
    pub stage: crate::ShaderStage,
    pub task_payload: Option<Handle<crate::GlobalVariable>>,
    pub mesh_info: Option<crate::MeshStageInfo>,
    pub workgroup_size: [u32; 3],
}

impl Function {
    pub(super) fn to_words(&self, sink: &mut impl Extend<Word>) {
        self.signature.as_ref().unwrap().to_words(sink);
        for argument in self.parameters.iter() {
            argument.instruction.to_words(sink);
        }
        for (index, block) in self.blocks.iter().enumerate() {
            Instruction::label(block.label_id).to_words(sink);
            if index == 0 {
                for local_var in self.variables.values() {
                    local_var.instruction.to_words(sink);
                }
                for local_var in self.ray_query_initialization_tracker_variables.values() {
                    local_var.instruction.to_words(sink);
                }
                for local_var in self.ray_query_t_max_tracker_variables.values() {
                    local_var.instruction.to_words(sink);
                }
                for local_var in self.force_loop_bounding_vars.iter() {
                    local_var.instruction.to_words(sink);
                }
                for internal_var in self.spilled_composites.values() {
                    internal_var.instruction.to_words(sink);
                }
            }
            for instruction in block.body.iter() {
                instruction.to_words(sink);
            }
        }
        Instruction::function_end().to_words(sink);
    }
}

impl Writer {
    pub fn new(options: &Options) -> Result<Self, Error> {
        let (major, minor) = options.lang_version;
        if major != 1 {
            return Err(Error::UnsupportedVersion(major, minor));
        }

        let mut capabilities_used = crate::FastIndexSet::default();
        capabilities_used.insert(spirv::Capability::Shader);

        let mut id_gen = IdGenerator::default();
        let gl450_ext_inst_id = id_gen.next();
        let void_type = id_gen.next();

        Ok(Writer {
            physical_layout: PhysicalLayout::new(major, minor),
            logical_layout: LogicalLayout::default(),
            id_gen,
            capabilities_available: options.capabilities.clone(),
            capabilities_used,
            extensions_used: crate::FastIndexSet::default(),
            debug_strings: vec![],
            debugs: vec![],
            annotations: vec![],
            flags: options.flags,
            bounds_check_policies: options.bounds_check_policies,
            zero_initialize_workgroup_memory: options.zero_initialize_workgroup_memory,
            force_loop_bounding: options.force_loop_bounding,
            ray_query_initialization_tracking: options.ray_query_initialization_tracking,
            trace_ray_argument_validation: options.trace_ray_argument_validation,
            use_storage_input_output_16: options.use_storage_input_output_16,
            emit_int_div_checks: options.emit_int_div_checks,
            void_type,
            tuple_of_u32s_ty_id: None,
            lookup_type: crate::FastHashMap::default(),
            lookup_function: crate::FastHashMap::default(),
            lookup_function_type: crate::FastHashMap::default(),
            wrapped_functions: crate::FastHashMap::default(),
            constant_ids: HandleVec::new(),
            cached_constants: crate::FastHashMap::default(),
            global_variables: HandleVec::new(),
            std140_compat_uniform_types: crate::FastHashMap::default(),
            fake_missing_bindings: options.fake_missing_bindings,
            binding_map: options.binding_map.clone(),
            saved_cached: CachedExpressions::default(),
            gl450_ext_inst_id,
            temp_list: Vec::new(),
            ray_query_functions: crate::FastHashMap::default(),
            ray_tracing_functions: crate::FastHashMap::default(),
            has_ray_tracing_pipeline: false,
            io_f16_polyfills: super::f16_polyfill::F16IoPolyfill::new(
                options.use_storage_input_output_16,
            ),
            debug_printf: None,
            task_dispatch_limits: options.task_dispatch_limits,
            mesh_shader_primitive_indices_clamp: options.mesh_shader_primitive_indices_clamp,
        })
    }

    pub fn set_options(&mut self, options: &Options) -> Result<(), Error> {
        let (major, minor) = options.lang_version;
        if major != 1 {
            return Err(Error::UnsupportedVersion(major, minor));
        }
        self.physical_layout = PhysicalLayout::new(major, minor);
        self.capabilities_available = options.capabilities.clone();
        self.flags = options.flags;
        self.bounds_check_policies = options.bounds_check_policies;
        self.zero_initialize_workgroup_memory = options.zero_initialize_workgroup_memory;
        self.force_loop_bounding = options.force_loop_bounding;
        self.use_storage_input_output_16 = options.use_storage_input_output_16;
        self.binding_map = options.binding_map.clone();
        self.io_f16_polyfills =
            super::f16_polyfill::F16IoPolyfill::new(options.use_storage_input_output_16);
        self.task_dispatch_limits = options.task_dispatch_limits;
        self.mesh_shader_primitive_indices_clamp = options.mesh_shader_primitive_indices_clamp;
        Ok(())
    }

    /// Returns `(major, minor)` of the SPIR-V language version.
    pub const fn lang_version(&self) -> (u8, u8) {
        self.physical_layout.lang_version()
    }

    /// Reset `Writer` to its initial state, retaining any allocations.
    ///
    /// Why not just implement `Reclaimable` for `Writer`? By design,
    /// `Reclaimable::reclaim` requires ownership of the value, not just
    /// `&mut`; see the trait documentation. But we need to use this method
    /// from functions like `Writer::write`, which only have `&mut Writer`.
    /// Workarounds include unsafe code (`core::ptr::read`, then `write`, ugh)
    /// or something like a `Default` impl that returns an oddly-initialized
    /// `Writer`, which is worse.
    fn reset(&mut self) {
        use super::reclaimable::Reclaimable;
        use core::mem::take;

        let mut id_gen = IdGenerator::default();
        let gl450_ext_inst_id = id_gen.next();
        let void_type = id_gen.next();

        // Every field of the old writer that is not determined by the `Options`
        // passed to `Writer::new` should be reset somehow.
        let fresh = Writer {
            // Copied from the old Writer:
            flags: self.flags,
            bounds_check_policies: self.bounds_check_policies,
            zero_initialize_workgroup_memory: self.zero_initialize_workgroup_memory,
            force_loop_bounding: self.force_loop_bounding,
            ray_query_initialization_tracking: self.ray_query_initialization_tracking,
            trace_ray_argument_validation: self.trace_ray_argument_validation,
            use_storage_input_output_16: self.use_storage_input_output_16,
            capabilities_available: take(&mut self.capabilities_available),
            fake_missing_bindings: self.fake_missing_bindings,
            binding_map: take(&mut self.binding_map),
            task_dispatch_limits: self.task_dispatch_limits,
            mesh_shader_primitive_indices_clamp: self.mesh_shader_primitive_indices_clamp,
            emit_int_div_checks: self.emit_int_div_checks,

            // Initialized afresh:
            id_gen,
            void_type,
            tuple_of_u32s_ty_id: None,
            gl450_ext_inst_id,

            // Reclaimed:
            capabilities_used: take(&mut self.capabilities_used).reclaim(),
            extensions_used: take(&mut self.extensions_used).reclaim(),
            physical_layout: self.physical_layout.clone().reclaim(),
            logical_layout: take(&mut self.logical_layout).reclaim(),
            debug_strings: take(&mut self.debug_strings).reclaim(),
            debugs: take(&mut self.debugs).reclaim(),
            annotations: take(&mut self.annotations).reclaim(),
            lookup_type: take(&mut self.lookup_type).reclaim(),
            lookup_function: take(&mut self.lookup_function).reclaim(),
            lookup_function_type: take(&mut self.lookup_function_type).reclaim(),
            wrapped_functions: take(&mut self.wrapped_functions).reclaim(),
            constant_ids: take(&mut self.constant_ids).reclaim(),
            cached_constants: take(&mut self.cached_constants).reclaim(),
            global_variables: take(&mut self.global_variables).reclaim(),
            std140_compat_uniform_types: take(&mut self.std140_compat_uniform_types).reclaim(),
            saved_cached: take(&mut self.saved_cached).reclaim(),
            temp_list: take(&mut self.temp_list).reclaim(),
            ray_query_functions: take(&mut self.ray_query_functions).reclaim(),
            ray_tracing_functions: take(&mut self.ray_tracing_functions).reclaim(),
            has_ray_tracing_pipeline: false,
            io_f16_polyfills: take(&mut self.io_f16_polyfills).reclaim(),
            debug_printf: None,
        };

        *self = fresh;

        self.capabilities_used.insert(spirv::Capability::Shader);
    }

    /// Indicate that the code requires any one of the listed capabilities.
    ///
    /// If nothing in `capabilities` appears in the available capabilities
    /// specified in the [`Options`] from which this `Writer` was created,
    /// return an error. The `what` string is used in the error message to
    /// explain what provoked the requirement. (If no available capabilities were
    /// given, assume everything is available.)
    ///
    /// The first acceptable capability will be added to this `Writer`'s
    /// [`capabilities_used`] table, and an `OpCapability` emitted for it in the
    /// result. For this reason, more specific capabilities should be listed
    /// before more general.
    ///
    /// [`capabilities_used`]: Writer::capabilities_used
    pub(super) fn require_any(
        &mut self,
        what: &'static str,
        capabilities: &[spirv::Capability],
    ) -> Result<(), Error> {
        match *capabilities {
            [] => Ok(()),
            [first, ..] => {
                // Find the first acceptable capability, or return an error if
                // there is none.
                let selected = match self.capabilities_available {
                    None => first,
                    Some(ref available) => {
                        match capabilities
                            .iter()
                            // need explicit type for hashbrown::HashSet::contains fn call to keep rustc happy
                            .find(|cap| available.contains::<spirv::Capability>(cap))
                        {
                            Some(&cap) => cap,
                            None => {
                                return Err(Error::MissingCapabilities(what, capabilities.to_vec()))
                            }
                        }
                    }
                };
                self.capabilities_used.insert(selected);
                Ok(())
            }
        }
    }

    /// Indicate that the code requires all of the listed capabilities.
    ///
    /// If all entries of `capabilities` appear in the available capabilities
    /// specified in the [`Options`] from which this `Writer` was created
    /// (including the case where [`Options::capabilities`] is `None`), add
    /// them all to this `Writer`'s [`capabilities_used`] table, and return
    /// `Ok(())`. If at least one of the listed capabilities is not available,
    /// do not add anything to the `capabilities_used` table, and return the
    /// first unavailable requested capability, wrapped in `Err()`.
    ///
    /// This method is does not return an [`enum@Error`] in case of failure
    /// because it may be used in cases where the caller can recover (e.g.,
    /// with a polyfill) if the requested capabilities are not available. In
    /// this case, it would be unnecessary work to find *all* the unavailable
    /// requested capabilities, and to allocate a `Vec` for them, just so we
    /// could return an [`Error::MissingCapabilities`]).
    ///
    /// [`capabilities_used`]: Writer::capabilities_used
    pub(super) fn require_all(
        &mut self,
        capabilities: &[spirv::Capability],
    ) -> Result<(), spirv::Capability> {
        if let Some(ref available) = self.capabilities_available {
            for requested in capabilities {
                if !available.contains(requested) {
                    return Err(*requested);
                }
            }
        }

        for requested in capabilities {
            self.capabilities_used.insert(*requested);
        }

        Ok(())
    }

    /// Indicate that the code uses the given extension.
    pub(super) fn use_extension(&mut self, extension: &'static str) {
        self.extensions_used.insert(extension);
    }

    pub(super) fn get_type_id(&mut self, lookup_ty: LookupType) -> Word {
        match self.lookup_type.entry(lookup_ty) {
            Entry::Occupied(e) => *e.get(),
            Entry::Vacant(e) => {
                let local = match lookup_ty {
                    LookupType::Handle(_handle) => unreachable!("Handles are populated at start"),
                    LookupType::Local(local) => local,
                };

                let id = self.id_gen.next();
                e.insert(id);
                self.write_type_declaration_local(id, local);
                id
            }
        }
    }

    pub(super) fn get_handle_type_id(&mut self, handle: Handle<crate::Type>) -> Word {
        self.get_type_id(LookupType::Handle(handle))
    }

    pub(super) fn get_expression_lookup_type(&mut self, tr: &TypeResolution) -> LookupType {
        match *tr {
            TypeResolution::Handle(ty_handle) => LookupType::Handle(ty_handle),
            TypeResolution::Value(ref inner) => {
                let inner_local_type = self.localtype_from_inner(inner).unwrap();
                LookupType::Local(inner_local_type)
            }
        }
    }

    pub(super) fn get_expression_type_id(&mut self, tr: &TypeResolution) -> Word {
        let lookup_ty = self.get_expression_lookup_type(tr);
        self.get_type_id(lookup_ty)
    }

    pub(super) fn get_localtype_id(&mut self, local: LocalType) -> Word {
        self.get_type_id(LookupType::Local(local))
    }

    pub(super) fn get_pointer_type_id(&mut self, base: Word, class: spirv::StorageClass) -> Word {
        self.get_type_id(LookupType::Local(LocalType::Pointer { base, class }))
    }

    pub(super) fn get_handle_pointer_type_id(
        &mut self,
        base: Handle<crate::Type>,
        class: spirv::StorageClass,
    ) -> Word {
        let base_id = self.get_handle_type_id(base);
        self.get_pointer_type_id(base_id, class)
    }

    pub(super) fn get_ray_query_pointer_id(&mut self) -> Word {
        let rq_id = self.get_type_id(LookupType::Local(LocalType::RayQuery));
        self.get_pointer_type_id(rq_id, spirv::StorageClass::Function)
    }

    /// Return a SPIR-V type for a pointer to `resolution`.
    ///
    /// The given `resolution` must be one that we can represent
    /// either as a `LocalType::Pointer` or `LocalType::LocalPointer`.
    pub(super) fn get_resolution_pointer_id(
        &mut self,
        resolution: &TypeResolution,
        class: spirv::StorageClass,
    ) -> Word {
        let resolution_type_id = self.get_expression_type_id(resolution);
        self.get_pointer_type_id(resolution_type_id, class)
    }

    pub(super) fn get_numeric_type_id(&mut self, numeric: NumericType) -> Word {
        self.get_type_id(LocalType::Numeric(numeric).into())
    }

    pub(super) fn get_u32_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::U32))
    }

    pub(super) fn get_f32_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::F32))
    }

    pub(super) fn get_vec2u_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Vector {
            size: crate::VectorSize::Bi,
            scalar: crate::Scalar::U32,
        })
    }

    pub(super) fn get_vec2f_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Vector {
            size: crate::VectorSize::Bi,
            scalar: crate::Scalar::F32,
        })
    }

    pub(super) fn get_vec3u_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Vector {
            size: crate::VectorSize::Tri,
            scalar: crate::Scalar::U32,
        })
    }

    pub(super) fn get_f32_pointer_type_id(&mut self, class: spirv::StorageClass) -> Word {
        let f32_id = self.get_f32_type_id();
        self.get_pointer_type_id(f32_id, class)
    }

    pub(super) fn get_vec2u_pointer_type_id(&mut self, class: spirv::StorageClass) -> Word {
        let vec2u_id = self.get_numeric_type_id(NumericType::Vector {
            size: crate::VectorSize::Bi,
            scalar: crate::Scalar::U32,
        });
        self.get_pointer_type_id(vec2u_id, class)
    }

    pub(super) fn get_bool_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::BOOL))
    }

    pub(super) fn get_vec2_bool_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Vector {
            size: crate::VectorSize::Bi,
            scalar: crate::Scalar::BOOL,
        })
    }

    pub(super) fn get_vec3_bool_type_id(&mut self) -> Word {
        self.get_numeric_type_id(NumericType::Vector {
            size: crate::VectorSize::Tri,
            scalar: crate::Scalar::BOOL,
        })
    }

    /// Used for "mulhi" to get the upper bits of multiplication.
    ///
    /// More specifically, `OpUMulExtended` multiplies 2 numbers and returns the lower and upper bits of the result
    /// as a user-defined struct type with 2 u32s. This defines that struct.
    pub(super) fn get_tuple_of_u32s_ty_id(&mut self) -> Word {
        if let Some(val) = self.tuple_of_u32s_ty_id {
            val
        } else {
            let id = self.id_gen.next();
            let u32_id = self.get_u32_type_id();
            let ins = Instruction::type_struct(id, &[u32_id, u32_id]);
            ins.to_words(&mut self.logical_layout.declarations);
            self.tuple_of_u32s_ty_id = Some(id);
            id
        }
    }

    pub(super) fn decorate(&mut self, id: Word, decoration: spirv::Decoration, operands: &[Word]) {
        self.annotations
            .push(Instruction::decorate(id, decoration, operands));
    }

    /// Return `inner` as a `LocalType`, if that's possible.
    ///
    /// If `inner` can be represented as a `LocalType`, return
    /// `Some(local_type)`.
    ///
    /// Otherwise, return `None`. In this case, the type must always be looked
    /// up using a `LookupType::Handle`.
    fn localtype_from_inner(&mut self, inner: &crate::TypeInner) -> Option<LocalType> {
        Some(match *inner {
            crate::TypeInner::Scalar(_)
            | crate::TypeInner::Atomic(_)
            | crate::TypeInner::Vector { .. }
            | crate::TypeInner::Matrix { .. } => {
                // We expect `NumericType::from_inner` to handle all
                // these cases, so unwrap.
                LocalType::Numeric(NumericType::from_inner(inner).unwrap())
            }
            crate::TypeInner::CooperativeMatrix { .. } => {
                LocalType::Cooperative(CooperativeType::from_inner(inner).unwrap())
            }
            crate::TypeInner::Pointer { base, space } => {
                let base_type_id = self.get_handle_type_id(base);
                LocalType::Pointer {
                    base: base_type_id,
                    class: map_storage_class(space),
                }
            }
            crate::TypeInner::ValuePointer {
                size,
                scalar,
                space,
            } => {
                let base_numeric_type = match size {
                    Some(size) => NumericType::Vector { size, scalar },
                    None => NumericType::Scalar(scalar),
                };
                LocalType::Pointer {
                    base: self.get_numeric_type_id(base_numeric_type),
                    class: map_storage_class(space),
                }
            }
            crate::TypeInner::Image {
                dim,
                arrayed,
                class,
            } => LocalType::Image(LocalImageType::from_inner(dim, arrayed, class)),
            crate::TypeInner::Sampler { comparison: _ } => LocalType::Sampler,
            crate::TypeInner::AccelerationStructure { .. } => LocalType::AccelerationStructure,
            crate::TypeInner::RayQuery { .. } => LocalType::RayQuery,
            crate::TypeInner::Array { .. }
            | crate::TypeInner::Struct { .. }
            | crate::TypeInner::BindingArray { .. } => return None,
        })
    }

    /// Resolve the [`BindingInfo`] for a [`crate::ResourceBinding`] from the
    /// provided [`Writer::binding_map`].
    ///
    /// If the specified resource is not present in the binding map this will
    /// return an error, unless [`Writer::fake_missing_bindings`] is set.
    pub(super) fn resolve_resource_binding(
        &self,
        res_binding: &crate::ResourceBinding,
    ) -> Result<BindingInfo, Error> {
        match self.binding_map.get(res_binding) {
            Some(target) => Ok(*target),
            None if self.fake_missing_bindings => Ok(BindingInfo {
                descriptor_set: res_binding.group,
                binding: res_binding.binding,
                binding_array_size: None,
                external_texture: false,
            }),
            None => Err(Error::MissingBinding(*res_binding)),
        }
    }

    /// Emits code for any wrapper functions required by the expressions in ir_function.
    /// The IDs of any emitted functions will be stored in [`Self::wrapped_functions`].
    fn write_wrapped_functions(
        &mut self,
        ir_function: &crate::Function,
        info: &FunctionInfo,
        ir_module: &crate::Module,
    ) -> Result<(), Error> {
        log::trace!("Generating wrapped functions for {:?}", ir_function.name);

        for (expr_handle, expr) in ir_function.expressions.iter() {
            match *expr {
                crate::Expression::Binary { op, left, right } => {
                    let expr_ty_inner = info[expr_handle].ty.inner_with(&ir_module.types);
                    if let Some(expr_ty) = NumericType::from_inner(expr_ty_inner) {
                        match (op, expr_ty.scalar().kind) {
                            // Division and modulo are undefined behaviour when the
                            // dividend is the minimum representable value and the divisor
                            // is negative one, or when the divisor is zero. These wrapped
                            // functions override the divisor to one in these cases,
                            // matching the WGSL spec.
                            //
                            // Signed `%` is additionally always wrapped (even without
                            // `emit_int_div_checks`) so it can be lowered to
                            // `a - b * (a / b)`: `OpSRem` produces a poison result for
                            // negative operands in the Vulkan SPIR-V environment unless
                            // `VK_KHR_maintenance8` is enabled. See
                            // <https://github.com/gfx-rs/wgpu/issues/8191>.
                            (
                                crate::BinaryOperator::Divide | crate::BinaryOperator::Modulo,
                                crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                            ) if self.emit_int_div_checks
                                || matches!(
                                    (op, expr_ty.scalar().kind),
                                    (crate::BinaryOperator::Modulo, crate::ScalarKind::Sint)
                                ) =>
                            {
                                self.write_wrapped_binary_op(
                                    op,
                                    expr_ty,
                                    &info[left].ty,
                                    &info[right].ty,
                                )?;
                            }
                            _ => {}
                        }
                    }
                }
                crate::Expression::Load { pointer } => {
                    if let crate::TypeInner::Pointer {
                        base: pointer_type,
                        space: crate::AddressSpace::Uniform,
                    } = *info[pointer].ty.inner_with(&ir_module.types)
                    {
                        if self.std140_compat_uniform_types.contains_key(&pointer_type) {
                            // Loading a std140 compat type requires the wrapper function
                            // to convert to the regular type.
                            self.write_wrapped_convert_from_std140_compat_type(
                                ir_module,
                                pointer_type,
                            )?;
                        }
                    }
                }
                crate::Expression::Access { base, .. } => {
                    if let crate::TypeInner::Pointer {
                        base: base_type,
                        space: crate::AddressSpace::Uniform,
                    } = *info[base].ty.inner_with(&ir_module.types)
                    {
                        // Dynamic accesses of a two-row matrix's columns require a
                        // wrapper function.
                        if let crate::TypeInner::Matrix {
                            rows: crate::VectorSize::Bi,
                            ..
                        } = ir_module.types[base_type].inner
                        {
                            self.write_wrapped_matcx2_get_column(ir_module, base_type)?;
                            // If the matrix is *not* directly a member of a struct, then
                            // we additionally require a wrapper function to convert from
                            // the std140 compat type to the regular type.
                            if !is_uniform_matcx2_struct_member_access(
                                ir_function,
                                info,
                                ir_module,
                                base,
                            ) {
                                self.write_wrapped_convert_from_std140_compat_type(
                                    ir_module, base_type,
                                )?;
                            }
                        }
                    }
                }
                _ => {}
            }
        }

        Ok(())
    }

    /// Write a SPIR-V function that performs the operator `op` with Naga IR semantics.
    ///
    /// Define a function that performs an integer division or modulo operation,
    /// except that using a divisor of zero or causing signed overflow with a
    /// divisor of -1 returns the numerator unchanged, rather than exhibiting
    /// undefined behavior.
    ///
    /// Store the generated function's id in the [`wrapped_functions`] table.
    ///
    /// The operator `op` must be either [`Divide`] or [`Modulo`].
    ///
    /// # Panics
    ///
    /// The `return_type`, `left_type` or `right_type` arguments must all be
    /// integer scalars or vectors. If not, this function panics.
    ///
    /// [`wrapped_functions`]: Writer::wrapped_functions
    /// [`Divide`]: crate::BinaryOperator::Divide
    /// [`Modulo`]: crate::BinaryOperator::Modulo
    fn write_wrapped_binary_op(
        &mut self,
        op: crate::BinaryOperator,
        return_type: NumericType,
        left_type: &TypeResolution,
        right_type: &TypeResolution,
    ) -> Result<(), Error> {
        let return_type_id = self.get_localtype_id(LocalType::Numeric(return_type));
        let left_type_id = self.get_expression_type_id(left_type);
        let right_type_id = self.get_expression_type_id(right_type);

        // Check if we've already emitted this function.
        let wrapped = WrappedFunction::BinaryOp {
            op,
            left_type_id,
            right_type_id,
        };
        let function_id = match self.wrapped_functions.entry(wrapped) {
            Entry::Occupied(_) => return Ok(()),
            Entry::Vacant(e) => *e.insert(self.id_gen.next()),
        };

        let scalar = return_type.scalar();

        if self.flags.contains(WriterFlags::DEBUG) {
            let function_name = match op {
                crate::BinaryOperator::Divide => "naga_div",
                crate::BinaryOperator::Modulo => "naga_mod",
                _ => unreachable!(),
            };
            self.debugs
                .push(Instruction::name(function_id, function_name));
        }
        let mut function = Function::default();

        let function_type_id = self.get_function_type(LookupFunctionType {
            parameter_type_ids: vec![left_type_id, right_type_id],
            return_type_id,
        });
        function.signature = Some(Instruction::function(
            return_type_id,
            function_id,
            spirv::FunctionControl::empty(),
            function_type_id,
        ));

        let lhs_id = self.id_gen.next();
        let rhs_id = self.id_gen.next();
        if self.flags.contains(WriterFlags::DEBUG) {
            self.debugs.push(Instruction::name(lhs_id, "lhs"));
            self.debugs.push(Instruction::name(rhs_id, "rhs"));
        }
        let left_par = Instruction::function_parameter(left_type_id, lhs_id);
        let right_par = Instruction::function_parameter(right_type_id, rhs_id);
        for instruction in [left_par, right_par] {
            function.parameters.push(FunctionArgument {
                instruction,
                handle_id: 0,
            });
        }

        let label_id = self.id_gen.next();
        let mut block = Block::new(label_id);

        let bool_type = return_type.with_scalar(crate::Scalar::BOOL);
        let bool_type_id = self.get_numeric_type_id(bool_type);

        let maybe_splat_const = |writer: &mut Self, const_id| match return_type {
            NumericType::Scalar(_) => const_id,
            NumericType::Vector { size, .. } => {
                let constituent_ids = [const_id; crate::VectorSize::MAX];
                writer.get_constant_composite(
                    LookupType::Local(LocalType::Numeric(return_type)),
                    &constituent_ids[..size as usize],
                )
            }
            NumericType::Matrix { .. } => unreachable!(),
        };

        let const_zero_id = self.get_constant_scalar_with(0, scalar)?;
        let composite_zero_id = maybe_splat_const(self, const_zero_id);
        let rhs_eq_zero_id = self.id_gen.next();
        block.body.push(Instruction::binary(
            spirv::Op::IEqual,
            bool_type_id,
            rhs_eq_zero_id,
            rhs_id,
            composite_zero_id,
        ));
        let divisor_selector_id = match scalar.kind {
            crate::ScalarKind::Sint => {
                let (const_min_id, const_neg_one_id) = match scalar.width {
                    2 => Ok((
                        self.get_constant_scalar(crate::Literal::I16(i16::MIN)),
                        self.get_constant_scalar(crate::Literal::I16(-1i16)),
                    )),
                    4 => Ok((
                        self.get_constant_scalar(crate::Literal::I32(i32::MIN)),
                        self.get_constant_scalar(crate::Literal::I32(-1i32)),
                    )),
                    8 => Ok((
                        self.get_constant_scalar(crate::Literal::I64(i64::MIN)),
                        self.get_constant_scalar(crate::Literal::I64(-1i64)),
                    )),
                    _ => Err(Error::Validation("Unexpected scalar width")),
                }?;
                let composite_min_id = maybe_splat_const(self, const_min_id);
                let composite_neg_one_id = maybe_splat_const(self, const_neg_one_id);

                let lhs_eq_int_min_id = self.id_gen.next();
                block.body.push(Instruction::binary(
                    spirv::Op::IEqual,
                    bool_type_id,
                    lhs_eq_int_min_id,
                    lhs_id,
                    composite_min_id,
                ));
                let rhs_eq_neg_one_id = self.id_gen.next();
                block.body.push(Instruction::binary(
                    spirv::Op::IEqual,
                    bool_type_id,
                    rhs_eq_neg_one_id,
                    rhs_id,
                    composite_neg_one_id,
                ));
                let lhs_eq_int_min_and_rhs_eq_neg_one_id = self.id_gen.next();
                block.body.push(Instruction::binary(
                    spirv::Op::LogicalAnd,
                    bool_type_id,
                    lhs_eq_int_min_and_rhs_eq_neg_one_id,
                    lhs_eq_int_min_id,
                    rhs_eq_neg_one_id,
                ));
                let rhs_eq_zero_or_lhs_eq_int_min_and_rhs_eq_neg_one_id = self.id_gen.next();
                block.body.push(Instruction::binary(
                    spirv::Op::LogicalOr,
                    bool_type_id,
                    rhs_eq_zero_or_lhs_eq_int_min_and_rhs_eq_neg_one_id,
                    rhs_eq_zero_id,
                    lhs_eq_int_min_and_rhs_eq_neg_one_id,
                ));
                rhs_eq_zero_or_lhs_eq_int_min_and_rhs_eq_neg_one_id
            }
            crate::ScalarKind::Uint => rhs_eq_zero_id,
            _ => unreachable!(),
        };

        let const_one_id = self.get_constant_scalar_with(1, scalar)?;
        let composite_one_id = maybe_splat_const(self, const_one_id);
        let divisor_id = self.id_gen.next();
        block.body.push(Instruction::select(
            right_type_id,
            divisor_id,
            divisor_selector_id,
            composite_one_id,
            rhs_id,
        ));
        let return_id = if matches!(op, crate::BinaryOperator::Modulo)
            && matches!(scalar.kind, crate::ScalarKind::Sint)
        {
            // `OpSRem` produces a poison result for negative operands in the Vulkan
            // environment without the `maintenance8` feature. `OpSDiv` is not poisoned, so
            // reconstruct the remainder as `a - b * (a / b)`, which is well-defined for
            // negative operands. `divisor_id` is the zero/overflow-guarded divisor selected
            // above, so the degenerate cases still match `OpSRem`'s guarded result (0).
            let quotient_id = self.id_gen.next();
            block.body.push(Instruction::binary(
                spirv::Op::SDiv,
                return_type_id,
                quotient_id,
                lhs_id,
                divisor_id,
            ));
            let product_id = self.id_gen.next();
            block.body.push(Instruction::binary(
                spirv::Op::IMul,
                return_type_id,
                product_id,
                quotient_id,
                divisor_id,
            ));
            let remainder_id = self.id_gen.next();
            block.body.push(Instruction::binary(
                spirv::Op::ISub,
                return_type_id,
                remainder_id,
                lhs_id,
                product_id,
            ));
            remainder_id
        } else {
            let spv_op = match (op, scalar.kind) {
                (crate::BinaryOperator::Divide, crate::ScalarKind::Sint) => spirv::Op::SDiv,
                (crate::BinaryOperator::Divide, crate::ScalarKind::Uint) => spirv::Op::UDiv,
                (crate::BinaryOperator::Modulo, crate::ScalarKind::Sint) => spirv::Op::SRem,
                (crate::BinaryOperator::Modulo, crate::ScalarKind::Uint) => spirv::Op::UMod,
                _ => unreachable!(),
            };
            let return_id = self.id_gen.next();
            block.body.push(Instruction::binary(
                spv_op,
                return_type_id,
                return_id,
                lhs_id,
                divisor_id,
            ));
            return_id
        };

        function.consume(block, Instruction::return_value(return_id));
        function.to_words(&mut self.logical_layout.function_definitions);
        Ok(())
    }

    /// Writes a wrapper function to convert from a std140 compat type to its
    /// corresponding regular type.
    ///
    /// See [`Self::write_std140_compat_type_declaration`] for more details.
    fn write_wrapped_convert_from_std140_compat_type(
        &mut self,
        ir_module: &crate::Module,
        r#type: Handle<crate::Type>,
    ) -> Result<(), Error> {
        if !self.std140_compat_uniform_types.contains_key(&r#type) {
            return Ok(());
        }
        // Check if we've already emitted this function.
        let wrapped = WrappedFunction::ConvertFromStd140CompatType { r#type };
        let function_id = match self.wrapped_functions.entry(wrapped) {
            Entry::Occupied(_) => return Ok(()),
            Entry::Vacant(e) => *e.insert(self.id_gen.next()),
        };
        if self.flags.contains(WriterFlags::DEBUG) {
            self.debugs.push(Instruction::name(
                function_id,
                &format!("{:?}_from_std140", r#type.for_debug(&ir_module.types)),
            ));
        }
        let param_type_id = self.std140_compat_uniform_types[&r#type].type_id;
        let return_type_id = self.get_handle_type_id(r#type);

        let mut function = Function::default();
        let function_type_id = self.get_function_type(LookupFunctionType {
            parameter_type_ids: vec![param_type_id],
            return_type_id,
        });
        function.signature = Some(Instruction::function(
            return_type_id,
            function_id,
            spirv::FunctionControl::empty(),
            function_type_id,
        ));
        let param_id = self.id_gen.next();
        function.parameters.push(FunctionArgument {
            instruction: Instruction::function_parameter(param_type_id, param_id),
            handle_id: 0,
        });

        let label_id = self.id_gen.next();
        let mut block = Block::new(label_id);

        let result_id = match ir_module.types[r#type].inner {
            // Param is struct containing a vector member for each of the
            // matrix's columns. Extract each column from the struct then
            // composite into a matrix.
            crate::TypeInner::Matrix {
                columns,
                rows: rows @ crate::VectorSize::Bi,
                scalar,
            } => {
                let column_type_id =
                    self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });

                let mut column_ids: ArrayVec<Word, 4> = ArrayVec::new();
                for column in 0..columns as u32 {
                    let column_id = self.id_gen.next();
                    block.body.push(Instruction::composite_extract(
                        column_type_id,
                        column_id,
                        param_id,
                        &[column],
                    ));
                    column_ids.push(column_id);
                }
                let result_id = self.id_gen.next();
                block.body.push(Instruction::composite_construct(
                    return_type_id,
                    result_id,
                    &column_ids,
                ));
                result_id
            }
            // Param is an array where the base type is the std140 compatible
            // type corresponding to `base`. Iterate through each element and
            // call its conversion function, then composite into a new array.
            crate::TypeInner::Array { base, size, .. } => {
                // Ensure the conversion function for the array's base type is
                // declared.
                self.write_wrapped_convert_from_std140_compat_type(ir_module, base)?;

                let element_type_id = self.get_handle_type_id(base);
                let std140_info = self.std140_compat_uniform_types.get(&base);
                let mut element_ids = Vec::new();
                let size = match size.resolve(ir_module.to_ctx())? {
                    crate::proc::IndexableLength::Known(size) => size,
                    crate::proc::IndexableLength::Dynamic => {
                        return Err(Error::Validation(
                            "Uniform buffers cannot contain dynamic arrays",
                        ))
                    }
                };
                for i in 0..size {
                    let std140_element_id = self.id_gen.next();
                    let std140_element_type_id =
                        std140_info.map_or(element_type_id, |info| info.type_id);
                    block.body.push(Instruction::composite_extract(
                        std140_element_type_id,
                        std140_element_id,
                        param_id,
                        &[i],
                    ));

                    // Only call the conversion function if a compatibility mapping actually exists.
                    let final_element_id = if std140_info.is_some() {
                        let conversion_fn_id = self.wrapped_functions
                            [&WrappedFunction::ConvertFromStd140CompatType { r#type: base }];
                        let id = self.id_gen.next();
                        block.body.push(Instruction::function_call(
                            element_type_id,
                            id,
                            conversion_fn_id,
                            &[std140_element_id],
                        ));
                        id
                    } else {
                        std140_element_type_id
                    };
                    element_ids.push(final_element_id);
                }
                let result_id = self.id_gen.next();
                block.body.push(Instruction::composite_construct(
                    return_type_id,
                    result_id,
                    &element_ids,
                ));
                result_id
            }
            // Param is a struct where each two-row matrix member has been
            // decomposed in to separate vector members for each column.
            // Other members use their std140 compatible type if one exists, or
            // else their regular type. Iterate through each member, converting
            // or composing any matrices if required, then finally compose into
            // the struct.
            crate::TypeInner::Struct { ref members, .. } => {
                let mut member_ids = Vec::new();
                let mut next_index = 0;
                for member in members {
                    let member_id = self.id_gen.next();
                    let member_type_id = self.get_handle_type_id(member.ty);
                    match ir_module.types[member.ty].inner {
                        crate::TypeInner::Matrix {
                            columns,
                            rows: rows @ crate::VectorSize::Bi,
                            scalar,
                        } => {
                            let mut column_ids: ArrayVec<Word, 4> = ArrayVec::new();
                            let column_type_id = self
                                .get_numeric_type_id(NumericType::Vector { size: rows, scalar });
                            for _ in 0..columns as u32 {
                                let column_id = self.id_gen.next();
                                block.body.push(Instruction::composite_extract(
                                    column_type_id,
                                    column_id,
                                    param_id,
                                    &[next_index],
                                ));
                                column_ids.push(column_id);
                                next_index += 1;
                            }
                            block.body.push(Instruction::composite_construct(
                                member_type_id,
                                member_id,
                                &column_ids,
                            ));
                        }
                        _ => {
                            // Ensure the conversion function for the member's
                            // type is declared.
                            self.write_wrapped_convert_from_std140_compat_type(
                                ir_module, member.ty,
                            )?;
                            match self.std140_compat_uniform_types.get(&member.ty) {
                                Some(std140_type_info) => {
                                    let std140_member_id = self.id_gen.next();
                                    block.body.push(Instruction::composite_extract(
                                        std140_type_info.type_id,
                                        std140_member_id,
                                        param_id,
                                        &[next_index],
                                    ));
                                    let function_id = self.wrapped_functions
                                        [&WrappedFunction::ConvertFromStd140CompatType {
                                            r#type: member.ty,
                                        }];
                                    block.body.push(Instruction::function_call(
                                        member_type_id,
                                        member_id,
                                        function_id,
                                        &[std140_member_id],
                                    ));
                                    next_index += 1;
                                }
                                None => {
                                    block.body.push(Instruction::composite_extract(
                                        member_type_id,
                                        member_id,
                                        param_id,
                                        &[next_index],
                                    ));
                                    next_index += 1;
                                }
                            }
                        }
                    }
                    member_ids.push(member_id);
                }
                let result_id = self.id_gen.next();
                block.body.push(Instruction::composite_construct(
                    return_type_id,
                    result_id,
                    &member_ids,
                ));
                result_id
            }
            _ => unreachable!(),
        };

        function.consume(block, Instruction::return_value(result_id));
        function.to_words(&mut self.logical_layout.function_definitions);
        Ok(())
    }

    /// Writes a wrapper function to get an `OpTypeVector` column from an
    /// `OpTypeMatrix` with a dynamic index.
    ///
    /// This is used when accessing a column of a [`TypeInner::Matrix`] through
    /// a [`Uniform`] address space pointer. In such cases, the matrix will have
    /// been declared in SPIR-V using an alternative type where each column is a
    /// member of a containing struct. SPIR-V is unable to dynamically access
    /// struct members, so instead we load the matrix then call this function to
    /// access a column from the loaded value.
    ///
    /// [`TypeInner::Matrix`]: crate::TypeInner::Matrix
    /// [`Uniform`]: crate::AddressSpace::Uniform
    fn write_wrapped_matcx2_get_column(
        &mut self,
        ir_module: &crate::Module,
        r#type: Handle<crate::Type>,
    ) -> Result<(), Error> {
        let wrapped = WrappedFunction::MatCx2GetColumn { r#type };
        let function_id = match self.wrapped_functions.entry(wrapped) {
            Entry::Occupied(_) => return Ok(()),
            Entry::Vacant(e) => *e.insert(self.id_gen.next()),
        };
        if self.flags.contains(WriterFlags::DEBUG) {
            self.debugs.push(Instruction::name(
                function_id,
                &format!("{:?}_get_column", r#type.for_debug(&ir_module.types)),
            ));
        }

        let crate::TypeInner::Matrix {
            columns,
            rows: rows @ crate::VectorSize::Bi,
            scalar,
        } = ir_module.types[r#type].inner
        else {
            unreachable!();
        };

        let mut function = Function::default();
        let matrix_type_id = self.get_handle_type_id(r#type);
        let column_index_type_id = self.get_u32_type_id();
        let column_type_id = self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
        let matrix_param_id = self.id_gen.next();
        let column_index_param_id = self.id_gen.next();
        function.parameters.push(FunctionArgument {
            instruction: Instruction::function_parameter(matrix_type_id, matrix_param_id),
            handle_id: 0,
        });
        function.parameters.push(FunctionArgument {
            instruction: Instruction::function_parameter(
                column_index_type_id,
                column_index_param_id,
            ),
            handle_id: 0,
        });
        let function_type_id = self.get_function_type(LookupFunctionType {
            parameter_type_ids: vec![matrix_type_id, column_index_type_id],
            return_type_id: column_type_id,
        });
        function.signature = Some(Instruction::function(
            column_type_id,
            function_id,
            spirv::FunctionControl::empty(),
            function_type_id,
        ));

        let label_id = self.id_gen.next();
        let mut block = Block::new(label_id);

        // Create a switch case for each column in the matrix, where each case
        // extracts its column from the matrix. Finally we use OpPhi to return
        // the correct column.
        let merge_id = self.id_gen.next();
        block.body.push(Instruction::selection_merge(
            merge_id,
            spirv::SelectionControl::NONE,
        ));
        let cases = (0..columns as u32)
            .map(|i| super::instructions::Case {
                value: i,
                label_id: self.id_gen.next(),
            })
            .collect::<ArrayVec<_, 4>>();

        // Which label we branch to in the default (column index out-of-bounds)
        // case depends on our bounds check policy.
        let default_id = match self.bounds_check_policies.index {
            // For `Restrict`, treat the same as the final column.
            crate::proc::BoundsCheckPolicy::Restrict => cases.last().unwrap().label_id,
            // For `ReadZeroSkipWrite`, branch directly to the merge block. This
            // will be handled in the `OpPhi` below to produce a zero value.
            crate::proc::BoundsCheckPolicy::ReadZeroSkipWrite => merge_id,
            // For `Unchecked` we create a new block containing an
            // `OpUnreachable`.
            crate::proc::BoundsCheckPolicy::Unchecked => self.id_gen.next(),
        };
        function.consume(
            block,
            Instruction::switch(column_index_param_id, default_id, &cases),
        );

        // Emit a block for each case, and produce a list of variable and parent
        // block IDs that will be used in an `OpPhi` below to select the right
        // value.
        let mut var_parent_pairs = cases
            .into_iter()
            .map(|case| {
                let mut block = Block::new(case.label_id);
                let column_id = self.id_gen.next();
                block.body.push(Instruction::composite_extract(
                    column_type_id,
                    column_id,
                    matrix_param_id,
                    &[case.value],
                ));
                function.consume(block, Instruction::branch(merge_id));
                (column_id, case.label_id)
            })
            // Need capacity for up to 4 columns plus possibly a default case.
            .collect::<ArrayVec<_, 5>>();

        // Emit a block or append the variable and parent `OpPhi` pair for the
        // column index out-of-bounds case, if required.
        match self.bounds_check_policies.index {
            // Don't need to do anything for `Restrict` as we have branched from
            // the final column case's block.
            crate::proc::BoundsCheckPolicy::Restrict => {}
            // For `ReadZeroSkipWrite` we have branched directly from the block
            // containing the `OpSwitch`. The `OpPhi` should produce a zero
            // value.
            crate::proc::BoundsCheckPolicy::ReadZeroSkipWrite => {
                var_parent_pairs.push((self.get_constant_null(column_type_id), label_id));
            }
            // For `Unchecked` create a new block containing `OpUnreachable`.
            // This does not need to be handled by the `OpPhi`.
            crate::proc::BoundsCheckPolicy::Unchecked => {
                function.consume(
                    Block::new(default_id),
                    Instruction::new(spirv::Op::Unreachable),
                );
            }
        }

        let mut block = Block::new(merge_id);
        let result_id = self.id_gen.next();
        block.body.push(Instruction::phi(
            column_type_id,
            result_id,
            &var_parent_pairs,
        ));

        function.consume(block, Instruction::return_value(result_id));
        function.to_words(&mut self.logical_layout.function_definitions);
        Ok(())
    }

    fn write_function(
        &mut self,
        ir_function: &crate::Function,
        info: &FunctionInfo,
        ir_module: &crate::Module,
        mut interface: Option<FunctionInterface>,
        debug_info: &Option<DebugInfoInner>,
    ) -> Result<Word, Error> {
        self.write_wrapped_functions(ir_function, info, ir_module)?;

        log::trace!("Generating code for {:?}", ir_function.name);
        let mut function = Function::default();

        let prelude_id = self.id_gen.next();
        let mut prelude = Block::new(prelude_id);
        let mut ep_context = EntryPointContext {
            argument_ids: Vec::new(),
            results: Vec::new(),
            task_payload_variable_id: if let Some(ref i) = interface {
                i.task_payload.map(|a| self.global_variables[a].var_id)
            } else {
                None
            },
            mesh_state: None,
        };

        let mut parameter_type_ids = Vec::with_capacity(ir_function.arguments.len());

        let mut local_invocation_index_var_id = None;
        let mut local_invocation_index_id = None;

        for argument in ir_function.arguments.iter() {
            let class = spirv::StorageClass::Input;
            let handle_ty = ir_module.types[argument.ty].inner.is_handle();
            let argument_type_id = if handle_ty {
                self.get_handle_pointer_type_id(argument.ty, spirv::StorageClass::UniformConstant)
            } else {
                self.get_handle_type_id(argument.ty)
            };

            if let Some(ref mut iface) = interface {
                let id = if let Some(ref binding) = argument.binding {
                    let name = argument.name.as_deref();

                    let varying_id = self.write_varying(
                        ir_module,
                        iface.stage,
                        class,
                        name,
                        argument.ty,
                        binding,
                    )?;
                    iface.varying_ids.push(varying_id);
                    let id = self.load_io_with_f16_polyfill(
                        &mut prelude.body,
                        varying_id,
                        argument_type_id,
                    );
                    if binding == &crate::Binding::BuiltIn(crate::BuiltIn::LocalInvocationIndex) {
                        local_invocation_index_id = Some(id);
                        local_invocation_index_var_id = Some(varying_id);
                    }

                    id
                } else if let crate::TypeInner::Struct { ref members, .. } =
                    ir_module.types[argument.ty].inner
                {
                    let struct_id = self.id_gen.next();
                    let mut constituent_ids = Vec::with_capacity(members.len());
                    for member in members {
                        let type_id = self.get_handle_type_id(member.ty);
                        let name = member.name.as_deref();
                        let binding = member.binding.as_ref().unwrap();
                        let varying_id = self.write_varying(
                            ir_module,
                            iface.stage,
                            class,
                            name,
                            member.ty,
                            binding,
                        )?;
                        iface.varying_ids.push(varying_id);
                        let id =
                            self.load_io_with_f16_polyfill(&mut prelude.body, varying_id, type_id);
                        constituent_ids.push(id);
                        if binding == &crate::Binding::BuiltIn(crate::BuiltIn::LocalInvocationIndex)
                        {
                            local_invocation_index_id = Some(id);
                            local_invocation_index_var_id = Some(varying_id);
                        }
                    }
                    prelude.body.push(Instruction::composite_construct(
                        argument_type_id,
                        struct_id,
                        &constituent_ids,
                    ));
                    struct_id
                } else {
                    unreachable!("Missing argument binding on an entry point");
                };
                ep_context.argument_ids.push(id);
            } else {
                let argument_id = self.id_gen.next();
                let instruction = Instruction::function_parameter(argument_type_id, argument_id);
                if self.flags.contains(WriterFlags::DEBUG) {
                    if let Some(ref name) = argument.name {
                        self.debugs.push(Instruction::name(argument_id, name));
                    }
                }
                function.parameters.push(FunctionArgument {
                    instruction,
                    handle_id: if handle_ty {
                        let id = self.id_gen.next();
                        prelude.body.push(Instruction::load(
                            self.get_handle_type_id(argument.ty),
                            id,
                            argument_id,
                            None,
                        ));
                        id
                    } else {
                        0
                    },
                });
                parameter_type_ids.push(argument_type_id);
            };
        }

        let return_type_id = match ir_function.result {
            Some(ref result) => {
                if let Some(ref mut iface) = interface {
                    let mut has_point_size = false;
                    let class = spirv::StorageClass::Output;
                    if let Some(ref binding) = result.binding {
                        has_point_size |=
                            *binding == crate::Binding::BuiltIn(crate::BuiltIn::PointSize);
                        let type_id = self.get_handle_type_id(result.ty);
                        let varying_id =
                            if *binding == crate::Binding::BuiltIn(crate::BuiltIn::MeshTaskSize) {
                                0
                            } else {
                                let varying_id = self.write_varying(
                                    ir_module,
                                    iface.stage,
                                    class,
                                    None,
                                    result.ty,
                                    binding,
                                )?;
                                iface.varying_ids.push(varying_id);
                                varying_id
                            };
                        ep_context.results.push(ResultMember {
                            id: varying_id,
                            type_id,
                            built_in: binding.to_built_in(),
                        });
                    } else if let crate::TypeInner::Struct { ref members, .. } =
                        ir_module.types[result.ty].inner
                    {
                        for member in members {
                            let type_id = self.get_handle_type_id(member.ty);
                            let name = member.name.as_deref();
                            let binding = member.binding.as_ref().unwrap();
                            has_point_size |=
                                *binding == crate::Binding::BuiltIn(crate::BuiltIn::PointSize);
                            // This isn't an actual builtin in SPIR-V. It can only appear as the
                            // output of a task shader and the output is used when writing the
                            // entry point return, in which case the id is ignored anyway.
                            let varying_id = if *binding
                                == crate::Binding::BuiltIn(crate::BuiltIn::MeshTaskSize)
                            {
                                0
                            } else {
                                let varying_id = self.write_varying(
                                    ir_module,
                                    iface.stage,
                                    class,
                                    name,
                                    member.ty,
                                    binding,
                                )?;
                                iface.varying_ids.push(varying_id);
                                varying_id
                            };
                            ep_context.results.push(ResultMember {
                                id: varying_id,
                                type_id,
                                built_in: binding.to_built_in(),
                            });
                        }
                    } else {
                        unreachable!("Missing result binding on an entry point");
                    }

                    if self.flags.contains(WriterFlags::FORCE_POINT_SIZE)
                        && iface.stage == crate::ShaderStage::Vertex
                        && !has_point_size
                    {
                        // add point size artificially
                        let varying_id = self.id_gen.next();
                        let pointer_type_id = self.get_f32_pointer_type_id(class);
                        Instruction::variable(pointer_type_id, varying_id, class, None)
                            .to_words(&mut self.logical_layout.declarations);
                        self.decorate(
                            varying_id,
                            spirv::Decoration::BuiltIn,
                            &[spirv::BuiltIn::PointSize as u32],
                        );
                        iface.varying_ids.push(varying_id);

                        let default_value_id = self.get_constant_scalar(crate::Literal::F32(1.0));
                        prelude
                            .body
                            .push(Instruction::store(varying_id, default_value_id, None));
                    }
                    if iface.stage == crate::ShaderStage::Task {
                        self.get_vec3u_type_id()
                    } else {
                        self.void_type
                    }
                } else {
                    self.get_handle_type_id(result.ty)
                }
            }
            None => self.void_type,
        };

        if let Some(ref mut iface) = interface {
            if let Some(task_payload) = iface.task_payload {
                iface
                    .varying_ids
                    .push(self.global_variables[task_payload].var_id);
            }
            // The mesh output copy and the workgroup memory polyfill both read
            // `LocalInvocationIndex`, which an entry point may only declare once.
            if iface.mesh_info.is_some() && local_invocation_index_var_id.is_none() {
                let u32_type_id = self.get_u32_type_id();
                let pointer_type_id =
                    self.get_pointer_type_id(u32_type_id, spirv::StorageClass::Input);
                let var_id = self.id_gen.next();
                Instruction::variable(pointer_type_id, var_id, spirv::StorageClass::Input, None)
                    .to_words(&mut self.logical_layout.declarations);
                self.decorate(
                    var_id,
                    spirv::Decoration::BuiltIn,
                    &[spirv::BuiltIn::LocalInvocationIndex as u32],
                );
                iface.varying_ids.push(var_id);

                let id = self.id_gen.next();
                prelude
                    .body
                    .push(Instruction::load(u32_type_id, id, var_id, None));
                local_invocation_index_var_id = Some(var_id);
                local_invocation_index_id = Some(id);
            }
            self.write_entry_point_mesh_shader_info(
                iface,
                local_invocation_index_var_id,
                ir_module,
                &mut ep_context,
            )?;
        }

        let lookup_function_type = LookupFunctionType {
            parameter_type_ids,
            return_type_id,
        };

        let function_id = self.id_gen.next();
        if self.flags.contains(WriterFlags::DEBUG) {
            if let Some(ref name) = ir_function.name {
                self.debugs.push(Instruction::name(function_id, name));
            }
        }

        let function_type = self.get_function_type(lookup_function_type);
        function.signature = Some(Instruction::function(
            return_type_id,
            function_id,
            spirv::FunctionControl::empty(),
            function_type,
        ));

        if interface.is_some() {
            function.entry_point_context = Some(ep_context);
        }

        // fill up the `GlobalVariable::access_id`
        for gv in self.global_variables.iter_mut() {
            gv.reset_for_function();
        }
        for (handle, var) in ir_module.global_variables.iter() {
            if info[handle].is_empty() {
                continue;
            }

            let mut gv = self.global_variables[handle].clone();
            if let Some(ref mut iface) = interface {
                // Have to include global variables in the interface
                if self.physical_layout.version >= 0x10400 && iface.task_payload != Some(handle) {
                    iface.varying_ids.push(gv.var_id);
                }
            }

            match ir_module.types[var.ty].inner {
                // Any that are binding arrays we skip as we cannot load the array, we must load the result after indexing.
                crate::TypeInner::BindingArray { .. } => {
                    gv.access_id = gv.var_id;
                }
                _ => {
                    // Handle globals are pre-emitted and should be loaded automatically.
                    if var.space == crate::AddressSpace::Handle {
                        let var_type_id = self.get_handle_type_id(var.ty);
                        let id = self.id_gen.next();
                        prelude
                            .body
                            .push(Instruction::load(var_type_id, id, gv.var_id, None));
                        gv.access_id = gv.var_id;
                        gv.handle_id = id;
                    } else if global_needs_wrapper(ir_module, var) {
                        let class = map_storage_class(var.space);
                        let pointer_type_id = match self.std140_compat_uniform_types.get(&var.ty) {
                            Some(std140_type_info) if var.space == crate::AddressSpace::Uniform => {
                                self.get_pointer_type_id(std140_type_info.type_id, class)
                            }
                            _ => self.get_handle_pointer_type_id(var.ty, class),
                        };
                        let index_id = self.get_index_constant(0);
                        let id = self.id_gen.next();
                        prelude.body.push(Instruction::access_chain(
                            pointer_type_id,
                            id,
                            gv.var_id,
                            &[index_id],
                        ));
                        gv.access_id = id;
                    } else {
                        // by default, the variable ID is accessed as is
                        gv.access_id = gv.var_id;
                    };
                }
            }

            // work around borrow checking in the presence of `self.xxx()` calls
            self.global_variables[handle] = gv;
        }

        // Create a `BlockContext` for generating SPIR-V for the function's
        // body.
        let mut context = BlockContext {
            ir_module,
            ir_function,
            fun_info: info,
            function: &mut function,
            // Re-use the cached expression table from prior functions.
            cached: core::mem::take(&mut self.saved_cached),

            // Steal the Writer's temp list for a bit.
            temp_list: core::mem::take(&mut self.temp_list),
            force_loop_bounding: self.force_loop_bounding,
            writer: self,
            expression_constness: super::ExpressionConstnessTracker::from_arena(
                &ir_function.expressions,
            ),
            ray_query_tracker_expr: crate::FastHashMap::default(),
        };

        // fill up the pre-emitted and const expressions
        context.cached.reset(ir_function.expressions.len());
        for (handle, expr) in ir_function.expressions.iter() {
            if (expr.needs_pre_emit() && !matches!(*expr, crate::Expression::LocalVariable(_)))
                || context.expression_constness.is_const(handle)
            {
                context.cache_expression_value(handle, &mut prelude)?;
            }
        }

        for (handle, variable) in ir_function.local_variables.iter() {
            let id = context.gen_id();

            if context.writer.flags.contains(WriterFlags::DEBUG) {
                if let Some(ref name) = variable.name {
                    context.writer.debugs.push(Instruction::name(id, name));
                }
            }

            let init_word = variable.init.map(|constant| context.cached[constant]);
            let pointer_type_id = context
                .writer
                .get_handle_pointer_type_id(variable.ty, spirv::StorageClass::Function);
            let instruction = Instruction::variable(
                pointer_type_id,
                id,
                spirv::StorageClass::Function,
                init_word.or_else(|| match ir_module.types[variable.ty].inner {
                    crate::TypeInner::RayQuery { .. } => None,
                    _ => {
                        let type_id = context.get_handle_type_id(variable.ty);
                        Some(context.writer.write_constant_null(type_id))
                    }
                }),
            );

            context
                .function
                .variables
                .insert(handle, LocalVariable { id, instruction });

            if let crate::TypeInner::RayQuery { .. } = ir_module.types[variable.ty].inner {
                // Don't refactor this into a struct: Although spirv itself allows opaque types in structs,
                // the vulkan environment for spirv does not. Putting ray queries into structs can cause
                // confusing bugs.
                let u32_type_id = context.writer.get_u32_type_id();
                let ptr_u32_type_id = context
                    .writer
                    .get_pointer_type_id(u32_type_id, spirv::StorageClass::Function);
                let tracker_id = context.gen_id();
                let tracker_init_id = context.writer.get_constant_scalar(crate::Literal::U32(
                    crate::back::RayQueryPoint::empty().bits(),
                ));
                let tracker_instruction = Instruction::variable(
                    ptr_u32_type_id,
                    tracker_id,
                    spirv::StorageClass::Function,
                    Some(tracker_init_id),
                );

                context
                    .function
                    .ray_query_initialization_tracker_variables
                    .insert(
                        handle,
                        LocalVariable {
                            id: tracker_id,
                            instruction: tracker_instruction,
                        },
                    );
                let f32_type_id = context.writer.get_f32_type_id();
                let ptr_f32_type_id = context
                    .writer
                    .get_pointer_type_id(f32_type_id, spirv::StorageClass::Function);
                let t_max_tracker_id = context.gen_id();
                let t_max_tracker_init_id =
                    context.writer.get_constant_scalar(crate::Literal::F32(0.0));
                let t_max_tracker_instruction = Instruction::variable(
                    ptr_f32_type_id,
                    t_max_tracker_id,
                    spirv::StorageClass::Function,
                    Some(t_max_tracker_init_id),
                );

                context.function.ray_query_t_max_tracker_variables.insert(
                    handle,
                    LocalVariable {
                        id: t_max_tracker_id,
                        instruction: t_max_tracker_instruction,
                    },
                );
            }
        }

        for (handle, expr) in ir_function.expressions.iter() {
            match *expr {
                crate::Expression::LocalVariable(_) => {
                    // Cache the `OpVariable` instruction we generated above as
                    // the value of this expression.
                    context.cache_expression_value(handle, &mut prelude)?;
                }
                crate::Expression::Access { base, .. }
                | crate::Expression::AccessIndex { base, .. } => {
                    // Count references to `base` by `Access` and `AccessIndex`
                    // instructions. See `access_uses` for details.
                    *context.function.access_uses.entry(base).or_insert(0) += 1;
                }
                _ => {}
            }
        }

        let next_id = context.gen_id();

        context
            .function
            .consume(prelude, Instruction::branch(next_id));

        let workgroup_vars_init_exit_block_id =
            match (context.writer.zero_initialize_workgroup_memory, interface) {
                (
                    super::ZeroInitializeWorkgroupMemoryMode::Polyfill,
                    Some(
                        ref mut interface @ FunctionInterface {
                            stage:
                                crate::ShaderStage::Compute
                                | crate::ShaderStage::Mesh
                                | crate::ShaderStage::Task,
                            ..
                        },
                    ),
                ) => context.writer.generate_workgroup_vars_init_block(
                    next_id,
                    ir_module,
                    info,
                    local_invocation_index_id,
                    interface,
                    context.function,
                ),
                _ => None,
            };

        let main_id = if let Some(exit_id) = workgroup_vars_init_exit_block_id {
            exit_id
        } else {
            next_id
        };

        context.write_function_body(main_id, debug_info.as_ref())?;

        // Consume the `BlockContext`, ending its borrows and letting the
        // `Writer` steal back its cached expression table and temp_list.
        let BlockContext {
            cached, temp_list, ..
        } = context;
        self.saved_cached = cached;
        self.temp_list = temp_list;

        function.to_words(&mut self.logical_layout.function_definitions);

        if let Some(EntryPointContext {
            mesh_state: Some(ref mesh_state),
            ..
        }) = function.entry_point_context
        {
            self.write_mesh_shader_wrapper(mesh_state, function_id)
        } else if let Some(EntryPointContext {
            task_payload_variable_id: Some(tp),
            ..
        }) = function.entry_point_context
        {
            self.write_task_shader_wrapper(tp, function_id)
        } else {
            Ok(function_id)
        }
    }

    fn write_execution_mode(
        &mut self,
        function_id: Word,
        mode: spirv::ExecutionMode,
    ) -> Result<(), Error> {
        //self.check(mode.required_capabilities())?;
        Instruction::execution_mode(function_id, mode, &[])
            .to_words(&mut self.logical_layout.execution_modes);
        Ok(())
    }

    // TODO Move to instructions module
    fn write_entry_point(
        &mut self,
        entry_point: &crate::EntryPoint,
        info: &FunctionInfo,
        ir_module: &crate::Module,
        debug_info: &Option<DebugInfoInner>,
    ) -> Result<Instruction, Error> {
        let mut interface_ids = Vec::new();

        let function_id = self.write_function(
            &entry_point.function,
            info,
            ir_module,
            Some(FunctionInterface {
                varying_ids: &mut interface_ids,
                stage: entry_point.stage,
                task_payload: entry_point.task_payload,
                mesh_info: entry_point.mesh_info.clone(),
                workgroup_size: entry_point.workgroup_size,
            }),
            debug_info,
        )?;

        let exec_model = match entry_point.stage {
            crate::ShaderStage::Vertex => spirv::ExecutionModel::Vertex,
            crate::ShaderStage::Fragment => {
                self.write_execution_mode(function_id, spirv::ExecutionMode::OriginUpperLeft)?;
                match entry_point.early_depth_test {
                    Some(crate::EarlyDepthTest::Force) => {
                        self.write_execution_mode(
                            function_id,
                            spirv::ExecutionMode::EarlyFragmentTests,
                        )?;
                    }
                    Some(crate::EarlyDepthTest::Allow { conservative }) => {
                        // TODO: Consider emitting EarlyAndLateFragmentTestsAMD here, if available.
                        // https://github.khronos.org/SPIRV-Registry/extensions/AMD/SPV_AMD_shader_early_and_late_fragment_tests.html
                        // This permits early depth tests even if the shader writes to a storage
                        // binding
                        match conservative {
                            crate::ConservativeDepth::GreaterEqual => self.write_execution_mode(
                                function_id,
                                spirv::ExecutionMode::DepthGreater,
                            )?,
                            crate::ConservativeDepth::LessEqual => self.write_execution_mode(
                                function_id,
                                spirv::ExecutionMode::DepthLess,
                            )?,
                            crate::ConservativeDepth::Unchanged => self.write_execution_mode(
                                function_id,
                                spirv::ExecutionMode::DepthUnchanged,
                            )?,
                        }
                    }
                    None => {}
                }
                if let Some(ref result) = entry_point.function.result {
                    if contains_builtin(
                        result.binding.as_ref(),
                        result.ty,
                        &ir_module.types,
                        crate::BuiltIn::FragDepth,
                    ) {
                        self.write_execution_mode(
                            function_id,
                            spirv::ExecutionMode::DepthReplacing,
                        )?;
                    }
                }
                spirv::ExecutionModel::Fragment
            }
            crate::ShaderStage::Compute => {
                let execution_mode = spirv::ExecutionMode::LocalSize;
                Instruction::execution_mode(
                    function_id,
                    execution_mode,
                    &entry_point.workgroup_size,
                )
                .to_words(&mut self.logical_layout.execution_modes);
                spirv::ExecutionModel::GLCompute
            }
            crate::ShaderStage::Task => {
                let execution_mode = spirv::ExecutionMode::LocalSize;
                Instruction::execution_mode(
                    function_id,
                    execution_mode,
                    &entry_point.workgroup_size,
                )
                .to_words(&mut self.logical_layout.execution_modes);
                spirv::ExecutionModel::TaskEXT
            }
            crate::ShaderStage::Mesh => {
                let execution_mode = spirv::ExecutionMode::LocalSize;
                Instruction::execution_mode(
                    function_id,
                    execution_mode,
                    &entry_point.workgroup_size,
                )
                .to_words(&mut self.logical_layout.execution_modes);
                let mesh_info = entry_point.mesh_info.as_ref().unwrap();
                Instruction::execution_mode(
                    function_id,
                    match mesh_info.topology {
                        crate::MeshOutputTopology::Points => spirv::ExecutionMode::OutputPoints,
                        crate::MeshOutputTopology::Lines => spirv::ExecutionMode::OutputLinesEXT,
                        crate::MeshOutputTopology::Triangles => {
                            spirv::ExecutionMode::OutputTrianglesEXT
                        }
                    },
                    &[],
                )
                .to_words(&mut self.logical_layout.execution_modes);
                Instruction::execution_mode(
                    function_id,
                    spirv::ExecutionMode::OutputVertices,
                    core::slice::from_ref(&mesh_info.max_vertices),
                )
                .to_words(&mut self.logical_layout.execution_modes);
                Instruction::execution_mode(
                    function_id,
                    spirv::ExecutionMode::OutputPrimitivesEXT,
                    core::slice::from_ref(&mesh_info.max_primitives),
                )
                .to_words(&mut self.logical_layout.execution_modes);
                spirv::ExecutionModel::MeshEXT
            }
            crate::ShaderStage::RayGeneration => {
                self.require_any("ray tracing pipelines", &[spirv::Capability::RayTracingKHR])?;
                spirv::ExecutionModel::RayGenerationKHR
            }
            crate::ShaderStage::AnyHit => {
                self.require_any("ray tracing pipelines", &[spirv::Capability::RayTracingKHR])?;
                spirv::ExecutionModel::AnyHitKHR
            }
            crate::ShaderStage::ClosestHit => {
                self.require_any("ray tracing pipelines", &[spirv::Capability::RayTracingKHR])?;
                spirv::ExecutionModel::ClosestHitKHR
            }
            crate::ShaderStage::Miss => {
                self.require_any("ray tracing pipelines", &[spirv::Capability::RayTracingKHR])?;
                spirv::ExecutionModel::MissKHR
            }
        };
        //self.check(exec_model.required_capabilities())?;

        Ok(Instruction::entry_point(
            exec_model,
            function_id,
            &entry_point.name,
            interface_ids.as_slice(),
        ))
    }

    fn make_scalar(&mut self, id: Word, scalar: crate::Scalar) -> Instruction {
        use crate::ScalarKind as Sk;

        let bits = (scalar.width * BITS_PER_BYTE) as u32;
        match scalar.kind {
            Sk::Sint | Sk::Uint => {
                let signedness = if scalar.kind == Sk::Sint {
                    super::instructions::Signedness::Signed
                } else {
                    super::instructions::Signedness::Unsigned
                };
                let cap = match bits {
                    8 => Some(spirv::Capability::Int8),
                    16 => Some(spirv::Capability::Int16),
                    64 => Some(spirv::Capability::Int64),
                    _ => None,
                };
                if let Some(cap) = cap {
                    self.capabilities_used.insert(cap);
                }
                if bits == 16 {
                    self.capabilities_used
                        .insert(spirv::Capability::StorageBuffer16BitAccess);
                    self.capabilities_used
                        .insert(spirv::Capability::UniformAndStorageBuffer16BitAccess);
                    if self.use_storage_input_output_16 {
                        self.capabilities_used
                            .insert(spirv::Capability::StorageInputOutput16);
                    }
                }
                Instruction::type_int(id, bits, signedness)
            }
            Sk::Float => {
                if bits == 64 {
                    self.capabilities_used.insert(spirv::Capability::Float64);
                }
                if bits == 16 {
                    self.capabilities_used.insert(spirv::Capability::Float16);
                    self.capabilities_used
                        .insert(spirv::Capability::StorageBuffer16BitAccess);
                    self.capabilities_used
                        .insert(spirv::Capability::UniformAndStorageBuffer16BitAccess);
                    if self.use_storage_input_output_16 {
                        self.capabilities_used
                            .insert(spirv::Capability::StorageInputOutput16);
                    }
                }
                Instruction::type_float(id, bits)
            }
            Sk::Bool => Instruction::type_bool(id),
            Sk::AbstractInt | Sk::AbstractFloat => {
                unreachable!("abstract types should never reach the backend");
            }
        }
    }

    fn request_type_capabilities(&mut self, inner: &crate::TypeInner) -> Result<(), Error> {
        match *inner {
            crate::TypeInner::Image {
                dim,
                arrayed,
                class,
            } => {
                let sampled = match class {
                    crate::ImageClass::Sampled { .. } => true,
                    crate::ImageClass::Depth { .. } => true,
                    crate::ImageClass::Storage { format, .. } => {
                        self.request_image_format_capabilities(format.into())?;
                        false
                    }
                    crate::ImageClass::External => unimplemented!(),
                };

                match dim {
                    crate::ImageDimension::D1 => {
                        if sampled {
                            self.require_any("sampled 1D images", &[spirv::Capability::Sampled1D])?;
                        } else {
                            self.require_any("1D storage images", &[spirv::Capability::Image1D])?;
                        }
                    }
                    crate::ImageDimension::Cube if arrayed => {
                        if sampled {
                            self.require_any(
                                "sampled cube array images",
                                &[spirv::Capability::SampledCubeArray],
                            )?;
                        } else {
                            self.require_any(
                                "cube array storage images",
                                &[spirv::Capability::ImageCubeArray],
                            )?;
                        }
                    }
                    _ => {}
                }
            }
            crate::TypeInner::AccelerationStructure { .. } => {
                self.require_any(
                    "Acceleration Structure",
                    // unless we use this conditional, the ray query snapshot
                    // tests pick the wrong capability
                    &[if self.has_ray_tracing_pipeline {
                        spirv::Capability::RayTracingKHR
                    } else {
                        spirv::Capability::RayQueryKHR
                    }],
                )?;
            }
            crate::TypeInner::RayQuery { .. } => {
                self.require_any("Ray Query", &[spirv::Capability::RayQueryKHR])?;
            }
            crate::TypeInner::Atomic(crate::Scalar { width: 8, kind: _ }) => {
                self.require_any("64 bit integer atomics", &[spirv::Capability::Int64Atomics])?;
            }
            crate::TypeInner::Atomic(crate::Scalar {
                width: 4,
                kind: crate::ScalarKind::Float,
            }) => {
                self.require_any(
                    "32 bit floating-point atomics",
                    &[spirv::Capability::AtomicFloat32AddEXT],
                )?;
                self.use_extension("SPV_EXT_shader_atomic_float_add");
            }
            // 16 bit floating-point support requires Float16 capability
            crate::TypeInner::Matrix {
                scalar: crate::Scalar::F16,
                ..
            }
            | crate::TypeInner::Vector {
                scalar: crate::Scalar::F16,
                ..
            }
            | crate::TypeInner::Scalar(crate::Scalar::F16) => {
                self.require_any("16 bit floating-point", &[spirv::Capability::Float16])?;
                self.use_extension("SPV_KHR_16bit_storage");
            }
            // 16 bit integer support requires Int16 capability
            crate::TypeInner::Vector {
                scalar:
                    crate::Scalar {
                        kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                        width: 2,
                    },
                ..
            }
            | crate::TypeInner::Scalar(crate::Scalar {
                kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                width: 2,
            }) => {
                self.require_any("16 bit integer", &[spirv::Capability::Int16])?;
                self.use_extension("SPV_KHR_16bit_storage");
            }
            // Cooperative types and ops
            crate::TypeInner::CooperativeMatrix { .. } => {
                self.require_any(
                    "cooperative matrix",
                    &[spirv::Capability::CooperativeMatrixKHR],
                )?;
                self.require_any("memory model", &[spirv::Capability::VulkanMemoryModel])?;
                self.use_extension("SPV_KHR_cooperative_matrix");
                self.use_extension("SPV_KHR_vulkan_memory_model");
            }
            _ => {}
        }
        Ok(())
    }

    fn write_numeric_type_declaration_local(&mut self, id: Word, numeric: NumericType) {
        let instruction = match numeric {
            NumericType::Scalar(scalar) => self.make_scalar(id, scalar),
            NumericType::Vector { size, scalar } => {
                let scalar_id = self.get_numeric_type_id(NumericType::Scalar(scalar));
                Instruction::type_vector(id, scalar_id, size)
            }
            NumericType::Matrix {
                columns,
                rows,
                scalar,
            } => {
                let column_id =
                    self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
                Instruction::type_matrix(id, column_id, columns)
            }
        };

        instruction.to_words(&mut self.logical_layout.declarations);
    }

    fn write_cooperative_type_declaration_local(&mut self, id: Word, coop: CooperativeType) {
        let instruction = match coop {
            CooperativeType::Matrix {
                columns,
                rows,
                scalar,
                role,
            } => {
                let scalar_id =
                    self.get_localtype_id(LocalType::Numeric(NumericType::Scalar(scalar)));
                let scope_id = self.get_index_constant(spirv::Scope::Subgroup as u32);
                let columns_id = self.get_index_constant(columns as u32);
                let rows_id = self.get_index_constant(rows as u32);
                let role_id =
                    self.get_index_constant(spirv::CooperativeMatrixUse::from(role) as u32);
                Instruction::type_coop_matrix(id, scalar_id, scope_id, rows_id, columns_id, role_id)
            }
        };

        instruction.to_words(&mut self.logical_layout.declarations);
    }

    fn write_type_declaration_local(&mut self, id: Word, local_ty: LocalType) {
        let instruction = match local_ty {
            LocalType::Numeric(numeric) => {
                self.write_numeric_type_declaration_local(id, numeric);
                return;
            }
            LocalType::Cooperative(coop) => {
                self.write_cooperative_type_declaration_local(id, coop);
                return;
            }
            LocalType::Pointer { base, class } => Instruction::type_pointer(id, class, base),
            LocalType::Image(image) => {
                let local_type = LocalType::Numeric(NumericType::Scalar(image.sampled_type));
                let type_id = self.get_localtype_id(local_type);
                Instruction::type_image(id, type_id, image.dim, image.flags, image.image_format)
            }
            LocalType::Sampler => Instruction::type_sampler(id),
            LocalType::SampledImage { image_type_id } => {
                Instruction::type_sampled_image(id, image_type_id)
            }
            LocalType::BindingArray { base, size } => {
                let inner_ty = self.get_handle_type_id(base);
                let scalar_id = self.get_constant_scalar(crate::Literal::U32(size));
                Instruction::type_array(id, inner_ty, scalar_id)
            }
            LocalType::AccelerationStructure => Instruction::type_acceleration_structure(id),
            LocalType::RayQuery => Instruction::type_ray_query(id),
        };

        instruction.to_words(&mut self.logical_layout.declarations);
    }

    fn write_type_declaration_arena(
        &mut self,
        module: &crate::Module,
        handle: Handle<crate::Type>,
    ) -> Result<Word, Error> {
        let ty = &module.types[handle];
        // If it's a type that needs SPIR-V capabilities, request them now.
        // This needs to happen regardless of the LocalType lookup succeeding,
        // because some types which map to the same LocalType have different
        // capability requirements. See https://github.com/gfx-rs/wgpu/issues/5569
        self.request_type_capabilities(&ty.inner)?;
        let id = if let Some(local) = self.localtype_from_inner(&ty.inner) {
            // This type can be represented as a `LocalType`, so check if we've
            // already written an instruction for it. If not, do so now, with
            // `write_type_declaration_local`.
            match self.lookup_type.entry(LookupType::Local(local)) {
                // We already have an id for this `LocalType`.
                Entry::Occupied(e) => *e.get(),

                // It's a type we haven't seen before.
                Entry::Vacant(e) => {
                    let id = self.id_gen.next();
                    e.insert(id);

                    self.write_type_declaration_local(id, local);

                    id
                }
            }
        } else {
            use spirv::Decoration;

            let id = self.id_gen.next();
            let instruction = match ty.inner {
                crate::TypeInner::Array { base, size, stride } => {
                    self.decorate(id, Decoration::ArrayStride, &[stride]);

                    let type_id = self.get_handle_type_id(base);
                    match size.resolve(module.to_ctx())? {
                        crate::proc::IndexableLength::Known(length) => {
                            let length_id = self.get_index_constant(length);
                            Instruction::type_array(id, type_id, length_id)
                        }
                        crate::proc::IndexableLength::Dynamic => {
                            Instruction::type_runtime_array(id, type_id)
                        }
                    }
                }
                crate::TypeInner::BindingArray { base, size } => {
                    let type_id = self.get_handle_type_id(base);
                    match size.resolve(module.to_ctx())? {
                        crate::proc::IndexableLength::Known(length) => {
                            let length_id = self.get_index_constant(length);
                            Instruction::type_array(id, type_id, length_id)
                        }
                        crate::proc::IndexableLength::Dynamic => {
                            Instruction::type_runtime_array(id, type_id)
                        }
                    }
                }
                crate::TypeInner::Struct {
                    ref members,
                    span: _,
                } => {
                    let mut has_runtime_array = false;
                    let mut member_ids = Vec::with_capacity(members.len());
                    for (index, member) in members.iter().enumerate() {
                        let member_ty = &module.types[member.ty];
                        match member_ty.inner {
                            crate::TypeInner::Array {
                                base: _,
                                size: crate::ArraySize::Dynamic,
                                stride: _,
                            } => {
                                has_runtime_array = true;
                            }
                            _ => (),
                        }
                        self.decorate_struct_member(id, index, member, &module.types)?;
                        let member_id = self.get_handle_type_id(member.ty);
                        member_ids.push(member_id);
                    }
                    if has_runtime_array {
                        self.decorate(id, Decoration::Block, &[]);
                    }
                    Instruction::type_struct(id, member_ids.as_slice())
                }

                // These all have TypeLocal representations, so they should have been
                // handled by `write_type_declaration_local` above.
                crate::TypeInner::Scalar(_)
                | crate::TypeInner::Atomic(_)
                | crate::TypeInner::Vector { .. }
                | crate::TypeInner::Matrix { .. }
                | crate::TypeInner::CooperativeMatrix { .. }
                | crate::TypeInner::Pointer { .. }
                | crate::TypeInner::ValuePointer { .. }
                | crate::TypeInner::Image { .. }
                | crate::TypeInner::Sampler { .. }
                | crate::TypeInner::AccelerationStructure { .. }
                | crate::TypeInner::RayQuery { .. } => unreachable!(),
            };

            instruction.to_words(&mut self.logical_layout.declarations);
            id
        };

        // Add this handle as a new alias for that type.
        self.lookup_type.insert(LookupType::Handle(handle), id);

        if self.flags.contains(WriterFlags::DEBUG) {
            if let Some(ref name) = ty.name {
                self.debugs.push(Instruction::name(id, name));
            }
        }

        Ok(id)
    }

    /// Writes a std140 layout compatible type declaration for a type. Returns
    /// the ID of the declared type, or None if no declaration is required.
    ///
    /// This should be called for any type for which there exists a
    /// [`GlobalVariable`] in the [`Uniform`] address space. If the type already
    /// adheres to std140 layout rules it will return without declaring any
    /// types. If the type contains another type which requires a std140
    /// compatible type declaration, it will recursively call itself.
    ///
    /// When `handle` refers to a [`TypeInner::Matrix`] with 2 rows, the
    /// declared type will be an `OpTypeStruct` containing an `OpVector` for
    /// each of the matrix's columns.
    ///
    /// When `handle` refers to a [`TypeInner::Array`] whose base type is a
    /// matrix with 2 rows, this will declare an `OpTypeArray` whose element
    /// type is the matrix's corresponding std140 compatible type.
    ///
    /// When `handle` refers to a [`TypeInner::Struct`] and any of its members
    /// require a std140 compatible type declaration, this will declare a new
    /// struct with the following rules:
    /// * Struct or array members will be declared with their std140 compatible
    ///   type declaration, if one is required.
    /// * Two-row matrix members will have each of their columns hoisted
    ///   directly into the struct as 2-component vector members.
    /// * All other members will be declared with their normal type.
    ///
    /// Note that this means the Naga IR index of a struct member may not match
    /// the index in the generated SPIR-V. The mapping can be obtained via
    /// `Std140TypeInfo::member_indices`.
    ///
    /// [`GlobalVariable`]: crate::GlobalVariable
    /// [`Uniform`]: crate::AddressSpace::Uniform
    /// [`TypeInner::Matrix`]: crate::TypeInner::Matrix
    /// [`TypeInner::Array`]: crate::TypeInner::Array
    /// [`TypeInner::Struct`]: crate::TypeInner::Struct
    fn write_std140_compat_type_declaration(
        &mut self,
        module: &crate::Module,
        handle: Handle<crate::Type>,
    ) -> Result<Option<Word>, Error> {
        if let Some(std140_type_info) = self.std140_compat_uniform_types.get(&handle) {
            return Ok(Some(std140_type_info.type_id));
        }

        let type_inner = &module.types[handle].inner;
        let std140_type_id = match *type_inner {
            crate::TypeInner::Matrix {
                columns,
                rows: rows @ crate::VectorSize::Bi,
                scalar,
            } => {
                let std140_type_id = self.id_gen.next();
                let mut member_type_ids: ArrayVec<Word, 4> = ArrayVec::new();
                let column_type_id =
                    self.get_numeric_type_id(NumericType::Vector { size: rows, scalar });
                for column in 0..columns as u32 {
                    member_type_ids.push(column_type_id);
                    self.annotations.push(Instruction::member_decorate(
                        std140_type_id,
                        column,
                        spirv::Decoration::Offset,
                        &[column * rows as u32 * scalar.width as u32],
                    ));
                    if self.flags.contains(WriterFlags::DEBUG) {
                        self.debugs.push(Instruction::member_name(
                            std140_type_id,
                            column,
                            &format!("col{column}"),
                        ));
                    }
                }
                Instruction::type_struct(std140_type_id, &member_type_ids)
                    .to_words(&mut self.logical_layout.declarations);
                self.std140_compat_uniform_types.insert(
                    handle,
                    Std140CompatTypeInfo {
                        type_id: std140_type_id,
                        member_indices: Vec::new(),
                    },
                );
                Some(std140_type_id)
            }
            crate::TypeInner::Array { base, size, stride } => {
                match self.write_std140_compat_type_declaration(module, base)? {
                    Some(std140_base_type_id) => {
                        let std140_type_id = self.id_gen.next();
                        self.decorate(std140_type_id, spirv::Decoration::ArrayStride, &[stride]);
                        let instruction = match size.resolve(module.to_ctx())? {
                            crate::proc::IndexableLength::Known(length) => {
                                let length_id = self.get_index_constant(length);
                                Instruction::type_array(
                                    std140_type_id,
                                    std140_base_type_id,
                                    length_id,
                                )
                            }
                            crate::proc::IndexableLength::Dynamic => {
                                unreachable!()
                            }
                        };
                        instruction.to_words(&mut self.logical_layout.declarations);
                        self.std140_compat_uniform_types.insert(
                            handle,
                            Std140CompatTypeInfo {
                                type_id: std140_type_id,
                                member_indices: Vec::new(),
                            },
                        );
                        Some(std140_type_id)
                    }
                    None => None,
                }
            }
            crate::TypeInner::Struct { ref members, .. } => {
                let mut needs_std140_type = false;
                for member in members {
                    match module.types[member.ty].inner {
                        // We don't need to write a std140 type for the matrix itself as
                        // it will be decomposed into the parent struct. As a result, the
                        // struct does need a std140 type, however.
                        crate::TypeInner::Matrix {
                            rows: crate::VectorSize::Bi,
                            ..
                        } => needs_std140_type = true,
                        // If an array member needs a std140 type, because it is an array
                        // (of an array, etc) of `matCx2`s, then the struct also needs
                        // a std140 type which uses the std140 type for this member.
                        crate::TypeInner::Array { .. }
                            if self
                                .write_std140_compat_type_declaration(module, member.ty)?
                                .is_some() =>
                        {
                            needs_std140_type = true;
                        }
                        _ => {}
                    }
                }

                if needs_std140_type {
                    let std140_type_id = self.id_gen.next();
                    let mut member_ids = Vec::new();
                    let mut member_indices = Vec::new();
                    let mut next_index = 0;

                    for member in members {
                        member_indices.push(next_index);
                        match module.types[member.ty].inner {
                            crate::TypeInner::Matrix {
                                columns,
                                rows: rows @ crate::VectorSize::Bi,
                                scalar,
                            } => {
                                let vector_type_id =
                                    self.get_numeric_type_id(NumericType::Vector {
                                        size: rows,
                                        scalar,
                                    });
                                for column in 0..columns as u32 {
                                    self.annotations.push(Instruction::member_decorate(
                                        std140_type_id,
                                        next_index,
                                        spirv::Decoration::Offset,
                                        &[member.offset
                                            + column * rows as u32 * scalar.width as u32],
                                    ));
                                    if self.flags.contains(WriterFlags::DEBUG) {
                                        if let Some(ref name) = member.name {
                                            self.debugs.push(Instruction::member_name(
                                                std140_type_id,
                                                next_index,
                                                &format!("{name}_col{column}"),
                                            ));
                                        }
                                    }
                                    member_ids.push(vector_type_id);
                                    next_index += 1;
                                }
                            }
                            _ => {
                                let member_id =
                                    match self.std140_compat_uniform_types.get(&member.ty) {
                                        Some(std140_member_type_info) => {
                                            self.annotations.push(Instruction::member_decorate(
                                                std140_type_id,
                                                next_index,
                                                spirv::Decoration::Offset,
                                                &[member.offset],
                                            ));
                                            if self.flags.contains(WriterFlags::DEBUG) {
                                                if let Some(ref name) = member.name {
                                                    self.debugs.push(Instruction::member_name(
                                                        std140_type_id,
                                                        next_index,
                                                        name,
                                                    ));
                                                }
                                            }
                                            std140_member_type_info.type_id
                                        }
                                        None => {
                                            self.decorate_struct_member(
                                                std140_type_id,
                                                next_index as usize,
                                                member,
                                                &module.types,
                                            )?;
                                            self.get_handle_type_id(member.ty)
                                        }
                                    };
                                member_ids.push(member_id);
                                next_index += 1;
                            }
                        }
                    }

                    Instruction::type_struct(std140_type_id, &member_ids)
                        .to_words(&mut self.logical_layout.declarations);
                    self.std140_compat_uniform_types.insert(
                        handle,
                        Std140CompatTypeInfo {
                            type_id: std140_type_id,
                            member_indices,
                        },
                    );
                    Some(std140_type_id)
                } else {
                    None
                }
            }
            _ => None,
        };

        if let Some(std140_type_id) = std140_type_id {
            if self.flags.contains(WriterFlags::DEBUG) {
                let name = format!("std140_{:?}", handle.for_debug(&module.types));
                self.debugs.push(Instruction::name(std140_type_id, &name));
            }
        }
        Ok(std140_type_id)
    }

    fn request_image_format_capabilities(
        &mut self,
        format: spirv::ImageFormat,
    ) -> Result<(), Error> {
        use spirv::ImageFormat as If;
        match format {
            If::Rg32f
            | If::Rg16f
            | If::R11fG11fB10f
            | If::R16f
            | If::Rgba16
            | If::Rgb10A2
            | If::Rg16
            | If::Rg8
            | If::R16
            | If::R8
            | If::Rgba16Snorm
            | If::Rg16Snorm
            | If::Rg8Snorm
            | If::R16Snorm
            | If::R8Snorm
            | If::Rg32i
            | If::Rg16i
            | If::Rg8i
            | If::R16i
            | If::R8i
            | If::Rgb10a2ui
            | If::Rg32ui
            | If::Rg16ui
            | If::Rg8ui
            | If::R16ui
            | If::R8ui => self.require_any(
                "storage image format",
                &[spirv::Capability::StorageImageExtendedFormats],
            ),
            If::R64ui | If::R64i => {
                self.use_extension("SPV_EXT_shader_image_int64");
                self.require_any(
                    "64-bit integer storage image format",
                    &[spirv::Capability::Int64ImageEXT],
                )
            }
            If::Unknown
            | If::Rgba32f
            | If::Rgba16f
            | If::R32f
            | If::Rgba8
            | If::Rgba8Snorm
            | If::Rgba32i
            | If::Rgba16i
            | If::Rgba8i
            | If::R32i
            | If::Rgba32ui
            | If::Rgba16ui
            | If::Rgba8ui
            | If::R32ui => Ok(()),
        }
    }

    pub(super) fn get_index_constant(&mut self, index: Word) -> Word {
        self.get_constant_scalar(crate::Literal::U32(index))
    }

    pub(super) fn get_constant_scalar_with(
        &mut self,
        value: u8,
        scalar: crate::Scalar,
    ) -> Result<Word, Error> {
        Ok(
            self.get_constant_scalar(crate::Literal::new(value, scalar).ok_or(
                Error::Validation("Unexpected kind and/or width for Literal"),
            )?),
        )
    }

    pub(super) fn get_constant_scalar(&mut self, value: crate::Literal) -> Word {
        let scalar = CachedConstant::Literal(value.into());
        if let Some(&id) = self.cached_constants.get(&scalar) {
            return id;
        }
        let id = self.id_gen.next();
        self.write_constant_scalar(id, &value, None);
        self.cached_constants.insert(scalar, id);
        id
    }

    fn write_constant_scalar(
        &mut self,
        id: Word,
        value: &crate::Literal,
        debug_name: Option<&String>,
    ) {
        if self.flags.contains(WriterFlags::DEBUG) {
            if let Some(name) = debug_name {
                self.debugs.push(Instruction::name(id, name));
            }
        }
        let type_id = self.get_numeric_type_id(NumericType::Scalar(value.scalar()));
        let instruction = match *value {
            crate::Literal::F64(value) => {
                let bits = value.to_bits();
                Instruction::constant_64bit(type_id, id, bits as u32, (bits >> 32) as u32)
            }
            crate::Literal::F32(value) => Instruction::constant_32bit(type_id, id, value.to_bits()),
            crate::Literal::F16(value) => {
                let low = value.to_bits();
                Instruction::constant_16bit(type_id, id, low as u32)
            }
            crate::Literal::U16(value) => Instruction::constant_16bit(type_id, id, value as u32),
            crate::Literal::I16(value) => {
                // Sign-extend into the 32-bit word so that `spirv-as` can
                // round-trip the disassembly (it expects signed values for
                // signed types).
                Instruction::constant_16bit(type_id, id, value as i32 as u32)
            }
            crate::Literal::U32(value) => Instruction::constant_32bit(type_id, id, value),
            crate::Literal::I32(value) => Instruction::constant_32bit(type_id, id, value as u32),
            crate::Literal::U64(value) => {
                Instruction::constant_64bit(type_id, id, value as u32, (value >> 32) as u32)
            }
            crate::Literal::I64(value) => {
                Instruction::constant_64bit(type_id, id, value as u32, (value >> 32) as u32)
            }
            crate::Literal::Bool(true) => Instruction::constant_true(type_id, id),
            crate::Literal::Bool(false) => Instruction::constant_false(type_id, id),
            crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
                unreachable!("Abstract types should not appear in IR presented to backends");
            }
        };

        instruction.to_words(&mut self.logical_layout.declarations);
    }

    pub(super) fn get_constant_composite(
        &mut self,
        ty: LookupType,
        constituent_ids: &[Word],
    ) -> Word {
        let composite = CachedConstant::Composite {
            ty,
            constituent_ids: constituent_ids.to_vec(),
        };
        if let Some(&id) = self.cached_constants.get(&composite) {
            return id;
        }
        let id = self.id_gen.next();
        self.write_constant_composite(id, ty, constituent_ids, None);
        self.cached_constants.insert(composite, id);
        id
    }

    fn write_constant_composite(
        &mut self,
        id: Word,
        ty: LookupType,
        constituent_ids: &[Word],
        debug_name: Option<&String>,
    ) {
        if self.flags.contains(WriterFlags::DEBUG) {
            if let Some(name) = debug_name {
                self.debugs.push(Instruction::name(id, name));
            }
        }
        let type_id = self.get_type_id(ty);
        Instruction::constant_composite(type_id, id, constituent_ids)
            .to_words(&mut self.logical_layout.declarations);
    }

    pub(super) fn get_constant_null(&mut self, type_id: Word) -> Word {
        let null = CachedConstant::ZeroValue(type_id);
        if let Some(&id) = self.cached_constants.get(&null) {
            return id;
        }
        let id = self.write_constant_null(type_id);
        self.cached_constants.insert(null, id);
        id
    }

    pub(super) fn write_constant_null(&mut self, type_id: Word) -> Word {
        let null_id = self.id_gen.next();
        Instruction::constant_null(type_id, null_id)
            .to_words(&mut self.logical_layout.declarations);
        null_id
    }

    fn write_constant_expr(
        &mut self,
        handle: Handle<crate::Expression>,
        ir_module: &crate::Module,
        mod_info: &ModuleInfo,
    ) -> Result<Word, Error> {
        let id = match ir_module.global_expressions[handle] {
            crate::Expression::Literal(literal) => self.get_constant_scalar(literal),
            crate::Expression::Constant(constant) => {
                let constant = &ir_module.constants[constant];
                self.constant_ids[constant.init]
            }
            crate::Expression::ZeroValue(ty) => {
                let type_id = self.get_handle_type_id(ty);
                self.get_constant_null(type_id)
            }
            crate::Expression::Compose { ty, ref components } => {
                let component_ids: Vec<_> = crate::proc::flatten_compose(
                    ty,
                    components,
                    &ir_module.global_expressions,
                    &ir_module.types,
                )
                .map(|component| self.constant_ids[component])
                .collect();
                self.get_constant_composite(LookupType::Handle(ty), component_ids.as_slice())
            }
            crate::Expression::Splat { size, value } => {
                let value_id = self.constant_ids[value];
                let component_ids = &[value_id; 4][..size as usize];

                let ty = self.get_expression_lookup_type(&mod_info[handle]);

                self.get_constant_composite(ty, component_ids)
            }
            _ => {
                return Err(Error::Override);
            }
        };

        self.constant_ids[handle] = id;

        Ok(id)
    }

    pub(super) fn write_control_barrier(
        &mut self,
        flags: crate::Barrier,
        body: &mut Vec<Instruction>,
    ) {
        let memory_scope = if flags.contains(crate::Barrier::STORAGE) {
            spirv::Scope::Device
        } else if flags.contains(crate::Barrier::SUB_GROUP) {
            spirv::Scope::Subgroup
        } else {
            spirv::Scope::Workgroup
        };
        let mut semantics = spirv::MemorySemantics::ACQUIRE_RELEASE;
        semantics.set(
            spirv::MemorySemantics::UNIFORM_MEMORY,
            flags.contains(crate::Barrier::STORAGE),
        );
        semantics.set(
            spirv::MemorySemantics::WORKGROUP_MEMORY,
            flags.contains(crate::Barrier::WORK_GROUP),
        );
        semantics.set(
            spirv::MemorySemantics::SUBGROUP_MEMORY,
            flags.contains(crate::Barrier::SUB_GROUP),
        );
        semantics.set(
            spirv::MemorySemantics::IMAGE_MEMORY,
            flags.contains(crate::Barrier::TEXTURE),
        );
        let exec_scope_id = if flags.contains(crate::Barrier::SUB_GROUP) {
            self.get_index_constant(spirv::Scope::Subgroup as u32)
        } else {
            self.get_index_constant(spirv::Scope::Workgroup as u32)
        };
        let mem_scope_id = self.get_index_constant(memory_scope as u32);
        let semantics_id = self.get_index_constant(semantics.bits());
        body.push(Instruction::control_barrier(
            exec_scope_id,
            mem_scope_id,
            semantics_id,
        ));
    }

    pub(super) fn write_memory_barrier(&mut self, flags: crate::Barrier, block: &mut Block) {
        let mut semantics = spirv::MemorySemantics::ACQUIRE_RELEASE;
        semantics.set(
            spirv::MemorySemantics::UNIFORM_MEMORY,
            flags.contains(crate::Barrier::STORAGE),
        );
        semantics.set(
            spirv::MemorySemantics::WORKGROUP_MEMORY,
            flags.contains(crate::Barrier::WORK_GROUP),
        );
        semantics.set(
            spirv::MemorySemantics::SUBGROUP_MEMORY,
            flags.contains(crate::Barrier::SUB_GROUP),
        );
        semantics.set(
            spirv::MemorySemantics::IMAGE_MEMORY,
            flags.contains(crate::Barrier::TEXTURE),
        );
        let mem_scope_id = if flags.contains(crate::Barrier::STORAGE) {
            self.get_index_constant(spirv::Scope::Device as u32)
        } else if flags.contains(crate::Barrier::SUB_GROUP) {
            self.get_index_constant(spirv::Scope::Subgroup as u32)
        } else {
            self.get_index_constant(spirv::Scope::Workgroup as u32)
        };
        let semantics_id = self.get_index_constant(semantics.bits());
        block
            .body
            .push(Instruction::memory_barrier(mem_scope_id, semantics_id));
    }

    fn generate_workgroup_vars_init_block(
        &mut self,
        entry_id: Word,
        ir_module: &crate::Module,
        info: &FunctionInfo,
        local_invocation_index: Option<Word>,
        interface: &mut FunctionInterface,
        function: &mut Function,
    ) -> Option<Word> {
        let body = ir_module
            .global_variables
            .iter()
            .filter(|&(handle, var)| {
                let task_exception = (var.space == crate::AddressSpace::TaskPayload)
                    && interface.stage == crate::ShaderStage::Task;
                !info[handle].is_empty()
                    && (var.space == crate::AddressSpace::WorkGroup || task_exception)
            })
            .map(|(handle, var)| {
                // It's safe to use `var_id` here, not `access_id`, because only
                // variables in the `Uniform` and `StorageBuffer` address spaces
                // get wrapped, and we're initializing `WorkGroup` variables.
                let var_id = self.global_variables[handle].var_id;
                let var_type_id = self.get_handle_type_id(var.ty);
                let init_word = self.get_constant_null(var_type_id);
                Instruction::store(var_id, init_word, None)
            })
            .collect::<Vec<_>>();

        if body.is_empty() {
            return None;
        }

        let mut pre_if_block = Block::new(entry_id);

        let local_invocation_index = if let Some(local_invocation_index) = local_invocation_index {
            local_invocation_index
        } else {
            let varying_id = self.id_gen.next();
            let class = spirv::StorageClass::Input;
            let u32_ty_id = self.get_u32_type_id();
            let pointer_type_id = self.get_pointer_type_id(u32_ty_id, class);

            Instruction::variable(pointer_type_id, varying_id, class, None)
                .to_words(&mut self.logical_layout.declarations);

            self.decorate(
                varying_id,
                spirv::Decoration::BuiltIn,
                &[spirv::BuiltIn::LocalInvocationIndex as u32],
            );

            interface.varying_ids.push(varying_id);
            let id = self.id_gen.next();
            pre_if_block
                .body
                .push(Instruction::load(u32_ty_id, id, varying_id, None));

            id
        };

        let zero_id = self.get_constant_scalar(crate::Literal::U32(0));

        let eq_id = self.id_gen.next();
        pre_if_block.body.push(Instruction::binary(
            spirv::Op::IEqual,
            self.get_bool_type_id(),
            eq_id,
            local_invocation_index,
            zero_id,
        ));

        let merge_id = self.id_gen.next();
        pre_if_block.body.push(Instruction::selection_merge(
            merge_id,
            spirv::SelectionControl::NONE,
        ));

        let accept_id = self.id_gen.next();
        function.consume(
            pre_if_block,
            Instruction::branch_conditional(eq_id, accept_id, merge_id),
        );

        let accept_block = Block {
            label_id: accept_id,
            body,
        };
        function.consume(accept_block, Instruction::branch(merge_id));

        let mut post_if_block = Block::new(merge_id);

        self.write_control_barrier(crate::Barrier::WORK_GROUP, &mut post_if_block.body);

        let next_id = self.id_gen.next();
        function.consume(post_if_block, Instruction::branch(next_id));
        Some(next_id)
    }

    /// Generate an `OpVariable` for one value in an [`EntryPoint`]'s IO interface.
    ///
    /// The [`Binding`]s of the arguments and result of an [`EntryPoint`]'s
    /// [`Function`] describe a SPIR-V shader interface. In SPIR-V, the
    /// interface is represented by global variables in the `Input` and `Output`
    /// storage classes, with decorations indicating which builtin or location
    /// each variable corresponds to.
    ///
    /// This function emits a single global `OpVariable` for a single value from
    /// the interface, and adds appropriate decorations to indicate which
    /// builtin or location it represents, how it should be interpolated, and so
    /// on. The `class` argument gives the variable's SPIR-V storage class,
    /// which should be either [`Input`] or [`Output`]. The one exception is
    /// `hit_barycentrics`, which overrides `class` to `HitAttributeKHR`.
    ///
    /// [`Binding`]: crate::Binding
    /// [`Function`]: crate::Function
    /// [`EntryPoint`]: crate::EntryPoint
    /// [`Input`]: spirv::StorageClass::Input
    /// [`Output`]: spirv::StorageClass::Output
    fn write_varying(
        &mut self,
        ir_module: &crate::Module,
        stage: crate::ShaderStage,
        class: spirv::StorageClass,
        debug_name: Option<&str>,
        ty: Handle<crate::Type>,
        binding: &crate::Binding,
    ) -> Result<Word, Error> {
        // Triangle barycentrics are supplied as a `HitAttributeKHR` variable, not as an `Input`
        // builtin: when a hit group has no intersection shader, the hit attribute an any-hit or
        // closest-hit shader reads is a two-component float vector holding the barycentric
        // coordinates of the hit, and hit attributes carry no `BuiltIn` decoration.
        //
        // SPEC: https://registry.khronos.org/vulkan/specs/latest/html/vkspec.html#interfaces-raypipeline
        //
        // Declaring it here, alongside the entry point's other varyings, works because ray tracing
        // requires SPIR-V 1.4, where the `OpEntryPoint` interface lists global variables from every
        // storage class rather than just `Input` and `Output`.
        //
        // SPEC: https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpEntryPoint
        let class = match *binding {
            crate::Binding::BuiltIn(crate::BuiltIn::HitBarycentrics) => {
                spirv::StorageClass::HitAttributeKHR
            }
            _ => class,
        };

        let id = self.id_gen.next();
        let ty_inner = &ir_module.types[ty].inner;
        let needs_polyfill = self.needs_f16_polyfill(ty_inner);

        let pointer_type_id = if needs_polyfill {
            let f32_value_local =
                super::f16_polyfill::F16IoPolyfill::create_polyfill_type(ty_inner)
                    .expect("needs_polyfill returned true but create_polyfill_type returned None");

            let f32_type_id = self.get_localtype_id(f32_value_local);
            let ptr_id = self.get_pointer_type_id(f32_type_id, class);
            self.io_f16_polyfills.register_io_var(id, f32_type_id);

            ptr_id
        } else {
            self.get_handle_pointer_type_id(ty, class)
        };

        Instruction::variable(pointer_type_id, id, class, None)
            .to_words(&mut self.logical_layout.declarations);

        if self
            .flags
            .contains(WriterFlags::DEBUG | WriterFlags::LABEL_VARYINGS)
        {
            if let Some(name) = debug_name {
                self.debugs.push(Instruction::name(id, name));
            }
        }

        let binding = self.map_binding(ir_module, stage, class, ty, binding)?;
        self.write_binding(id, binding);

        Ok(id)
    }

    pub fn write_binding(&mut self, id: Word, binding: BindingDecorations) {
        match binding {
            BindingDecorations::None => (),
            BindingDecorations::BuiltIn(bi, others) => {
                self.decorate(id, spirv::Decoration::BuiltIn, &[bi as u32]);
                for other in others {
                    self.decorate(id, other, &[]);
                }
            }
            BindingDecorations::Location {
                location,
                others,
                blend_src,
            } => {
                self.decorate(id, spirv::Decoration::Location, &[location]);
                for other in others {
                    self.decorate(id, other, &[]);
                }
                if let Some(blend_src) = blend_src {
                    self.decorate(id, spirv::Decoration::Index, &[blend_src]);
                }
            }
        }
    }

    pub fn write_binding_struct_member(
        &mut self,
        struct_id: Word,
        member_idx: Word,
        binding_info: BindingDecorations,
    ) {
        match binding_info {
            BindingDecorations::None => (),
            BindingDecorations::BuiltIn(bi, others) => {
                self.annotations.push(Instruction::member_decorate(
                    struct_id,
                    member_idx,
                    spirv::Decoration::BuiltIn,
                    &[bi as Word],
                ));
                for other in others {
                    self.annotations.push(Instruction::member_decorate(
                        struct_id,
                        member_idx,
                        other,
                        &[],
                    ));
                }
            }
            BindingDecorations::Location {
                location,
                others,
                blend_src,
            } => {
                self.annotations.push(Instruction::member_decorate(
                    struct_id,
                    member_idx,
                    spirv::Decoration::Location,
                    &[location],
                ));
                for other in others {
                    self.annotations.push(Instruction::member_decorate(
                        struct_id,
                        member_idx,
                        other,
                        &[],
                    ));
                }
                if let Some(blend_src) = blend_src {
                    self.annotations.push(Instruction::member_decorate(
                        struct_id,
                        member_idx,
                        spirv::Decoration::Index,
                        &[blend_src],
                    ));
                }
            }
        }
    }

    pub fn map_binding(
        &mut self,
        ir_module: &crate::Module,
        stage: crate::ShaderStage,
        class: spirv::StorageClass,
        ty: Handle<crate::Type>,
        binding: &crate::Binding,
    ) -> Result<BindingDecorations, Error> {
        use spirv::BuiltIn;
        use spirv::Decoration;
        match *binding {
            crate::Binding::Location {
                location,
                interpolation,
                sampling,
                blend_src,
                per_primitive,
            } => {
                let mut others = ArrayVec::new();

                let no_decorations =
                    // VUID-StandaloneSpirv-Flat-06202
                    // > The Flat, NoPerspective, Sample, and Centroid decorations
                    // > must not be used on variables with the Input storage class in a vertex shader
                    (class == spirv::StorageClass::Input && stage == crate::ShaderStage::Vertex) ||
                    // VUID-StandaloneSpirv-Flat-06201
                    // > The Flat, NoPerspective, Sample, and Centroid decorations
                    // > must not be used on variables with the Output storage class in a fragment shader
                    (class == spirv::StorageClass::Output && stage == crate::ShaderStage::Fragment);

                if !no_decorations {
                    match interpolation {
                        // Perspective-correct interpolation is the default in SPIR-V.
                        None | Some(crate::Interpolation::Perspective) => (),
                        Some(crate::Interpolation::Flat) => {
                            others.push(Decoration::Flat);
                        }
                        Some(crate::Interpolation::Linear) => {
                            others.push(Decoration::NoPerspective);
                        }
                        Some(crate::Interpolation::PerVertex) => {
                            others.push(Decoration::PerVertexKHR);
                            self.require_any(
                                "`per_vertex` interpolation",
                                &[spirv::Capability::FragmentBarycentricKHR],
                            )?;
                            self.use_extension("SPV_KHR_fragment_shader_barycentric");
                        }
                    }
                    match sampling {
                        // Center sampling is the default in SPIR-V.
                        None
                        | Some(
                            crate::Sampling::Center
                            | crate::Sampling::First
                            | crate::Sampling::Either,
                        ) => (),
                        Some(crate::Sampling::Centroid) => {
                            others.push(Decoration::Centroid);
                        }
                        Some(crate::Sampling::Sample) => {
                            self.require_any(
                                "per-sample interpolation",
                                &[spirv::Capability::SampleRateShading],
                            )?;
                            others.push(Decoration::Sample);
                        }
                    }
                }
                if per_primitive && stage == crate::ShaderStage::Fragment {
                    others.push(Decoration::PerPrimitiveEXT);
                }
                Ok(BindingDecorations::Location {
                    location,
                    others,
                    blend_src,
                })
            }
            crate::Binding::BuiltIn(built_in) => {
                use crate::BuiltIn as Bi;
                let mut others = ArrayVec::new();

                let built_in = match built_in {
                    Bi::Position { invariant } => {
                        if invariant {
                            others.push(Decoration::Invariant);
                        }

                        if class == spirv::StorageClass::Output {
                            BuiltIn::Position
                        } else {
                            BuiltIn::FragCoord
                        }
                    }
                    Bi::ViewIndex => {
                        self.require_any("`view_index` built-in", &[spirv::Capability::MultiView])?;
                        BuiltIn::ViewIndex
                    }
                    // vertex
                    Bi::BaseInstance => BuiltIn::BaseInstance,
                    Bi::BaseVertex => BuiltIn::BaseVertex,
                    Bi::ClipDistances => {
                        self.require_any(
                            "`clip_distances` built-in",
                            &[spirv::Capability::ClipDistance],
                        )?;
                        BuiltIn::ClipDistance
                    }
                    Bi::CullDistance => {
                        self.require_any(
                            "`cull_distance` built-in",
                            &[spirv::Capability::CullDistance],
                        )?;
                        BuiltIn::CullDistance
                    }
                    Bi::InstanceIndex => match stage {
                        crate::ShaderStage::AnyHit | crate::ShaderStage::ClosestHit => {
                            BuiltIn::InstanceId
                        }
                        _ => BuiltIn::InstanceIndex,
                    },
                    Bi::PointSize => BuiltIn::PointSize,
                    Bi::VertexIndex => BuiltIn::VertexIndex,
                    Bi::DrawIndex => {
                        self.use_extension("SPV_KHR_shader_draw_parameters");
                        self.require_any(
                            "`draw_index built-in",
                            &[spirv::Capability::DrawParameters],
                        )?;
                        BuiltIn::DrawIndex
                    }
                    // fragment
                    Bi::FragDepth => BuiltIn::FragDepth,
                    Bi::PointCoord => BuiltIn::PointCoord,
                    Bi::FrontFacing => BuiltIn::FrontFacing,
                    Bi::PrimitiveIndex => {
                        // `PrimitiveId` is enabled by any of `Geometry`, `Tessellation`,
                        // `RayTracingKHR` or `MeshShadingEXT`. `require_any` picks the first one the target
                        // allows.
                        //
                        // SPEC: https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#:~:text=PrimitiveId
                        let enabled_by: &[spirv::Capability] = match stage {
                            // The stage itself already requires `RayTracingKHR`.
                            crate::ShaderStage::AnyHit | crate::ShaderStage::ClosestHit => &[],
                            // The stage itself already requires `MeshShadingEXT`.
                            crate::ShaderStage::Mesh => &[],
                            // A fragment shader can be fed primitive IDs by a geometry,
                            // tessellation or mesh pipeline.
                            crate::ShaderStage::Fragment => &[
                                spirv::Capability::Geometry,
                                spirv::Capability::Tessellation,
                                spirv::Capability::MeshShadingEXT,
                            ],
                            // `PrimitiveId` isn't permitted in these execution models.
                            _ => return Err(Error::Validation(
                                "`primitive_index` built-in is not allowed in this shader stage",
                            )),
                        };
                        self.require_any("`primitive_index` built-in", enabled_by)?;

                        if stage == crate::ShaderStage::Mesh {
                            others.push(Decoration::PerPrimitiveEXT);
                        }

                        BuiltIn::PrimitiveId
                    }
                    Bi::Barycentric { perspective } => {
                        self.require_any(
                            "`barycentric` built-in",
                            &[spirv::Capability::FragmentBarycentricKHR],
                        )?;
                        self.use_extension("SPV_KHR_fragment_shader_barycentric");
                        if perspective {
                            BuiltIn::BaryCoordKHR
                        } else {
                            BuiltIn::BaryCoordNoPerspKHR
                        }
                    }
                    Bi::SampleIndex => {
                        self.require_any(
                            "`sample_index` built-in",
                            &[spirv::Capability::SampleRateShading],
                        )?;

                        BuiltIn::SampleId
                    }
                    Bi::SampleMask => BuiltIn::SampleMask,
                    // compute
                    Bi::GlobalInvocationId => BuiltIn::GlobalInvocationId,
                    Bi::LocalInvocationId => BuiltIn::LocalInvocationId,
                    Bi::LocalInvocationIndex => BuiltIn::LocalInvocationIndex,
                    Bi::WorkGroupId => BuiltIn::WorkgroupId,
                    Bi::WorkGroupSize => BuiltIn::WorkgroupSize,
                    Bi::NumWorkGroups => BuiltIn::NumWorkgroups,
                    // Subgroup
                    Bi::NumSubgroups => {
                        self.require_any(
                            "`num_subgroups` built-in",
                            &[spirv::Capability::GroupNonUniform],
                        )?;
                        BuiltIn::NumSubgroups
                    }
                    Bi::SubgroupId => {
                        self.require_any(
                            "`subgroup_id` built-in",
                            &[spirv::Capability::GroupNonUniform],
                        )?;
                        BuiltIn::SubgroupId
                    }
                    Bi::SubgroupSize => {
                        self.require_any(
                            "`subgroup_size` built-in",
                            &[
                                spirv::Capability::GroupNonUniform,
                                spirv::Capability::SubgroupBallotKHR,
                            ],
                        )?;
                        BuiltIn::SubgroupSize
                    }
                    Bi::SubgroupInvocationId => {
                        self.require_any(
                            "`subgroup_invocation_id` built-in",
                            &[
                                spirv::Capability::GroupNonUniform,
                                spirv::Capability::SubgroupBallotKHR,
                            ],
                        )?;
                        BuiltIn::SubgroupLocalInvocationId
                    }
                    Bi::CullPrimitive => {
                        others.push(Decoration::PerPrimitiveEXT);
                        BuiltIn::CullPrimitiveEXT
                    }
                    Bi::PointIndex => BuiltIn::PrimitivePointIndicesEXT,
                    Bi::LineIndices => BuiltIn::PrimitiveLineIndicesEXT,
                    Bi::TriangleIndices => BuiltIn::PrimitiveTriangleIndicesEXT,
                    // No decoration, this EmitMeshTasksEXT is called at function return
                    Bi::MeshTaskSize => return Ok(BindingDecorations::None),
                    // These aren't normal builtins and don't occur in function output
                    Bi::VertexCount | Bi::Vertices | Bi::PrimitiveCount | Bi::Primitives => {
                        unreachable!()
                    }
                    // ray tracing pipeline
                    Bi::RayInvocationId => BuiltIn::LaunchIdKHR,
                    Bi::NumRayInvocations => BuiltIn::LaunchSizeKHR,
                    Bi::InstanceCustomData => BuiltIn::InstanceCustomIndexKHR,
                    Bi::GeometryIndex => BuiltIn::RayGeometryIndexKHR,
                    Bi::WorldRayOrigin => BuiltIn::WorldRayOriginKHR,
                    Bi::WorldRayDirection => BuiltIn::WorldRayDirectionKHR,
                    Bi::ObjectRayOrigin => BuiltIn::ObjectRayOriginKHR,
                    Bi::ObjectRayDirection => BuiltIn::ObjectRayDirectionKHR,
                    Bi::RayTmin => BuiltIn::RayTminKHR,
                    Bi::RayTCurrentMax => BuiltIn::RayTmaxKHR,
                    Bi::ObjectToWorld => BuiltIn::ObjectToWorldKHR,
                    Bi::WorldToObject => BuiltIn::WorldToObjectKHR,
                    Bi::HitKind => BuiltIn::HitKindKHR,
                    // Read from the entry point's `HitAttributeKHR` variable instead.
                    Bi::HitBarycentrics => return Ok(BindingDecorations::None),
                };

                use crate::ScalarKind as Sk;

                // Per the Vulkan spec, `VUID-StandaloneSpirv-Flat-04744`:
                //
                // > Any variable with integer or double-precision floating-
                // > point type and with Input storage class in a fragment
                // > shader, must be decorated Flat
                if class == spirv::StorageClass::Input && stage == crate::ShaderStage::Fragment {
                    let is_flat = match ir_module.types[ty].inner {
                        crate::TypeInner::Scalar(scalar)
                        | crate::TypeInner::Vector { scalar, .. } => match scalar.kind {
                            Sk::Uint | Sk::Sint | Sk::Bool => true,
                            Sk::Float => false,
                            Sk::AbstractInt | Sk::AbstractFloat => {
                                return Err(Error::Validation(
                                    "Abstract types should not appear in IR presented to backends",
                                ))
                            }
                        },
                        _ => false,
                    };

                    if is_flat {
                        others.push(Decoration::Flat);
                    }
                }
                Ok(BindingDecorations::BuiltIn(built_in, others))
            }
        }
    }

    /// Load an IO variable, converting from `f32` to `f16` if polyfill is active.
    /// Returns the id of the loaded value matching `target_type_id`.
    pub(super) fn load_io_with_f16_polyfill(
        &mut self,
        body: &mut Vec<Instruction>,
        varying_id: Word,
        target_type_id: Word,
    ) -> Word {
        let tmp = self.id_gen.next();
        if let Some(f32_ty) = self.io_f16_polyfills.get_f32_io_type(varying_id) {
            body.push(Instruction::load(f32_ty, tmp, varying_id, None));
            let converted = self.id_gen.next();
            super::f16_polyfill::F16IoPolyfill::emit_f32_to_f16_conversion(
                tmp,
                target_type_id,
                converted,
                body,
            );
            converted
        } else {
            body.push(Instruction::load(target_type_id, tmp, varying_id, None));
            tmp
        }
    }

    /// Store an IO variable, converting from `f16` to `f32` if polyfill is active.
    pub(super) fn store_io_with_f16_polyfill(
        &mut self,
        body: &mut Vec<Instruction>,
        varying_id: Word,
        value_id: Word,
    ) {
        if let Some(f32_ty) = self.io_f16_polyfills.get_f32_io_type(varying_id) {
            let converted = self.id_gen.next();
            super::f16_polyfill::F16IoPolyfill::emit_f16_to_f32_conversion(
                value_id, f32_ty, converted, body,
            );
            body.push(Instruction::store(varying_id, converted, None));
        } else {
            body.push(Instruction::store(varying_id, value_id, None));
        }
    }

    fn write_global_variable(
        &mut self,
        ir_module: &crate::Module,
        global_variable: &crate::GlobalVariable,
    ) -> Result<Word, Error> {
        use spirv::Decoration;

        let id = self.id_gen.next();
        let class = map_storage_class(global_variable.space);

        if let crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload =
            global_variable.space
        {
            self.require_any("ray tracing pipelines", &[spirv::Capability::RayTracingKHR])?;
        }

        //self.check(class.required_capabilities())?;

        if global_variable
            .memory_decorations
            .contains(crate::MemoryDecorations::COHERENT)
        {
            self.decorate(id, Decoration::Coherent, &[]);
        }
        if global_variable
            .memory_decorations
            .contains(crate::MemoryDecorations::VOLATILE)
        {
            self.decorate(id, Decoration::Volatile, &[]);
        }

        if self.flags.contains(WriterFlags::DEBUG) {
            if let Some(ref name) = global_variable.name {
                self.debugs.push(Instruction::name(id, name));
            }
        }

        let storage_access = match global_variable.space {
            crate::AddressSpace::Storage { access } => Some(access),
            _ => match ir_module.types[global_variable.ty].inner {
                crate::TypeInner::Image {
                    class: crate::ImageClass::Storage { access, .. },
                    ..
                } => Some(access),
                _ => None,
            },
        };
        if let Some(storage_access) = storage_access {
            if !storage_access.contains(crate::StorageAccess::LOAD) {
                self.decorate(id, Decoration::NonReadable, &[]);
            }
            if !storage_access.contains(crate::StorageAccess::STORE) {
                self.decorate(id, Decoration::NonWritable, &[]);
            }
        }

        // Note: we should be able to substitute `binding_array<Foo, 0>`,
        // but there is still code that tries to register the pre-substituted type,
        // and it is failing on 0.
        let mut substitute_inner_type_lookup = None;
        if let Some(ref res_binding) = global_variable.binding {
            let bind_target = self.resolve_resource_binding(res_binding)?;
            self.decorate(id, Decoration::DescriptorSet, &[bind_target.descriptor_set]);
            self.decorate(id, Decoration::Binding, &[bind_target.binding]);

            if let Some(remapped_binding_array_size) = bind_target.binding_array_size {
                if let crate::TypeInner::BindingArray { base, .. } =
                    ir_module.types[global_variable.ty].inner
                {
                    let binding_array_type_id =
                        self.get_type_id(LookupType::Local(LocalType::BindingArray {
                            base,
                            size: remapped_binding_array_size,
                        }));
                    substitute_inner_type_lookup = Some(LookupType::Local(LocalType::Pointer {
                        base: binding_array_type_id,
                        class,
                    }));
                }
            }
        };

        let init_word = global_variable
            .init
            .map(|constant| self.constant_ids[constant]);
        let inner_type_id = self.get_type_id(
            substitute_inner_type_lookup.unwrap_or(LookupType::Handle(global_variable.ty)),
        );

        // generate the wrapping structure if needed
        let pointer_type_id = if global_needs_wrapper(ir_module, global_variable) {
            let wrapper_type_id = self.id_gen.next();

            self.decorate(wrapper_type_id, Decoration::Block, &[]);

            match self.std140_compat_uniform_types.get(&global_variable.ty) {
                Some(std140_type_info) if global_variable.space == crate::AddressSpace::Uniform => {
                    self.annotations.push(Instruction::member_decorate(
                        wrapper_type_id,
                        0,
                        Decoration::Offset,
                        &[0],
                    ));
                    Instruction::type_struct(wrapper_type_id, &[std140_type_info.type_id])
                        .to_words(&mut self.logical_layout.declarations);
                }
                _ => {
                    let member = crate::StructMember {
                        name: None,
                        ty: global_variable.ty,
                        binding: None,
                        offset: 0,
                    };
                    self.decorate_struct_member(wrapper_type_id, 0, &member, &ir_module.types)?;

                    Instruction::type_struct(wrapper_type_id, &[inner_type_id])
                        .to_words(&mut self.logical_layout.declarations);
                }
            }

            let pointer_type_id = self.id_gen.next();
            Instruction::type_pointer(pointer_type_id, class, wrapper_type_id)
                .to_words(&mut self.logical_layout.declarations);

            pointer_type_id
        } else {
            // This is a global variable in the Storage address space. The only
            // way it could have `global_needs_wrapper() == false` is if it has
            // a runtime-sized or binding array.
            // Runtime-sized arrays were decorated when iterating through struct content.
            // Now binding arrays require Block decorating.
            if let crate::AddressSpace::Storage { .. } = global_variable.space {
                match ir_module.types[global_variable.ty].inner {
                    crate::TypeInner::BindingArray { base, .. } => {
                        let ty = &ir_module.types[base];
                        let mut should_decorate = true;
                        // Check if the type has a runtime array.
                        // A normal runtime array gets validated out,
                        // so only structs can be with runtime arrays
                        if let crate::TypeInner::Struct { ref members, .. } = ty.inner {
                            // only the last member in a struct can be dynamically sized
                            if let Some(last_member) = members.last() {
                                if let &crate::TypeInner::Array {
                                    size: crate::ArraySize::Dynamic,
                                    ..
                                } = &ir_module.types[last_member.ty].inner
                                {
                                    should_decorate = false;
                                }
                            }
                        }
                        if should_decorate {
                            let decorated_id = self.get_handle_type_id(base);
                            self.decorate(decorated_id, Decoration::Block, &[]);
                        }
                    }
                    _ => (),
                };
            }
            if substitute_inner_type_lookup.is_some() {
                inner_type_id
            } else {
                self.get_handle_pointer_type_id(global_variable.ty, class)
            }
        };

        let init_word = match (global_variable.space, self.zero_initialize_workgroup_memory) {
            (crate::AddressSpace::Private, _)
            | (crate::AddressSpace::WorkGroup, super::ZeroInitializeWorkgroupMemoryMode::Native) => {
                init_word.or_else(|| Some(self.get_constant_null(inner_type_id)))
            }
            _ => init_word,
        };

        Instruction::variable(pointer_type_id, id, class, init_word)
            .to_words(&mut self.logical_layout.declarations);
        Ok(id)
    }

    /// Write the necessary decorations for a struct member.
    ///
    /// Emit decorations for the `index`'th member of the struct type
    /// designated by `struct_id`, described by `member`.
    fn decorate_struct_member(
        &mut self,
        struct_id: Word,
        index: usize,
        member: &crate::StructMember,
        arena: &UniqueArena<crate::Type>,
    ) -> Result<(), Error> {
        use spirv::Decoration;

        self.annotations.push(Instruction::member_decorate(
            struct_id,
            index as u32,
            Decoration::Offset,
            &[member.offset],
        ));

        if self.flags.contains(WriterFlags::DEBUG) {
            if let Some(ref name) = member.name {
                self.debugs
                    .push(Instruction::member_name(struct_id, index as u32, name));
            }
        }

        // Matrices and (potentially nested) arrays of matrices both require decorations,
        // so "see through" any arrays to determine if they're needed.
        let mut member_array_subty_inner = &arena[member.ty].inner;
        while let crate::TypeInner::Array { base, .. } = *member_array_subty_inner {
            member_array_subty_inner = &arena[base].inner;
        }

        if let crate::TypeInner::Matrix {
            columns: _,
            rows,
            scalar,
        } = *member_array_subty_inner
        {
            let byte_stride = Alignment::from(rows) * scalar.width as u32;
            self.annotations.push(Instruction::member_decorate(
                struct_id,
                index as u32,
                Decoration::ColMajor,
                &[],
            ));
            self.annotations.push(Instruction::member_decorate(
                struct_id,
                index as u32,
                Decoration::MatrixStride,
                &[byte_stride],
            ));
        }

        Ok(())
    }

    pub(super) fn get_function_type(&mut self, lookup_function_type: LookupFunctionType) -> Word {
        match self
            .lookup_function_type
            .entry(lookup_function_type.clone())
        {
            Entry::Occupied(e) => *e.get(),
            Entry::Vacant(_) => {
                let id = self.id_gen.next();
                let instruction = Instruction::type_function(
                    id,
                    lookup_function_type.return_type_id,
                    &lookup_function_type.parameter_type_ids,
                );
                instruction.to_words(&mut self.logical_layout.declarations);
                self.lookup_function_type.insert(lookup_function_type, id);
                id
            }
        }
    }

    const fn write_physical_layout(&mut self) {
        self.physical_layout.bound = self.id_gen.0 + 1;
    }

    fn write_logical_layout(
        &mut self,
        ir_module: &crate::Module,
        mod_info: &ModuleInfo,
        ep_index: Option<usize>,
        debug_info: &Option<DebugInfo>,
    ) -> Result<(), Error> {
        fn has_view_index_check(
            ir_module: &crate::Module,
            binding: Option<&crate::Binding>,
            ty: Handle<crate::Type>,
        ) -> bool {
            match ir_module.types[ty].inner {
                crate::TypeInner::Struct { ref members, .. } => members.iter().any(|member| {
                    has_view_index_check(ir_module, member.binding.as_ref(), member.ty)
                }),
                _ => binding == Some(&crate::Binding::BuiltIn(crate::BuiltIn::ViewIndex)),
            }
        }

        let has_storage_buffers =
            ir_module
                .global_variables
                .iter()
                .any(|(_, var)| match var.space {
                    crate::AddressSpace::Storage { .. } => true,
                    _ => false,
                });
        let has_view_index = ir_module
            .entry_points
            .iter()
            .flat_map(|entry| entry.function.arguments.iter())
            .any(|arg| has_view_index_check(ir_module, arg.binding.as_ref(), arg.ty));
        let has_vertex_return = ir_module.special_types.ray_vertex_return.is_some();

        let rt_uses = ir_module.uses_ray_tracing(ep_index);
        let has_ray_query = rt_uses.queries;
        let has_ray_tracing_pipeline = rt_uses.pipelines;

        self.has_ray_tracing_pipeline = has_ray_tracing_pipeline;

        if self.physical_layout.version < 0x10300 && has_storage_buffers {
            // enable the storage buffer class on < SPV-1.3
            Instruction::extension("SPV_KHR_storage_buffer_storage_class")
                .to_words(&mut self.logical_layout.extensions);
        }
        if has_view_index {
            Instruction::extension("SPV_KHR_multiview")
                .to_words(&mut self.logical_layout.extensions)
        }
        if has_ray_query {
            Instruction::extension("SPV_KHR_ray_query")
                .to_words(&mut self.logical_layout.extensions)
        }
        if has_vertex_return {
            Instruction::extension("SPV_KHR_ray_tracing_position_fetch")
                .to_words(&mut self.logical_layout.extensions);
        }
        if ir_module.uses_mesh_shaders() {
            self.use_extension("SPV_EXT_mesh_shader");
            self.require_any("Mesh Shaders", &[spirv::Capability::MeshShadingEXT])?;
            let lang_version = self.lang_version();
            if lang_version.0 <= 1 && lang_version.1 < 4 {
                return Err(Error::SpirvVersionTooLow(1, 4));
            }
        }
        if has_ray_tracing_pipeline {
            // `SPV_KHR_ray_tracing` requires SPIR-V 1.4. That is also what lets `write_varying`
            // emit `hit_barycentrics` as a `HitAttributeKHR` variable, since only from 1.4 on does
            // the `OpEntryPoint` interface cover storage classes other than `Input` and `Output`.
            //
            // SPEC: https://github.khronos.org/SPIRV-Registry/extensions/KHR/SPV_KHR_ray_tracing.html
            let lang_version = self.lang_version();
            if lang_version.0 <= 1 && lang_version.1 < 4 {
                return Err(Error::SpirvVersionTooLow(1, 4));
            }
            Instruction::extension("SPV_KHR_ray_tracing")
                .to_words(&mut self.logical_layout.extensions)
        }
        Instruction::type_void(self.void_type).to_words(&mut self.logical_layout.declarations);
        Instruction::ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450")
            .to_words(&mut self.logical_layout.ext_inst_imports);

        let mut debug_info_inner = None;
        if self.flags.contains(WriterFlags::DEBUG) {
            if let Some(debug_info) = debug_info.as_ref() {
                let source_file_id = self.id_gen.next();
                self.debugs
                    .push(Instruction::string(debug_info.file_name, source_file_id));

                debug_info_inner = Some(DebugInfoInner {
                    source_code: debug_info.source_code,
                    source_file_id,
                });
                self.debugs.append(&mut Instruction::source_auto_continued(
                    debug_info.language,
                    0,
                    &debug_info_inner,
                ));
            }
        }

        // write all types
        for (handle, _) in ir_module.types.iter() {
            self.write_type_declaration_arena(ir_module, handle)?;
        }

        // write std140 layout compatible types required by uniforms
        for (_, var) in ir_module.global_variables.iter() {
            if var.space == crate::AddressSpace::Uniform {
                self.write_std140_compat_type_declaration(ir_module, var.ty)?;
            }
        }

        // write all const-expressions as constants
        self.constant_ids
            .resize(ir_module.global_expressions.len(), 0);
        for (handle, _) in ir_module.global_expressions.iter() {
            self.write_constant_expr(handle, ir_module, mod_info)?;
        }
        debug_assert!(self.constant_ids.iter().all(|&id| id != 0));

        // write the name of constants on their respective const-expression initializer
        if self.flags.contains(WriterFlags::DEBUG) {
            for (_, constant) in ir_module.constants.iter() {
                if let Some(ref name) = constant.name {
                    let id = self.constant_ids[constant.init];
                    self.debugs.push(Instruction::name(id, name));
                }
            }
        }

        // write all global variables
        for (handle, var) in ir_module.global_variables.iter() {
            // If a single entry point was specified, only write `OpVariable` instructions
            // for the globals it actually uses. Emit dummies for the others,
            // to preserve the indices in `global_variables`.
            let gvar = match ep_index {
                Some(index) if mod_info.get_entry_point(index)[handle].is_empty() => {
                    GlobalVariable::dummy()
                }
                _ => {
                    let id = self.write_global_variable(ir_module, var)?;
                    GlobalVariable::new(id)
                }
            };
            self.global_variables.insert(handle, gvar);
        }

        // write all functions
        for (handle, ir_function) in ir_module.functions.iter() {
            let info = &mod_info[handle];
            if let Some(index) = ep_index {
                let ep_info = mod_info.get_entry_point(index);
                // If this function uses globals that we omitted from the SPIR-V
                // because the entry point and its callees didn't use them,
                // then we must skip it.
                if !ep_info.dominates_global_use(info) {
                    log::debug!("Skip function {:?}", ir_function.name);
                    continue;
                }

                // Skip functions that that are not compatible with this entry point's stage.
                //
                // When validation is enabled, it rejects modules whose entry points try to call
                // incompatible functions, so if we got this far, then any functions incompatible
                // with our selected entry point must not be used.
                //
                // When validation is disabled, `fun_info.available_stages` is always just
                // `ShaderStages::all()`, so this will write all functions in the module, and
                // the downstream GLSL compiler will catch any problems.
                if !info.available_stages.contains(ep_info.available_stages) {
                    continue;
                }
            }
            let id = self.write_function(ir_function, info, ir_module, None, &debug_info_inner)?;
            self.lookup_function.insert(handle, id);
        }

        // write all or one entry points
        for (index, ir_ep) in ir_module.entry_points.iter().enumerate() {
            if ep_index.is_some() && ep_index != Some(index) {
                continue;
            }
            let info = mod_info.get_entry_point(index);
            let ep_instruction =
                self.write_entry_point(ir_ep, info, ir_module, &debug_info_inner)?;
            ep_instruction.to_words(&mut self.logical_layout.entry_points);
        }

        for capability in self.capabilities_used.iter() {
            Instruction::capability(*capability).to_words(&mut self.logical_layout.capabilities);
        }
        for extension in self.extensions_used.iter() {
            Instruction::extension(extension).to_words(&mut self.logical_layout.extensions);
        }
        if ir_module.entry_points.is_empty() {
            // SPIR-V doesn't like modules without entry points
            Instruction::capability(spirv::Capability::Linkage)
                .to_words(&mut self.logical_layout.capabilities);
        }

        let addressing_model = spirv::AddressingModel::Logical;
        let memory_model = if self
            .capabilities_used
            .contains(&spirv::Capability::VulkanMemoryModel)
        {
            spirv::MemoryModel::Vulkan
        } else {
            spirv::MemoryModel::GLSL450
        };
        //self.check(addressing_model.required_capabilities())?;
        //self.check(memory_model.required_capabilities())?;

        Instruction::memory_model(addressing_model, memory_model)
            .to_words(&mut self.logical_layout.memory_model);

        for debug_string in self.debug_strings.iter() {
            debug_string.to_words(&mut self.logical_layout.debugs);
        }

        if self.flags.contains(WriterFlags::DEBUG) {
            for debug in self.debugs.iter() {
                debug.to_words(&mut self.logical_layout.debugs);
            }
        }

        for annotation in self.annotations.iter() {
            annotation.to_words(&mut self.logical_layout.annotations);
        }

        Ok(())
    }

    pub fn write(
        &mut self,
        ir_module: &crate::Module,
        info: &ModuleInfo,
        pipeline_options: Option<&PipelineOptions>,
        debug_info: &Option<DebugInfo>,
        words: &mut Vec<Word>,
    ) -> Result<(), Error> {
        self.reset();

        // Try to find the entry point and corresponding index
        let ep_index = match pipeline_options {
            Some(po) => {
                let index = ir_module
                    .entry_points
                    .iter()
                    .position(|ep| po.shader_stage == ep.stage && po.entry_point == ep.name)
                    .ok_or(Error::EntryPointNotFound)?;
                Some(index)
            }
            None => None,
        };

        self.write_logical_layout(ir_module, info, ep_index, debug_info)?;
        self.write_physical_layout();

        self.physical_layout.in_words(words);
        self.logical_layout.in_words(words);
        Ok(())
    }

    /// Return the set of capabilities the last module written used.
    pub const fn get_capabilities_used(&self) -> &crate::FastIndexSet<spirv::Capability> {
        &self.capabilities_used
    }

    pub fn decorate_non_uniform_binding_array_access(&mut self, id: Word) -> Result<(), Error> {
        self.require_any("NonUniformEXT", &[spirv::Capability::ShaderNonUniform])?;
        self.use_extension("SPV_EXT_descriptor_indexing");
        self.decorate(id, spirv::Decoration::NonUniform, &[]);
        Ok(())
    }

    pub(super) fn needs_f16_polyfill(&self, ty_inner: &crate::TypeInner) -> bool {
        self.io_f16_polyfills.needs_polyfill(ty_inner)
    }

    pub(super) fn write_debug_printf(
        &mut self,
        block: &mut Block,
        string: &str,
        format_params: &[Word],
    ) {
        if self.debug_printf.is_none() {
            self.use_extension("SPV_KHR_non_semantic_info");
            let import_id = self.id_gen.next();
            Instruction::ext_inst_import(import_id, "NonSemantic.DebugPrintf")
                .to_words(&mut self.logical_layout.ext_inst_imports);
            self.debug_printf = Some(import_id)
        }

        let import_id = self.debug_printf.unwrap();

        let string_id = self.id_gen.next();
        self.debug_strings
            .push(Instruction::string(string, string_id));

        let mut operands = Vec::with_capacity(1 + format_params.len());
        operands.push(string_id);
        operands.extend(format_params.iter());

        let print_id = self.id_gen.next();
        block.body.push(Instruction::ext_inst(
            import_id,
            1,
            self.void_type,
            print_id,
            &operands,
        ));
    }
}

#[test]
fn test_write_physical_layout() {
    let mut writer = Writer::new(&Options::default()).unwrap();
    assert_eq!(writer.physical_layout.bound, 0);
    writer.write_physical_layout();
    assert_eq!(writer.physical_layout.bound, 3);
}