naga/back/glsl/
mod.rs

1/*!
2Backend for [GLSL][glsl] (OpenGL Shading Language).
3
4The main structure is [`Writer`], it maintains internal state that is used
5to output a [`Module`](crate::Module) into glsl
6
7# Supported versions
8### Core
9- 330
10- 400
11- 410
12- 420
13- 430
14- 450
15
16### ES
17- 300
18- 310
19
20[glsl]: https://www.khronos.org/registry/OpenGL/index_gl.php
21*/
22
23// GLSL is mostly a superset of C but it also removes some parts of it this is a list of relevant
24// aspects for this backend.
25//
26// The most notable change is the introduction of the version preprocessor directive that must
27// always be the first line of a glsl file and is written as
28// `#version number profile`
29// `number` is the version itself (i.e. 300) and `profile` is the
30// shader profile we only support "core" and "es", the former is used in desktop applications and
31// the later is used in embedded contexts, mobile devices and browsers. Each one as it's own
32// versions (at the time of writing this the latest version for "core" is 460 and for "es" is 320)
33//
34// Other important preprocessor addition is the extension directive which is written as
35// `#extension name: behaviour`
36// Extensions provide increased features in a plugin fashion but they aren't required to be
37// supported hence why they are called extensions, that's why `behaviour` is used it specifies
38// whether the extension is strictly required or if it should only be enabled if needed. In our case
39// when we use extensions we set behaviour to `require` always.
40//
41// The only thing that glsl removes that makes a difference are pointers.
42//
43// Additions that are relevant for the backend are the discard keyword, the introduction of
44// vector, matrices, samplers, image types and functions that provide common shader operations
45
46pub use features::Features;
47
48use crate::{
49    back,
50    proc::{self, NameKey},
51    valid, Handle, ShaderStage, TypeInner,
52};
53use features::FeaturesManager;
54use std::{
55    cmp::Ordering,
56    fmt,
57    fmt::{Error as FmtError, Write},
58};
59use thiserror::Error;
60
61/// Contains the features related code and the features querying method
62mod features;
63/// Contains a constant with a slice of all the reserved keywords RESERVED_KEYWORDS
64mod keywords;
65
66/// List of supported `core` GLSL versions.
67pub const SUPPORTED_CORE_VERSIONS: &[u16] = &[330, 400, 410, 420, 430, 440, 450];
68/// List of supported `es` GLSL versions.
69pub const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310, 320];
70
71/// The suffix of the variable that will hold the calculated clamped level
72/// of detail for bounds checking in `ImageLoad`
73const CLAMPED_LOD_SUFFIX: &str = "_clamped_lod";
74
75/// Mapping between resources and bindings.
76pub type BindingMap = std::collections::BTreeMap<crate::ResourceBinding, u8>;
77
78impl crate::AtomicFunction {
79    const fn to_glsl(self) -> &'static str {
80        match self {
81            Self::Add | Self::Subtract => "Add",
82            Self::And => "And",
83            Self::InclusiveOr => "Or",
84            Self::ExclusiveOr => "Xor",
85            Self::Min => "Min",
86            Self::Max => "Max",
87            Self::Exchange { compare: None } => "Exchange",
88            Self::Exchange { compare: Some(_) } => "", //TODO
89        }
90    }
91}
92
93impl crate::AddressSpace {
94    const fn is_buffer(&self) -> bool {
95        match *self {
96            crate::AddressSpace::Uniform | crate::AddressSpace::Storage { .. } => true,
97            _ => false,
98        }
99    }
100
101    /// Whether a variable with this address space can be initialized
102    const fn initializable(&self) -> bool {
103        match *self {
104            crate::AddressSpace::Function | crate::AddressSpace::Private => true,
105            crate::AddressSpace::WorkGroup
106            | crate::AddressSpace::Uniform
107            | crate::AddressSpace::Storage { .. }
108            | crate::AddressSpace::Handle
109            | crate::AddressSpace::PushConstant => false,
110        }
111    }
112}
113
114/// A GLSL version.
115#[derive(Debug, Copy, Clone, PartialEq)]
116#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
117#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
118pub enum Version {
119    /// `core` GLSL.
120    Desktop(u16),
121    /// `es` GLSL.
122    Embedded { version: u16, is_webgl: bool },
123}
124
125impl Version {
126    /// Create a new gles version
127    pub const fn new_gles(version: u16) -> Self {
128        Self::Embedded {
129            version,
130            is_webgl: false,
131        }
132    }
133
134    /// Returns true if self is `Version::Embedded` (i.e. is a es version)
135    const fn is_es(&self) -> bool {
136        match *self {
137            Version::Desktop(_) => false,
138            Version::Embedded { .. } => true,
139        }
140    }
141
142    /// Returns true if targetting WebGL
143    const fn is_webgl(&self) -> bool {
144        match *self {
145            Version::Desktop(_) => false,
146            Version::Embedded { is_webgl, .. } => is_webgl,
147        }
148    }
149
150    /// Checks the list of currently supported versions and returns true if it contains the
151    /// specified version
152    ///
153    /// # Notes
154    /// As an invalid version number will never be added to the supported version list
155    /// so this also checks for version validity
156    fn is_supported(&self) -> bool {
157        match *self {
158            Version::Desktop(v) => SUPPORTED_CORE_VERSIONS.contains(&v),
159            Version::Embedded { version: v, .. } => SUPPORTED_ES_VERSIONS.contains(&v),
160        }
161    }
162
163    /// Checks if the version supports all of the explicit layouts:
164    /// - `location=` qualifiers for bindings
165    /// - `binding=` qualifiers for resources
166    ///
167    /// Note: `location=` for vertex inputs and fragment outputs is supported
168    /// unconditionally for GLES 300.
169    fn supports_explicit_locations(&self) -> bool {
170        *self >= Version::Desktop(410) || *self >= Version::new_gles(310)
171    }
172
173    fn supports_early_depth_test(&self) -> bool {
174        *self >= Version::Desktop(130) || *self >= Version::new_gles(310)
175    }
176
177    fn supports_std430_layout(&self) -> bool {
178        *self >= Version::Desktop(430) || *self >= Version::new_gles(310)
179    }
180
181    fn supports_fma_function(&self) -> bool {
182        *self >= Version::Desktop(400) || *self >= Version::new_gles(320)
183    }
184
185    fn supports_integer_functions(&self) -> bool {
186        *self >= Version::Desktop(400) || *self >= Version::new_gles(310)
187    }
188
189    fn supports_derivative_control(&self) -> bool {
190        *self >= Version::Desktop(450)
191    }
192}
193
194impl PartialOrd for Version {
195    fn partial_cmp(&self, other: &Self) -> Option<Ordering> {
196        match (*self, *other) {
197            (Version::Desktop(x), Version::Desktop(y)) => Some(x.cmp(&y)),
198            (Version::Embedded { version: x, .. }, Version::Embedded { version: y, .. }) => {
199                Some(x.cmp(&y))
200            }
201            _ => None,
202        }
203    }
204}
205
206impl fmt::Display for Version {
207    fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result {
208        match *self {
209            Version::Desktop(v) => write!(f, "{v} core"),
210            Version::Embedded { version: v, .. } => write!(f, "{v} es"),
211        }
212    }
213}
214
215bitflags::bitflags! {
216    /// Configuration flags for the [`Writer`].
217    #[cfg_attr(feature = "serialize", derive(serde::Serialize))]
218    #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
219    #[derive(Clone, Copy, Debug, Eq, PartialEq)]
220    pub struct WriterFlags: u32 {
221        /// Flip output Y and extend Z from (0, 1) to (-1, 1).
222        const ADJUST_COORDINATE_SPACE = 0x1;
223        /// Supports GL_EXT_texture_shadow_lod on the host, which provides
224        /// additional functions on shadows and arrays of shadows.
225        const TEXTURE_SHADOW_LOD = 0x2;
226        /// Include unused global variables, constants and functions. By default the output will exclude
227        /// global variables that are not used in the specified entrypoint (including indirect use),
228        /// all constant declarations, and functions that use excluded global variables.
229        const INCLUDE_UNUSED_ITEMS = 0x4;
230        /// Emit `PointSize` output builtin to vertex shaders, which is
231        /// required for drawing with `PointList` topology.
232        ///
233        /// https://registry.khronos.org/OpenGL/specs/es/3.2/GLSL_ES_Specification_3.20.html#built-in-language-variables
234        /// The variable gl_PointSize is intended for a shader to write the size of the point to be rasterized. It is measured in pixels.
235        /// If gl_PointSize is not written to, its value is undefined in subsequent pipe stages.
236        const FORCE_POINT_SIZE = 0x10;
237    }
238}
239
240/// Configuration used in the [`Writer`].
241#[derive(Debug, Clone)]
242#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
243#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
244pub struct Options {
245    /// The GLSL version to be used.
246    pub version: Version,
247    /// Configuration flags for the [`Writer`].
248    pub writer_flags: WriterFlags,
249    /// Map of resources association to binding locations.
250    pub binding_map: BindingMap,
251    /// Should workgroup variables be zero initialized (by polyfilling)?
252    pub zero_initialize_workgroup_memory: bool,
253}
254
255impl Default for Options {
256    fn default() -> Self {
257        Options {
258            version: Version::new_gles(310),
259            writer_flags: WriterFlags::ADJUST_COORDINATE_SPACE,
260            binding_map: BindingMap::default(),
261            zero_initialize_workgroup_memory: true,
262        }
263    }
264}
265
266/// A subset of options meant to be changed per pipeline.
267#[derive(Debug, Clone, PartialEq, Eq, Hash)]
268#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
269#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
270pub struct PipelineOptions {
271    /// The stage of the entry point.
272    pub shader_stage: ShaderStage,
273    /// The name of the entry point.
274    ///
275    /// If no entry point that matches is found while creating a [`Writer`], a error will be thrown.
276    pub entry_point: String,
277    /// How many views to render to, if doing multiview rendering.
278    pub multiview: Option<std::num::NonZeroU32>,
279}
280
281/// Reflection info for texture mappings and uniforms.
282pub struct ReflectionInfo {
283    /// Mapping between texture names and variables/samplers.
284    pub texture_mapping: crate::FastHashMap<String, TextureMapping>,
285    /// Mapping between uniform variables and names.
286    pub uniforms: crate::FastHashMap<Handle<crate::GlobalVariable>, String>,
287}
288
289/// Mapping between a texture and its sampler, if it exists.
290///
291/// GLSL pre-Vulkan has no concept of separate textures and samplers. Instead, everything is a
292/// `gsamplerN` where `g` is the scalar type and `N` is the dimension. But naga uses separate textures
293/// and samplers in the IR, so the backend produces a [`FastHashMap`](crate::FastHashMap) with the texture name
294/// as a key and a [`TextureMapping`] as a value. This way, the user knows where to bind.
295///
296/// [`Storage`](crate::ImageClass::Storage) images produce `gimageN` and don't have an associated sampler,
297/// so the [`sampler`](Self::sampler) field will be [`None`].
298#[derive(Debug, Clone)]
299pub struct TextureMapping {
300    /// Handle to the image global variable.
301    pub texture: Handle<crate::GlobalVariable>,
302    /// Handle to the associated sampler global variable, if it exists.
303    pub sampler: Option<Handle<crate::GlobalVariable>>,
304}
305
306/// Helper structure that generates a number
307#[derive(Default)]
308struct IdGenerator(u32);
309
310impl IdGenerator {
311    /// Generates a number that's guaranteed to be unique for this `IdGenerator`
312    fn generate(&mut self) -> u32 {
313        // It's just an increasing number but it does the job
314        let ret = self.0;
315        self.0 += 1;
316        ret
317    }
318}
319
320/// Helper wrapper used to get a name for a varying
321///
322/// Varying have different naming schemes depending on their binding:
323/// - Varyings with builtin bindings get the from [`glsl_built_in`](glsl_built_in).
324/// - Varyings with location bindings are named `_S_location_X` where `S` is a
325///   prefix identifying which pipeline stage the varying connects, and `X` is
326///   the location.
327struct VaryingName<'a> {
328    binding: &'a crate::Binding,
329    stage: ShaderStage,
330    output: bool,
331    targetting_webgl: bool,
332}
333impl fmt::Display for VaryingName<'_> {
334    fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result {
335        match *self.binding {
336            crate::Binding::Location { location, .. } => {
337                let prefix = match (self.stage, self.output) {
338                    (ShaderStage::Compute, _) => unreachable!(),
339                    // pipeline to vertex
340                    (ShaderStage::Vertex, false) => "p2vs",
341                    // vertex to fragment
342                    (ShaderStage::Vertex, true) | (ShaderStage::Fragment, false) => "vs2fs",
343                    // fragment to pipeline
344                    (ShaderStage::Fragment, true) => "fs2p",
345                };
346                write!(f, "_{prefix}_location{location}",)
347            }
348            crate::Binding::BuiltIn(built_in) => {
349                write!(
350                    f,
351                    "{}",
352                    glsl_built_in(built_in, self.output, self.targetting_webgl)
353                )
354            }
355        }
356    }
357}
358
359impl ShaderStage {
360    const fn to_str(self) -> &'static str {
361        match self {
362            ShaderStage::Compute => "cs",
363            ShaderStage::Fragment => "fs",
364            ShaderStage::Vertex => "vs",
365        }
366    }
367}
368
369/// Shorthand result used internally by the backend
370type BackendResult<T = ()> = Result<T, Error>;
371
372/// A GLSL compilation error.
373#[derive(Debug, Error)]
374pub enum Error {
375    /// A error occurred while writing to the output.
376    #[error("Format error")]
377    FmtError(#[from] FmtError),
378    /// The specified [`Version`] doesn't have all required [`Features`].
379    ///
380    /// Contains the missing [`Features`].
381    #[error("The selected version doesn't support {0:?}")]
382    MissingFeatures(Features),
383    /// [`AddressSpace::PushConstant`](crate::AddressSpace::PushConstant) was used more than
384    /// once in the entry point, which isn't supported.
385    #[error("Multiple push constants aren't supported")]
386    MultiplePushConstants,
387    /// The specified [`Version`] isn't supported.
388    #[error("The specified version isn't supported")]
389    VersionNotSupported,
390    /// The entry point couldn't be found.
391    #[error("The requested entry point couldn't be found")]
392    EntryPointNotFound,
393    /// A call was made to an unsupported external.
394    #[error("A call was made to an unsupported external: {0}")]
395    UnsupportedExternal(String),
396    /// A scalar with an unsupported width was requested.
397    #[error("A scalar with an unsupported width was requested: {0:?} {1:?}")]
398    UnsupportedScalar(crate::ScalarKind, crate::Bytes),
399    /// A image was used with multiple samplers, which isn't supported.
400    #[error("A image was used with multiple samplers")]
401    ImageMultipleSamplers,
402    #[error("{0}")]
403    Custom(String),
404}
405
406/// Binary operation with a different logic on the GLSL side.
407enum BinaryOperation {
408    /// Vector comparison should use the function like `greaterThan()`, etc.
409    VectorCompare,
410    /// Vector component wise operation; used to polyfill unsupported ops like `|` and `&` for `bvecN`'s
411    VectorComponentWise,
412    /// GLSL `%` is SPIR-V `OpUMod/OpSMod` and `mod()` is `OpFMod`, but [`BinaryOperator::Modulo`](crate::BinaryOperator::Modulo) is `OpFRem`.
413    Modulo,
414    /// Any plain operation. No additional logic required.
415    Other,
416}
417
418/// Writer responsible for all code generation.
419pub struct Writer<'a, W> {
420    // Inputs
421    /// The module being written.
422    module: &'a crate::Module,
423    /// The module analysis.
424    info: &'a valid::ModuleInfo,
425    /// The output writer.
426    out: W,
427    /// User defined configuration to be used.
428    options: &'a Options,
429    /// The bound checking policies to be used
430    policies: proc::BoundsCheckPolicies,
431
432    // Internal State
433    /// Features manager used to store all the needed features and write them.
434    features: FeaturesManager,
435    namer: proc::Namer,
436    /// A map with all the names needed for writing the module
437    /// (generated by a [`Namer`](crate::proc::Namer)).
438    names: crate::FastHashMap<NameKey, String>,
439    /// A map with the names of global variables needed for reflections.
440    reflection_names_globals: crate::FastHashMap<Handle<crate::GlobalVariable>, String>,
441    /// The selected entry point.
442    entry_point: &'a crate::EntryPoint,
443    /// The index of the selected entry point.
444    entry_point_idx: proc::EntryPointIndex,
445    /// A generator for unique block numbers.
446    block_id: IdGenerator,
447    /// Set of expressions that have associated temporary variables.
448    named_expressions: crate::NamedExpressions,
449    /// Set of expressions that need to be baked to avoid unnecessary repetition in output
450    need_bake_expressions: back::NeedBakeExpressions,
451    /// How many views to render to, if doing multiview rendering.
452    multiview: Option<std::num::NonZeroU32>,
453}
454
455impl<'a, W: Write> Writer<'a, W> {
456    /// Creates a new [`Writer`] instance.
457    ///
458    /// # Errors
459    /// - If the version specified is invalid or supported.
460    /// - If the entry point couldn't be found in the module.
461    /// - If the version specified doesn't support some used features.
462    pub fn new(
463        out: W,
464        module: &'a crate::Module,
465        info: &'a valid::ModuleInfo,
466        options: &'a Options,
467        pipeline_options: &'a PipelineOptions,
468        policies: proc::BoundsCheckPolicies,
469    ) -> Result<Self, Error> {
470        // Check if the requested version is supported
471        if !options.version.is_supported() {
472            log::error!("Version {}", options.version);
473            return Err(Error::VersionNotSupported);
474        }
475
476        // Try to find the entry point and corresponding index
477        let ep_idx = module
478            .entry_points
479            .iter()
480            .position(|ep| {
481                pipeline_options.shader_stage == ep.stage && pipeline_options.entry_point == ep.name
482            })
483            .ok_or(Error::EntryPointNotFound)?;
484
485        // Generate a map with names required to write the module
486        let mut names = crate::FastHashMap::default();
487        let mut namer = proc::Namer::default();
488        namer.reset(
489            module,
490            keywords::RESERVED_KEYWORDS,
491            &[],
492            &[],
493            &["gl_"],
494            &mut names,
495        );
496
497        // Build the instance
498        let mut this = Self {
499            module,
500            info,
501            out,
502            options,
503            policies,
504
505            namer,
506            features: FeaturesManager::new(),
507            names,
508            reflection_names_globals: crate::FastHashMap::default(),
509            entry_point: &module.entry_points[ep_idx],
510            entry_point_idx: ep_idx as u16,
511            multiview: pipeline_options.multiview,
512            block_id: IdGenerator::default(),
513            named_expressions: Default::default(),
514            need_bake_expressions: Default::default(),
515        };
516
517        // Find all features required to print this module
518        this.collect_required_features()?;
519
520        Ok(this)
521    }
522
523    /// Writes the [`Module`](crate::Module) as glsl to the output
524    ///
525    /// # Notes
526    /// If an error occurs while writing, the output might have been written partially
527    ///
528    /// # Panics
529    /// Might panic if the module is invalid
530    pub fn write(&mut self) -> Result<ReflectionInfo, Error> {
531        // We use `writeln!(self.out)` throughout the write to add newlines
532        // to make the output more readable
533
534        let es = self.options.version.is_es();
535
536        // Write the version (It must be the first thing or it isn't a valid glsl output)
537        writeln!(self.out, "#version {}", self.options.version)?;
538        // Write all the needed extensions
539        //
540        // This used to be the last thing being written as it allowed to search for features while
541        // writing the module saving some loops but some older versions (420 or less) required the
542        // extensions to appear before being used, even though extensions are part of the
543        // preprocessor not the processor ¯\_(ツ)_/¯
544        self.features.write(self.options.version, &mut self.out)?;
545
546        // Write the additional extensions
547        if self
548            .options
549            .writer_flags
550            .contains(WriterFlags::TEXTURE_SHADOW_LOD)
551        {
552            // https://www.khronos.org/registry/OpenGL/extensions/EXT/EXT_texture_shadow_lod.txt
553            writeln!(self.out, "#extension GL_EXT_texture_shadow_lod : require")?;
554        }
555
556        // glsl es requires a precision to be specified for floats and ints
557        // TODO: Should this be user configurable?
558        if es {
559            writeln!(self.out)?;
560            writeln!(self.out, "precision highp float;")?;
561            writeln!(self.out, "precision highp int;")?;
562            writeln!(self.out)?;
563        }
564
565        if self.entry_point.stage == ShaderStage::Compute {
566            let workgroup_size = self.entry_point.workgroup_size;
567            writeln!(
568                self.out,
569                "layout(local_size_x = {}, local_size_y = {}, local_size_z = {}) in;",
570                workgroup_size[0], workgroup_size[1], workgroup_size[2]
571            )?;
572            writeln!(self.out)?;
573        }
574
575        // Enable early depth tests if needed
576        if let Some(depth_test) = self.entry_point.early_depth_test {
577            // If early depth test is supported for this version of GLSL
578            if self.options.version.supports_early_depth_test() {
579                writeln!(self.out, "layout(early_fragment_tests) in;")?;
580
581                if let Some(conservative) = depth_test.conservative {
582                    use crate::ConservativeDepth as Cd;
583
584                    let depth = match conservative {
585                        Cd::GreaterEqual => "greater",
586                        Cd::LessEqual => "less",
587                        Cd::Unchanged => "unchanged",
588                    };
589                    writeln!(self.out, "layout (depth_{depth}) out float gl_FragDepth;")?;
590                }
591                writeln!(self.out)?;
592            } else {
593                log::warn!(
594                    "Early depth testing is not supported for this version of GLSL: {}",
595                    self.options.version
596                );
597            }
598        }
599
600        if self.entry_point.stage == ShaderStage::Vertex && self.options.version.is_webgl() {
601            if let Some(multiview) = self.multiview.as_ref() {
602                writeln!(self.out, "layout(num_views = {multiview}) in;")?;
603                writeln!(self.out)?;
604            }
605        }
606
607        // Write struct types.
608        //
609        // This are always ordered because the IR is structured in a way that
610        // you can't make a struct without adding all of its members first.
611        for (handle, ty) in self.module.types.iter() {
612            if let TypeInner::Struct { ref members, .. } = ty.inner {
613                // Structures ending with runtime-sized arrays can only be
614                // rendered as shader storage blocks in GLSL, not stand-alone
615                // struct types.
616                if !self.module.types[members.last().unwrap().ty]
617                    .inner
618                    .is_dynamically_sized(&self.module.types)
619                {
620                    let name = &self.names[&NameKey::Type(handle)];
621                    write!(self.out, "struct {name} ")?;
622                    self.write_struct_body(handle, members)?;
623                    writeln!(self.out, ";")?;
624                }
625            }
626        }
627
628        // Write all named constants
629        let mut constants = self
630            .module
631            .constants
632            .iter()
633            .filter(|&(_, c)| c.name.is_some())
634            .peekable();
635        while let Some((handle, _)) = constants.next() {
636            self.write_global_constant(handle)?;
637            // Add extra newline for readability on last iteration
638            if constants.peek().is_none() {
639                writeln!(self.out)?;
640            }
641        }
642
643        let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);
644
645        // Write the globals
646        //
647        // Unless explicitly disabled with WriterFlags::INCLUDE_UNUSED_ITEMS,
648        // we filter all globals that aren't used by the selected entry point as they might be
649        // interfere with each other (i.e. two globals with the same location but different with
650        // different classes)
651        let include_unused = self
652            .options
653            .writer_flags
654            .contains(WriterFlags::INCLUDE_UNUSED_ITEMS);
655        for (handle, global) in self.module.global_variables.iter() {
656            let is_unused = ep_info[handle].is_empty();
657            if !include_unused && is_unused {
658                continue;
659            }
660
661            match self.module.types[global.ty].inner {
662                // We treat images separately because they might require
663                // writing the storage format
664                TypeInner::Image {
665                    mut dim,
666                    arrayed,
667                    class,
668                } => {
669                    // Gather the storage format if needed
670                    let storage_format_access = match self.module.types[global.ty].inner {
671                        TypeInner::Image {
672                            class: crate::ImageClass::Storage { format, access },
673                            ..
674                        } => Some((format, access)),
675                        _ => None,
676                    };
677
678                    if dim == crate::ImageDimension::D1 && es {
679                        dim = crate::ImageDimension::D2
680                    }
681
682                    // Gether the location if needed
683                    let layout_binding = if self.options.version.supports_explicit_locations() {
684                        let br = global.binding.as_ref().unwrap();
685                        self.options.binding_map.get(br).cloned()
686                    } else {
687                        None
688                    };
689
690                    // Write all the layout qualifiers
691                    if layout_binding.is_some() || storage_format_access.is_some() {
692                        write!(self.out, "layout(")?;
693                        if let Some(binding) = layout_binding {
694                            write!(self.out, "binding = {binding}")?;
695                        }
696                        if let Some((format, _)) = storage_format_access {
697                            let format_str = glsl_storage_format(format);
698                            let separator = match layout_binding {
699                                Some(_) => ",",
700                                None => "",
701                            };
702                            write!(self.out, "{separator}{format_str}")?;
703                        }
704                        write!(self.out, ") ")?;
705                    }
706
707                    if let Some((_, access)) = storage_format_access {
708                        self.write_storage_access(access)?;
709                    }
710
711                    // All images in glsl are `uniform`
712                    // The trailing space is important
713                    write!(self.out, "uniform ")?;
714
715                    // write the type
716                    //
717                    // This is way we need the leading space because `write_image_type` doesn't add
718                    // any spaces at the beginning or end
719                    self.write_image_type(dim, arrayed, class)?;
720
721                    // Finally write the name and end the global with a `;`
722                    // The leading space is important
723                    let global_name = self.get_global_name(handle, global);
724                    writeln!(self.out, " {global_name};")?;
725                    writeln!(self.out)?;
726
727                    self.reflection_names_globals.insert(handle, global_name);
728                }
729                // glsl has no concept of samplers so we just ignore it
730                TypeInner::Sampler { .. } => continue,
731                // All other globals are written by `write_global`
732                _ => {
733                    self.write_global(handle, global)?;
734                    // Add a newline (only for readability)
735                    writeln!(self.out)?;
736                }
737            }
738        }
739
740        for arg in self.entry_point.function.arguments.iter() {
741            self.write_varying(arg.binding.as_ref(), arg.ty, false)?;
742        }
743        if let Some(ref result) = self.entry_point.function.result {
744            self.write_varying(result.binding.as_ref(), result.ty, true)?;
745        }
746        writeln!(self.out)?;
747
748        // Write all regular functions
749        for (handle, function) in self.module.functions.iter() {
750            // Check that the function doesn't use globals that aren't supported
751            // by the current entry point
752            if !include_unused && !ep_info.dominates_global_use(&self.info[handle]) {
753                continue;
754            }
755
756            let fun_info = &self.info[handle];
757
758            // Write the function
759            self.write_function(back::FunctionType::Function(handle), function, fun_info)?;
760
761            writeln!(self.out)?;
762        }
763
764        self.write_function(
765            back::FunctionType::EntryPoint(self.entry_point_idx),
766            &self.entry_point.function,
767            ep_info,
768        )?;
769
770        // Add newline at the end of file
771        writeln!(self.out)?;
772
773        // Collect all reflection info and return it to the user
774        self.collect_reflection_info()
775    }
776
777    fn write_array_size(
778        &mut self,
779        base: Handle<crate::Type>,
780        size: crate::ArraySize,
781    ) -> BackendResult {
782        write!(self.out, "[")?;
783
784        // Write the array size
785        // Writes nothing if `ArraySize::Dynamic`
786        match size {
787            crate::ArraySize::Constant(size) => {
788                write!(self.out, "{size}")?;
789            }
790            crate::ArraySize::Dynamic => (),
791        }
792
793        write!(self.out, "]")?;
794
795        if let TypeInner::Array {
796            base: next_base,
797            size: next_size,
798            ..
799        } = self.module.types[base].inner
800        {
801            self.write_array_size(next_base, next_size)?;
802        }
803
804        Ok(())
805    }
806
807    /// Helper method used to write value types
808    ///
809    /// # Notes
810    /// Adds no trailing or leading whitespace
811    fn write_value_type(&mut self, inner: &TypeInner) -> BackendResult {
812        match *inner {
813            // Scalars are simple we just get the full name from `glsl_scalar`
814            TypeInner::Scalar { kind, width }
815            | TypeInner::Atomic { kind, width }
816            | TypeInner::ValuePointer {
817                size: None,
818                kind,
819                width,
820                space: _,
821            } => write!(self.out, "{}", glsl_scalar(kind, width)?.full)?,
822            // Vectors are just `gvecN` where `g` is the scalar prefix and `N` is the vector size
823            TypeInner::Vector { size, kind, width }
824            | TypeInner::ValuePointer {
825                size: Some(size),
826                kind,
827                width,
828                space: _,
829            } => write!(
830                self.out,
831                "{}vec{}",
832                glsl_scalar(kind, width)?.prefix,
833                size as u8
834            )?,
835            // Matrices are written with `gmatMxN` where `g` is the scalar prefix (only floats and
836            // doubles are allowed), `M` is the columns count and `N` is the rows count
837            //
838            // glsl supports a matrix shorthand `gmatN` where `N` = `M` but it doesn't justify the
839            // extra branch to write matrices this way
840            TypeInner::Matrix {
841                columns,
842                rows,
843                width,
844            } => write!(
845                self.out,
846                "{}mat{}x{}",
847                glsl_scalar(crate::ScalarKind::Float, width)?.prefix,
848                columns as u8,
849                rows as u8
850            )?,
851            // GLSL arrays are written as `type name[size]`
852            // Here we only write the size of the array i.e. `[size]`
853            // Base `type` and `name` should be written outside
854            TypeInner::Array { base, size, .. } => self.write_array_size(base, size)?,
855            // Write all variants instead of `_` so that if new variants are added a
856            // no exhaustiveness error is thrown
857            TypeInner::Pointer { .. }
858            | TypeInner::Struct { .. }
859            | TypeInner::Image { .. }
860            | TypeInner::Sampler { .. }
861            | TypeInner::AccelerationStructure
862            | TypeInner::RayQuery
863            | TypeInner::BindingArray { .. } => {
864                return Err(Error::Custom(format!("Unable to write type {inner:?}")))
865            }
866        }
867
868        Ok(())
869    }
870
871    /// Helper method used to write non image/sampler types
872    ///
873    /// # Notes
874    /// Adds no trailing or leading whitespace
875    fn write_type(&mut self, ty: Handle<crate::Type>) -> BackendResult {
876        match self.module.types[ty].inner {
877            // glsl has no pointer types so just write types as normal and loads are skipped
878            TypeInner::Pointer { base, .. } => self.write_type(base),
879            // glsl structs are written as just the struct name
880            TypeInner::Struct { .. } => {
881                // Get the struct name
882                let name = &self.names[&NameKey::Type(ty)];
883                write!(self.out, "{name}")?;
884                Ok(())
885            }
886            // glsl array has the size separated from the base type
887            TypeInner::Array { base, .. } => self.write_type(base),
888            ref other => self.write_value_type(other),
889        }
890    }
891
892    /// Helper method to write a image type
893    ///
894    /// # Notes
895    /// Adds no leading or trailing whitespace
896    fn write_image_type(
897        &mut self,
898        dim: crate::ImageDimension,
899        arrayed: bool,
900        class: crate::ImageClass,
901    ) -> BackendResult {
902        // glsl images consist of four parts the scalar prefix, the image "type", the dimensions
903        // and modifiers
904        //
905        // There exists two image types
906        // - sampler - for sampled images
907        // - image - for storage images
908        //
909        // There are three possible modifiers that can be used together and must be written in
910        // this order to be valid
911        // - MS - used if it's a multisampled image
912        // - Array - used if it's an image array
913        // - Shadow - used if it's a depth image
914        use crate::ImageClass as Ic;
915
916        let (base, kind, ms, comparison) = match class {
917            Ic::Sampled { kind, multi: true } => ("sampler", kind, "MS", ""),
918            Ic::Sampled { kind, multi: false } => ("sampler", kind, "", ""),
919            Ic::Depth { multi: true } => ("sampler", crate::ScalarKind::Float, "MS", ""),
920            Ic::Depth { multi: false } => ("sampler", crate::ScalarKind::Float, "", "Shadow"),
921            Ic::Storage { format, .. } => ("image", format.into(), "", ""),
922        };
923
924        write!(
925            self.out,
926            "highp {}{}{}{}{}{}",
927            glsl_scalar(kind, 4)?.prefix,
928            base,
929            glsl_dimension(dim),
930            ms,
931            if arrayed { "Array" } else { "" },
932            comparison
933        )?;
934
935        Ok(())
936    }
937
938    /// Helper method used to write non images/sampler globals
939    ///
940    /// # Notes
941    /// Adds a newline
942    ///
943    /// # Panics
944    /// If the global has type sampler
945    fn write_global(
946        &mut self,
947        handle: Handle<crate::GlobalVariable>,
948        global: &crate::GlobalVariable,
949    ) -> BackendResult {
950        if self.options.version.supports_explicit_locations() {
951            if let Some(ref br) = global.binding {
952                match self.options.binding_map.get(br) {
953                    Some(binding) => {
954                        let layout = match global.space {
955                            crate::AddressSpace::Storage { .. } => {
956                                if self.options.version.supports_std430_layout() {
957                                    "std430, "
958                                } else {
959                                    "std140, "
960                                }
961                            }
962                            crate::AddressSpace::Uniform => "std140, ",
963                            _ => "",
964                        };
965                        write!(self.out, "layout({layout}binding = {binding}) ")?
966                    }
967                    None => {
968                        log::debug!("unassigned binding for {:?}", global.name);
969                        if let crate::AddressSpace::Storage { .. } = global.space {
970                            if self.options.version.supports_std430_layout() {
971                                write!(self.out, "layout(std430) ")?
972                            }
973                        }
974                    }
975                }
976            }
977        }
978
979        if let crate::AddressSpace::Storage { access } = global.space {
980            self.write_storage_access(access)?;
981        }
982
983        if let Some(storage_qualifier) = glsl_storage_qualifier(global.space) {
984            write!(self.out, "{storage_qualifier} ")?;
985        }
986
987        match global.space {
988            crate::AddressSpace::Private => {
989                self.write_simple_global(handle, global)?;
990            }
991            crate::AddressSpace::WorkGroup => {
992                self.write_simple_global(handle, global)?;
993            }
994            crate::AddressSpace::PushConstant => {
995                self.write_simple_global(handle, global)?;
996            }
997            crate::AddressSpace::Uniform => {
998                self.write_interface_block(handle, global)?;
999            }
1000            crate::AddressSpace::Storage { .. } => {
1001                self.write_interface_block(handle, global)?;
1002            }
1003            // A global variable in the `Function` address space is a
1004            // contradiction in terms.
1005            crate::AddressSpace::Function => unreachable!(),
1006            // Textures and samplers are handled directly in `Writer::write`.
1007            crate::AddressSpace::Handle => unreachable!(),
1008        }
1009
1010        Ok(())
1011    }
1012
1013    fn write_simple_global(
1014        &mut self,
1015        handle: Handle<crate::GlobalVariable>,
1016        global: &crate::GlobalVariable,
1017    ) -> BackendResult {
1018        self.write_type(global.ty)?;
1019        write!(self.out, " ")?;
1020        self.write_global_name(handle, global)?;
1021
1022        if let TypeInner::Array { base, size, .. } = self.module.types[global.ty].inner {
1023            self.write_array_size(base, size)?;
1024        }
1025
1026        if global.space.initializable() && is_value_init_supported(self.module, global.ty) {
1027            write!(self.out, " = ")?;
1028            if let Some(init) = global.init {
1029                self.write_const_expr(init)?;
1030            } else {
1031                self.write_zero_init_value(global.ty)?;
1032            }
1033        }
1034
1035        writeln!(self.out, ";")?;
1036
1037        if let crate::AddressSpace::PushConstant = global.space {
1038            let global_name = self.get_global_name(handle, global);
1039            self.reflection_names_globals.insert(handle, global_name);
1040        }
1041
1042        Ok(())
1043    }
1044
1045    /// Write an interface block for a single Naga global.
1046    ///
1047    /// Write `block_name { members }`. Since `block_name` must be unique
1048    /// between blocks and structs, we add `_block_ID` where `ID` is a
1049    /// `IdGenerator` generated number. Write `members` in the same way we write
1050    /// a struct's members.
1051    fn write_interface_block(
1052        &mut self,
1053        handle: Handle<crate::GlobalVariable>,
1054        global: &crate::GlobalVariable,
1055    ) -> BackendResult {
1056        // Write the block name, it's just the struct name appended with `_block_ID`
1057        let ty_name = &self.names[&NameKey::Type(global.ty)];
1058        let block_name = format!(
1059            "{}_block_{}{:?}",
1060            ty_name,
1061            self.block_id.generate(),
1062            self.entry_point.stage,
1063        );
1064        write!(self.out, "{block_name} ")?;
1065        self.reflection_names_globals.insert(handle, block_name);
1066
1067        match self.module.types[global.ty].inner {
1068            crate::TypeInner::Struct { ref members, .. }
1069                if self.module.types[members.last().unwrap().ty]
1070                    .inner
1071                    .is_dynamically_sized(&self.module.types) =>
1072            {
1073                // Structs with dynamically sized arrays must have their
1074                // members lifted up as members of the interface block. GLSL
1075                // can't write such struct types anyway.
1076                self.write_struct_body(global.ty, members)?;
1077                write!(self.out, " ")?;
1078                self.write_global_name(handle, global)?;
1079            }
1080            _ => {
1081                // A global of any other type is written as the sole member
1082                // of the interface block. Since the interface block is
1083                // anonymous, this becomes visible in the global scope.
1084                write!(self.out, "{{ ")?;
1085                self.write_type(global.ty)?;
1086                write!(self.out, " ")?;
1087                self.write_global_name(handle, global)?;
1088                if let TypeInner::Array { base, size, .. } = self.module.types[global.ty].inner {
1089                    self.write_array_size(base, size)?;
1090                }
1091                write!(self.out, "; }}")?;
1092            }
1093        }
1094
1095        writeln!(self.out, ";")?;
1096
1097        Ok(())
1098    }
1099
1100    /// Helper method used to find which expressions of a given function require baking
1101    ///
1102    /// # Notes
1103    /// Clears `need_bake_expressions` set before adding to it
1104    fn update_expressions_to_bake(&mut self, func: &crate::Function, info: &valid::FunctionInfo) {
1105        use crate::Expression;
1106        self.need_bake_expressions.clear();
1107        for (fun_handle, expr) in func.expressions.iter() {
1108            let expr_info = &info[fun_handle];
1109            let min_ref_count = func.expressions[fun_handle].bake_ref_count();
1110            if min_ref_count <= expr_info.ref_count {
1111                self.need_bake_expressions.insert(fun_handle);
1112            }
1113
1114            let inner = expr_info.ty.inner_with(&self.module.types);
1115
1116            if let Expression::Math { fun, arg, arg1, .. } = *expr {
1117                match fun {
1118                    crate::MathFunction::Dot => {
1119                        // if the expression is a Dot product with integer arguments,
1120                        // then the args needs baking as well
1121                        if let TypeInner::Scalar { kind, .. } = *inner {
1122                            match kind {
1123                                crate::ScalarKind::Sint | crate::ScalarKind::Uint => {
1124                                    self.need_bake_expressions.insert(arg);
1125                                    self.need_bake_expressions.insert(arg1.unwrap());
1126                                }
1127                                _ => {}
1128                            }
1129                        }
1130                    }
1131                    crate::MathFunction::CountLeadingZeros => {
1132                        if let Some(crate::ScalarKind::Sint) = inner.scalar_kind() {
1133                            self.need_bake_expressions.insert(arg);
1134                        }
1135                    }
1136                    _ => {}
1137                }
1138            }
1139        }
1140    }
1141
1142    /// Helper method used to get a name for a global
1143    ///
1144    /// Globals have different naming schemes depending on their binding:
1145    /// - Globals without bindings use the name from the [`Namer`](crate::proc::Namer)
1146    /// - Globals with resource binding are named `_group_X_binding_Y` where `X`
1147    ///   is the group and `Y` is the binding
1148    fn get_global_name(
1149        &self,
1150        handle: Handle<crate::GlobalVariable>,
1151        global: &crate::GlobalVariable,
1152    ) -> String {
1153        match global.binding {
1154            Some(ref br) => {
1155                format!(
1156                    "_group_{}_binding_{}_{}",
1157                    br.group,
1158                    br.binding,
1159                    self.entry_point.stage.to_str()
1160                )
1161            }
1162            None => self.names[&NameKey::GlobalVariable(handle)].clone(),
1163        }
1164    }
1165
1166    /// Helper method used to write a name for a global without additional heap allocation
1167    fn write_global_name(
1168        &mut self,
1169        handle: Handle<crate::GlobalVariable>,
1170        global: &crate::GlobalVariable,
1171    ) -> BackendResult {
1172        match global.binding {
1173            Some(ref br) => write!(
1174                self.out,
1175                "_group_{}_binding_{}_{}",
1176                br.group,
1177                br.binding,
1178                self.entry_point.stage.to_str()
1179            )?,
1180            None => write!(
1181                self.out,
1182                "{}",
1183                &self.names[&NameKey::GlobalVariable(handle)]
1184            )?,
1185        }
1186
1187        Ok(())
1188    }
1189
1190    /// Write a GLSL global that will carry a Naga entry point's argument or return value.
1191    ///
1192    /// A Naga entry point's arguments and return value are rendered in GLSL as
1193    /// variables at global scope with the `in` and `out` storage qualifiers.
1194    /// The code we generate for `main` loads from all the `in` globals into
1195    /// appropriately named locals. Before it returns, `main` assigns the
1196    /// components of its return value into all the `out` globals.
1197    ///
1198    /// This function writes a declaration for one such GLSL global,
1199    /// representing a value passed into or returned from [`self.entry_point`]
1200    /// that has a [`Location`] binding. The global's name is generated based on
1201    /// the location index and the shader stages being connected; see
1202    /// [`VaryingName`]. This means we don't need to know the names of
1203    /// arguments, just their types and bindings.
1204    ///
1205    /// Emit nothing for entry point arguments or return values with [`BuiltIn`]
1206    /// bindings; `main` will read from or assign to the appropriate GLSL
1207    /// special variable; these are pre-declared. As an exception, we do declare
1208    /// `gl_Position` or `gl_FragCoord` with the `invariant` qualifier if
1209    /// needed.
1210    ///
1211    /// Use `output` together with [`self.entry_point.stage`] to determine which
1212    /// shader stages are being connected, and choose the `in` or `out` storage
1213    /// qualifier.
1214    ///
1215    /// [`self.entry_point`]: Writer::entry_point
1216    /// [`self.entry_point.stage`]: crate::EntryPoint::stage
1217    /// [`Location`]: crate::Binding::Location
1218    /// [`BuiltIn`]: crate::Binding::BuiltIn
1219    fn write_varying(
1220        &mut self,
1221        binding: Option<&crate::Binding>,
1222        ty: Handle<crate::Type>,
1223        output: bool,
1224    ) -> Result<(), Error> {
1225        // For a struct, emit a separate global for each member with a binding.
1226        if let crate::TypeInner::Struct { ref members, .. } = self.module.types[ty].inner {
1227            for member in members {
1228                self.write_varying(member.binding.as_ref(), member.ty, output)?;
1229            }
1230            return Ok(());
1231        }
1232
1233        let binding = match binding {
1234            None => return Ok(()),
1235            Some(binding) => binding,
1236        };
1237
1238        let (location, interpolation, sampling) = match *binding {
1239            crate::Binding::Location {
1240                location,
1241                interpolation,
1242                sampling,
1243            } => (location, interpolation, sampling),
1244            crate::Binding::BuiltIn(built_in) => {
1245                if let crate::BuiltIn::Position { invariant: true } = built_in {
1246                    match (self.options.version, self.entry_point.stage) {
1247                        (
1248                            Version::Embedded {
1249                                version: 300,
1250                                is_webgl: true,
1251                            },
1252                            ShaderStage::Fragment,
1253                        ) => {
1254                            // `invariant gl_FragCoord` is not allowed in WebGL2 and possibly
1255                            // OpenGL ES in general (waiting on confirmation).
1256                            //
1257                            // See https://github.com/KhronosGroup/WebGL/issues/3518
1258                        }
1259                        _ => {
1260                            writeln!(
1261                                self.out,
1262                                "invariant {};",
1263                                glsl_built_in(built_in, output, self.options.version.is_webgl())
1264                            )?;
1265                        }
1266                    }
1267                }
1268                return Ok(());
1269            }
1270        };
1271
1272        // Write the interpolation modifier if needed
1273        //
1274        // We ignore all interpolation and auxiliary modifiers that aren't used in fragment
1275        // shaders' input globals or vertex shaders' output globals.
1276        let emit_interpolation_and_auxiliary = match self.entry_point.stage {
1277            ShaderStage::Vertex => output,
1278            ShaderStage::Fragment => !output,
1279            ShaderStage::Compute => false,
1280        };
1281
1282        // Write the I/O locations, if allowed
1283        if self.options.version.supports_explicit_locations() || !emit_interpolation_and_auxiliary {
1284            write!(self.out, "layout(location = {location}) ")?;
1285        }
1286
1287        // Write the interpolation qualifier.
1288        if let Some(interp) = interpolation {
1289            if emit_interpolation_and_auxiliary {
1290                write!(self.out, "{} ", glsl_interpolation(interp))?;
1291            }
1292        }
1293
1294        // Write the sampling auxiliary qualifier.
1295        //
1296        // Before GLSL 4.2, the `centroid` and `sample` qualifiers were required to appear
1297        // immediately before the `in` / `out` qualifier, so we'll just follow that rule
1298        // here, regardless of the version.
1299        if let Some(sampling) = sampling {
1300            if emit_interpolation_and_auxiliary {
1301                if let Some(qualifier) = glsl_sampling(sampling) {
1302                    write!(self.out, "{qualifier} ")?;
1303                }
1304            }
1305        }
1306
1307        // Write the input/output qualifier.
1308        write!(self.out, "{} ", if output { "out" } else { "in" })?;
1309
1310        // Write the type
1311        // `write_type` adds no leading or trailing spaces
1312        self.write_type(ty)?;
1313
1314        // Finally write the global name and end the global with a `;` and a newline
1315        // Leading space is important
1316        let vname = VaryingName {
1317            binding: &crate::Binding::Location {
1318                location,
1319                interpolation: None,
1320                sampling: None,
1321            },
1322            stage: self.entry_point.stage,
1323            output,
1324            targetting_webgl: self.options.version.is_webgl(),
1325        };
1326        writeln!(self.out, " {vname};")?;
1327
1328        Ok(())
1329    }
1330
1331    /// Helper method used to write functions (both entry points and regular functions)
1332    ///
1333    /// # Notes
1334    /// Adds a newline
1335    fn write_function(
1336        &mut self,
1337        ty: back::FunctionType,
1338        func: &crate::Function,
1339        info: &valid::FunctionInfo,
1340    ) -> BackendResult {
1341        // Create a function context for the function being written
1342        let ctx = back::FunctionCtx {
1343            ty,
1344            info,
1345            expressions: &func.expressions,
1346            named_expressions: &func.named_expressions,
1347        };
1348
1349        self.named_expressions.clear();
1350        self.update_expressions_to_bake(func, info);
1351
1352        // Write the function header
1353        //
1354        // glsl headers are the same as in c:
1355        // `ret_type name(args)`
1356        // `ret_type` is the return type
1357        // `name` is the function name
1358        // `args` is a comma separated list of `type name`
1359        //  | - `type` is the argument type
1360        //  | - `name` is the argument name
1361
1362        // Start by writing the return type if any otherwise write void
1363        // This is the only place where `void` is a valid type
1364        // (though it's more a keyword than a type)
1365        if let back::FunctionType::EntryPoint(_) = ctx.ty {
1366            write!(self.out, "void")?;
1367        } else if let Some(ref result) = func.result {
1368            self.write_type(result.ty)?;
1369            if let TypeInner::Array { base, size, .. } = self.module.types[result.ty].inner {
1370                self.write_array_size(base, size)?
1371            }
1372        } else {
1373            write!(self.out, "void")?;
1374        }
1375
1376        // Write the function name and open parentheses for the argument list
1377        let function_name = match ctx.ty {
1378            back::FunctionType::Function(handle) => &self.names[&NameKey::Function(handle)],
1379            back::FunctionType::EntryPoint(_) => "main",
1380        };
1381        write!(self.out, " {function_name}(")?;
1382
1383        // Write the comma separated argument list
1384        //
1385        // We need access to `Self` here so we use the reference passed to the closure as an
1386        // argument instead of capturing as that would cause a borrow checker error
1387        let arguments = match ctx.ty {
1388            back::FunctionType::EntryPoint(_) => &[][..],
1389            back::FunctionType::Function(_) => &func.arguments,
1390        };
1391        let arguments: Vec<_> = arguments
1392            .iter()
1393            .enumerate()
1394            .filter(|&(_, arg)| match self.module.types[arg.ty].inner {
1395                TypeInner::Sampler { .. } => false,
1396                _ => true,
1397            })
1398            .collect();
1399        self.write_slice(&arguments, |this, _, &(i, arg)| {
1400            // Write the argument type
1401            match this.module.types[arg.ty].inner {
1402                // We treat images separately because they might require
1403                // writing the storage format
1404                TypeInner::Image {
1405                    dim,
1406                    arrayed,
1407                    class,
1408                } => {
1409                    // Write the storage format if needed
1410                    if let TypeInner::Image {
1411                        class: crate::ImageClass::Storage { format, .. },
1412                        ..
1413                    } = this.module.types[arg.ty].inner
1414                    {
1415                        write!(this.out, "layout({}) ", glsl_storage_format(format))?;
1416                    }
1417
1418                    // write the type
1419                    //
1420                    // This is way we need the leading space because `write_image_type` doesn't add
1421                    // any spaces at the beginning or end
1422                    this.write_image_type(dim, arrayed, class)?;
1423                }
1424                TypeInner::Pointer { base, .. } => {
1425                    // write parameter qualifiers
1426                    write!(this.out, "inout ")?;
1427                    this.write_type(base)?;
1428                }
1429                // All other types are written by `write_type`
1430                _ => {
1431                    this.write_type(arg.ty)?;
1432                }
1433            }
1434
1435            // Write the argument name
1436            // The leading space is important
1437            write!(this.out, " {}", &this.names[&ctx.argument_key(i as u32)])?;
1438
1439            // Write array size
1440            match this.module.types[arg.ty].inner {
1441                TypeInner::Array { base, size, .. } => {
1442                    this.write_array_size(base, size)?;
1443                }
1444                TypeInner::Pointer { base, .. } => {
1445                    if let TypeInner::Array { base, size, .. } = this.module.types[base].inner {
1446                        this.write_array_size(base, size)?;
1447                    }
1448                }
1449                _ => {}
1450            }
1451
1452            Ok(())
1453        })?;
1454
1455        // Close the parentheses and open braces to start the function body
1456        writeln!(self.out, ") {{")?;
1457
1458        if self.options.zero_initialize_workgroup_memory
1459            && ctx.ty.is_compute_entry_point(self.module)
1460        {
1461            self.write_workgroup_variables_initialization(&ctx)?;
1462        }
1463
1464        // Compose the function arguments from globals, in case of an entry point.
1465        if let back::FunctionType::EntryPoint(ep_index) = ctx.ty {
1466            let stage = self.module.entry_points[ep_index as usize].stage;
1467            for (index, arg) in func.arguments.iter().enumerate() {
1468                write!(self.out, "{}", back::INDENT)?;
1469                self.write_type(arg.ty)?;
1470                let name = &self.names[&NameKey::EntryPointArgument(ep_index, index as u32)];
1471                write!(self.out, " {name}")?;
1472                write!(self.out, " = ")?;
1473                match self.module.types[arg.ty].inner {
1474                    crate::TypeInner::Struct { ref members, .. } => {
1475                        self.write_type(arg.ty)?;
1476                        write!(self.out, "(")?;
1477                        for (index, member) in members.iter().enumerate() {
1478                            let varying_name = VaryingName {
1479                                binding: member.binding.as_ref().unwrap(),
1480                                stage,
1481                                output: false,
1482                                targetting_webgl: self.options.version.is_webgl(),
1483                            };
1484                            if index != 0 {
1485                                write!(self.out, ", ")?;
1486                            }
1487                            write!(self.out, "{varying_name}")?;
1488                        }
1489                        writeln!(self.out, ");")?;
1490                    }
1491                    _ => {
1492                        let varying_name = VaryingName {
1493                            binding: arg.binding.as_ref().unwrap(),
1494                            stage,
1495                            output: false,
1496                            targetting_webgl: self.options.version.is_webgl(),
1497                        };
1498                        writeln!(self.out, "{varying_name};")?;
1499                    }
1500                }
1501            }
1502        }
1503
1504        // Write all function locals
1505        // Locals are `type name (= init)?;` where the init part (including the =) are optional
1506        //
1507        // Always adds a newline
1508        for (handle, local) in func.local_variables.iter() {
1509            // Write indentation (only for readability) and the type
1510            // `write_type` adds no trailing space
1511            write!(self.out, "{}", back::INDENT)?;
1512            self.write_type(local.ty)?;
1513
1514            // Write the local name
1515            // The leading space is important
1516            write!(self.out, " {}", self.names[&ctx.name_key(handle)])?;
1517            // Write size for array type
1518            if let TypeInner::Array { base, size, .. } = self.module.types[local.ty].inner {
1519                self.write_array_size(base, size)?;
1520            }
1521            // Write the local initializer if needed
1522            if let Some(init) = local.init {
1523                // Put the equal signal only if there's a initializer
1524                // The leading and trailing spaces aren't needed but help with readability
1525                write!(self.out, " = ")?;
1526
1527                // Write the constant
1528                // `write_constant` adds no trailing or leading space/newline
1529                self.write_const_expr(init)?;
1530            } else if is_value_init_supported(self.module, local.ty) {
1531                write!(self.out, " = ")?;
1532                self.write_zero_init_value(local.ty)?;
1533            }
1534
1535            // Finish the local with `;` and add a newline (only for readability)
1536            writeln!(self.out, ";")?
1537        }
1538
1539        // Write the function body (statement list)
1540        for sta in func.body.iter() {
1541            // Write a statement, the indentation should always be 1 when writing the function body
1542            // `write_stmt` adds a newline
1543            self.write_stmt(sta, &ctx, back::Level(1))?;
1544        }
1545
1546        // Close braces and add a newline
1547        writeln!(self.out, "}}")?;
1548
1549        Ok(())
1550    }
1551
1552    fn write_workgroup_variables_initialization(
1553        &mut self,
1554        ctx: &back::FunctionCtx,
1555    ) -> BackendResult {
1556        let mut vars = self
1557            .module
1558            .global_variables
1559            .iter()
1560            .filter(|&(handle, var)| {
1561                !ctx.info[handle].is_empty() && var.space == crate::AddressSpace::WorkGroup
1562            })
1563            .peekable();
1564
1565        if vars.peek().is_some() {
1566            let level = back::Level(1);
1567
1568            writeln!(self.out, "{level}if (gl_LocalInvocationID == uvec3(0u)) {{")?;
1569
1570            for (handle, var) in vars {
1571                let name = &self.names[&NameKey::GlobalVariable(handle)];
1572                write!(self.out, "{}{} = ", level.next(), name)?;
1573                self.write_zero_init_value(var.ty)?;
1574                writeln!(self.out, ";")?;
1575            }
1576
1577            writeln!(self.out, "{level}}}")?;
1578            self.write_barrier(crate::Barrier::WORK_GROUP, level)?;
1579        }
1580
1581        Ok(())
1582    }
1583
1584    /// Write a list of comma separated `T` values using a writer function `F`.
1585    ///
1586    /// The writer function `F` receives a mutable reference to `self` that if needed won't cause
1587    /// borrow checker issues (using for example a closure with `self` will cause issues), the
1588    /// second argument is the 0 based index of the element on the list, and the last element is
1589    /// a reference to the element `T` being written
1590    ///
1591    /// # Notes
1592    /// - Adds no newlines or leading/trailing whitespace
1593    /// - The last element won't have a trailing `,`
1594    fn write_slice<T, F: FnMut(&mut Self, u32, &T) -> BackendResult>(
1595        &mut self,
1596        data: &[T],
1597        mut f: F,
1598    ) -> BackendResult {
1599        // Loop through `data` invoking `f` for each element
1600        for (index, item) in data.iter().enumerate() {
1601            if index != 0 {
1602                write!(self.out, ", ")?;
1603            }
1604            f(self, index as u32, item)?;
1605        }
1606
1607        Ok(())
1608    }
1609
1610    /// Helper method used to write global constants
1611    fn write_global_constant(&mut self, handle: Handle<crate::Constant>) -> BackendResult {
1612        write!(self.out, "const ")?;
1613        let constant = &self.module.constants[handle];
1614        self.write_type(constant.ty)?;
1615        let name = &self.names[&NameKey::Constant(handle)];
1616        write!(self.out, " {name}")?;
1617        if let TypeInner::Array { base, size, .. } = self.module.types[constant.ty].inner {
1618            self.write_array_size(base, size)?;
1619        }
1620        write!(self.out, " = ")?;
1621        self.write_const_expr(constant.init)?;
1622        writeln!(self.out, ";")?;
1623        Ok(())
1624    }
1625
1626    /// Helper method used to output a dot product as an arithmetic expression
1627    ///
1628    fn write_dot_product(
1629        &mut self,
1630        arg: Handle<crate::Expression>,
1631        arg1: Handle<crate::Expression>,
1632        size: usize,
1633        ctx: &back::FunctionCtx<'_>,
1634    ) -> BackendResult {
1635        // Write parantheses around the dot product expression to prevent operators
1636        // with different precedences from applying earlier.
1637        write!(self.out, "(")?;
1638
1639        // Cycle trough all the components of the vector
1640        for index in 0..size {
1641            let component = back::COMPONENTS[index];
1642            // Write the addition to the previous product
1643            // This will print an extra '+' at the beginning but that is fine in glsl
1644            write!(self.out, " + ")?;
1645            // Write the first vector expression, this expression is marked to be
1646            // cached so unless it can't be cached (for example, it's a Constant)
1647            // it shouldn't produce large expressions.
1648            self.write_expr(arg, ctx)?;
1649            // Access the current component on the first vector
1650            write!(self.out, ".{component} * ")?;
1651            // Write the second vector expression, this expression is marked to be
1652            // cached so unless it can't be cached (for example, it's a Constant)
1653            // it shouldn't produce large expressions.
1654            self.write_expr(arg1, ctx)?;
1655            // Access the current component on the second vector
1656            write!(self.out, ".{component}")?;
1657        }
1658
1659        write!(self.out, ")")?;
1660        Ok(())
1661    }
1662
1663    /// Helper method used to write structs
1664    ///
1665    /// # Notes
1666    /// Ends in a newline
1667    fn write_struct_body(
1668        &mut self,
1669        handle: Handle<crate::Type>,
1670        members: &[crate::StructMember],
1671    ) -> BackendResult {
1672        // glsl structs are written as in C
1673        // `struct name() { members };`
1674        //  | `struct` is a keyword
1675        //  | `name` is the struct name
1676        //  | `members` is a semicolon separated list of `type name`
1677        //      | `type` is the member type
1678        //      | `name` is the member name
1679        writeln!(self.out, "{{")?;
1680
1681        for (idx, member) in members.iter().enumerate() {
1682            // The indentation is only for readability
1683            write!(self.out, "{}", back::INDENT)?;
1684
1685            match self.module.types[member.ty].inner {
1686                TypeInner::Array {
1687                    base,
1688                    size,
1689                    stride: _,
1690                } => {
1691                    self.write_type(base)?;
1692                    write!(
1693                        self.out,
1694                        " {}",
1695                        &self.names[&NameKey::StructMember(handle, idx as u32)]
1696                    )?;
1697                    // Write [size]
1698                    self.write_array_size(base, size)?;
1699                    // Newline is important
1700                    writeln!(self.out, ";")?;
1701                }
1702                _ => {
1703                    // Write the member type
1704                    // Adds no trailing space
1705                    self.write_type(member.ty)?;
1706
1707                    // Write the member name and put a semicolon
1708                    // The leading space is important
1709                    // All members must have a semicolon even the last one
1710                    writeln!(
1711                        self.out,
1712                        " {};",
1713                        &self.names[&NameKey::StructMember(handle, idx as u32)]
1714                    )?;
1715                }
1716            }
1717        }
1718
1719        write!(self.out, "}}")?;
1720        Ok(())
1721    }
1722
1723    /// Helper method used to write statements
1724    ///
1725    /// # Notes
1726    /// Always adds a newline
1727    fn write_stmt(
1728        &mut self,
1729        sta: &crate::Statement,
1730        ctx: &back::FunctionCtx,
1731        level: back::Level,
1732    ) -> BackendResult {
1733        use crate::Statement;
1734
1735        match *sta {
1736            // This is where we can generate intermediate constants for some expression types.
1737            Statement::Emit(ref range) => {
1738                for handle in range.clone() {
1739                    let info = &ctx.info[handle];
1740                    let ptr_class = info.ty.inner_with(&self.module.types).pointer_space();
1741                    let expr_name = if ptr_class.is_some() {
1742                        // GLSL can't save a pointer-valued expression in a variable,
1743                        // but we shouldn't ever need to: they should never be named expressions,
1744                        // and none of the expression types flagged by bake_ref_count can be pointer-valued.
1745                        None
1746                    } else if let Some(name) = ctx.named_expressions.get(&handle) {
1747                        // Front end provides names for all variables at the start of writing.
1748                        // But we write them to step by step. We need to recache them
1749                        // Otherwise, we could accidentally write variable name instead of full expression.
1750                        // Also, we use sanitized names! It defense backend from generating variable with name from reserved keywords.
1751                        Some(self.namer.call(name))
1752                    } else if self.need_bake_expressions.contains(&handle) {
1753                        Some(format!("{}{}", back::BAKE_PREFIX, handle.index()))
1754                    } else {
1755                        None
1756                    };
1757
1758                    // If we are going to write an `ImageLoad` next and the target image
1759                    // is sampled and we are using the `Restrict` policy for bounds
1760                    // checking images we need to write a local holding the clamped lod.
1761                    if let crate::Expression::ImageLoad {
1762                        image,
1763                        level: Some(level_expr),
1764                        ..
1765                    } = ctx.expressions[handle]
1766                    {
1767                        if let TypeInner::Image {
1768                            class: crate::ImageClass::Sampled { .. },
1769                            ..
1770                        } = *ctx.info[image].ty.inner_with(&self.module.types)
1771                        {
1772                            if let proc::BoundsCheckPolicy::Restrict = self.policies.image_load {
1773                                write!(self.out, "{level}")?;
1774                                self.write_clamped_lod(ctx, handle, image, level_expr)?
1775                            }
1776                        }
1777                    }
1778
1779                    if let Some(name) = expr_name {
1780                        write!(self.out, "{level}")?;
1781                        self.write_named_expr(handle, name, handle, ctx)?;
1782                    }
1783                }
1784            }
1785            // Blocks are simple we just need to write the block statements between braces
1786            // We could also just print the statements but this is more readable and maps more
1787            // closely to the IR
1788            Statement::Block(ref block) => {
1789                write!(self.out, "{level}")?;
1790                writeln!(self.out, "{{")?;
1791                for sta in block.iter() {
1792                    // Increase the indentation to help with readability
1793                    self.write_stmt(sta, ctx, level.next())?
1794                }
1795                writeln!(self.out, "{level}}}")?
1796            }
1797            // Ifs are written as in C:
1798            // ```
1799            // if(condition) {
1800            //  accept
1801            // } else {
1802            //  reject
1803            // }
1804            // ```
1805            Statement::If {
1806                condition,
1807                ref accept,
1808                ref reject,
1809            } => {
1810                write!(self.out, "{level}")?;
1811                write!(self.out, "if (")?;
1812                self.write_expr(condition, ctx)?;
1813                writeln!(self.out, ") {{")?;
1814
1815                for sta in accept {
1816                    // Increase indentation to help with readability
1817                    self.write_stmt(sta, ctx, level.next())?;
1818                }
1819
1820                // If there are no statements in the reject block we skip writing it
1821                // This is only for readability
1822                if !reject.is_empty() {
1823                    writeln!(self.out, "{level}}} else {{")?;
1824
1825                    for sta in reject {
1826                        // Increase indentation to help with readability
1827                        self.write_stmt(sta, ctx, level.next())?;
1828                    }
1829                }
1830
1831                writeln!(self.out, "{level}}}")?
1832            }
1833            // Switch are written as in C:
1834            // ```
1835            // switch (selector) {
1836            //      // Fallthrough
1837            //      case label:
1838            //          block
1839            //      // Non fallthrough
1840            //      case label:
1841            //          block
1842            //          break;
1843            //      default:
1844            //          block
1845            //  }
1846            //  ```
1847            //  Where the `default` case happens isn't important but we put it last
1848            //  so that we don't need to print a `break` for it
1849            Statement::Switch {
1850                selector,
1851                ref cases,
1852            } => {
1853                // Start the switch
1854                write!(self.out, "{level}")?;
1855                write!(self.out, "switch(")?;
1856                self.write_expr(selector, ctx)?;
1857                writeln!(self.out, ") {{")?;
1858
1859                // Write all cases
1860                let l2 = level.next();
1861                for case in cases {
1862                    match case.value {
1863                        crate::SwitchValue::I32(value) => write!(self.out, "{l2}case {value}:")?,
1864                        crate::SwitchValue::U32(value) => write!(self.out, "{l2}case {value}u:")?,
1865                        crate::SwitchValue::Default => write!(self.out, "{l2}default:")?,
1866                    }
1867
1868                    let write_block_braces = !(case.fall_through && case.body.is_empty());
1869                    if write_block_braces {
1870                        writeln!(self.out, " {{")?;
1871                    } else {
1872                        writeln!(self.out)?;
1873                    }
1874
1875                    for sta in case.body.iter() {
1876                        self.write_stmt(sta, ctx, l2.next())?;
1877                    }
1878
1879                    if !case.fall_through && case.body.last().map_or(true, |s| !s.is_terminator()) {
1880                        writeln!(self.out, "{}break;", l2.next())?;
1881                    }
1882
1883                    if write_block_braces {
1884                        writeln!(self.out, "{l2}}}")?;
1885                    }
1886                }
1887
1888                writeln!(self.out, "{level}}}")?
1889            }
1890            // Loops in naga IR are based on wgsl loops, glsl can emulate the behaviour by using a
1891            // while true loop and appending the continuing block to the body resulting on:
1892            // ```
1893            // bool loop_init = true;
1894            // while(true) {
1895            //  if (!loop_init) { <continuing> }
1896            //  loop_init = false;
1897            //  <body>
1898            // }
1899            // ```
1900            Statement::Loop {
1901                ref body,
1902                ref continuing,
1903                break_if,
1904            } => {
1905                if !continuing.is_empty() || break_if.is_some() {
1906                    let gate_name = self.namer.call("loop_init");
1907                    writeln!(self.out, "{level}bool {gate_name} = true;")?;
1908                    writeln!(self.out, "{level}while(true) {{")?;
1909                    let l2 = level.next();
1910                    let l3 = l2.next();
1911                    writeln!(self.out, "{l2}if (!{gate_name}) {{")?;
1912                    for sta in continuing {
1913                        self.write_stmt(sta, ctx, l3)?;
1914                    }
1915                    if let Some(condition) = break_if {
1916                        write!(self.out, "{l3}if (")?;
1917                        self.write_expr(condition, ctx)?;
1918                        writeln!(self.out, ") {{")?;
1919                        writeln!(self.out, "{}break;", l3.next())?;
1920                        writeln!(self.out, "{l3}}}")?;
1921                    }
1922                    writeln!(self.out, "{l2}}}")?;
1923                    writeln!(self.out, "{}{} = false;", level.next(), gate_name)?;
1924                } else {
1925                    writeln!(self.out, "{level}while(true) {{")?;
1926                }
1927                for sta in body {
1928                    self.write_stmt(sta, ctx, level.next())?;
1929                }
1930                writeln!(self.out, "{level}}}")?
1931            }
1932            // Break, continue and return as written as in C
1933            // `break;`
1934            Statement::Break => {
1935                write!(self.out, "{level}")?;
1936                writeln!(self.out, "break;")?
1937            }
1938            // `continue;`
1939            Statement::Continue => {
1940                write!(self.out, "{level}")?;
1941                writeln!(self.out, "continue;")?
1942            }
1943            // `return expr;`, `expr` is optional
1944            Statement::Return { value } => {
1945                write!(self.out, "{level}")?;
1946                match ctx.ty {
1947                    back::FunctionType::Function(_) => {
1948                        write!(self.out, "return")?;
1949                        // Write the expression to be returned if needed
1950                        if let Some(expr) = value {
1951                            write!(self.out, " ")?;
1952                            self.write_expr(expr, ctx)?;
1953                        }
1954                        writeln!(self.out, ";")?;
1955                    }
1956                    back::FunctionType::EntryPoint(ep_index) => {
1957                        let mut has_point_size = false;
1958                        let ep = &self.module.entry_points[ep_index as usize];
1959                        if let Some(ref result) = ep.function.result {
1960                            let value = value.unwrap();
1961                            match self.module.types[result.ty].inner {
1962                                crate::TypeInner::Struct { ref members, .. } => {
1963                                    let temp_struct_name = match ctx.expressions[value] {
1964                                        crate::Expression::Compose { .. } => {
1965                                            let return_struct = "_tmp_return";
1966                                            write!(
1967                                                self.out,
1968                                                "{} {} = ",
1969                                                &self.names[&NameKey::Type(result.ty)],
1970                                                return_struct
1971                                            )?;
1972                                            self.write_expr(value, ctx)?;
1973                                            writeln!(self.out, ";")?;
1974                                            write!(self.out, "{level}")?;
1975                                            Some(return_struct)
1976                                        }
1977                                        _ => None,
1978                                    };
1979
1980                                    for (index, member) in members.iter().enumerate() {
1981                                        if let Some(crate::Binding::BuiltIn(
1982                                            crate::BuiltIn::PointSize,
1983                                        )) = member.binding
1984                                        {
1985                                            has_point_size = true;
1986                                        }
1987
1988                                        let varying_name = VaryingName {
1989                                            binding: member.binding.as_ref().unwrap(),
1990                                            stage: ep.stage,
1991                                            output: true,
1992                                            targetting_webgl: self.options.version.is_webgl(),
1993                                        };
1994                                        write!(self.out, "{varying_name} = ")?;
1995
1996                                        if let Some(struct_name) = temp_struct_name {
1997                                            write!(self.out, "{struct_name}")?;
1998                                        } else {
1999                                            self.write_expr(value, ctx)?;
2000                                        }
2001
2002                                        // Write field name
2003                                        writeln!(
2004                                            self.out,
2005                                            ".{};",
2006                                            &self.names
2007                                                [&NameKey::StructMember(result.ty, index as u32)]
2008                                        )?;
2009                                        write!(self.out, "{level}")?;
2010                                    }
2011                                }
2012                                _ => {
2013                                    let name = VaryingName {
2014                                        binding: result.binding.as_ref().unwrap(),
2015                                        stage: ep.stage,
2016                                        output: true,
2017                                        targetting_webgl: self.options.version.is_webgl(),
2018                                    };
2019                                    write!(self.out, "{name} = ")?;
2020                                    self.write_expr(value, ctx)?;
2021                                    writeln!(self.out, ";")?;
2022                                    write!(self.out, "{level}")?;
2023                                }
2024                            }
2025                        }
2026
2027                        let is_vertex_stage = self.module.entry_points[ep_index as usize].stage
2028                            == ShaderStage::Vertex;
2029                        if is_vertex_stage
2030                            && self
2031                                .options
2032                                .writer_flags
2033                                .contains(WriterFlags::ADJUST_COORDINATE_SPACE)
2034                        {
2035                            writeln!(
2036                                self.out,
2037                                "gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);",
2038                            )?;
2039                            write!(self.out, "{level}")?;
2040                        }
2041
2042                        if is_vertex_stage
2043                            && self
2044                                .options
2045                                .writer_flags
2046                                .contains(WriterFlags::FORCE_POINT_SIZE)
2047                            && !has_point_size
2048                        {
2049                            writeln!(self.out, "gl_PointSize = 1.0;")?;
2050                            write!(self.out, "{level}")?;
2051                        }
2052                        writeln!(self.out, "return;")?;
2053                    }
2054                }
2055            }
2056            // This is one of the places were glsl adds to the syntax of C in this case the discard
2057            // keyword which ceases all further processing in a fragment shader, it's called OpKill
2058            // in spir-v that's why it's called `Statement::Kill`
2059            Statement::Kill => writeln!(self.out, "{level}discard;")?,
2060            Statement::Barrier(flags) => {
2061                self.write_barrier(flags, level)?;
2062            }
2063            // Stores in glsl are just variable assignments written as `pointer = value;`
2064            Statement::Store { pointer, value } => {
2065                write!(self.out, "{level}")?;
2066                self.write_expr(pointer, ctx)?;
2067                write!(self.out, " = ")?;
2068                self.write_expr(value, ctx)?;
2069                writeln!(self.out, ";")?
2070            }
2071            Statement::WorkGroupUniformLoad { pointer, result } => {
2072                // GLSL doesn't have pointers, which means that this backend needs to ensure that
2073                // the actual "loading" is happening between the two barriers.
2074                // This is done in `Emit` by never emitting a variable name for pointer variables
2075                self.write_barrier(crate::Barrier::WORK_GROUP, level)?;
2076
2077                let result_name = format!("{}{}", back::BAKE_PREFIX, result.index());
2078                write!(self.out, "{level}")?;
2079                // Expressions cannot have side effects, so just writing the expression here is fine.
2080                self.write_named_expr(pointer, result_name, result, ctx)?;
2081
2082                self.write_barrier(crate::Barrier::WORK_GROUP, level)?;
2083            }
2084            // Stores a value into an image.
2085            Statement::ImageStore {
2086                image,
2087                coordinate,
2088                array_index,
2089                value,
2090            } => {
2091                write!(self.out, "{level}")?;
2092                self.write_image_store(ctx, image, coordinate, array_index, value)?
2093            }
2094            // A `Call` is written `name(arguments)` where `arguments` is a comma separated expressions list
2095            Statement::Call {
2096                function,
2097                ref arguments,
2098                result,
2099            } => {
2100                write!(self.out, "{level}")?;
2101                if let Some(expr) = result {
2102                    let name = format!("{}{}", back::BAKE_PREFIX, expr.index());
2103                    let result = self.module.functions[function].result.as_ref().unwrap();
2104                    self.write_type(result.ty)?;
2105                    write!(self.out, " {name}")?;
2106                    if let TypeInner::Array { base, size, .. } = self.module.types[result.ty].inner
2107                    {
2108                        self.write_array_size(base, size)?
2109                    }
2110                    write!(self.out, " = ")?;
2111                    self.named_expressions.insert(expr, name);
2112                }
2113                write!(self.out, "{}(", &self.names[&NameKey::Function(function)])?;
2114                let arguments: Vec<_> = arguments
2115                    .iter()
2116                    .enumerate()
2117                    .filter_map(|(i, arg)| {
2118                        let arg_ty = self.module.functions[function].arguments[i].ty;
2119                        match self.module.types[arg_ty].inner {
2120                            TypeInner::Sampler { .. } => None,
2121                            _ => Some(*arg),
2122                        }
2123                    })
2124                    .collect();
2125                self.write_slice(&arguments, |this, _, arg| this.write_expr(*arg, ctx))?;
2126                writeln!(self.out, ");")?
2127            }
2128            Statement::Atomic {
2129                pointer,
2130                ref fun,
2131                value,
2132                result,
2133            } => {
2134                write!(self.out, "{level}")?;
2135                let res_name = format!("{}{}", back::BAKE_PREFIX, result.index());
2136                let res_ty = ctx.info[result].ty.inner_with(&self.module.types);
2137                self.write_value_type(res_ty)?;
2138                write!(self.out, " {res_name} = ")?;
2139                self.named_expressions.insert(result, res_name);
2140
2141                let fun_str = fun.to_glsl();
2142                write!(self.out, "atomic{fun_str}(")?;
2143                self.write_expr(pointer, ctx)?;
2144                write!(self.out, ", ")?;
2145                // handle the special cases
2146                match *fun {
2147                    crate::AtomicFunction::Subtract => {
2148                        // we just wrote `InterlockedAdd`, so negate the argument
2149                        write!(self.out, "-")?;
2150                    }
2151                    crate::AtomicFunction::Exchange { compare: Some(_) } => {
2152                        return Err(Error::Custom(
2153                            "atomic CompareExchange is not implemented".to_string(),
2154                        ));
2155                    }
2156                    _ => {}
2157                }
2158                self.write_expr(value, ctx)?;
2159                writeln!(self.out, ");")?;
2160            }
2161            Statement::RayQuery { .. } => unreachable!(),
2162        }
2163
2164        Ok(())
2165    }
2166
2167    /// Write a const expression.
2168    ///
2169    /// Write `expr`, a handle to an [`Expression`] in the current [`Module`]'s
2170    /// constant expression arena, as GLSL expression.
2171    ///
2172    /// # Notes
2173    /// Adds no newlines or leading/trailing whitespace
2174    ///
2175    /// [`Expression`]: crate::Expression
2176    /// [`Module`]: crate::Module
2177    fn write_const_expr(&mut self, expr: Handle<crate::Expression>) -> BackendResult {
2178        self.write_possibly_const_expr(expr, &self.module.const_expressions, |writer, expr| {
2179            writer.write_const_expr(expr)
2180        })
2181    }
2182
2183    /// Write [`Expression`] variants that can occur in both runtime and const expressions.
2184    ///
2185    /// Write `expr`, a handle to an [`Expression`] in the arena `expressions`,
2186    /// as as GLSL expression. This must be one of the [`Expression`] variants
2187    /// that is allowed to occur in constant expressions.
2188    ///
2189    /// Use `write_expression` to write subexpressions.
2190    ///
2191    /// This is the common code for `write_expr`, which handles arbitrary
2192    /// runtime expressions, and `write_const_expr`, which only handles
2193    /// const-expressions. Each of those callers passes itself (essentially) as
2194    /// the `write_expression` callback, so that subexpressions are restricted
2195    /// to the appropriate variants.
2196    ///
2197    /// # Notes
2198    /// Adds no newlines or leading/trailing whitespace
2199    ///
2200    /// [`Expression`]: crate::Expression
2201    fn write_possibly_const_expr<E>(
2202        &mut self,
2203        expr: Handle<crate::Expression>,
2204        expressions: &crate::Arena<crate::Expression>,
2205        write_expression: E,
2206    ) -> BackendResult
2207    where
2208        E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
2209    {
2210        use crate::Expression;
2211
2212        match expressions[expr] {
2213            Expression::Literal(literal) => {
2214                match literal {
2215                    // Floats are written using `Debug` instead of `Display` because it always appends the
2216                    // decimal part even it's zero which is needed for a valid glsl float constant
2217                    crate::Literal::F64(value) => write!(self.out, "{:?}LF", value)?,
2218                    crate::Literal::F32(value) => write!(self.out, "{:?}", value)?,
2219                    // Unsigned integers need a `u` at the end
2220                    //
2221                    // While `core` doesn't necessarily need it, it's allowed and since `es` needs it we
2222                    // always write it as the extra branch wouldn't have any benefit in readability
2223                    crate::Literal::U32(value) => write!(self.out, "{}u", value)?,
2224                    crate::Literal::I32(value) => write!(self.out, "{}", value)?,
2225                    crate::Literal::Bool(value) => write!(self.out, "{}", value)?,
2226                }
2227            }
2228            Expression::Constant(handle) => {
2229                let constant = &self.module.constants[handle];
2230                if constant.name.is_some() {
2231                    write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
2232                } else {
2233                    self.write_const_expr(constant.init)?;
2234                }
2235            }
2236            Expression::ZeroValue(ty) => {
2237                self.write_zero_init_value(ty)?;
2238            }
2239            Expression::Compose { ty, ref components } => {
2240                self.write_type(ty)?;
2241
2242                if let TypeInner::Array { base, size, .. } = self.module.types[ty].inner {
2243                    self.write_array_size(base, size)?;
2244                }
2245
2246                write!(self.out, "(")?;
2247                for (index, component) in components.iter().enumerate() {
2248                    if index != 0 {
2249                        write!(self.out, ", ")?;
2250                    }
2251                    write_expression(self, *component)?;
2252                }
2253                write!(self.out, ")")?
2254            }
2255            _ => unreachable!(),
2256        }
2257
2258        Ok(())
2259    }
2260
2261    /// Helper method to write expressions
2262    ///
2263    /// # Notes
2264    /// Doesn't add any newlines or leading/trailing spaces
2265    fn write_expr(
2266        &mut self,
2267        expr: Handle<crate::Expression>,
2268        ctx: &back::FunctionCtx<'_>,
2269    ) -> BackendResult {
2270        use crate::Expression;
2271
2272        if let Some(name) = self.named_expressions.get(&expr) {
2273            write!(self.out, "{name}")?;
2274            return Ok(());
2275        }
2276
2277        match ctx.expressions[expr] {
2278            Expression::Literal(_)
2279            | Expression::Constant(_)
2280            | Expression::ZeroValue(_)
2281            | Expression::Compose { .. } => {
2282                self.write_possibly_const_expr(expr, ctx.expressions, |writer, expr| {
2283                    writer.write_expr(expr, ctx)
2284                })?;
2285            }
2286            // `Access` is applied to arrays, vectors and matrices and is written as indexing
2287            Expression::Access { base, index } => {
2288                self.write_expr(base, ctx)?;
2289                write!(self.out, "[")?;
2290                self.write_expr(index, ctx)?;
2291                write!(self.out, "]")?
2292            }
2293            // `AccessIndex` is the same as `Access` except that the index is a constant and it can
2294            // be applied to structs, in this case we need to find the name of the field at that
2295            // index and write `base.field_name`
2296            Expression::AccessIndex { base, index } => {
2297                self.write_expr(base, ctx)?;
2298
2299                let base_ty_res = &ctx.info[base].ty;
2300                let mut resolved = base_ty_res.inner_with(&self.module.types);
2301                let base_ty_handle = match *resolved {
2302                    TypeInner::Pointer { base, space: _ } => {
2303                        resolved = &self.module.types[base].inner;
2304                        Some(base)
2305                    }
2306                    _ => base_ty_res.handle(),
2307                };
2308
2309                match *resolved {
2310                    TypeInner::Vector { .. } => {
2311                        // Write vector access as a swizzle
2312                        write!(self.out, ".{}", back::COMPONENTS[index as usize])?
2313                    }
2314                    TypeInner::Matrix { .. }
2315                    | TypeInner::Array { .. }
2316                    | TypeInner::ValuePointer { .. } => write!(self.out, "[{index}]")?,
2317                    TypeInner::Struct { .. } => {
2318                        // This will never panic in case the type is a `Struct`, this is not true
2319                        // for other types so we can only check while inside this match arm
2320                        let ty = base_ty_handle.unwrap();
2321
2322                        write!(
2323                            self.out,
2324                            ".{}",
2325                            &self.names[&NameKey::StructMember(ty, index)]
2326                        )?
2327                    }
2328                    ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))),
2329                }
2330            }
2331            // `Splat` needs to actually write down a vector, it's not always inferred in GLSL.
2332            Expression::Splat { size: _, value } => {
2333                let resolved = ctx.info[expr].ty.inner_with(&self.module.types);
2334                self.write_value_type(resolved)?;
2335                write!(self.out, "(")?;
2336                self.write_expr(value, ctx)?;
2337                write!(self.out, ")")?
2338            }
2339            // `Swizzle` adds a few letters behind the dot.
2340            Expression::Swizzle {
2341                size,
2342                vector,
2343                pattern,
2344            } => {
2345                self.write_expr(vector, ctx)?;
2346                write!(self.out, ".")?;
2347                for &sc in pattern[..size as usize].iter() {
2348                    self.out.write_char(back::COMPONENTS[sc as usize])?;
2349                }
2350            }
2351            // Function arguments are written as the argument name
2352            Expression::FunctionArgument(pos) => {
2353                write!(self.out, "{}", &self.names[&ctx.argument_key(pos)])?
2354            }
2355            // Global variables need some special work for their name but
2356            // `get_global_name` does the work for us
2357            Expression::GlobalVariable(handle) => {
2358                let global = &self.module.global_variables[handle];
2359                self.write_global_name(handle, global)?
2360            }
2361            // A local is written as it's name
2362            Expression::LocalVariable(handle) => {
2363                write!(self.out, "{}", self.names[&ctx.name_key(handle)])?
2364            }
2365            // glsl has no pointers so there's no load operation, just write the pointer expression
2366            Expression::Load { pointer } => self.write_expr(pointer, ctx)?,
2367            // `ImageSample` is a bit complicated compared to the rest of the IR.
2368            //
2369            // First there are three variations depending whether the sample level is explicitly set,
2370            // if it's automatic or it it's bias:
2371            // `texture(image, coordinate)` - Automatic sample level
2372            // `texture(image, coordinate, bias)` - Bias sample level
2373            // `textureLod(image, coordinate, level)` - Zero or Exact sample level
2374            //
2375            // Furthermore if `depth_ref` is some we need to append it to the coordinate vector
2376            Expression::ImageSample {
2377                image,
2378                sampler: _, //TODO?
2379                gather,
2380                coordinate,
2381                array_index,
2382                offset,
2383                level,
2384                depth_ref,
2385            } => {
2386                let dim = match *ctx.info[image].ty.inner_with(&self.module.types) {
2387                    TypeInner::Image { dim, .. } => dim,
2388                    _ => unreachable!(),
2389                };
2390
2391                if dim == crate::ImageDimension::Cube
2392                    && array_index.is_some()
2393                    && depth_ref.is_some()
2394                {
2395                    match level {
2396                        crate::SampleLevel::Zero
2397                        | crate::SampleLevel::Exact(_)
2398                        | crate::SampleLevel::Gradient { .. }
2399                        | crate::SampleLevel::Bias(_) => {
2400                            return Err(Error::Custom(String::from(
2401                                "gsamplerCubeArrayShadow isn't supported in textureGrad, \
2402                                 textureLod or texture with bias",
2403                            )))
2404                        }
2405                        crate::SampleLevel::Auto => {}
2406                    }
2407                }
2408
2409                // textureLod on sampler2DArrayShadow and samplerCubeShadow does not exist in GLSL.
2410                // To emulate this, we will have to use textureGrad with a constant gradient of 0.
2411                let workaround_lod_array_shadow_as_grad = (array_index.is_some()
2412                    || dim == crate::ImageDimension::Cube)
2413                    && depth_ref.is_some()
2414                    && gather.is_none()
2415                    && !self
2416                        .options
2417                        .writer_flags
2418                        .contains(WriterFlags::TEXTURE_SHADOW_LOD);
2419
2420                //Write the function to be used depending on the sample level
2421                let fun_name = match level {
2422                    crate::SampleLevel::Zero if gather.is_some() => "textureGather",
2423                    crate::SampleLevel::Auto | crate::SampleLevel::Bias(_) => "texture",
2424                    crate::SampleLevel::Zero | crate::SampleLevel::Exact(_) => {
2425                        if workaround_lod_array_shadow_as_grad {
2426                            "textureGrad"
2427                        } else {
2428                            "textureLod"
2429                        }
2430                    }
2431                    crate::SampleLevel::Gradient { .. } => "textureGrad",
2432                };
2433                let offset_name = match offset {
2434                    Some(_) => "Offset",
2435                    None => "",
2436                };
2437
2438                write!(self.out, "{fun_name}{offset_name}(")?;
2439
2440                // Write the image that will be used
2441                self.write_expr(image, ctx)?;
2442                // The space here isn't required but it helps with readability
2443                write!(self.out, ", ")?;
2444
2445                // We need to get the coordinates vector size to later build a vector that's `size + 1`
2446                // if `depth_ref` is some, if it isn't a vector we panic as that's not a valid expression
2447                let mut coord_dim = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
2448                    TypeInner::Vector { size, .. } => size as u8,
2449                    TypeInner::Scalar { .. } => 1,
2450                    _ => unreachable!(),
2451                };
2452
2453                if array_index.is_some() {
2454                    coord_dim += 1;
2455                }
2456                let merge_depth_ref = depth_ref.is_some() && gather.is_none() && coord_dim < 4;
2457                if merge_depth_ref {
2458                    coord_dim += 1;
2459                }
2460
2461                let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es();
2462                let is_vec = tex_1d_hack || coord_dim != 1;
2463                // Compose a new texture coordinates vector
2464                if is_vec {
2465                    write!(self.out, "vec{}(", coord_dim + tex_1d_hack as u8)?;
2466                }
2467                self.write_expr(coordinate, ctx)?;
2468                if tex_1d_hack {
2469                    write!(self.out, ", 0.0")?;
2470                }
2471                if let Some(expr) = array_index {
2472                    write!(self.out, ", ")?;
2473                    self.write_expr(expr, ctx)?;
2474                }
2475                if merge_depth_ref {
2476                    write!(self.out, ", ")?;
2477                    self.write_expr(depth_ref.unwrap(), ctx)?;
2478                }
2479                if is_vec {
2480                    write!(self.out, ")")?;
2481                }
2482
2483                if let (Some(expr), false) = (depth_ref, merge_depth_ref) {
2484                    write!(self.out, ", ")?;
2485                    self.write_expr(expr, ctx)?;
2486                }
2487
2488                match level {
2489                    // Auto needs no more arguments
2490                    crate::SampleLevel::Auto => (),
2491                    // Zero needs level set to 0
2492                    crate::SampleLevel::Zero => {
2493                        if workaround_lod_array_shadow_as_grad {
2494                            let vec_dim = match dim {
2495                                crate::ImageDimension::Cube => 3,
2496                                _ => 2,
2497                            };
2498                            write!(self.out, ", vec{vec_dim}(0.0), vec{vec_dim}(0.0)")?;
2499                        } else if gather.is_none() {
2500                            write!(self.out, ", 0.0")?;
2501                        }
2502                    }
2503                    // Exact and bias require another argument
2504                    crate::SampleLevel::Exact(expr) => {
2505                        if workaround_lod_array_shadow_as_grad {
2506                            log::warn!("Unable to `textureLod` a shadow array, ignoring the LOD");
2507                            write!(self.out, ", vec2(0,0), vec2(0,0)")?;
2508                        } else {
2509                            write!(self.out, ", ")?;
2510                            self.write_expr(expr, ctx)?;
2511                        }
2512                    }
2513                    crate::SampleLevel::Bias(_) => {
2514                        // This needs to be done after the offset writing
2515                    }
2516                    crate::SampleLevel::Gradient { x, y } => {
2517                        // If we are using sampler2D to replace sampler1D, we also
2518                        // need to make sure to use vec2 gradients
2519                        if tex_1d_hack {
2520                            write!(self.out, ", vec2(")?;
2521                            self.write_expr(x, ctx)?;
2522                            write!(self.out, ", 0.0)")?;
2523                            write!(self.out, ", vec2(")?;
2524                            self.write_expr(y, ctx)?;
2525                            write!(self.out, ", 0.0)")?;
2526                        } else {
2527                            write!(self.out, ", ")?;
2528                            self.write_expr(x, ctx)?;
2529                            write!(self.out, ", ")?;
2530                            self.write_expr(y, ctx)?;
2531                        }
2532                    }
2533                }
2534
2535                if let Some(constant) = offset {
2536                    write!(self.out, ", ")?;
2537                    if tex_1d_hack {
2538                        write!(self.out, "ivec2(")?;
2539                    }
2540                    self.write_const_expr(constant)?;
2541                    if tex_1d_hack {
2542                        write!(self.out, ", 0)")?;
2543                    }
2544                }
2545
2546                // Bias is always the last argument
2547                if let crate::SampleLevel::Bias(expr) = level {
2548                    write!(self.out, ", ")?;
2549                    self.write_expr(expr, ctx)?;
2550                }
2551
2552                if let (Some(component), None) = (gather, depth_ref) {
2553                    write!(self.out, ", {}", component as usize)?;
2554                }
2555
2556                // End the function
2557                write!(self.out, ")")?
2558            }
2559            Expression::ImageLoad {
2560                image,
2561                coordinate,
2562                array_index,
2563                sample,
2564                level,
2565            } => self.write_image_load(expr, ctx, image, coordinate, array_index, sample, level)?,
2566            // Query translates into one of the:
2567            // - textureSize/imageSize
2568            // - textureQueryLevels
2569            // - textureSamples/imageSamples
2570            Expression::ImageQuery { image, query } => {
2571                use crate::ImageClass;
2572
2573                // This will only panic if the module is invalid
2574                let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
2575                    TypeInner::Image {
2576                        dim,
2577                        arrayed: _,
2578                        class,
2579                    } => (dim, class),
2580                    _ => unreachable!(),
2581                };
2582                let components = match dim {
2583                    crate::ImageDimension::D1 => 1,
2584                    crate::ImageDimension::D2 => 2,
2585                    crate::ImageDimension::D3 => 3,
2586                    crate::ImageDimension::Cube => 2,
2587                };
2588
2589                if let crate::ImageQuery::Size { .. } = query {
2590                    match components {
2591                        1 => write!(self.out, "uint(")?,
2592                        _ => write!(self.out, "uvec{components}(")?,
2593                    }
2594                } else {
2595                    write!(self.out, "uint(")?;
2596                }
2597
2598                match query {
2599                    crate::ImageQuery::Size { level } => {
2600                        match class {
2601                            ImageClass::Sampled { multi, .. } | ImageClass::Depth { multi } => {
2602                                write!(self.out, "textureSize(")?;
2603                                self.write_expr(image, ctx)?;
2604                                if let Some(expr) = level {
2605                                    let cast_to_int = matches!(
2606                                        *ctx.info[expr].ty.inner_with(&self.module.types),
2607                                        crate::TypeInner::Scalar {
2608                                            kind: crate::ScalarKind::Uint,
2609                                            ..
2610                                        }
2611                                    );
2612
2613                                    write!(self.out, ", ")?;
2614
2615                                    if cast_to_int {
2616                                        write!(self.out, "int(")?;
2617                                    }
2618
2619                                    self.write_expr(expr, ctx)?;
2620
2621                                    if cast_to_int {
2622                                        write!(self.out, ")")?;
2623                                    }
2624                                } else if !multi {
2625                                    // All textureSize calls requires an lod argument
2626                                    // except for multisampled samplers
2627                                    write!(self.out, ", 0")?;
2628                                }
2629                            }
2630                            ImageClass::Storage { .. } => {
2631                                write!(self.out, "imageSize(")?;
2632                                self.write_expr(image, ctx)?;
2633                            }
2634                        }
2635                        write!(self.out, ")")?;
2636                        if components != 1 || self.options.version.is_es() {
2637                            write!(self.out, ".{}", &"xyz"[..components])?;
2638                        }
2639                    }
2640                    crate::ImageQuery::NumLevels => {
2641                        write!(self.out, "textureQueryLevels(",)?;
2642                        self.write_expr(image, ctx)?;
2643                        write!(self.out, ")",)?;
2644                    }
2645                    crate::ImageQuery::NumLayers => {
2646                        let fun_name = match class {
2647                            ImageClass::Sampled { .. } | ImageClass::Depth { .. } => "textureSize",
2648                            ImageClass::Storage { .. } => "imageSize",
2649                        };
2650                        write!(self.out, "{fun_name}(")?;
2651                        self.write_expr(image, ctx)?;
2652                        // All textureSize calls requires an lod argument
2653                        // except for multisampled samplers
2654                        if class.is_multisampled() {
2655                            write!(self.out, ", 0")?;
2656                        }
2657                        write!(self.out, ")")?;
2658                        if components != 1 || self.options.version.is_es() {
2659                            write!(self.out, ".{}", back::COMPONENTS[components])?;
2660                        }
2661                    }
2662                    crate::ImageQuery::NumSamples => {
2663                        let fun_name = match class {
2664                            ImageClass::Sampled { .. } | ImageClass::Depth { .. } => {
2665                                "textureSamples"
2666                            }
2667                            ImageClass::Storage { .. } => "imageSamples",
2668                        };
2669                        write!(self.out, "{fun_name}(")?;
2670                        self.write_expr(image, ctx)?;
2671                        write!(self.out, ")",)?;
2672                    }
2673                }
2674
2675                write!(self.out, ")")?;
2676            }
2677            // `Unary` is pretty straightforward
2678            // "-" - for `Negate`
2679            // "~" - for `Not` if it's an integer
2680            // "!" - for `Not` if it's a boolean
2681            //
2682            // We also wrap the everything in parentheses to avoid precedence issues
2683            Expression::Unary { op, expr } => {
2684                use crate::{ScalarKind as Sk, UnaryOperator as Uo};
2685
2686                let ty = ctx.info[expr].ty.inner_with(&self.module.types);
2687
2688                match *ty {
2689                    TypeInner::Vector { kind: Sk::Bool, .. } => {
2690                        write!(self.out, "not(")?;
2691                    }
2692                    _ => {
2693                        let operator = match op {
2694                            Uo::Negate => "-",
2695                            Uo::Not => match ty.scalar_kind() {
2696                                Some(Sk::Sint) | Some(Sk::Uint) => "~",
2697                                Some(Sk::Bool) => "!",
2698                                ref other => {
2699                                    return Err(Error::Custom(format!(
2700                                        "Cannot apply not to type {other:?}"
2701                                    )))
2702                                }
2703                            },
2704                        };
2705
2706                        write!(self.out, "{operator}(")?;
2707                    }
2708                }
2709
2710                self.write_expr(expr, ctx)?;
2711
2712                write!(self.out, ")")?
2713            }
2714            // `Binary` we just write `left op right`, except when dealing with
2715            // comparison operations on vectors as they are implemented with
2716            // builtin functions.
2717            // Once again we wrap everything in parentheses to avoid precedence issues
2718            Expression::Binary {
2719                mut op,
2720                left,
2721                right,
2722            } => {
2723                // Holds `Some(function_name)` if the binary operation is
2724                // implemented as a function call
2725                use crate::{BinaryOperator as Bo, ScalarKind as Sk, TypeInner as Ti};
2726
2727                let left_inner = ctx.info[left].ty.inner_with(&self.module.types);
2728                let right_inner = ctx.info[right].ty.inner_with(&self.module.types);
2729
2730                let function = match (left_inner, right_inner) {
2731                    (&Ti::Vector { kind, .. }, &Ti::Vector { .. }) => match op {
2732                        Bo::Less
2733                        | Bo::LessEqual
2734                        | Bo::Greater
2735                        | Bo::GreaterEqual
2736                        | Bo::Equal
2737                        | Bo::NotEqual => BinaryOperation::VectorCompare,
2738                        Bo::Modulo if kind == Sk::Float => BinaryOperation::Modulo,
2739                        Bo::And if kind == Sk::Bool => {
2740                            op = crate::BinaryOperator::LogicalAnd;
2741                            BinaryOperation::VectorComponentWise
2742                        }
2743                        Bo::InclusiveOr if kind == Sk::Bool => {
2744                            op = crate::BinaryOperator::LogicalOr;
2745                            BinaryOperation::VectorComponentWise
2746                        }
2747                        _ => BinaryOperation::Other,
2748                    },
2749                    _ => match (left_inner.scalar_kind(), right_inner.scalar_kind()) {
2750                        (Some(Sk::Float), _) | (_, Some(Sk::Float)) => match op {
2751                            Bo::Modulo => BinaryOperation::Modulo,
2752                            _ => BinaryOperation::Other,
2753                        },
2754                        (Some(Sk::Bool), Some(Sk::Bool)) => match op {
2755                            Bo::InclusiveOr => {
2756                                op = crate::BinaryOperator::LogicalOr;
2757                                BinaryOperation::Other
2758                            }
2759                            Bo::And => {
2760                                op = crate::BinaryOperator::LogicalAnd;
2761                                BinaryOperation::Other
2762                            }
2763                            _ => BinaryOperation::Other,
2764                        },
2765                        _ => BinaryOperation::Other,
2766                    },
2767                };
2768
2769                match function {
2770                    BinaryOperation::VectorCompare => {
2771                        let op_str = match op {
2772                            Bo::Less => "lessThan(",
2773                            Bo::LessEqual => "lessThanEqual(",
2774                            Bo::Greater => "greaterThan(",
2775                            Bo::GreaterEqual => "greaterThanEqual(",
2776                            Bo::Equal => "equal(",
2777                            Bo::NotEqual => "notEqual(",
2778                            _ => unreachable!(),
2779                        };
2780                        write!(self.out, "{op_str}")?;
2781                        self.write_expr(left, ctx)?;
2782                        write!(self.out, ", ")?;
2783                        self.write_expr(right, ctx)?;
2784                        write!(self.out, ")")?;
2785                    }
2786                    BinaryOperation::VectorComponentWise => {
2787                        self.write_value_type(left_inner)?;
2788                        write!(self.out, "(")?;
2789
2790                        let size = match *left_inner {
2791                            Ti::Vector { size, .. } => size,
2792                            _ => unreachable!(),
2793                        };
2794
2795                        for i in 0..size as usize {
2796                            if i != 0 {
2797                                write!(self.out, ", ")?;
2798                            }
2799
2800                            self.write_expr(left, ctx)?;
2801                            write!(self.out, ".{}", back::COMPONENTS[i])?;
2802
2803                            write!(self.out, " {} ", back::binary_operation_str(op))?;
2804
2805                            self.write_expr(right, ctx)?;
2806                            write!(self.out, ".{}", back::COMPONENTS[i])?;
2807                        }
2808
2809                        write!(self.out, ")")?;
2810                    }
2811                    // TODO: handle undefined behavior of BinaryOperator::Modulo
2812                    //
2813                    // sint:
2814                    // if right == 0 return 0
2815                    // if left == min(type_of(left)) && right == -1 return 0
2816                    // if sign(left) == -1 || sign(right) == -1 return result as defined by WGSL
2817                    //
2818                    // uint:
2819                    // if right == 0 return 0
2820                    //
2821                    // float:
2822                    // if right == 0 return ? see https://github.com/gpuweb/gpuweb/issues/2798
2823                    BinaryOperation::Modulo => {
2824                        write!(self.out, "(")?;
2825
2826                        // write `e1 - e2 * trunc(e1 / e2)`
2827                        self.write_expr(left, ctx)?;
2828                        write!(self.out, " - ")?;
2829                        self.write_expr(right, ctx)?;
2830                        write!(self.out, " * ")?;
2831                        write!(self.out, "trunc(")?;
2832                        self.write_expr(left, ctx)?;
2833                        write!(self.out, " / ")?;
2834                        self.write_expr(right, ctx)?;
2835                        write!(self.out, ")")?;
2836
2837                        write!(self.out, ")")?;
2838                    }
2839                    BinaryOperation::Other => {
2840                        write!(self.out, "(")?;
2841
2842                        self.write_expr(left, ctx)?;
2843                        write!(self.out, " {} ", back::binary_operation_str(op))?;
2844                        self.write_expr(right, ctx)?;
2845
2846                        write!(self.out, ")")?;
2847                    }
2848                }
2849            }
2850            // `Select` is written as `condition ? accept : reject`
2851            // We wrap everything in parentheses to avoid precedence issues
2852            Expression::Select {
2853                condition,
2854                accept,
2855                reject,
2856            } => {
2857                let cond_ty = ctx.info[condition].ty.inner_with(&self.module.types);
2858                let vec_select = if let TypeInner::Vector { .. } = *cond_ty {
2859                    true
2860                } else {
2861                    false
2862                };
2863
2864                // TODO: Boolean mix on desktop required GL_EXT_shader_integer_mix
2865                if vec_select {
2866                    // Glsl defines that for mix when the condition is a boolean the first element
2867                    // is picked if condition is false and the second if condition is true
2868                    write!(self.out, "mix(")?;
2869                    self.write_expr(reject, ctx)?;
2870                    write!(self.out, ", ")?;
2871                    self.write_expr(accept, ctx)?;
2872                    write!(self.out, ", ")?;
2873                    self.write_expr(condition, ctx)?;
2874                } else {
2875                    write!(self.out, "(")?;
2876                    self.write_expr(condition, ctx)?;
2877                    write!(self.out, " ? ")?;
2878                    self.write_expr(accept, ctx)?;
2879                    write!(self.out, " : ")?;
2880                    self.write_expr(reject, ctx)?;
2881                }
2882
2883                write!(self.out, ")")?
2884            }
2885            // `Derivative` is a function call to a glsl provided function
2886            Expression::Derivative { axis, ctrl, expr } => {
2887                use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl};
2888                let fun_name = if self.options.version.supports_derivative_control() {
2889                    match (axis, ctrl) {
2890                        (Axis::X, Ctrl::Coarse) => "dFdxCoarse",
2891                        (Axis::X, Ctrl::Fine) => "dFdxFine",
2892                        (Axis::X, Ctrl::None) => "dFdx",
2893                        (Axis::Y, Ctrl::Coarse) => "dFdyCoarse",
2894                        (Axis::Y, Ctrl::Fine) => "dFdyFine",
2895                        (Axis::Y, Ctrl::None) => "dFdy",
2896                        (Axis::Width, Ctrl::Coarse) => "fwidthCoarse",
2897                        (Axis::Width, Ctrl::Fine) => "fwidthFine",
2898                        (Axis::Width, Ctrl::None) => "fwidth",
2899                    }
2900                } else {
2901                    match axis {
2902                        Axis::X => "dFdx",
2903                        Axis::Y => "dFdy",
2904                        Axis::Width => "fwidth",
2905                    }
2906                };
2907                write!(self.out, "{fun_name}(")?;
2908                self.write_expr(expr, ctx)?;
2909                write!(self.out, ")")?
2910            }
2911            // `Relational` is a normal function call to some glsl provided functions
2912            Expression::Relational { fun, argument } => {
2913                use crate::RelationalFunction as Rf;
2914
2915                let fun_name = match fun {
2916                    // There's no specific function for this but we can invert the result of `isinf`
2917                    Rf::IsFinite => "!isinf",
2918                    Rf::IsInf => "isinf",
2919                    Rf::IsNan => "isnan",
2920                    // There's also no function for this but we can invert `isnan`
2921                    Rf::IsNormal => "!isnan",
2922                    Rf::All => "all",
2923                    Rf::Any => "any",
2924                };
2925                write!(self.out, "{fun_name}(")?;
2926
2927                self.write_expr(argument, ctx)?;
2928
2929                write!(self.out, ")")?
2930            }
2931            Expression::Math {
2932                fun,
2933                arg,
2934                arg1,
2935                arg2,
2936                arg3,
2937            } => {
2938                use crate::MathFunction as Mf;
2939
2940                let fun_name = match fun {
2941                    // comparison
2942                    Mf::Abs => "abs",
2943                    Mf::Min => "min",
2944                    Mf::Max => "max",
2945                    Mf::Clamp => "clamp",
2946                    Mf::Saturate => {
2947                        write!(self.out, "clamp(")?;
2948
2949                        self.write_expr(arg, ctx)?;
2950
2951                        match *ctx.info[arg].ty.inner_with(&self.module.types) {
2952                            crate::TypeInner::Vector { size, .. } => write!(
2953                                self.out,
2954                                ", vec{}(0.0), vec{0}(1.0)",
2955                                back::vector_size_str(size)
2956                            )?,
2957                            _ => write!(self.out, ", 0.0, 1.0")?,
2958                        }
2959
2960                        write!(self.out, ")")?;
2961
2962                        return Ok(());
2963                    }
2964                    // trigonometry
2965                    Mf::Cos => "cos",
2966                    Mf::Cosh => "cosh",
2967                    Mf::Sin => "sin",
2968                    Mf::Sinh => "sinh",
2969                    Mf::Tan => "tan",
2970                    Mf::Tanh => "tanh",
2971                    Mf::Acos => "acos",
2972                    Mf::Asin => "asin",
2973                    Mf::Atan => "atan",
2974                    Mf::Asinh => "asinh",
2975                    Mf::Acosh => "acosh",
2976                    Mf::Atanh => "atanh",
2977                    Mf::Radians => "radians",
2978                    Mf::Degrees => "degrees",
2979                    // glsl doesn't have atan2 function
2980                    // use two-argument variation of the atan function
2981                    Mf::Atan2 => "atan",
2982                    // decomposition
2983                    Mf::Ceil => "ceil",
2984                    Mf::Floor => "floor",
2985                    Mf::Round => "roundEven",
2986                    Mf::Fract => "fract",
2987                    Mf::Trunc => "trunc",
2988                    Mf::Modf => "modf",
2989                    Mf::Frexp => "frexp",
2990                    Mf::Ldexp => "ldexp",
2991                    // exponent
2992                    Mf::Exp => "exp",
2993                    Mf::Exp2 => "exp2",
2994                    Mf::Log => "log",
2995                    Mf::Log2 => "log2",
2996                    Mf::Pow => "pow",
2997                    // geometry
2998                    Mf::Dot => match *ctx.info[arg].ty.inner_with(&self.module.types) {
2999                        crate::TypeInner::Vector {
3000                            kind: crate::ScalarKind::Float,
3001                            ..
3002                        } => "dot",
3003                        crate::TypeInner::Vector { size, .. } => {
3004                            return self.write_dot_product(arg, arg1.unwrap(), size as usize, ctx)
3005                        }
3006                        _ => unreachable!(
3007                            "Correct TypeInner for dot product should be already validated"
3008                        ),
3009                    },
3010                    Mf::Outer => "outerProduct",
3011                    Mf::Cross => "cross",
3012                    Mf::Distance => "distance",
3013                    Mf::Length => "length",
3014                    Mf::Normalize => "normalize",
3015                    Mf::FaceForward => "faceforward",
3016                    Mf::Reflect => "reflect",
3017                    Mf::Refract => "refract",
3018                    // computational
3019                    Mf::Sign => "sign",
3020                    Mf::Fma => {
3021                        if self.options.version.supports_fma_function() {
3022                            // Use the fma function when available
3023                            "fma"
3024                        } else {
3025                            // No fma support. Transform the function call into an arithmetic expression
3026                            write!(self.out, "(")?;
3027
3028                            self.write_expr(arg, ctx)?;
3029                            write!(self.out, " * ")?;
3030
3031                            let arg1 =
3032                                arg1.ok_or_else(|| Error::Custom("Missing fma arg1".to_owned()))?;
3033                            self.write_expr(arg1, ctx)?;
3034                            write!(self.out, " + ")?;
3035
3036                            let arg2 =
3037                                arg2.ok_or_else(|| Error::Custom("Missing fma arg2".to_owned()))?;
3038                            self.write_expr(arg2, ctx)?;
3039                            write!(self.out, ")")?;
3040
3041                            return Ok(());
3042                        }
3043                    }
3044                    Mf::Mix => "mix",
3045                    Mf::Step => "step",
3046                    Mf::SmoothStep => "smoothstep",
3047                    Mf::Sqrt => "sqrt",
3048                    Mf::InverseSqrt => "inversesqrt",
3049                    Mf::Inverse => "inverse",
3050                    Mf::Transpose => "transpose",
3051                    Mf::Determinant => "determinant",
3052                    // bits
3053                    Mf::CountTrailingZeros => {
3054                        match *ctx.info[arg].ty.inner_with(&self.module.types) {
3055                            crate::TypeInner::Vector { size, kind, .. } => {
3056                                let s = back::vector_size_str(size);
3057                                if let crate::ScalarKind::Uint = kind {
3058                                    write!(self.out, "min(uvec{s}(findLSB(")?;
3059                                    self.write_expr(arg, ctx)?;
3060                                    write!(self.out, ")), uvec{s}(32u))")?;
3061                                } else {
3062                                    write!(self.out, "ivec{s}(min(uvec{s}(findLSB(")?;
3063                                    self.write_expr(arg, ctx)?;
3064                                    write!(self.out, ")), uvec{s}(32u)))")?;
3065                                }
3066                            }
3067                            crate::TypeInner::Scalar { kind, .. } => {
3068                                if let crate::ScalarKind::Uint = kind {
3069                                    write!(self.out, "min(uint(findLSB(")?;
3070                                    self.write_expr(arg, ctx)?;
3071                                    write!(self.out, ")), 32u)")?;
3072                                } else {
3073                                    write!(self.out, "int(min(uint(findLSB(")?;
3074                                    self.write_expr(arg, ctx)?;
3075                                    write!(self.out, ")), 32u))")?;
3076                                }
3077                            }
3078                            _ => unreachable!(),
3079                        };
3080                        return Ok(());
3081                    }
3082                    Mf::CountLeadingZeros => {
3083                        if self.options.version.supports_integer_functions() {
3084                            match *ctx.info[arg].ty.inner_with(&self.module.types) {
3085                                crate::TypeInner::Vector { size, kind, .. } => {
3086                                    let s = back::vector_size_str(size);
3087
3088                                    if let crate::ScalarKind::Uint = kind {
3089                                        write!(self.out, "uvec{s}(ivec{s}(31) - findMSB(")?;
3090                                        self.write_expr(arg, ctx)?;
3091                                        write!(self.out, "))")?;
3092                                    } else {
3093                                        write!(self.out, "mix(ivec{s}(31) - findMSB(")?;
3094                                        self.write_expr(arg, ctx)?;
3095                                        write!(self.out, "), ivec{s}(0), lessThan(")?;
3096                                        self.write_expr(arg, ctx)?;
3097                                        write!(self.out, ", ivec{s}(0)))")?;
3098                                    }
3099                                }
3100                                crate::TypeInner::Scalar { kind, .. } => {
3101                                    if let crate::ScalarKind::Uint = kind {
3102                                        write!(self.out, "uint(31 - findMSB(")?;
3103                                    } else {
3104                                        write!(self.out, "(")?;
3105                                        self.write_expr(arg, ctx)?;
3106                                        write!(self.out, " < 0 ? 0 : 31 - findMSB(")?;
3107                                    }
3108
3109                                    self.write_expr(arg, ctx)?;
3110                                    write!(self.out, "))")?;
3111                                }
3112                                _ => unreachable!(),
3113                            };
3114                        } else {
3115                            match *ctx.info[arg].ty.inner_with(&self.module.types) {
3116                                crate::TypeInner::Vector { size, kind, .. } => {
3117                                    let s = back::vector_size_str(size);
3118
3119                                    if let crate::ScalarKind::Uint = kind {
3120                                        write!(self.out, "uvec{s}(")?;
3121                                        write!(self.out, "vec{s}(31.0) - floor(log2(vec{s}(")?;
3122                                        self.write_expr(arg, ctx)?;
3123                                        write!(self.out, ") + 0.5)))")?;
3124                                    } else {
3125                                        write!(self.out, "ivec{s}(")?;
3126                                        write!(self.out, "mix(vec{s}(31.0) - floor(log2(vec{s}(")?;
3127                                        self.write_expr(arg, ctx)?;
3128                                        write!(self.out, ") + 0.5)), ")?;
3129                                        write!(self.out, "vec{s}(0.0), lessThan(")?;
3130                                        self.write_expr(arg, ctx)?;
3131                                        write!(self.out, ", ivec{s}(0u))))")?;
3132                                    }
3133                                }
3134                                crate::TypeInner::Scalar { kind, .. } => {
3135                                    if let crate::ScalarKind::Uint = kind {
3136                                        write!(self.out, "uint(31.0 - floor(log2(float(")?;
3137                                        self.write_expr(arg, ctx)?;
3138                                        write!(self.out, ") + 0.5)))")?;
3139                                    } else {
3140                                        write!(self.out, "(")?;
3141                                        self.write_expr(arg, ctx)?;
3142                                        write!(self.out, " < 0 ? 0 : int(")?;
3143                                        write!(self.out, "31.0 - floor(log2(float(")?;
3144                                        self.write_expr(arg, ctx)?;
3145                                        write!(self.out, ") + 0.5))))")?;
3146                                    }
3147                                }
3148                                _ => unreachable!(),
3149                            };
3150                        }
3151
3152                        return Ok(());
3153                    }
3154                    Mf::CountOneBits => "bitCount",
3155                    Mf::ReverseBits => "bitfieldReverse",
3156                    Mf::ExtractBits => "bitfieldExtract",
3157                    Mf::InsertBits => "bitfieldInsert",
3158                    Mf::FindLsb => "findLSB",
3159                    Mf::FindMsb => "findMSB",
3160                    // data packing
3161                    Mf::Pack4x8snorm => "packSnorm4x8",
3162                    Mf::Pack4x8unorm => "packUnorm4x8",
3163                    Mf::Pack2x16snorm => "packSnorm2x16",
3164                    Mf::Pack2x16unorm => "packUnorm2x16",
3165                    Mf::Pack2x16float => "packHalf2x16",
3166                    // data unpacking
3167                    Mf::Unpack4x8snorm => "unpackSnorm4x8",
3168                    Mf::Unpack4x8unorm => "unpackUnorm4x8",
3169                    Mf::Unpack2x16snorm => "unpackSnorm2x16",
3170                    Mf::Unpack2x16unorm => "unpackUnorm2x16",
3171                    Mf::Unpack2x16float => "unpackHalf2x16",
3172                };
3173
3174                let extract_bits = fun == Mf::ExtractBits;
3175                let insert_bits = fun == Mf::InsertBits;
3176
3177                // Some GLSL functions always return signed integers (like findMSB),
3178                // so they need to be cast to uint if the argument is also an uint.
3179                let ret_might_need_int_to_uint =
3180                    matches!(fun, Mf::FindLsb | Mf::FindMsb | Mf::CountOneBits | Mf::Abs);
3181
3182                // Some GLSL functions only accept signed integers (like abs),
3183                // so they need their argument cast from uint to int.
3184                let arg_might_need_uint_to_int = matches!(fun, Mf::Abs);
3185
3186                // Check if the argument is an unsigned integer and return the vector size
3187                // in case it's a vector
3188                let maybe_uint_size = match *ctx.info[arg].ty.inner_with(&self.module.types) {
3189                    crate::TypeInner::Scalar {
3190                        kind: crate::ScalarKind::Uint,
3191                        ..
3192                    } => Some(None),
3193                    crate::TypeInner::Vector {
3194                        kind: crate::ScalarKind::Uint,
3195                        size,
3196                        ..
3197                    } => Some(Some(size)),
3198                    _ => None,
3199                };
3200
3201                // Cast to uint if the function needs it
3202                if ret_might_need_int_to_uint {
3203                    if let Some(maybe_size) = maybe_uint_size {
3204                        match maybe_size {
3205                            Some(size) => write!(self.out, "uvec{}(", size as u8)?,
3206                            None => write!(self.out, "uint(")?,
3207                        }
3208                    }
3209                }
3210
3211                write!(self.out, "{fun_name}(")?;
3212
3213                // Cast to int if the function needs it
3214                if arg_might_need_uint_to_int {
3215                    if let Some(maybe_size) = maybe_uint_size {
3216                        match maybe_size {
3217                            Some(size) => write!(self.out, "ivec{}(", size as u8)?,
3218                            None => write!(self.out, "int(")?,
3219                        }
3220                    }
3221                }
3222
3223                self.write_expr(arg, ctx)?;
3224
3225                // Close the cast from uint to int
3226                if arg_might_need_uint_to_int && maybe_uint_size.is_some() {
3227                    write!(self.out, ")")?
3228                }
3229
3230                if let Some(arg) = arg1 {
3231                    write!(self.out, ", ")?;
3232                    if extract_bits {
3233                        write!(self.out, "int(")?;
3234                        self.write_expr(arg, ctx)?;
3235                        write!(self.out, ")")?;
3236                    } else {
3237                        self.write_expr(arg, ctx)?;
3238                    }
3239                }
3240                if let Some(arg) = arg2 {
3241                    write!(self.out, ", ")?;
3242                    if extract_bits || insert_bits {
3243                        write!(self.out, "int(")?;
3244                        self.write_expr(arg, ctx)?;
3245                        write!(self.out, ")")?;
3246                    } else {
3247                        self.write_expr(arg, ctx)?;
3248                    }
3249                }
3250                if let Some(arg) = arg3 {
3251                    write!(self.out, ", ")?;
3252                    if insert_bits {
3253                        write!(self.out, "int(")?;
3254                        self.write_expr(arg, ctx)?;
3255                        write!(self.out, ")")?;
3256                    } else {
3257                        self.write_expr(arg, ctx)?;
3258                    }
3259                }
3260                write!(self.out, ")")?;
3261
3262                // Close the cast from int to uint
3263                if ret_might_need_int_to_uint && maybe_uint_size.is_some() {
3264                    write!(self.out, ")")?
3265                }
3266            }
3267            // `As` is always a call.
3268            // If `convert` is true the function name is the type
3269            // Else the function name is one of the glsl provided bitcast functions
3270            Expression::As {
3271                expr,
3272                kind: target_kind,
3273                convert,
3274            } => {
3275                let inner = ctx.info[expr].ty.inner_with(&self.module.types);
3276                match convert {
3277                    Some(width) => {
3278                        // this is similar to `write_type`, but with the target kind
3279                        let scalar = glsl_scalar(target_kind, width)?;
3280                        match *inner {
3281                            TypeInner::Matrix { columns, rows, .. } => write!(
3282                                self.out,
3283                                "{}mat{}x{}",
3284                                scalar.prefix, columns as u8, rows as u8
3285                            )?,
3286                            TypeInner::Vector { size, .. } => {
3287                                write!(self.out, "{}vec{}", scalar.prefix, size as u8)?
3288                            }
3289                            _ => write!(self.out, "{}", scalar.full)?,
3290                        }
3291
3292                        write!(self.out, "(")?;
3293                        self.write_expr(expr, ctx)?;
3294                        write!(self.out, ")")?
3295                    }
3296                    None => {
3297                        use crate::ScalarKind as Sk;
3298
3299                        let target_vector_type = match *inner {
3300                            TypeInner::Vector { size, width, .. } => Some(TypeInner::Vector {
3301                                size,
3302                                width,
3303                                kind: target_kind,
3304                            }),
3305                            _ => None,
3306                        };
3307
3308                        let source_kind = inner.scalar_kind().unwrap();
3309
3310                        match (source_kind, target_kind, target_vector_type) {
3311                            // No conversion needed
3312                            (Sk::Sint, Sk::Sint, _)
3313                            | (Sk::Uint, Sk::Uint, _)
3314                            | (Sk::Float, Sk::Float, _)
3315                            | (Sk::Bool, Sk::Bool, _) => {
3316                                self.write_expr(expr, ctx)?;
3317                                return Ok(());
3318                            }
3319
3320                            // Cast to/from floats
3321                            (Sk::Float, Sk::Sint, _) => write!(self.out, "floatBitsToInt")?,
3322                            (Sk::Float, Sk::Uint, _) => write!(self.out, "floatBitsToUint")?,
3323                            (Sk::Sint, Sk::Float, _) => write!(self.out, "intBitsToFloat")?,
3324                            (Sk::Uint, Sk::Float, _) => write!(self.out, "uintBitsToFloat")?,
3325
3326                            // Cast between vector types
3327                            (_, _, Some(vector)) => {
3328                                self.write_value_type(&vector)?;
3329                            }
3330
3331                            // There is no way to bitcast between Uint/Sint in glsl. Use constructor conversion
3332                            (Sk::Uint | Sk::Bool, Sk::Sint, None) => write!(self.out, "int")?,
3333                            (Sk::Sint | Sk::Bool, Sk::Uint, None) => write!(self.out, "uint")?,
3334                            (Sk::Bool, Sk::Float, None) => write!(self.out, "float")?,
3335                            (Sk::Sint | Sk::Uint | Sk::Float, Sk::Bool, None) => {
3336                                write!(self.out, "bool")?
3337                            }
3338                        };
3339
3340                        write!(self.out, "(")?;
3341                        self.write_expr(expr, ctx)?;
3342                        write!(self.out, ")")?;
3343                    }
3344                }
3345            }
3346            // These expressions never show up in `Emit`.
3347            Expression::CallResult(_)
3348            | Expression::AtomicResult { .. }
3349            | Expression::RayQueryProceedResult
3350            | Expression::WorkGroupUniformLoadResult { .. } => unreachable!(),
3351            // `ArrayLength` is written as `expr.length()` and we convert it to a uint
3352            Expression::ArrayLength(expr) => {
3353                write!(self.out, "uint(")?;
3354                self.write_expr(expr, ctx)?;
3355                write!(self.out, ".length())")?
3356            }
3357            // not supported yet
3358            Expression::RayQueryGetIntersection { .. } => unreachable!(),
3359        }
3360
3361        Ok(())
3362    }
3363
3364    /// Helper function to write the local holding the clamped lod
3365    fn write_clamped_lod(
3366        &mut self,
3367        ctx: &back::FunctionCtx,
3368        expr: Handle<crate::Expression>,
3369        image: Handle<crate::Expression>,
3370        level_expr: Handle<crate::Expression>,
3371    ) -> Result<(), Error> {
3372        // Define our local and start a call to `clamp`
3373        write!(
3374            self.out,
3375            "int {}{}{} = clamp(",
3376            back::BAKE_PREFIX,
3377            expr.index(),
3378            CLAMPED_LOD_SUFFIX
3379        )?;
3380        // Write the lod that will be clamped
3381        self.write_expr(level_expr, ctx)?;
3382        // Set the min value to 0 and start a call to `textureQueryLevels` to get
3383        // the maximum value
3384        write!(self.out, ", 0, textureQueryLevels(")?;
3385        // Write the target image as an argument to `textureQueryLevels`
3386        self.write_expr(image, ctx)?;
3387        // Close the call to `textureQueryLevels` subtract 1 from it since
3388        // the lod argument is 0 based, close the `clamp` call and end the
3389        // local declaration statement.
3390        writeln!(self.out, ") - 1);")?;
3391
3392        Ok(())
3393    }
3394
3395    // Helper method used to retrieve how many elements a coordinate vector
3396    // for the images operations need.
3397    fn get_coordinate_vector_size(&self, dim: crate::ImageDimension, arrayed: bool) -> u8 {
3398        // openGL es doesn't have 1D images so we need workaround it
3399        let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es();
3400        // Get how many components the coordinate vector needs for the dimensions only
3401        let tex_coord_size = match dim {
3402            crate::ImageDimension::D1 => 1,
3403            crate::ImageDimension::D2 => 2,
3404            crate::ImageDimension::D3 => 3,
3405            crate::ImageDimension::Cube => 2,
3406        };
3407        // Calculate the true size of the coordinate vector by adding 1 for arrayed images
3408        // and another 1 if we need to workaround 1D images by making them 2D
3409        tex_coord_size + tex_1d_hack as u8 + arrayed as u8
3410    }
3411
3412    /// Helper method to write the coordinate vector for image operations
3413    fn write_texture_coord(
3414        &mut self,
3415        ctx: &back::FunctionCtx,
3416        vector_size: u8,
3417        coordinate: Handle<crate::Expression>,
3418        array_index: Option<Handle<crate::Expression>>,
3419        // Emulate 1D images as 2D for profiles that don't support it (glsl es)
3420        tex_1d_hack: bool,
3421    ) -> Result<(), Error> {
3422        match array_index {
3423            // If the image needs an array indice we need to add it to the end of our
3424            // coordinate vector, to do so we will use the `ivec(ivec, scalar)`
3425            // constructor notation (NOTE: the inner `ivec` can also be a scalar, this
3426            // is important for 1D arrayed images).
3427            Some(layer_expr) => {
3428                write!(self.out, "ivec{vector_size}(")?;
3429                self.write_expr(coordinate, ctx)?;
3430                write!(self.out, ", ")?;
3431                // If we are replacing sampler1D with sampler2D we also need
3432                // to add another zero to the coordinates vector for the y component
3433                if tex_1d_hack {
3434                    write!(self.out, "0, ")?;
3435                }
3436                self.write_expr(layer_expr, ctx)?;
3437                write!(self.out, ")")?;
3438            }
3439            // Otherwise write just the expression (and the 1D hack if needed)
3440            None => {
3441                let uvec_size = match *ctx.info[coordinate].ty.inner_with(&self.module.types) {
3442                    TypeInner::Scalar {
3443                        kind: crate::ScalarKind::Uint,
3444                        ..
3445                    } => Some(None),
3446                    TypeInner::Vector {
3447                        size,
3448                        kind: crate::ScalarKind::Uint,
3449                        ..
3450                    } => Some(Some(size as u32)),
3451                    _ => None,
3452                };
3453                if tex_1d_hack {
3454                    write!(self.out, "ivec2(")?;
3455                } else if uvec_size.is_some() {
3456                    match uvec_size {
3457                        Some(None) => write!(self.out, "int(")?,
3458                        Some(Some(size)) => write!(self.out, "ivec{size}(")?,
3459                        _ => {}
3460                    }
3461                }
3462                self.write_expr(coordinate, ctx)?;
3463                if tex_1d_hack {
3464                    write!(self.out, ", 0)")?;
3465                } else if uvec_size.is_some() {
3466                    write!(self.out, ")")?;
3467                }
3468            }
3469        }
3470
3471        Ok(())
3472    }
3473
3474    /// Helper method to write the `ImageStore` statement
3475    fn write_image_store(
3476        &mut self,
3477        ctx: &back::FunctionCtx,
3478        image: Handle<crate::Expression>,
3479        coordinate: Handle<crate::Expression>,
3480        array_index: Option<Handle<crate::Expression>>,
3481        value: Handle<crate::Expression>,
3482    ) -> Result<(), Error> {
3483        use crate::ImageDimension as IDim;
3484
3485        // NOTE: openGL requires that `imageStore`s have no effets when the texel is invalid
3486        // so we don't need to generate bounds checks (OpenGL 4.2 Core §3.9.20)
3487
3488        // This will only panic if the module is invalid
3489        let dim = match *ctx.info[image].ty.inner_with(&self.module.types) {
3490            TypeInner::Image { dim, .. } => dim,
3491            _ => unreachable!(),
3492        };
3493
3494        // Begin our call to `imageStore`
3495        write!(self.out, "imageStore(")?;
3496        self.write_expr(image, ctx)?;
3497        // Separate the image argument from the coordinates
3498        write!(self.out, ", ")?;
3499
3500        // openGL es doesn't have 1D images so we need workaround it
3501        let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
3502        // Write the coordinate vector
3503        self.write_texture_coord(
3504            ctx,
3505            // Get the size of the coordinate vector
3506            self.get_coordinate_vector_size(dim, array_index.is_some()),
3507            coordinate,
3508            array_index,
3509            tex_1d_hack,
3510        )?;
3511
3512        // Separate the coordinate from the value to write and write the expression
3513        // of the value to write.
3514        write!(self.out, ", ")?;
3515        self.write_expr(value, ctx)?;
3516        // End the call to `imageStore` and the statement.
3517        writeln!(self.out, ");")?;
3518
3519        Ok(())
3520    }
3521
3522    /// Helper method for writing an `ImageLoad` expression.
3523    #[allow(clippy::too_many_arguments)]
3524    fn write_image_load(
3525        &mut self,
3526        handle: Handle<crate::Expression>,
3527        ctx: &back::FunctionCtx,
3528        image: Handle<crate::Expression>,
3529        coordinate: Handle<crate::Expression>,
3530        array_index: Option<Handle<crate::Expression>>,
3531        sample: Option<Handle<crate::Expression>>,
3532        level: Option<Handle<crate::Expression>>,
3533    ) -> Result<(), Error> {
3534        use crate::ImageDimension as IDim;
3535
3536        // `ImageLoad` is a bit complicated.
3537        // There are two functions one for sampled
3538        // images another for storage images, the former uses `texelFetch` and the
3539        // latter uses `imageLoad`.
3540        //
3541        // Furthermore we have `level` which is always `Some` for sampled images
3542        // and `None` for storage images, so we end up with two functions:
3543        // - `texelFetch(image, coordinate, level)` for sampled images
3544        // - `imageLoad(image, coordinate)` for storage images
3545        //
3546        // Finally we also have to consider bounds checking, for storage images
3547        // this is easy since openGL requires that invalid texels always return
3548        // 0, for sampled images we need to either verify that all arguments are
3549        // in bounds (`ReadZeroSkipWrite`) or make them a valid texel (`Restrict`).
3550
3551        // This will only panic if the module is invalid
3552        let (dim, class) = match *ctx.info[image].ty.inner_with(&self.module.types) {
3553            TypeInner::Image {
3554                dim,
3555                arrayed: _,
3556                class,
3557            } => (dim, class),
3558            _ => unreachable!(),
3559        };
3560
3561        // Get the name of the function to be used for the load operation
3562        // and the policy to be used with it.
3563        let (fun_name, policy) = match class {
3564            // Sampled images inherit the policy from the user passed policies
3565            crate::ImageClass::Sampled { .. } => ("texelFetch", self.policies.image_load),
3566            crate::ImageClass::Storage { .. } => {
3567                // OpenGL ES 3.1 mentiones in Chapter "8.22 Texture Image Loads and Stores" that:
3568                // "Invalid image loads will return a vector where the value of R, G, and B components
3569                // is 0 and the value of the A component is undefined."
3570                //
3571                // OpenGL 4.2 Core mentiones in Chapter "3.9.20 Texture Image Loads and Stores" that:
3572                // "Invalid image loads will return zero."
3573                //
3574                // So, we only inject bounds checks for ES
3575                let policy = if self.options.version.is_es() {
3576                    self.policies.image_load
3577                } else {
3578                    proc::BoundsCheckPolicy::Unchecked
3579                };
3580                ("imageLoad", policy)
3581            }
3582            // TODO: Is there even a function for this?
3583            crate::ImageClass::Depth { multi: _ } => {
3584                return Err(Error::Custom(
3585                    "WGSL `textureLoad` from depth textures is not supported in GLSL".to_string(),
3586                ))
3587            }
3588        };
3589
3590        // openGL es doesn't have 1D images so we need workaround it
3591        let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
3592        // Get the size of the coordinate vector
3593        let vector_size = self.get_coordinate_vector_size(dim, array_index.is_some());
3594
3595        if let proc::BoundsCheckPolicy::ReadZeroSkipWrite = policy {
3596            // To write the bounds checks for `ReadZeroSkipWrite` we will use a
3597            // ternary operator since we are in the middle of an expression and
3598            // need to return a value.
3599            //
3600            // NOTE: glsl does short circuit when evaluating logical
3601            // expressions so we can be sure that after we test a
3602            // condition it will be true for the next ones
3603
3604            // Write parantheses around the ternary operator to prevent problems with
3605            // expressions emitted before or after it having more precedence
3606            write!(self.out, "(",)?;
3607
3608            // The lod check needs to precede the size check since we need
3609            // to use the lod to get the size of the image at that level.
3610            if let Some(level_expr) = level {
3611                self.write_expr(level_expr, ctx)?;
3612                write!(self.out, " < textureQueryLevels(",)?;
3613                self.write_expr(image, ctx)?;
3614                // Chain the next check
3615                write!(self.out, ") && ")?;
3616            }
3617
3618            // Check that the sample arguments doesn't exceed the number of samples
3619            if let Some(sample_expr) = sample {
3620                self.write_expr(sample_expr, ctx)?;
3621                write!(self.out, " < textureSamples(",)?;
3622                self.write_expr(image, ctx)?;
3623                // Chain the next check
3624                write!(self.out, ") && ")?;
3625            }
3626
3627            // We now need to write the size checks for the coordinates and array index
3628            // first we write the comparation function in case the image is 1D non arrayed
3629            // (and no 1D to 2D hack was needed) we are comparing scalars so the less than
3630            // operator will suffice, but otherwise we'll be comparing two vectors so we'll
3631            // need to use the `lessThan` function but it returns a vector of booleans (one
3632            // for each comparison) so we need to fold it all in one scalar boolean, since
3633            // we want all comparisons to pass we use the `all` function which will only
3634            // return `true` if all the elements of the boolean vector are also `true`.
3635            //
3636            // So we'll end with one of the following forms
3637            // - `coord < textureSize(image, lod)` for 1D images
3638            // - `all(lessThan(coord, textureSize(image, lod)))` for normal images
3639            // - `all(lessThan(ivec(coord, array_index), textureSize(image, lod)))`
3640            //    for arrayed images
3641            // - `all(lessThan(coord, textureSize(image)))` for multi sampled images
3642
3643            if vector_size != 1 {
3644                write!(self.out, "all(lessThan(")?;
3645            }
3646
3647            // Write the coordinate vector
3648            self.write_texture_coord(ctx, vector_size, coordinate, array_index, tex_1d_hack)?;
3649
3650            if vector_size != 1 {
3651                // If we used the `lessThan` function we need to separate the
3652                // coordinates from the image size.
3653                write!(self.out, ", ")?;
3654            } else {
3655                // If we didn't use it (ie. 1D images) we perform the comparsion
3656                // using the less than operator.
3657                write!(self.out, " < ")?;
3658            }
3659
3660            // Call `textureSize` to get our image size
3661            write!(self.out, "textureSize(")?;
3662            self.write_expr(image, ctx)?;
3663            // `textureSize` uses the lod as a second argument for mipmapped images
3664            if let Some(level_expr) = level {
3665                // Separate the image from the lod
3666                write!(self.out, ", ")?;
3667                self.write_expr(level_expr, ctx)?;
3668            }
3669            // Close the `textureSize` call
3670            write!(self.out, ")")?;
3671
3672            if vector_size != 1 {
3673                // Close the `all` and `lessThan` calls
3674                write!(self.out, "))")?;
3675            }
3676
3677            // Finally end the condition part of the ternary operator
3678            write!(self.out, " ? ")?;
3679        }
3680
3681        // Begin the call to the function used to load the texel
3682        write!(self.out, "{fun_name}(")?;
3683        self.write_expr(image, ctx)?;
3684        write!(self.out, ", ")?;
3685
3686        // If we are using `Restrict` bounds checking we need to pass valid texel
3687        // coordinates, to do so we use the `clamp` function to get a value between
3688        // 0 and the image size - 1 (indexing begins at 0)
3689        if let proc::BoundsCheckPolicy::Restrict = policy {
3690            write!(self.out, "clamp(")?;
3691        }
3692
3693        // Write the coordinate vector
3694        self.write_texture_coord(ctx, vector_size, coordinate, array_index, tex_1d_hack)?;
3695
3696        // If we are using `Restrict` bounds checking we need to write the rest of the
3697        // clamp we initiated before writing the coordinates.
3698        if let proc::BoundsCheckPolicy::Restrict = policy {
3699            // Write the min value 0
3700            if vector_size == 1 {
3701                write!(self.out, ", 0")?;
3702            } else {
3703                write!(self.out, ", ivec{vector_size}(0)")?;
3704            }
3705            // Start the `textureSize` call to use as the max value.
3706            write!(self.out, ", textureSize(")?;
3707            self.write_expr(image, ctx)?;
3708            // If the image is mipmapped we need to add the lod argument to the
3709            // `textureSize` call, but this needs to be the clamped lod, this should
3710            // have been generated earlier and put in a local.
3711            if class.is_mipmapped() {
3712                write!(
3713                    self.out,
3714                    ", {}{}{}",
3715                    back::BAKE_PREFIX,
3716                    handle.index(),
3717                    CLAMPED_LOD_SUFFIX
3718                )?;
3719            }
3720            // Close the `textureSize` call
3721            write!(self.out, ")")?;
3722
3723            // Subtract 1 from the `textureSize` call since the coordinates are zero based.
3724            if vector_size == 1 {
3725                write!(self.out, " - 1")?;
3726            } else {
3727                write!(self.out, " - ivec{vector_size}(1)")?;
3728            }
3729
3730            // Close the `clamp` call
3731            write!(self.out, ")")?;
3732
3733            // Add the clamped lod (if present) as the second argument to the
3734            // image load function.
3735            if level.is_some() {
3736                write!(
3737                    self.out,
3738                    ", {}{}{}",
3739                    back::BAKE_PREFIX,
3740                    handle.index(),
3741                    CLAMPED_LOD_SUFFIX
3742                )?;
3743            }
3744
3745            // If a sample argument is needed we need to clamp it between 0 and
3746            // the number of samples the image has.
3747            if let Some(sample_expr) = sample {
3748                write!(self.out, ", clamp(")?;
3749                self.write_expr(sample_expr, ctx)?;
3750                // Set the min value to 0 and start the call to `textureSamples`
3751                write!(self.out, ", 0, textureSamples(")?;
3752                self.write_expr(image, ctx)?;
3753                // Close the `textureSamples` call, subtract 1 from it since the sample
3754                // argument is zero based, and close the `clamp` call
3755                writeln!(self.out, ") - 1)")?;
3756            }
3757        } else if let Some(sample_or_level) = sample.or(level) {
3758            // If no bounds checking is need just add the sample or level argument
3759            // after the coordinates
3760            write!(self.out, ", ")?;
3761            self.write_expr(sample_or_level, ctx)?;
3762        }
3763
3764        // Close the image load function.
3765        write!(self.out, ")")?;
3766
3767        // If we were using the `ReadZeroSkipWrite` policy we need to end the first branch
3768        // (which is taken if the condition is `true`) with a colon (`:`) and write the
3769        // second branch which is just a 0 value.
3770        if let proc::BoundsCheckPolicy::ReadZeroSkipWrite = policy {
3771            // Get the kind of the output value.
3772            let kind = match class {
3773                // Only sampled images can reach here since storage images
3774                // don't need bounds checks and depth images aren't implmented
3775                crate::ImageClass::Sampled { kind, .. } => kind,
3776                _ => unreachable!(),
3777            };
3778
3779            // End the first branch
3780            write!(self.out, " : ")?;
3781            // Write the 0 value
3782            write!(self.out, "{}vec4(", glsl_scalar(kind, 4)?.prefix,)?;
3783            self.write_zero_init_scalar(kind)?;
3784            // Close the zero value constructor
3785            write!(self.out, ")")?;
3786            // Close the parantheses surrounding our ternary
3787            write!(self.out, ")")?;
3788        }
3789
3790        Ok(())
3791    }
3792
3793    fn write_named_expr(
3794        &mut self,
3795        handle: Handle<crate::Expression>,
3796        name: String,
3797        // The expression which is being named.
3798        // Generally, this is the same as handle, except in WorkGroupUniformLoad
3799        named: Handle<crate::Expression>,
3800        ctx: &back::FunctionCtx,
3801    ) -> BackendResult {
3802        match ctx.info[named].ty {
3803            proc::TypeResolution::Handle(ty_handle) => match self.module.types[ty_handle].inner {
3804                TypeInner::Struct { .. } => {
3805                    let ty_name = &self.names[&NameKey::Type(ty_handle)];
3806                    write!(self.out, "{ty_name}")?;
3807                }
3808                _ => {
3809                    self.write_type(ty_handle)?;
3810                }
3811            },
3812            proc::TypeResolution::Value(ref inner) => {
3813                self.write_value_type(inner)?;
3814            }
3815        }
3816
3817        let base_ty_res = &ctx.info[named].ty;
3818        let resolved = base_ty_res.inner_with(&self.module.types);
3819
3820        write!(self.out, " {name}")?;
3821        if let TypeInner::Array { base, size, .. } = *resolved {
3822            self.write_array_size(base, size)?;
3823        }
3824        write!(self.out, " = ")?;
3825        self.write_expr(handle, ctx)?;
3826        writeln!(self.out, ";")?;
3827        self.named_expressions.insert(named, name);
3828
3829        Ok(())
3830    }
3831
3832    /// Helper function that write string with default zero initialization for supported types
3833    fn write_zero_init_value(&mut self, ty: Handle<crate::Type>) -> BackendResult {
3834        let inner = &self.module.types[ty].inner;
3835        match *inner {
3836            TypeInner::Scalar { kind, .. } | TypeInner::Atomic { kind, .. } => {
3837                self.write_zero_init_scalar(kind)?;
3838            }
3839            TypeInner::Vector { kind, .. } => {
3840                self.write_value_type(inner)?;
3841                write!(self.out, "(")?;
3842                self.write_zero_init_scalar(kind)?;
3843                write!(self.out, ")")?;
3844            }
3845            TypeInner::Matrix { .. } => {
3846                self.write_value_type(inner)?;
3847                write!(self.out, "(")?;
3848                self.write_zero_init_scalar(crate::ScalarKind::Float)?;
3849                write!(self.out, ")")?;
3850            }
3851            TypeInner::Array { base, size, .. } => {
3852                let count = match size
3853                    .to_indexable_length(self.module)
3854                    .expect("Bad array size")
3855                {
3856                    proc::IndexableLength::Known(count) => count,
3857                    proc::IndexableLength::Dynamic => return Ok(()),
3858                };
3859                self.write_type(base)?;
3860                self.write_array_size(base, size)?;
3861                write!(self.out, "(")?;
3862                for _ in 1..count {
3863                    self.write_zero_init_value(base)?;
3864                    write!(self.out, ", ")?;
3865                }
3866                // write last parameter without comma and space
3867                self.write_zero_init_value(base)?;
3868                write!(self.out, ")")?;
3869            }
3870            TypeInner::Struct { ref members, .. } => {
3871                let name = &self.names[&NameKey::Type(ty)];
3872                write!(self.out, "{name}(")?;
3873                for (index, member) in members.iter().enumerate() {
3874                    if index != 0 {
3875                        write!(self.out, ", ")?;
3876                    }
3877                    self.write_zero_init_value(member.ty)?;
3878                }
3879                write!(self.out, ")")?;
3880            }
3881            _ => unreachable!(),
3882        }
3883
3884        Ok(())
3885    }
3886
3887    /// Helper function that write string with zero initialization for scalar
3888    fn write_zero_init_scalar(&mut self, kind: crate::ScalarKind) -> BackendResult {
3889        match kind {
3890            crate::ScalarKind::Bool => write!(self.out, "false")?,
3891            crate::ScalarKind::Uint => write!(self.out, "0u")?,
3892            crate::ScalarKind::Float => write!(self.out, "0.0")?,
3893            crate::ScalarKind::Sint => write!(self.out, "0")?,
3894        }
3895
3896        Ok(())
3897    }
3898
3899    /// Issue a memory barrier. Please note that to ensure visibility,
3900    /// OpenGL always requires a call to the `barrier()` function after a `memoryBarrier*()`
3901    fn write_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult {
3902        if flags.contains(crate::Barrier::STORAGE) {
3903            writeln!(self.out, "{level}memoryBarrierBuffer();")?;
3904        }
3905        if flags.contains(crate::Barrier::WORK_GROUP) {
3906            writeln!(self.out, "{level}memoryBarrierShared();")?;
3907        }
3908        writeln!(self.out, "{level}barrier();")?;
3909        Ok(())
3910    }
3911
3912    /// Helper function that return the glsl storage access string of [`StorageAccess`](crate::StorageAccess)
3913    ///
3914    /// glsl allows adding both `readonly` and `writeonly` but this means that
3915    /// they can only be used to query information about the resource which isn't what
3916    /// we want here so when storage access is both `LOAD` and `STORE` add no modifiers
3917    fn write_storage_access(&mut self, storage_access: crate::StorageAccess) -> BackendResult {
3918        if !storage_access.contains(crate::StorageAccess::STORE) {
3919            write!(self.out, "readonly ")?;
3920        }
3921        if !storage_access.contains(crate::StorageAccess::LOAD) {
3922            write!(self.out, "writeonly ")?;
3923        }
3924        Ok(())
3925    }
3926
3927    /// Helper method used to produce the reflection info that's returned to the user
3928    fn collect_reflection_info(&self) -> Result<ReflectionInfo, Error> {
3929        use std::collections::hash_map::Entry;
3930        let info = self.info.get_entry_point(self.entry_point_idx as usize);
3931        let mut texture_mapping = crate::FastHashMap::default();
3932        let mut uniforms = crate::FastHashMap::default();
3933
3934        for sampling in info.sampling_set.iter() {
3935            let tex_name = self.reflection_names_globals[&sampling.image].clone();
3936
3937            match texture_mapping.entry(tex_name) {
3938                Entry::Vacant(v) => {
3939                    v.insert(TextureMapping {
3940                        texture: sampling.image,
3941                        sampler: Some(sampling.sampler),
3942                    });
3943                }
3944                Entry::Occupied(e) => {
3945                    if e.get().sampler != Some(sampling.sampler) {
3946                        log::error!("Conflicting samplers for {}", e.key());
3947                        return Err(Error::ImageMultipleSamplers);
3948                    }
3949                }
3950            }
3951        }
3952
3953        for (handle, var) in self.module.global_variables.iter() {
3954            if info[handle].is_empty() {
3955                continue;
3956            }
3957            match self.module.types[var.ty].inner {
3958                crate::TypeInner::Image { .. } => {
3959                    let tex_name = self.reflection_names_globals[&handle].clone();
3960                    match texture_mapping.entry(tex_name) {
3961                        Entry::Vacant(v) => {
3962                            v.insert(TextureMapping {
3963                                texture: handle,
3964                                sampler: None,
3965                            });
3966                        }
3967                        Entry::Occupied(_) => {
3968                            // already used with a sampler, do nothing
3969                        }
3970                    }
3971                }
3972                _ => match var.space {
3973                    crate::AddressSpace::Uniform | crate::AddressSpace::Storage { .. } => {
3974                        let name = self.reflection_names_globals[&handle].clone();
3975                        uniforms.insert(handle, name);
3976                    }
3977                    _ => (),
3978                },
3979            }
3980        }
3981
3982        Ok(ReflectionInfo {
3983            texture_mapping,
3984            uniforms,
3985        })
3986    }
3987}
3988
3989/// Structure returned by [`glsl_scalar`](glsl_scalar)
3990///
3991/// It contains both a prefix used in other types and the full type name
3992struct ScalarString<'a> {
3993    /// The prefix used to compose other types
3994    prefix: &'a str,
3995    /// The name of the scalar type
3996    full: &'a str,
3997}
3998
3999/// Helper function that returns scalar related strings
4000///
4001/// Check [`ScalarString`](ScalarString) for the information provided
4002///
4003/// # Errors
4004/// If a [`Float`](crate::ScalarKind::Float) with an width that isn't 4 or 8
4005const fn glsl_scalar(
4006    kind: crate::ScalarKind,
4007    width: crate::Bytes,
4008) -> Result<ScalarString<'static>, Error> {
4009    use crate::ScalarKind as Sk;
4010
4011    Ok(match kind {
4012        Sk::Sint => ScalarString {
4013            prefix: "i",
4014            full: "int",
4015        },
4016        Sk::Uint => ScalarString {
4017            prefix: "u",
4018            full: "uint",
4019        },
4020        Sk::Float => match width {
4021            4 => ScalarString {
4022                prefix: "",
4023                full: "float",
4024            },
4025            8 => ScalarString {
4026                prefix: "d",
4027                full: "double",
4028            },
4029            _ => return Err(Error::UnsupportedScalar(kind, width)),
4030        },
4031        Sk::Bool => ScalarString {
4032            prefix: "b",
4033            full: "bool",
4034        },
4035    })
4036}
4037
4038/// Helper function that returns the glsl variable name for a builtin
4039const fn glsl_built_in(
4040    built_in: crate::BuiltIn,
4041    output: bool,
4042    targetting_webgl: bool,
4043) -> &'static str {
4044    use crate::BuiltIn as Bi;
4045
4046    match built_in {
4047        Bi::Position { .. } => {
4048            if output {
4049                "gl_Position"
4050            } else {
4051                "gl_FragCoord"
4052            }
4053        }
4054        Bi::ViewIndex if targetting_webgl => "int(gl_ViewID_OVR)",
4055        Bi::ViewIndex => "gl_ViewIndex",
4056        // vertex
4057        Bi::BaseInstance => "uint(gl_BaseInstance)",
4058        Bi::BaseVertex => "uint(gl_BaseVertex)",
4059        Bi::ClipDistance => "gl_ClipDistance",
4060        Bi::CullDistance => "gl_CullDistance",
4061        Bi::InstanceIndex => "uint(gl_InstanceID)",
4062        Bi::PointSize => "gl_PointSize",
4063        Bi::VertexIndex => "uint(gl_VertexID)",
4064        // fragment
4065        Bi::FragDepth => "gl_FragDepth",
4066        Bi::PointCoord => "gl_PointCoord",
4067        Bi::FrontFacing => "gl_FrontFacing",
4068        Bi::PrimitiveIndex => "uint(gl_PrimitiveID)",
4069        Bi::SampleIndex => "gl_SampleID",
4070        Bi::SampleMask => {
4071            if output {
4072                "gl_SampleMask"
4073            } else {
4074                "gl_SampleMaskIn"
4075            }
4076        }
4077        // compute
4078        Bi::GlobalInvocationId => "gl_GlobalInvocationID",
4079        Bi::LocalInvocationId => "gl_LocalInvocationID",
4080        Bi::LocalInvocationIndex => "gl_LocalInvocationIndex",
4081        Bi::WorkGroupId => "gl_WorkGroupID",
4082        Bi::WorkGroupSize => "gl_WorkGroupSize",
4083        Bi::NumWorkGroups => "gl_NumWorkGroups",
4084    }
4085}
4086
4087/// Helper function that returns the string corresponding to the address space
4088const fn glsl_storage_qualifier(space: crate::AddressSpace) -> Option<&'static str> {
4089    use crate::AddressSpace as As;
4090
4091    match space {
4092        As::Function => None,
4093        As::Private => None,
4094        As::Storage { .. } => Some("buffer"),
4095        As::Uniform => Some("uniform"),
4096        As::Handle => Some("uniform"),
4097        As::WorkGroup => Some("shared"),
4098        As::PushConstant => Some("uniform"),
4099    }
4100}
4101
4102/// Helper function that returns the string corresponding to the glsl interpolation qualifier
4103const fn glsl_interpolation(interpolation: crate::Interpolation) -> &'static str {
4104    use crate::Interpolation as I;
4105
4106    match interpolation {
4107        I::Perspective => "smooth",
4108        I::Linear => "noperspective",
4109        I::Flat => "flat",
4110    }
4111}
4112
4113/// Return the GLSL auxiliary qualifier for the given sampling value.
4114const fn glsl_sampling(sampling: crate::Sampling) -> Option<&'static str> {
4115    use crate::Sampling as S;
4116
4117    match sampling {
4118        S::Center => None,
4119        S::Centroid => Some("centroid"),
4120        S::Sample => Some("sample"),
4121    }
4122}
4123
4124/// Helper function that returns the glsl dimension string of [`ImageDimension`](crate::ImageDimension)
4125const fn glsl_dimension(dim: crate::ImageDimension) -> &'static str {
4126    use crate::ImageDimension as IDim;
4127
4128    match dim {
4129        IDim::D1 => "1D",
4130        IDim::D2 => "2D",
4131        IDim::D3 => "3D",
4132        IDim::Cube => "Cube",
4133    }
4134}
4135
4136/// Helper function that returns the glsl storage format string of [`StorageFormat`](crate::StorageFormat)
4137const fn glsl_storage_format(format: crate::StorageFormat) -> &'static str {
4138    use crate::StorageFormat as Sf;
4139
4140    match format {
4141        Sf::R8Unorm => "r8",
4142        Sf::R8Snorm => "r8_snorm",
4143        Sf::R8Uint => "r8ui",
4144        Sf::R8Sint => "r8i",
4145        Sf::R16Uint => "r16ui",
4146        Sf::R16Sint => "r16i",
4147        Sf::R16Float => "r16f",
4148        Sf::Rg8Unorm => "rg8",
4149        Sf::Rg8Snorm => "rg8_snorm",
4150        Sf::Rg8Uint => "rg8ui",
4151        Sf::Rg8Sint => "rg8i",
4152        Sf::R32Uint => "r32ui",
4153        Sf::R32Sint => "r32i",
4154        Sf::R32Float => "r32f",
4155        Sf::Rg16Uint => "rg16ui",
4156        Sf::Rg16Sint => "rg16i",
4157        Sf::Rg16Float => "rg16f",
4158        Sf::Rgba8Unorm => "rgba8",
4159        Sf::Rgba8Snorm => "rgba8_snorm",
4160        Sf::Rgba8Uint => "rgba8ui",
4161        Sf::Rgba8Sint => "rgba8i",
4162        Sf::Rgb10a2Unorm => "rgb10_a2ui",
4163        Sf::Rg11b10Float => "r11f_g11f_b10f",
4164        Sf::Rg32Uint => "rg32ui",
4165        Sf::Rg32Sint => "rg32i",
4166        Sf::Rg32Float => "rg32f",
4167        Sf::Rgba16Uint => "rgba16ui",
4168        Sf::Rgba16Sint => "rgba16i",
4169        Sf::Rgba16Float => "rgba16f",
4170        Sf::Rgba32Uint => "rgba32ui",
4171        Sf::Rgba32Sint => "rgba32i",
4172        Sf::Rgba32Float => "rgba32f",
4173        Sf::R16Unorm => "r16",
4174        Sf::R16Snorm => "r16_snorm",
4175        Sf::Rg16Unorm => "rg16",
4176        Sf::Rg16Snorm => "rg16_snorm",
4177        Sf::Rgba16Unorm => "rgba16",
4178        Sf::Rgba16Snorm => "rgba16_snorm",
4179    }
4180}
4181
4182fn is_value_init_supported(module: &crate::Module, ty: Handle<crate::Type>) -> bool {
4183    match module.types[ty].inner {
4184        TypeInner::Scalar { .. } | TypeInner::Vector { .. } | TypeInner::Matrix { .. } => true,
4185        TypeInner::Array { base, size, .. } => {
4186            size != crate::ArraySize::Dynamic && is_value_init_supported(module, base)
4187        }
4188        TypeInner::Struct { ref members, .. } => members
4189            .iter()
4190            .all(|member| is_value_init_supported(module, member.ty)),
4191        _ => false,
4192    }
4193}