Quellcode-Bibliothek mod.rs   Sprache: unbekannt

/*!
Backend for [GLSL][glsl] (OpenGL Shading Language).

The main structure is [`Writer`], it maintains internal state that is used
to output a [`Module`](crate::Module) into glsl

# Supported versions
### Core
- 330
- 400
- 410
- 420
- 430
- 450

### ES
- 300
- 310

[glsl]: https://www.khronos.org/registry/OpenGL/index_gl.php
*/

// GLSL is mostly a superset of C but it also removes some parts of it this is a list of relevant
// aspects for this backend.
//
// The most notable change is the introduction of the version preprocessor directive that must
// always be the first line of a glsl file and is written as
// `#version number profile`
// `number` is the version itself (i.e. 300) and `profile` is the
// shader profile we only support "core" and "es", the former is used in desktop applications and
// the later is used in embedded contexts, mobile devices and browsers. Each one as it's own
// versions (at the time of writing this the latest version for "core" is 460 and for "es" is 320)
//
// Other important preprocessor addition is the extension directive which is written as
// `#extension name: behaviour`
// Extensions provide increased features in a plugin fashion but they aren't required to be
// supported hence why they are called extensions, that's why `behaviour` is used it specifies
// whether the extension is strictly required or if it should only be enabled if needed. In our case
// when we use extensions we set behaviour to `require` always.
//
// The only thing that glsl removes that makes a difference are pointers.
//
// Additions that are relevant for the backend are the discard keyword, the introduction of
// vector, matrices, samplers, image types and functions that provide common shader operations

pub use features::Features;

use crate::{
    back::{self, Baked},
    proc::{self, ExpressionKindTracker, NameKey},
    valid, Handle, ShaderStage, TypeInner,
};
use features::FeaturesManager;
use std::{
    cmp::Ordering,
    fmt::{self, Error as FmtError, Write},
    mem,
};
use thiserror::Error;

/// Contains the features related code and the features querying method
mod features;
/// Contains a constant with a slice of all the reserved keywords RESERVED_KEYWORDS
mod keywords;

/// List of supported `core` GLSL versions.
pub const SUPPORTED_CORE_VERSIONS: &[u16] = &[140, 150, 330, 400, 410, 420, 430, 440, 450, 460];
/// List of supported `es` GLSL versions.
pub const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310, 320];

/// The suffix of the variable that will hold the calculated clamped level
/// of detail for bounds checking in `ImageLoad`
const CLAMPED_LOD_SUFFIX: &str = "_clamped_lod";

pub(crate) const MODF_FUNCTION: &str = "naga_modf";
pub(crate) const FREXP_FUNCTION: &str = "naga_frexp";

// Must match code in glsl_built_in
pub const FIRST_INSTANCE_BINDING: &str = "naga_vs_first_instance";

/// Mapping between resources and bindings.
pub type BindingMap = std::collections::BTreeMap<crate::ResourceBinding, u8>;

impl crate::AtomicFunction {
    const fn to_glsl(self) -> &'static str {
        match self {
            Self::Add | Self::Subtract => "Add",
            Self::And => "And",
            Self::InclusiveOr => "Or",
            Self::ExclusiveOr => "Xor",
            Self::Min => "Min",
            Self::Max => "Max",
            Self::Exchange { compare: None } => "Exchange",
            Self::Exchange { compare: Some(_) } => "", //TODO
        }
    }
}

impl crate::AddressSpace {
    const fn is_buffer(&self) -> bool {
        match *self {
            crate::AddressSpace::Uniform | crate::AddressSpace::Storage { .. } => true,
            _ => false,
        }
    }

    /// Whether a variable with this address space can be initialized
    const fn initializable(&self) -> bool {
        match *self {
            crate::AddressSpace::Function | crate::AddressSpace::Private => true,
            crate::AddressSpace::WorkGroup
            | crate::AddressSpace::Uniform
            | crate::AddressSpace::Storage { .. }
            | crate::AddressSpace::Handle
            | crate::AddressSpace::PushConstant => false,
        }
    }
}

/// A GLSL version.
#[derive(Debug, Copy, Clone, PartialEq)]
#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
pub enum Version {
    /// `core` GLSL.
    Desktop(u16),
    /// `es` GLSL.
    Embedded { version: u16, is_webgl: bool },
}

impl Version {
    /// Create a new gles version
    pub const fn new_gles(version: u16) -> Self {
        Self::Embedded {
            version,
            is_webgl: false,
        }
    }

    /// Returns true if self is `Version::Embedded` (i.e. is a es version)
    const fn is_es(&self) -> bool {
        match *self {
            Version::Desktop(_) => false,
            Version::Embedded { .. } => true,
        }
    }

    /// Returns true if targeting WebGL
    const fn is_webgl(&self) -> bool {
        match *self {
            Version::Desktop(_) => false,
            Version::Embedded { is_webgl, .. } => is_webgl,
        }
    }

    /// Checks the list of currently supported versions and returns true if it contains the
    /// specified version
    ///
    /// # Notes
    /// As an invalid version number will never be added to the supported version list
    /// so this also checks for version validity
    fn is_supported(&self) -> bool {
        match *self {
            Version::Desktop(v) => SUPPORTED_CORE_VERSIONS.contains(&v),
            Version::Embedded { version: v, .. } => SUPPORTED_ES_VERSIONS.contains(&v),
        }
    }

    fn supports_io_locations(&self) -> bool {
        *self >= Version::Desktop(330) || *self >= Version::new_gles(300)
    }

    /// Checks if the version supports all of the explicit layouts:
    /// - `location=` qualifiers for bindings
    /// - `binding=` qualifiers for resources
    ///
    /// Note: `location=` for vertex inputs and fragment outputs is supported
    /// unconditionally for GLES 300.
    fn supports_explicit_locations(&self) -> bool {
        *self >= Version::Desktop(420) || *self >= Version::new_gles(310)
    }

    fn supports_early_depth_test(&self) -> bool {
        *self >= Version::Desktop(130) || *self >= Version::new_gles(310)
    }

    fn supports_std430_layout(&self) -> bool {
        *self >= Version::Desktop(430) || *self >= Version::new_gles(310)
    }

    fn supports_fma_function(&self) -> bool {
        *self >= Version::Desktop(400) || *self >= Version::new_gles(320)
    }

    fn supports_integer_functions(&self) -> bool {
        *self >= Version::Desktop(400) || *self >= Version::new_gles(310)
    }

    fn supports_frexp_function(&self) -> bool {
        *self >= Version::Desktop(400) || *self >= Version::new_gles(310)
    }

    fn supports_derivative_control(&self) -> bool {
        *self >= Version::Desktop(450)
    }
}

impl PartialOrd for Version {
    fn partial_cmp(&self, other: &Self) -> Option<Ordering> {
        match (*self, *other) {
            (Version::Desktop(x), Version::Desktop(y)) => Some(x.cmp(&y)),
            (Version::Embedded { version: x, .. }, Version::Embedded { version: y, .. }) => {
                Some(x.cmp(&y))
            }
            _ => None,
        }
    }
}

impl fmt::Display for Version {
    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
        match *self {
            Version::Desktop(v) => write!(f, "{v} core"),
            Version::Embedded { version: v, .. } => write!(f, "{v} es"),
        }
    }
}

bitflags::bitflags! {
    /// Configuration flags for the [`Writer`].
    #[cfg_attr(feature = "serialize", derive(serde::Serialize))]
    #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
    #[derive(Clone, Copy, Debug, Eq, PartialEq)]
    pub struct WriterFlags: u32 {
        /// Flip output Y and extend Z from (0, 1) to (-1, 1).
        const ADJUST_COORDINATE_SPACE = 0x1;
        /// Supports GL_EXT_texture_shadow_lod on the host, which provides
        /// additional functions on shadows and arrays of shadows.
        const TEXTURE_SHADOW_LOD = 0x2;
        /// Supports ARB_shader_draw_parameters on the host, which provides
        /// support for `gl_BaseInstanceARB`, `gl_BaseVertexARB`, `gl_DrawIDARB`, and `gl_DrawID`.
        const DRAW_PARAMETERS = 0x4;
        /// Include unused global variables, constants and functions. By default the output will exclude
        /// global variables that are not used in the specified entrypoint (including indirect use),
        /// all constant declarations, and functions that use excluded global variables.
        const INCLUDE_UNUSED_ITEMS = 0x10;
        /// Emit `PointSize` output builtin to vertex shaders, which is
        /// required for drawing with `PointList` topology.
        ///
        /// https://registry.khronos.org/OpenGL/specs/es/3.2/GLSL_ES_Specification_3.20.html#built-in-language-variables
        /// The variable gl_PointSize is intended for a shader to write the size of the point to be rasterized. It is measured in pixels.
        /// If gl_PointSize is not written to, its value is undefined in subsequent pipe stages.
        const FORCE_POINT_SIZE = 0x20;
    }
}

/// Configuration used in the [`Writer`].
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
#[cfg_attr(feature = "deserialize", serde(default))]
pub struct Options {
    /// The GLSL version to be used.
    pub version: Version,
    /// Configuration flags for the [`Writer`].
    pub writer_flags: WriterFlags,
    /// Map of resources association to binding locations.
    pub binding_map: BindingMap,
    /// Should workgroup variables be zero initialized (by polyfilling)?
    pub zero_initialize_workgroup_memory: bool,
}

impl Default for Options {
    fn default() -> Self {
        Options {
            version: Version::new_gles(310),
            writer_flags: WriterFlags::ADJUST_COORDINATE_SPACE,
            binding_map: BindingMap::default(),
            zero_initialize_workgroup_memory: true,
        }
    }
}

/// A subset of options meant to be changed per pipeline.
#[derive(Debug, Clone)]
#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
pub struct PipelineOptions {
    /// The stage of the entry point.
    pub shader_stage: ShaderStage,
    /// The name of the entry point.
    ///
    /// If no entry point that matches is found while creating a [`Writer`], a error will be thrown.
    pub entry_point: String,
    /// How many views to render to, if doing multiview rendering.
    pub multiview: Option<std::num::NonZeroU32>,
}

#[derive(Debug)]
pub struct VaryingLocation {
    /// The location of the global.
    /// This corresponds to `layout(location = ..)` in GLSL.
    pub location: u32,
    /// The index which can be used for dual source blending.
    /// This corresponds to `layout(index = ..)` in GLSL.
    pub index: u32,
}

/// Reflection info for texture mappings and uniforms.
#[derive(Debug)]
pub struct ReflectionInfo {
    /// Mapping between texture names and variables/samplers.
    pub texture_mapping: crate::FastHashMap<String, TextureMapping>,
    /// Mapping between uniform variables and names.
    pub uniforms: crate::FastHashMap<Handle<crate::GlobalVariable>, String>,
    /// Mapping between names and attribute locations.
    pub varying: crate::FastHashMap<String, VaryingLocation>,
    /// List of push constant items in the shader.
    pub push_constant_items: Vec<PushConstantItem>,
}

/// Mapping between a texture and its sampler, if it exists.
///
/// GLSL pre-Vulkan has no concept of separate textures and samplers. Instead, everything is a
/// `gsamplerN` where `g` is the scalar type and `N` is the dimension. But naga uses separate textures
/// and samplers in the IR, so the backend produces a [`FastHashMap`](crate::FastHashMap) with the texture name
/// as a key and a [`TextureMapping`] as a value. This way, the user knows where to bind.
///
/// [`Storage`](crate::ImageClass::Storage) images produce `gimageN` and don't have an associated sampler,
/// so the [`sampler`](Self::sampler) field will be [`None`].
#[derive(Debug, Clone)]
pub struct TextureMapping {
    /// Handle to the image global variable.
    pub texture: Handle<crate::GlobalVariable>,
    /// Handle to the associated sampler global variable, if it exists.
    pub sampler: Option<Handle<crate::GlobalVariable>>,
}

/// All information to bind a single uniform value to the shader.
///
/// Push constants are emulated using traditional uniforms in OpenGL.
///
/// These are composed of a set of primitives (scalar, vector, matrix) that
/// are given names. Because they are not backed by the concept of a buffer,
/// we must do the work of calculating the offset of each primitive in the
/// push constant block.
#[derive(Debug, Clone)]
pub struct PushConstantItem {
    /// GL uniform name for the item. This name is the same as if you were
    /// to access it directly from a GLSL shader.
    ///
    /// The with the following example, the following names will be generated,
    /// one name per GLSL uniform.
    ///
    /// ```glsl
    /// struct InnerStruct {
    ///     value: f32,
    /// }
    ///
    /// struct PushConstant {
    ///     InnerStruct inner;
    ///     vec4 array[2];
    /// }
    ///
    /// uniform PushConstants _push_constant_binding_cs;
    /// ```
    ///
    /// ```text
    /// - _push_constant_binding_cs.inner.value
    /// - _push_constant_binding_cs.array[0]
    /// - _push_constant_binding_cs.array[1]
    /// ```
    ///
    pub access_path: String,
    /// Type of the uniform. This will only ever be a scalar, vector, or matrix.
    pub ty: Handle<crate::Type>,
    /// The offset in the push constant memory block this uniform maps to.
    ///
    /// The size of the uniform can be derived from the type.
    pub offset: u32,
}

/// Helper structure that generates a number
#[derive(Default)]
struct IdGenerator(u32);

impl IdGenerator {
    /// Generates a number that's guaranteed to be unique for this `IdGenerator`
    fn generate(&mut self) -> u32 {
        // It's just an increasing number but it does the job
        let ret = self.0;
        self.0 += 1;
        ret
    }
}

/// Assorted options needed for generating varyings.
#[derive(Clone, Copy)]
struct VaryingOptions {
    output: bool,
    targeting_webgl: bool,
    draw_parameters: bool,
}

impl VaryingOptions {
    const fn from_writer_options(options: &Options, output: bool) -> Self {
        Self {
            output,
            targeting_webgl: options.version.is_webgl(),
            draw_parameters: options.writer_flags.contains(WriterFlags::DRAW_PARAMETERS),
        }
    }
}

/// Helper wrapper used to get a name for a varying
///
/// Varying have different naming schemes depending on their binding:
/// - Varyings with builtin bindings get the from [`glsl_built_in`].
/// - Varyings with location bindings are named `_S_location_X` where `S` is a
///   prefix identifying which pipeline stage the varying connects, and `X` is
///   the location.
struct VaryingName<'a> {
    binding: &'a crate::Binding,
    stage: ShaderStage,
    options: VaryingOptions,
}
impl fmt::Display for VaryingName<'_> {
    fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result {
        match *self.binding {
            crate::Binding::Location {
                second_blend_source: true,
                ..
            } => {
                write!(f, "_fs2p_location1",)
            }
            crate::Binding::Location { location, .. } => {
                let prefix = match (self.stage, self.options.output) {
                    (ShaderStage::Compute, _) => unreachable!(),
                    // pipeline to vertex
                    (ShaderStage::Vertex, false) => "p2vs",
                    // vertex to fragment
                    (ShaderStage::Vertex, true) | (ShaderStage::Fragment, false) => "vs2fs",
                    // fragment to pipeline
                    (ShaderStage::Fragment, true) => "fs2p",
                };
                write!(f, "_{prefix}_location{location}",)
            }
            crate::Binding::BuiltIn(built_in) => {
                write!(f, "{}", glsl_built_in(built_in, self.options))
            }
        }
    }
}

impl ShaderStage {
    const fn to_str(self) -> &'static str {
        match self {
            ShaderStage::Compute => "cs",
            ShaderStage::Fragment => "fs",
            ShaderStage::Vertex => "vs",
        }
    }
}

/// Shorthand result used internally by the backend
type BackendResult<T = ()> = Result<T, Error>;

/// A GLSL compilation error.
#[derive(Debug, Error)]
pub enum Error {
    /// A error occurred while writing to the output.
    #[error("Format error")]
    FmtError(#[from] FmtError),
    /// The specified [`Version`] doesn't have all required [`Features`].
    ///
    /// Contains the missing [`Features`].
    #[error("The selected version doesn't support {0:?}")]
    MissingFeatures(Features),
    /// [`AddressSpace::PushConstant`](crate::AddressSpace::PushConstant) was used more than
    /// once in the entry point, which isn't supported.
    #[error("Multiple push constants aren't supported")]
    MultiplePushConstants,
    /// The specified [`Version`] isn't supported.
    #[error("The specified version isn't supported")]
    VersionNotSupported,
    /// The entry point couldn't be found.
    #[error("The requested entry point couldn't be found")]
    EntryPointNotFound,
    /// A call was made to an unsupported external.
    #[error("A call was made to an unsupported external: {0}")]
    UnsupportedExternal(String),
    /// A scalar with an unsupported width was requested.
    #[error("A scalar with an unsupported width was requested: {0:?}")]
    UnsupportedScalar(crate::Scalar),
    /// A image was used with multiple samplers, which isn't supported.
    #[error("A image was used with multiple samplers")]
    ImageMultipleSamplers,
    #[error("{0}")]
    Custom(String),
    #[error("overrides should not be present at this stage")]
    Override,
    /// [`crate::Sampling::First`] is unsupported.
    #[error("`{:?}` sampling is unsupported", crate::Sampling::First)]
    FirstSamplingNotSupported,
}

/// Binary operation with a different logic on the GLSL side.
enum BinaryOperation {
    /// Vector comparison should use the function like `greaterThan()`, etc.
    VectorCompare,
    /// Vector component wise operation; used to polyfill unsupported ops like `|` and `&` for `bvecN`'s
    VectorComponentWise,
    /// GLSL `%` is SPIR-V `OpUMod/OpSMod` and `mod()` is `OpFMod`, but [`BinaryOperator::Modulo`](crate::BinaryOperator::Modulo) is `OpFRem`.
    Modulo,
    /// Any plain operation. No additional logic required.
    Other,
}

/// Writer responsible for all code generation.
pub struct Writer<'a, W> {
    // Inputs
    /// The module being written.
    module: &'a crate::Module,
    /// The module analysis.
    info: &'a valid::ModuleInfo,
    /// The output writer.
    out: W,
    /// User defined configuration to be used.
    options: &'a Options,
    /// The bound checking policies to be used
    policies: proc::BoundsCheckPolicies,

    // Internal State
    /// Features manager used to store all the needed features and write them.
    features: FeaturesManager,
    namer: proc::Namer,
    /// A map with all the names needed for writing the module
    /// (generated by a [`Namer`](crate::proc::Namer)).
    names: crate::FastHashMap<NameKey, String>,
    /// A map with the names of global variables needed for reflections.
    reflection_names_globals: crate::FastHashMap<Handle<crate::GlobalVariable>, String>,
    /// The selected entry point.
    entry_point: &'a crate::EntryPoint,
    /// The index of the selected entry point.
    entry_point_idx: proc::EntryPointIndex,
    /// A generator for unique block numbers.
    block_id: IdGenerator,
    /// Set of expressions that have associated temporary variables.
    named_expressions: crate::NamedExpressions,
    /// Set of expressions that need to be baked to avoid unnecessary repetition in output
    need_bake_expressions: back::NeedBakeExpressions,
    /// Information about nesting of loops and switches.
    ///
    /// Used for forwarding continue statements in switches that have been
    /// transformed to `do {} while(false);` loops.
    continue_ctx: back::continue_forward::ContinueCtx,
    /// How many views to render to, if doing multiview rendering.
    multiview: Option<std::num::NonZeroU32>,
    /// Mapping of varying variables to their location. Needed for reflections.
    varying: crate::FastHashMap<String, VaryingLocation>,
}

impl<'a, W: Write> Writer<'a, W> {
    /// Creates a new [`Writer`] instance.
    ///
    /// # Errors
    /// - If the version specified is invalid or supported.
    /// - If the entry point couldn't be found in the module.
    /// - If the version specified doesn't support some used features.
    pub fn new(
        out: W,
        module: &'a crate::Module,
        info: &'a valid::ModuleInfo,
        options: &'a Options,
        pipeline_options: &'a PipelineOptions,
        policies: proc::BoundsCheckPolicies,
    ) -> Result<Self, Error> {
        if !module.overrides.is_empty() {
            return Err(Error::Override);
        }

        // Check if the requested version is supported
        if !options.version.is_supported() {
            log::error!("Version {}", options.version);
            return Err(Error::VersionNotSupported);
        }

        // Try to find the entry point and corresponding index
        let ep_idx = module
            .entry_points
            .iter()
            .position(|ep| {
                pipeline_options.shader_stage == ep.stage && pipeline_options.entry_point == ep.name
            })
            .ok_or(Error::EntryPointNotFound)?;

        // Generate a map with names required to write the module
        let mut names = crate::FastHashMap::default();
        let mut namer = proc::Namer::default();
        namer.reset(
            module,
            keywords::RESERVED_KEYWORDS,
            &[],
            &[],
            &[
                "gl_",                     // all GL built-in variables
                "_group",                  // all normal bindings
                "_push_constant_binding_", // all push constant bindings
            ],
            &mut names,
        );

        // Build the instance
        let mut this = Self {
            module,
            info,
            out,
            options,
            policies,

            namer,
            features: FeaturesManager::new(),
            names,
            reflection_names_globals: crate::FastHashMap::default(),
            entry_point: &module.entry_points[ep_idx],
            entry_point_idx: ep_idx as u16,
            multiview: pipeline_options.multiview,
            block_id: IdGenerator::default(),
            named_expressions: Default::default(),
            need_bake_expressions: Default::default(),
            continue_ctx: back::continue_forward::ContinueCtx::default(),
            varying: Default::default(),
        };

        // Find all features required to print this module
        this.collect_required_features()?;

        Ok(this)
    }

    /// Writes the [`Module`](crate::Module) as glsl to the output
    ///
    /// # Notes
    /// If an error occurs while writing, the output might have been written partially
    ///
    /// # Panics
    /// Might panic if the module is invalid
    pub fn write(&mut self) -> Result<ReflectionInfo, Error> {
        // We use `writeln!(self.out)` throughout the write to add newlines
        // to make the output more readable

        let es = self.options.version.is_es();

        // Write the version (It must be the first thing or it isn't a valid glsl output)
        writeln!(self.out, "#version {}", self.options.version)?;
        // Write all the needed extensions
        //
        // This used to be the last thing being written as it allowed to search for features while
        // writing the module saving some loops but some older versions (420 or less) required the
        // extensions to appear before being used, even though extensions are part of the
        // preprocessor not the processor ¯\_(ツ)_/¯
        self.features.write(self.options, &mut self.out)?;

        // glsl es requires a precision to be specified for floats and ints
        // TODO: Should this be user configurable?
        if es {
            writeln!(self.out)?;
            writeln!(self.out, "precision highp float;")?;
            writeln!(self.out, "precision highp int;")?;
            writeln!(self.out)?;
        }

        if self.entry_point.stage == ShaderStage::Compute {
            let workgroup_size = self.entry_point.workgroup_size;
            writeln!(
                self.out,
                "layout(local_size_x = {}, local_size_y = {}, local_size_z = {}) in;",
                workgroup_size[0], workgroup_size[1], workgroup_size[2]
            )?;
            writeln!(self.out)?;
        }

        if self.entry_point.stage == ShaderStage::Vertex
            && !self
                .options
                .writer_flags
                .contains(WriterFlags::DRAW_PARAMETERS)
            && self.features.contains(Features::INSTANCE_INDEX)
        {
            writeln!(self.out, "uniform uint {FIRST_INSTANCE_BINDING};")?;
            writeln!(self.out)?;
        }

        // Enable early depth tests if needed
        if let Some(depth_test) = self.entry_point.early_depth_test {
            // If early depth test is supported for this version of GLSL
            if self.options.version.supports_early_depth_test() {
                writeln!(self.out, "layout(early_fragment_tests) in;")?;

                if let Some(conservative) = depth_test.conservative {
                    use crate::ConservativeDepth as Cd;

                    let depth = match conservative {
                        Cd::GreaterEqual => "greater",
                        Cd::LessEqual => "less",
                        Cd::Unchanged => "unchanged",
                    };
                    writeln!(self.out, "layout (depth_{depth}) out float gl_FragDepth;")?;
                }
                writeln!(self.out)?;
            } else {
                log::warn!(
                    "Early depth testing is not supported for this version of GLSL: {}",
                    self.options.version
                );
            }
        }

        if self.entry_point.stage == ShaderStage::Vertex && self.options.version.is_webgl() {
            if let Some(multiview) = self.multiview.as_ref() {
                writeln!(self.out, "layout(num_views = {multiview}) in;")?;
                writeln!(self.out)?;
            }
        }

        // Write struct types.
        //
        // This are always ordered because the IR is structured in a way that
        // you can't make a struct without adding all of its members first.
        for (handle, ty) in self.module.types.iter() {
            if let TypeInner::Struct { ref members, .. } = ty.inner {
                // Structures ending with runtime-sized arrays can only be
                // rendered as shader storage blocks in GLSL, not stand-alone
                // struct types.
                if !self.module.types[members.last().unwrap().ty]
                    .inner
                    .is_dynamically_sized(&self.module.types)
                {
                    let name = &self.names[&NameKey::Type(handle)];
                    write!(self.out, "struct {name} ")?;
                    self.write_struct_body(handle, members)?;
                    writeln!(self.out, ";")?;
                }
            }
        }

        // Write functions to create special types.
        for (type_key, struct_ty) in self.module.special_types.predeclared_types.iter() {
            match type_key {
                &crate::PredeclaredType::ModfResult { size, scalar }
                | &crate::PredeclaredType::FrexpResult { size, scalar } => {
                    let arg_type_name_owner;
                    let arg_type_name = if let Some(size) = size {
                        arg_type_name_owner = format!(
                            "{}vec{}",
                            if scalar.width == 8 { "d" } else { "" },
                            size as u8
                        );
                        &arg_type_name_owner
                    } else if scalar.width == 8 {
                        "double"
                    } else {
                        "float"
                    };

                    let other_type_name_owner;
                    let (defined_func_name, called_func_name, other_type_name) =
                        if matches!(type_key, &crate::PredeclaredType::ModfResult { .. }) {
                            (MODF_FUNCTION, "modf", arg_type_name)
                        } else {
                            let other_type_name = if let Some(size) = size {
                                other_type_name_owner = format!("ivec{}", size as u8);
                                &other_type_name_owner
                            } else {
                                "int"
                            };
                            (FREXP_FUNCTION, "frexp", other_type_name)
                        };

                    let struct_name = &self.names[&NameKey::Type(*struct_ty)];

                    writeln!(self.out)?;
                    if !self.options.version.supports_frexp_function()
                        && matches!(type_key, &crate::PredeclaredType::FrexpResult { .. })
                    {
                        writeln!(
                            self.out,
                            "{struct_name} {defined_func_name}({arg_type_name} arg) {{
    {other_type_name} other = arg == {arg_type_name}(0) ? {other_type_name}(0) : {other_type_name}({arg_type_name}(1) + log2(arg));
    {arg_type_name} fract = arg * exp2({arg_type_name}(-other));
    return {struct_name}(fract, other);
}}",
                        )?;
                    } else {
                        writeln!(
                            self.out,
                            "{struct_name} {defined_func_name}({arg_type_name} arg) {{
    {other_type_name} other;
    {arg_type_name} fract = {called_func_name}(arg, other);
    return {struct_name}(fract, other);
}}",
                        )?;
                    }
                }
                &crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}
            }
        }

        // Write all named constants
        let mut constants = self
            .module
            .constants
            .iter()
            .filter(|&(_, c)| c.name.is_some())
            .peekable();
        while let Some((handle, _)) = constants.next() {
            self.write_global_constant(handle)?;
            // Add extra newline for readability on last iteration
            if constants.peek().is_none() {
                writeln!(self.out)?;
            }
        }

        let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);

        // Write the globals
        //
        // Unless explicitly disabled with WriterFlags::INCLUDE_UNUSED_ITEMS,
        // we filter all globals that aren't used by the selected entry point as they might be
        // interfere with each other (i.e. two globals with the same location but different with
        // different classes)
        let include_unused = self
            .options
            .writer_flags
            .contains(WriterFlags::INCLUDE_UNUSED_ITEMS);
        for (handle, global) in self.module.global_variables.iter() {
            let is_unused = ep_info[handle].is_empty();
            if !include_unused && is_unused {
                continue;
            }

            match self.module.types[global.ty].inner {
                // We treat images separately because they might require
                // writing the storage format
                TypeInner::Image {
                    mut dim,
                    arrayed,
                    class,
                } => {
                    // Gather the storage format if needed
                    let storage_format_access = match self.module.types[global.ty].inner {
                        TypeInner::Image {
                            class: crate::ImageClass::Storage { format, access },
                            ..
                        } => Some((format, access)),
                        _ => None,
                    };

                    if dim == crate::ImageDimension::D1 && es {
                        dim = crate::ImageDimension::D2
                    }

                    // Gether the location if needed
                    let layout_binding = if self.options.version.supports_explicit_locations() {
                        let br = global.binding.as_ref().unwrap();
                        self.options.binding_map.get(br).cloned()
                    } else {
                        None
                    };

                    // Write all the layout qualifiers
                    if layout_binding.is_some() || storage_format_access.is_some() {
                        write!(self.out, "layout(")?;
                        if let Some(binding) = layout_binding {
                            write!(self.out, "binding = {binding}")?;
                        }
                        if let Some((format, _)) = storage_format_access {
                            let format_str = glsl_storage_format(format)?;
                            let separator = match layout_binding {
                                Some(_) => ",",
                                None => "",
                            };
                            write!(self.out, "{separator}{format_str}")?;
                        }
                        write!(self.out, ") ")?;
                    }

                    if let Some((_, access)) = storage_format_access {
                        self.write_storage_access(access)?;
                    }

                    // All images in glsl are `uniform`
                    // The trailing space is important
                    write!(self.out, "uniform ")?;

                    // write the type
                    //
                    // This is way we need the leading space because `write_image_type` doesn't add
                    // any spaces at the beginning or end
                    self.write_image_type(dim, arrayed, class)?;

                    // Finally write the name and end the global with a `;`
                    // The leading space is important
                    let global_name = self.get_global_name(handle, global);
                    writeln!(self.out, " {global_name};")?;
                    writeln!(self.out)?;

                    self.reflection_names_globals.insert(handle, global_name);
                }
                // glsl has no concept of samplers so we just ignore it
                TypeInner::Sampler { .. } => continue,
                // All other globals are written by `write_global`
                _ => {
                    self.write_global(handle, global)?;
                    // Add a newline (only for readability)
                    writeln!(self.out)?;
                }
            }
        }

        for arg in self.entry_point.function.arguments.iter() {
            self.write_varying(arg.binding.as_ref(), arg.ty, false)?;
        }
        if let Some(ref result) = self.entry_point.function.result {
            self.write_varying(result.binding.as_ref(), result.ty, true)?;
        }
        writeln!(self.out)?;

        // Write all regular functions
        for (handle, function) in self.module.functions.iter() {
            // Check that the function doesn't use globals that aren't supported
            // by the current entry point
            if !include_unused && !ep_info.dominates_global_use(&self.info[handle]) {
                continue;
            }

            let fun_info = &self.info[handle];

            // 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 !fun_info.available_stages.contains(ep_info.available_stages) {
                continue;
            }

            // Write the function
            self.write_function(back::FunctionType::Function(handle), function, fun_info)?;

            writeln!(self.out)?;
        }

        self.write_function(
            back::FunctionType::EntryPoint(self.entry_point_idx),
            &self.entry_point.function,
            ep_info,
        )?;

        // Add newline at the end of file
        writeln!(self.out)?;

        // Collect all reflection info and return it to the user
        self.collect_reflection_info()
    }

    fn write_array_size(
        &mut self,
        base: Handle<crate::Type>,
        size: crate::ArraySize,
    ) -> BackendResult {
        write!(self.out, "[")?;

        // Write the array size
        // Writes nothing if `ArraySize::Dynamic`
        match size {
            crate::ArraySize::Constant(size) => {
                write!(self.out, "{size}")?;
            }
            crate::ArraySize::Pending(_) => unreachable!(),
            crate::ArraySize::Dynamic => (),
        }

        write!(self.out, "]")?;

        if let TypeInner::Array {
            base: next_base,
            size: next_size,
            ..
        } = self.module.types[base].inner
        {
            self.write_array_size(next_base, next_size)?;
        }

        Ok(())
    }

    /// Helper method used to write value types
    ///
    /// # Notes
    /// Adds no trailing or leading whitespace
    fn write_value_type(&mut self, inner: &TypeInner) -> BackendResult {
        match *inner {
            // Scalars are simple we just get the full name from `glsl_scalar`
            TypeInner::Scalar(scalar)
            | TypeInner::Atomic(scalar)
            | TypeInner::ValuePointer {
                size: None,
                scalar,
                space: _,
            } => write!(self.out, "{}", glsl_scalar(scalar)?.full)?,
            // Vectors are just `gvecN` where `g` is the scalar prefix and `N` is the vector size
            TypeInner::Vector { size, scalar }
            | TypeInner::ValuePointer {
                size: Some(size),
                scalar,
                space: _,
            } => write!(self.out, "{}vec{}", glsl_scalar(scalar)?.prefix, size as u8)?,
            // Matrices are written with `gmatMxN` where `g` is the scalar prefix (only floats and
            // doubles are allowed), `M` is the columns count and `N` is the rows count
            //
            // glsl supports a matrix shorthand `gmatN` where `N` = `M` but it doesn't justify the
            // extra branch to write matrices this way
            TypeInner::Matrix {
                columns,
                rows,
                scalar,
            } => write!(
                self.out,
                "{}mat{}x{}",
                glsl_scalar(scalar)?.prefix,
                columns as u8,
                rows as u8
            )?,
            // GLSL arrays are written as `type name[size]`
            // Here we only write the size of the array i.e. `[size]`
            // Base `type` and `name` should be written outside
            TypeInner::Array { base, size, .. } => self.write_array_size(base, size)?,
            // Write all variants instead of `_` so that if new variants are added a
            // no exhaustiveness error is thrown
            TypeInner::Pointer { .. }
            | TypeInner::Struct { .. }
            | TypeInner::Image { .. }
            | TypeInner::Sampler { .. }
            | TypeInner::AccelerationStructure
            | TypeInner::RayQuery
            | TypeInner::BindingArray { .. } => {
                return Err(Error::Custom(format!("Unable to write type {inner:?}")))
            }
        }

        Ok(())
    }

    /// Helper method used to write non image/sampler types
    ///
    /// # Notes
    /// Adds no trailing or leading whitespace
    fn write_type(&mut self, ty: Handle<crate::Type>) -> BackendResult {
        match self.module.types[ty].inner {
            // glsl has no pointer types so just write types as normal and loads are skipped
            TypeInner::Pointer { base, .. } => self.write_type(base),
            // glsl structs are written as just the struct name
            TypeInner::Struct { .. } => {
                // Get the struct name
                let name = &self.names[&NameKey::Type(ty)];
                write!(self.out, "{name}")?;
                Ok(())
            }
            // glsl array has the size separated from the base type
            TypeInner::Array { base, .. } => self.write_type(base),
            ref other => self.write_value_type(other),
        }
    }

    /// Helper method to write a image type
    ///
    /// # Notes
    /// Adds no leading or trailing whitespace
    fn write_image_type(
        &mut self,
        dim: crate::ImageDimension,
        arrayed: bool,
        class: crate::ImageClass,
    ) -> BackendResult {
        // glsl images consist of four parts the scalar prefix, the image "type", the dimensions
        // and modifiers
        //
        // There exists two image types
        // - sampler - for sampled images
        // - image - for storage images
        //
        // There are three possible modifiers that can be used together and must be written in
        // this order to be valid
        // - MS - used if it's a multisampled image
        // - Array - used if it's an image array
        // - Shadow - used if it's a depth image
        use crate::ImageClass as Ic;
        use crate::Scalar as S;
        let float = S {
            kind: crate::ScalarKind::Float,
            width: 4,
        };
        let (base, scalar, ms, comparison) = match class {
            Ic::Sampled { kind, multi: true } => ("sampler", S { kind, width: 4 }, "MS", ""),
            Ic::Sampled { kind, multi: false } => ("sampler", S { kind, width: 4 }, "", ""),
            Ic::Depth { multi: true } => ("sampler", float, "MS", ""),
            Ic::Depth { multi: false } => ("sampler", float, "", "Shadow"),
            Ic::Storage { format, .. } => ("image", format.into(), "", ""),
        };

        let precision = if self.options.version.is_es() {
            "highp "
        } else {
            ""
        };

        write!(
            self.out,
            "{}{}{}{}{}{}{}",
            precision,
            glsl_scalar(scalar)?.prefix,
            base,
            glsl_dimension(dim),
            ms,
            if arrayed { "Array" } else { "" },
            comparison
        )?;

        Ok(())
    }

    /// Helper method used to write non images/sampler globals
    ///
    /// # Notes
    /// Adds a newline
    ///
    /// # Panics
    /// If the global has type sampler
    fn write_global(
        &mut self,
        handle: Handle<crate::GlobalVariable>,
        global: &crate::GlobalVariable,
    ) -> BackendResult {
        if self.options.version.supports_explicit_locations() {
            if let Some(ref br) = global.binding {
                match self.options.binding_map.get(br) {
                    Some(binding) => {
                        let layout = match global.space {
                            crate::AddressSpace::Storage { .. } => {
                                if self.options.version.supports_std430_layout() {
                                    "std430, "
                                } else {
                                    "std140, "
                                }
                            }
                            crate::AddressSpace::Uniform => "std140, ",
                            _ => "",
                        };
                        write!(self.out, "layout({layout}binding = {binding}) ")?
                    }
                    None => {
                        log::debug!("unassigned binding for {:?}", global.name);
                        if let crate::AddressSpace::Storage { .. } = global.space {
                            if self.options.version.supports_std430_layout() {
                                write!(self.out, "layout(std430) ")?
                            }
                        }
                    }
                }
            }
        }

        if let crate::AddressSpace::Storage { access } = global.space {
            self.write_storage_access(access)?;
        }

        if let Some(storage_qualifier) = glsl_storage_qualifier(global.space) {
            write!(self.out, "{storage_qualifier} ")?;
        }

        match global.space {
            crate::AddressSpace::Private => {
                self.write_simple_global(handle, global)?;
            }
            crate::AddressSpace::WorkGroup => {
                self.write_simple_global(handle, global)?;
            }
            crate::AddressSpace::PushConstant => {
                self.write_simple_global(handle, global)?;
            }
            crate::AddressSpace::Uniform => {
                self.write_interface_block(handle, global)?;
            }
            crate::AddressSpace::Storage { .. } => {
                self.write_interface_block(handle, global)?;
            }
            // A global variable in the `Function` address space is a
            // contradiction in terms.
            crate::AddressSpace::Function => unreachable!(),
            // Textures and samplers are handled directly in `Writer::write`.
            crate::AddressSpace::Handle => unreachable!(),
        }

        Ok(())
    }

    fn write_simple_global(
        &mut self,
        handle: Handle<crate::GlobalVariable>,
        global: &crate::GlobalVariable,
    ) -> BackendResult {
        self.write_type(global.ty)?;
        write!(self.out, " ")?;
        self.write_global_name(handle, global)?;

        if let TypeInner::Array { base, size, .. } = self.module.types[global.ty].inner {
            self.write_array_size(base, size)?;
        }

        if global.space.initializable() && is_value_init_supported(self.module, global.ty) {
            write!(self.out, " = ")?;
            if let Some(init) = global.init {
                self.write_const_expr(init)?;
            } else {
                self.write_zero_init_value(global.ty)?;
            }
        }

        writeln!(self.out, ";")?;

        if let crate::AddressSpace::PushConstant = global.space {
            let global_name = self.get_global_name(handle, global);
            self.reflection_names_globals.insert(handle, global_name);
        }

        Ok(())
    }

    /// Write an interface block for a single Naga global.
    ///
    /// Write `block_name { members }`. Since `block_name` must be unique
    /// between blocks and structs, we add `_block_ID` where `ID` is a
    /// `IdGenerator` generated number. Write `members` in the same way we write
    /// a struct's members.
    fn write_interface_block(
        &mut self,
        handle: Handle<crate::GlobalVariable>,
        global: &crate::GlobalVariable,
    ) -> BackendResult {
        // Write the block name, it's just the struct name appended with `_block_ID`
        let ty_name = &self.names[&NameKey::Type(global.ty)];
        let block_name = format!(
            "{}_block_{}{:?}",
            // avoid double underscores as they are reserved in GLSL
            ty_name.trim_end_matches('_'),
            self.block_id.generate(),
            self.entry_point.stage,
        );
        write!(self.out, "{block_name} ")?;
        self.reflection_names_globals.insert(handle, block_name);

        match self.module.types[global.ty].inner {
            TypeInner::Struct { ref members, .. }
                if self.module.types[members.last().unwrap().ty]
                    .inner
                    .is_dynamically_sized(&self.module.types) =>
            {
                // Structs with dynamically sized arrays must have their
                // members lifted up as members of the interface block. GLSL
                // can't write such struct types anyway.
                self.write_struct_body(global.ty, members)?;
                write!(self.out, " ")?;
                self.write_global_name(handle, global)?;
            }
            _ => {
                // A global of any other type is written as the sole member
                // of the interface block. Since the interface block is
                // anonymous, this becomes visible in the global scope.
                write!(self.out, "{{ ")?;
                self.write_type(global.ty)?;
                write!(self.out, " ")?;
                self.write_global_name(handle, global)?;
                if let TypeInner::Array { base, size, .. } = self.module.types[global.ty].inner {
                    self.write_array_size(base, size)?;
                }
                write!(self.out, "; }}")?;
            }
        }

        writeln!(self.out, ";")?;

        Ok(())
    }

    /// Helper method used to find which expressions of a given function require baking
    ///
    /// # Notes
    /// Clears `need_bake_expressions` set before adding to it
    fn update_expressions_to_bake(&mut self, func: &crate::Function, info: &valid::FunctionInfo) {
        use crate::Expression;
        self.need_bake_expressions.clear();
        for (fun_handle, expr) in func.expressions.iter() {
            let expr_info = &info[fun_handle];
            let min_ref_count = func.expressions[fun_handle].bake_ref_count();
            if min_ref_count <= expr_info.ref_count {
                self.need_bake_expressions.insert(fun_handle);
            }

            let inner = expr_info.ty.inner_with(&self.module.types);

            if let Expression::Math {
                fun,
                arg,
                arg1,
                arg2,
                ..
            } = *expr
            {
                match fun {
                    crate::MathFunction::Dot => {
                        // if the expression is a Dot product with integer arguments,
                        // then the args needs baking as well
                        if let TypeInner::Scalar(crate::Scalar {
                            kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint,
                            ..
                        }) = *inner
                        {
                            self.need_bake_expressions.insert(arg);
                            self.need_bake_expressions.insert(arg1.unwrap());
                        }
                    }
                    crate::MathFunction::Pack4xI8
                    | crate::MathFunction::Pack4xU8
                    | crate::MathFunction::Unpack4xI8
                    | crate::MathFunction::Unpack4xU8
                    | crate::MathFunction::QuantizeToF16 => {
                        self.need_bake_expressions.insert(arg);
                    }
                    crate::MathFunction::ExtractBits => {
                        // Only argument 1 is re-used.
                        self.need_bake_expressions.insert(arg1.unwrap());
                    }
                    crate::MathFunction::InsertBits => {
                        // Only argument 2 is re-used.
                        self.need_bake_expressions.insert(arg2.unwrap());
                    }
                    crate::MathFunction::CountLeadingZeros => {
                        if let Some(crate::ScalarKind::Sint) = inner.scalar_kind() {
                            self.need_bake_expressions.insert(arg);
                        }
                    }
                    _ => {}
                }
            }
        }
    }

    /// Helper method used to get a name for a global
    ///
    /// Globals have different naming schemes depending on their binding:
    /// - Globals without bindings use the name from the [`Namer`](crate::proc::Namer)
    /// - Globals with resource binding are named `_group_X_binding_Y` where `X`
    ///   is the group and `Y` is the binding
    fn get_global_name(
        &self,
        handle: Handle<crate::GlobalVariable>,
        global: &crate::GlobalVariable,
    ) -> String {
        match (&global.binding, global.space) {
            (&Some(ref br), _) => {
                format!(
                    "_group_{}_binding_{}_{}",
                    br.group,
                    br.binding,
                    self.entry_point.stage.to_str()
                )
            }
            (&None, crate::AddressSpace::PushConstant) => {
                format!("_push_constant_binding_{}", self.entry_point.stage.to_str())
            }
            (&None, _) => self.names[&NameKey::GlobalVariable(handle)].clone(),
        }
    }

    /// Helper method used to write a name for a global without additional heap allocation
    fn write_global_name(
        &mut self,
        handle: Handle<crate::GlobalVariable>,
        global: &crate::GlobalVariable,
    ) -> BackendResult {
        match (&global.binding, global.space) {
            (&Some(ref br), _) => write!(
                self.out,
                "_group_{}_binding_{}_{}",
                br.group,
                br.binding,
                self.entry_point.stage.to_str()
            )?,
            (&None, crate::AddressSpace::PushConstant) => write!(
                self.out,
                "_push_constant_binding_{}",
                self.entry_point.stage.to_str()
            )?,
            (&None, _) => write!(
                self.out,
                "{}",
                &self.names[&NameKey::GlobalVariable(handle)]
            )?,
        }

        Ok(())
    }

    /// Write a GLSL global that will carry a Naga entry point's argument or return value.
    ///
    /// A Naga entry point's arguments and return value are rendered in GLSL as
    /// variables at global scope with the `in` and `out` storage qualifiers.
    /// The code we generate for `main` loads from all the `in` globals into
    /// appropriately named locals. Before it returns, `main` assigns the
    /// components of its return value into all the `out` globals.
    ///
    /// This function writes a declaration for one such GLSL global,
    /// representing a value passed into or returned from [`self.entry_point`]
    /// that has a [`Location`] binding. The global's name is generated based on
    /// the location index and the shader stages being connected; see
    /// [`VaryingName`]. This means we don't need to know the names of
    /// arguments, just their types and bindings.
    ///
    /// Emit nothing for entry point arguments or return values with [`BuiltIn`]
    /// bindings; `main` will read from or assign to the appropriate GLSL
    /// special variable; these are pre-declared. As an exception, we do declare
    /// `gl_Position` or `gl_FragCoord` with the `invariant` qualifier if
    /// needed.
    ///
    /// Use `output` together with [`self.entry_point.stage`] to determine which
    /// shader stages are being connected, and choose the `in` or `out` storage
    /// qualifier.
    ///
    /// [`self.entry_point`]: Writer::entry_point
    /// [`self.entry_point.stage`]: crate::EntryPoint::stage
    /// [`Location`]: crate::Binding::Location
    /// [`BuiltIn`]: crate::Binding::BuiltIn
    fn write_varying(
        &mut self,
        binding: Option<&crate::Binding>,
        ty: Handle<crate::Type>,
        output: bool,
    ) -> Result<(), Error> {
        // For a struct, emit a separate global for each member with a binding.
        if let TypeInner::Struct { ref members, .. } = self.module.types[ty].inner {
            for member in members {
                self.write_varying(member.binding.as_ref(), member.ty, output)?;
            }
            return Ok(());
        }

        let binding = match binding {
            None => return Ok(()),
            Some(binding) => binding,
        };

        let (location, interpolation, sampling, second_blend_source) = match *binding {
            crate::Binding::Location {
                location,
                interpolation,
                sampling,
                second_blend_source,
            } => (location, interpolation, sampling, second_blend_source),
            crate::Binding::BuiltIn(built_in) => {
                if let crate::BuiltIn::Position { invariant: true } = built_in {
                    match (self.options.version, self.entry_point.stage) {
                        (
                            Version::Embedded {
                                version: 300,
                                is_webgl: true,
                            },
                            ShaderStage::Fragment,
                        ) => {
                            // `invariant gl_FragCoord` is not allowed in WebGL2 and possibly
                            // OpenGL ES in general (waiting on confirmation).
                            //
                            // See https://github.com/KhronosGroup/WebGL/issues/3518
                        }
                        _ => {
                            writeln!(
                                self.out,
                                "invariant {};",
                                glsl_built_in(
                                    built_in,
                                    VaryingOptions::from_writer_options(self.options, output)
                                )
                            )?;
                        }
                    }
                }
                return Ok(());
            }
        };

        // Write the interpolation modifier if needed
        //
        // We ignore all interpolation and auxiliary modifiers that aren't used in fragment
        // shaders' input globals or vertex shaders' output globals.
        let emit_interpolation_and_auxiliary = match self.entry_point.stage {
            ShaderStage::Vertex => output,
            ShaderStage::Fragment => !output,
            ShaderStage::Compute => false,
        };

        // Write the I/O locations, if allowed
        let io_location = if self.options.version.supports_explicit_locations()
            || !emit_interpolation_and_auxiliary
        {
            if self.options.version.supports_io_locations() {
                if second_blend_source {
                    write!(self.out, "layout(location = {location}, index = 1) ")?;
                } else {
                    write!(self.out, "layout(location = {location}) ")?;
                }
                None
            } else {
                Some(VaryingLocation {
                    location,
                    index: second_blend_source as u32,
                })
            }
        } else {
            None
        };

        // Write the interpolation qualifier.
        if let Some(interp) = interpolation {
            if emit_interpolation_and_auxiliary {
                write!(self.out, "{} ", glsl_interpolation(interp))?;
            }
        }

        // Write the sampling auxiliary qualifier.
        //
        // Before GLSL 4.2, the `centroid` and `sample` qualifiers were required to appear
        // immediately before the `in` / `out` qualifier, so we'll just follow that rule
        // here, regardless of the version.
        if let Some(sampling) = sampling {
            if emit_interpolation_and_auxiliary {
                if let Some(qualifier) = glsl_sampling(sampling)? {
                    write!(self.out, "{qualifier} ")?;
                }
            }
        }

        // Write the input/output qualifier.
        write!(self.out, "{} ", if output { "out" } else { "in" })?;

        // Write the type
        // `write_type` adds no leading or trailing spaces
        self.write_type(ty)?;

        // Finally write the global name and end the global with a `;` and a newline
        // Leading space is important
        let vname = VaryingName {
            binding: &crate::Binding::Location {
                location,
                interpolation: None,
                sampling: None,
                second_blend_source,
            },
            stage: self.entry_point.stage,
            options: VaryingOptions::from_writer_options(self.options, output),
        };
        writeln!(self.out, " {vname};")?;

        if let Some(location) = io_location {
            self.varying.insert(vname.to_string(), location);
        }

        Ok(())
    }

    /// Helper method used to write functions (both entry points and regular functions)
    ///
    /// # Notes
    /// Adds a newline
    fn write_function(
        &mut self,
        ty: back::FunctionType,
        func: &crate::Function,
        info: &valid::FunctionInfo,
    ) -> BackendResult {
        // Create a function context for the function being written
        let ctx = back::FunctionCtx {
            ty,
            info,
            expressions: &func.expressions,
            named_expressions: &func.named_expressions,
            expr_kind_tracker: ExpressionKindTracker::from_arena(&func.expressions),
        };

        self.named_expressions.clear();
        self.update_expressions_to_bake(func, info);

        // Write the function header
        //
        // glsl headers are the same as in c:
        // `ret_type name(args)`
        // `ret_type` is the return type
        // `name` is the function name
        // `args` is a comma separated list of `type name`
        //  | - `type` is the argument type
        //  | - `name` is the argument name

        // Start by writing the return type if any otherwise write void
        // This is the only place where `void` is a valid type
        // (though it's more a keyword than a type)
        if let back::FunctionType::EntryPoint(_) = ctx.ty {
            write!(self.out, "void")?;
        } else if let Some(ref result) = func.result {
            self.write_type(result.ty)?;
            if let TypeInner::Array { base, size, .. } = self.module.types[result.ty].inner {
                self.write_array_size(base, size)?
            }
        } else {
            write!(self.out, "void")?;
        }

        // Write the function name and open parentheses for the argument list
        let function_name = match ctx.ty {
            back::FunctionType::Function(handle) => &self.names[&NameKey::Function(handle)],
            back::FunctionType::EntryPoint(_) => "main",
        };
        write!(self.out, " {function_name}(")?;

        // Write the comma separated argument list
        //
        // We need access to `Self` here so we use the reference passed to the closure as an
        // argument instead of capturing as that would cause a borrow checker error
        let arguments = match ctx.ty {
            back::FunctionType::EntryPoint(_) => &[][..],
            back::FunctionType::Function(_) => &func.arguments,
        };
        let arguments: Vec<_> = arguments
            .iter()
            .enumerate()
            .filter(|&(_, arg)| match self.module.types[arg.ty].inner {
                TypeInner::Sampler { .. } => false,
                _ => true,
            })
            .collect();
        self.write_slice(&arguments, |this, _, &(i, arg)| {
            // Write the argument type
            match this.module.types[arg.ty].inner {
                // We treat images separately because they might require
                // writing the storage format
                TypeInner::Image {
                    dim,
                    arrayed,
                    class,
                } => {
                    // Write the storage format if needed
                    if let TypeInner::Image {
                        class: crate::ImageClass::Storage { format, .. },
                        ..
                    } = this.module.types[arg.ty].inner
                    {
                        write!(this.out, "layout({}) ", glsl_storage_format(format)?)?;
                    }

                    // write the type
                    //
                    // This is way we need the leading space because `write_image_type` doesn't add
                    // any spaces at the beginning or end
                    this.write_image_type(dim, arrayed, class)?;
                }
                TypeInner::Pointer { base, .. } => {
                    // write parameter qualifiers
                    write!(this.out, "inout ")?;
                    this.write_type(base)?;
                }
                // All other types are written by `write_type`
                _ => {
--> --------------------

--> maximum size reached

--> --------------------