naga/back/wgsl/
writer.rs

1use alloc::{
2    format,
3    string::{String, ToString},
4    vec,
5    vec::Vec,
6};
7use core::fmt::Write;
8
9use super::Error;
10use super::ToWgslIfImplemented as _;
11use crate::{back::wgsl::polyfill::InversePolyfill, common::wgsl::TypeContext};
12use crate::{
13    back::{self, Baked},
14    common::{
15        self,
16        wgsl::{address_space_str, ToWgsl, TryToWgsl},
17    },
18    proc::{self, NameKey},
19    valid, Handle, Module, ShaderStage, TypeInner,
20};
21
22/// Shorthand result used internally by the backend
23type BackendResult = Result<(), Error>;
24
25/// WGSL [attribute](https://gpuweb.github.io/gpuweb/wgsl/#attributes)
26enum Attribute {
27    Binding(u32),
28    BuiltIn(crate::BuiltIn),
29    Group(u32),
30    Invariant,
31    Interpolate(Option<crate::Interpolation>, Option<crate::Sampling>),
32    Location(u32),
33    BlendSrc(u32),
34    Stage(ShaderStage),
35    WorkGroupSize([u32; 3]),
36}
37
38/// The WGSL form that `write_expr_with_indirection` should use to render a Naga
39/// expression.
40///
41/// Sometimes a Naga `Expression` alone doesn't provide enough information to
42/// choose the right rendering for it in WGSL. For example, one natural WGSL
43/// rendering of a Naga `LocalVariable(x)` expression might be `&x`, since
44/// `LocalVariable` produces a pointer to the local variable's storage. But when
45/// rendering a `Store` statement, the `pointer` operand must be the left hand
46/// side of a WGSL assignment, so the proper rendering is `x`.
47///
48/// The caller of `write_expr_with_indirection` must provide an `Expected` value
49/// to indicate how ambiguous expressions should be rendered.
50#[derive(Clone, Copy, Debug)]
51enum Indirection {
52    /// Render pointer-construction expressions as WGSL `ptr`-typed expressions.
53    ///
54    /// This is the right choice for most cases. Whenever a Naga pointer
55    /// expression is not the `pointer` operand of a `Load` or `Store`, it
56    /// must be a WGSL pointer expression.
57    Ordinary,
58
59    /// Render pointer-construction expressions as WGSL reference-typed
60    /// expressions.
61    ///
62    /// For example, this is the right choice for the `pointer` operand when
63    /// rendering a `Store` statement as a WGSL assignment.
64    Reference,
65}
66
67bitflags::bitflags! {
68    #[cfg_attr(feature = "serialize", derive(serde::Serialize))]
69    #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
70    #[derive(Clone, Copy, Debug, Eq, PartialEq)]
71    pub struct WriterFlags: u32 {
72        /// Always annotate the type information instead of inferring.
73        const EXPLICIT_TYPES = 0x1;
74    }
75}
76
77pub struct Writer<W> {
78    out: W,
79    flags: WriterFlags,
80    names: crate::FastHashMap<NameKey, String>,
81    namer: proc::Namer,
82    named_expressions: crate::NamedExpressions,
83    required_polyfills: crate::FastIndexSet<InversePolyfill>,
84}
85
86impl<W: Write> Writer<W> {
87    pub fn new(out: W, flags: WriterFlags) -> Self {
88        Writer {
89            out,
90            flags,
91            names: crate::FastHashMap::default(),
92            namer: proc::Namer::default(),
93            named_expressions: crate::NamedExpressions::default(),
94            required_polyfills: crate::FastIndexSet::default(),
95        }
96    }
97
98    fn reset(&mut self, module: &Module) {
99        self.names.clear();
100        self.namer.reset(
101            module,
102            &crate::keywords::wgsl::RESERVED_SET,
103            // an identifier must not start with two underscore
104            &[],
105            &["__", "_naga"],
106            &mut self.names,
107        );
108        self.named_expressions.clear();
109        self.required_polyfills.clear();
110    }
111
112    fn is_builtin_wgsl_struct(&self, module: &Module, handle: Handle<crate::Type>) -> bool {
113        module
114            .special_types
115            .predeclared_types
116            .values()
117            .any(|t| *t == handle)
118    }
119
120    pub fn write(&mut self, module: &Module, info: &valid::ModuleInfo) -> BackendResult {
121        if !module.overrides.is_empty() {
122            return Err(Error::Unimplemented(
123                "Pipeline constants are not yet supported for this back-end".to_string(),
124            ));
125        }
126
127        self.reset(module);
128
129        // Write all `enable` declarations
130        self.write_enable_declarations(module)?;
131
132        // Write all structs
133        for (handle, ty) in module.types.iter() {
134            if let TypeInner::Struct { ref members, .. } = ty.inner {
135                {
136                    if !self.is_builtin_wgsl_struct(module, handle) {
137                        self.write_struct(module, handle, members)?;
138                        writeln!(self.out)?;
139                    }
140                }
141            }
142        }
143
144        // Write all named constants
145        let mut constants = module
146            .constants
147            .iter()
148            .filter(|&(_, c)| c.name.is_some())
149            .peekable();
150        while let Some((handle, _)) = constants.next() {
151            self.write_global_constant(module, handle)?;
152            // Add extra newline for readability on last iteration
153            if constants.peek().is_none() {
154                writeln!(self.out)?;
155            }
156        }
157
158        // Write all globals
159        for (ty, global) in module.global_variables.iter() {
160            self.write_global(module, global, ty)?;
161        }
162
163        if !module.global_variables.is_empty() {
164            // Add extra newline for readability
165            writeln!(self.out)?;
166        }
167
168        // Write all regular functions
169        for (handle, function) in module.functions.iter() {
170            let fun_info = &info[handle];
171
172            let func_ctx = back::FunctionCtx {
173                ty: back::FunctionType::Function(handle),
174                info: fun_info,
175                expressions: &function.expressions,
176                named_expressions: &function.named_expressions,
177            };
178
179            // Write the function
180            self.write_function(module, function, &func_ctx)?;
181
182            writeln!(self.out)?;
183        }
184
185        // Write all entry points
186        for (index, ep) in module.entry_points.iter().enumerate() {
187            let attributes = match ep.stage {
188                ShaderStage::Vertex | ShaderStage::Fragment => vec![Attribute::Stage(ep.stage)],
189                ShaderStage::Compute => vec![
190                    Attribute::Stage(ShaderStage::Compute),
191                    Attribute::WorkGroupSize(ep.workgroup_size),
192                ],
193                ShaderStage::Task | ShaderStage::Mesh => unreachable!(),
194            };
195
196            self.write_attributes(&attributes)?;
197            // Add a newline after attribute
198            writeln!(self.out)?;
199
200            let func_ctx = back::FunctionCtx {
201                ty: back::FunctionType::EntryPoint(index as u16),
202                info: info.get_entry_point(index),
203                expressions: &ep.function.expressions,
204                named_expressions: &ep.function.named_expressions,
205            };
206            self.write_function(module, &ep.function, &func_ctx)?;
207
208            if index < module.entry_points.len() - 1 {
209                writeln!(self.out)?;
210            }
211        }
212
213        // Write any polyfills that were required.
214        for polyfill in &self.required_polyfills {
215            writeln!(self.out)?;
216            write!(self.out, "{}", polyfill.source)?;
217            writeln!(self.out)?;
218        }
219
220        Ok(())
221    }
222
223    /// Helper method which writes all the `enable` declarations
224    /// needed for a module.
225    fn write_enable_declarations(&mut self, module: &Module) -> BackendResult {
226        let mut needs_f16 = false;
227        let mut needs_dual_source_blending = false;
228        let mut needs_clip_distances = false;
229
230        // Determine which `enable` declarations are needed
231        for (_, ty) in module.types.iter() {
232            match ty.inner {
233                TypeInner::Scalar(scalar)
234                | TypeInner::Vector { scalar, .. }
235                | TypeInner::Matrix { scalar, .. } => {
236                    needs_f16 |= scalar == crate::Scalar::F16;
237                }
238                TypeInner::Struct { ref members, .. } => {
239                    for binding in members.iter().filter_map(|m| m.binding.as_ref()) {
240                        match *binding {
241                            crate::Binding::Location {
242                                blend_src: Some(_), ..
243                            } => {
244                                needs_dual_source_blending = true;
245                            }
246                            crate::Binding::BuiltIn(crate::BuiltIn::ClipDistance) => {
247                                needs_clip_distances = true;
248                            }
249                            _ => {}
250                        }
251                    }
252                }
253                _ => {}
254            }
255        }
256
257        // Write required declarations
258        let mut any_written = false;
259        if needs_f16 {
260            writeln!(self.out, "enable f16;")?;
261            any_written = true;
262        }
263        if needs_dual_source_blending {
264            writeln!(self.out, "enable dual_source_blending;")?;
265            any_written = true;
266        }
267        if needs_clip_distances {
268            writeln!(self.out, "enable clip_distances;")?;
269            any_written = true;
270        }
271        if any_written {
272            // Empty line for readability
273            writeln!(self.out)?;
274        }
275
276        Ok(())
277    }
278
279    /// Helper method used to write
280    /// [functions](https://gpuweb.github.io/gpuweb/wgsl/#functions)
281    ///
282    /// # Notes
283    /// Ends in a newline
284    fn write_function(
285        &mut self,
286        module: &Module,
287        func: &crate::Function,
288        func_ctx: &back::FunctionCtx<'_>,
289    ) -> BackendResult {
290        let func_name = match func_ctx.ty {
291            back::FunctionType::EntryPoint(index) => &self.names[&NameKey::EntryPoint(index)],
292            back::FunctionType::Function(handle) => &self.names[&NameKey::Function(handle)],
293        };
294
295        // Write function name
296        write!(self.out, "fn {func_name}(")?;
297
298        // Write function arguments
299        for (index, arg) in func.arguments.iter().enumerate() {
300            // Write argument attribute if a binding is present
301            if let Some(ref binding) = arg.binding {
302                self.write_attributes(&map_binding_to_attribute(binding))?;
303            }
304            // Write argument name
305            let argument_name = &self.names[&func_ctx.argument_key(index as u32)];
306
307            write!(self.out, "{argument_name}: ")?;
308            // Write argument type
309            self.write_type(module, arg.ty)?;
310            if index < func.arguments.len() - 1 {
311                // Add a separator between args
312                write!(self.out, ", ")?;
313            }
314        }
315
316        write!(self.out, ")")?;
317
318        // Write function return type
319        if let Some(ref result) = func.result {
320            write!(self.out, " -> ")?;
321            if let Some(ref binding) = result.binding {
322                self.write_attributes(&map_binding_to_attribute(binding))?;
323            }
324            self.write_type(module, result.ty)?;
325        }
326
327        write!(self.out, " {{")?;
328        writeln!(self.out)?;
329
330        // Write function local variables
331        for (handle, local) in func.local_variables.iter() {
332            // Write indentation (only for readability)
333            write!(self.out, "{}", back::INDENT)?;
334
335            // Write the local name
336            // The leading space is important
337            write!(self.out, "var {}: ", self.names[&func_ctx.name_key(handle)])?;
338
339            // Write the local type
340            self.write_type(module, local.ty)?;
341
342            // Write the local initializer if needed
343            if let Some(init) = local.init {
344                // Put the equal signal only if there's a initializer
345                // The leading and trailing spaces aren't needed but help with readability
346                write!(self.out, " = ")?;
347
348                // Write the constant
349                // `write_constant` adds no trailing or leading space/newline
350                self.write_expr(module, init, func_ctx)?;
351            }
352
353            // Finish the local with `;` and add a newline (only for readability)
354            writeln!(self.out, ";")?
355        }
356
357        if !func.local_variables.is_empty() {
358            writeln!(self.out)?;
359        }
360
361        // Write the function body (statement list)
362        for sta in func.body.iter() {
363            // The indentation should always be 1 when writing the function body
364            self.write_stmt(module, sta, func_ctx, back::Level(1))?;
365        }
366
367        writeln!(self.out, "}}")?;
368
369        self.named_expressions.clear();
370
371        Ok(())
372    }
373
374    /// Helper method to write a attribute
375    fn write_attributes(&mut self, attributes: &[Attribute]) -> BackendResult {
376        for attribute in attributes {
377            match *attribute {
378                Attribute::Location(id) => write!(self.out, "@location({id}) ")?,
379                Attribute::BlendSrc(blend_src) => write!(self.out, "@blend_src({blend_src}) ")?,
380                Attribute::BuiltIn(builtin_attrib) => {
381                    let builtin = builtin_attrib.to_wgsl_if_implemented()?;
382                    write!(self.out, "@builtin({builtin}) ")?;
383                }
384                Attribute::Stage(shader_stage) => {
385                    let stage_str = match shader_stage {
386                        ShaderStage::Vertex => "vertex",
387                        ShaderStage::Fragment => "fragment",
388                        ShaderStage::Compute => "compute",
389                        ShaderStage::Task | ShaderStage::Mesh => unreachable!(),
390                    };
391                    write!(self.out, "@{stage_str} ")?;
392                }
393                Attribute::WorkGroupSize(size) => {
394                    write!(
395                        self.out,
396                        "@workgroup_size({}, {}, {}) ",
397                        size[0], size[1], size[2]
398                    )?;
399                }
400                Attribute::Binding(id) => write!(self.out, "@binding({id}) ")?,
401                Attribute::Group(id) => write!(self.out, "@group({id}) ")?,
402                Attribute::Invariant => write!(self.out, "@invariant ")?,
403                Attribute::Interpolate(interpolation, sampling) => {
404                    if sampling.is_some() && sampling != Some(crate::Sampling::Center) {
405                        let interpolation = interpolation
406                            .unwrap_or(crate::Interpolation::Perspective)
407                            .to_wgsl();
408                        let sampling = sampling.unwrap_or(crate::Sampling::Center).to_wgsl();
409                        write!(self.out, "@interpolate({interpolation}, {sampling}) ")?;
410                    } else if interpolation.is_some()
411                        && interpolation != Some(crate::Interpolation::Perspective)
412                    {
413                        let interpolation = interpolation
414                            .unwrap_or(crate::Interpolation::Perspective)
415                            .to_wgsl();
416                        write!(self.out, "@interpolate({interpolation}) ")?;
417                    }
418                }
419            };
420        }
421        Ok(())
422    }
423
424    /// Helper method used to write structs
425    /// Write the full declaration of a struct type.
426    ///
427    /// Write out a definition of the struct type referred to by
428    /// `handle` in `module`. The output will be an instance of the
429    /// `struct_decl` production in the WGSL grammar.
430    ///
431    /// Use `members` as the list of `handle`'s members. (This
432    /// function is usually called after matching a `TypeInner`, so
433    /// the callers already have the members at hand.)
434    fn write_struct(
435        &mut self,
436        module: &Module,
437        handle: Handle<crate::Type>,
438        members: &[crate::StructMember],
439    ) -> BackendResult {
440        write!(self.out, "struct {}", self.names[&NameKey::Type(handle)])?;
441        write!(self.out, " {{")?;
442        writeln!(self.out)?;
443        for (index, member) in members.iter().enumerate() {
444            // The indentation is only for readability
445            write!(self.out, "{}", back::INDENT)?;
446            if let Some(ref binding) = member.binding {
447                self.write_attributes(&map_binding_to_attribute(binding))?;
448            }
449            // Write struct member name and type
450            let member_name = &self.names[&NameKey::StructMember(handle, index as u32)];
451            write!(self.out, "{member_name}: ")?;
452            self.write_type(module, member.ty)?;
453            write!(self.out, ",")?;
454            writeln!(self.out)?;
455        }
456
457        writeln!(self.out, "}}")?;
458
459        Ok(())
460    }
461
462    fn write_type(&mut self, module: &Module, ty: Handle<crate::Type>) -> BackendResult {
463        // This actually can't be factored out into a nice constructor method,
464        // because the borrow checker needs to be able to see that the borrows
465        // of `self.names` and `self.out` are disjoint.
466        let type_context = WriterTypeContext {
467            module,
468            names: &self.names,
469        };
470        type_context.write_type(ty, &mut self.out)?;
471
472        Ok(())
473    }
474
475    fn write_type_resolution(
476        &mut self,
477        module: &Module,
478        resolution: &proc::TypeResolution,
479    ) -> BackendResult {
480        // This actually can't be factored out into a nice constructor method,
481        // because the borrow checker needs to be able to see that the borrows
482        // of `self.names` and `self.out` are disjoint.
483        let type_context = WriterTypeContext {
484            module,
485            names: &self.names,
486        };
487        type_context.write_type_resolution(resolution, &mut self.out)?;
488
489        Ok(())
490    }
491
492    /// Helper method used to write statements
493    ///
494    /// # Notes
495    /// Always adds a newline
496    fn write_stmt(
497        &mut self,
498        module: &Module,
499        stmt: &crate::Statement,
500        func_ctx: &back::FunctionCtx<'_>,
501        level: back::Level,
502    ) -> BackendResult {
503        use crate::{Expression, Statement};
504
505        match *stmt {
506            Statement::Emit(ref range) => {
507                for handle in range.clone() {
508                    let info = &func_ctx.info[handle];
509                    let expr_name = if let Some(name) = func_ctx.named_expressions.get(&handle) {
510                        // Front end provides names for all variables at the start of writing.
511                        // But we write them to step by step. We need to recache them
512                        // Otherwise, we could accidentally write variable name instead of full expression.
513                        // Also, we use sanitized names! It defense backend from generating variable with name from reserved keywords.
514                        Some(self.namer.call(name))
515                    } else {
516                        let expr = &func_ctx.expressions[handle];
517                        let min_ref_count = expr.bake_ref_count();
518                        // Forcefully creating baking expressions in some cases to help with readability
519                        let required_baking_expr = match *expr {
520                            Expression::ImageLoad { .. }
521                            | Expression::ImageQuery { .. }
522                            | Expression::ImageSample { .. } => true,
523                            _ => false,
524                        };
525                        if min_ref_count <= info.ref_count || required_baking_expr {
526                            Some(Baked(handle).to_string())
527                        } else {
528                            None
529                        }
530                    };
531
532                    if let Some(name) = expr_name {
533                        write!(self.out, "{level}")?;
534                        self.start_named_expr(module, handle, func_ctx, &name)?;
535                        self.write_expr(module, handle, func_ctx)?;
536                        self.named_expressions.insert(handle, name);
537                        writeln!(self.out, ";")?;
538                    }
539                }
540            }
541            // TODO: copy-paste from glsl-out
542            Statement::If {
543                condition,
544                ref accept,
545                ref reject,
546            } => {
547                write!(self.out, "{level}")?;
548                write!(self.out, "if ")?;
549                self.write_expr(module, condition, func_ctx)?;
550                writeln!(self.out, " {{")?;
551
552                let l2 = level.next();
553                for sta in accept {
554                    // Increase indentation to help with readability
555                    self.write_stmt(module, sta, func_ctx, l2)?;
556                }
557
558                // If there are no statements in the reject block we skip writing it
559                // This is only for readability
560                if !reject.is_empty() {
561                    writeln!(self.out, "{level}}} else {{")?;
562
563                    for sta in reject {
564                        // Increase indentation to help with readability
565                        self.write_stmt(module, sta, func_ctx, l2)?;
566                    }
567                }
568
569                writeln!(self.out, "{level}}}")?
570            }
571            Statement::Return { value } => {
572                write!(self.out, "{level}")?;
573                write!(self.out, "return")?;
574                if let Some(return_value) = value {
575                    // The leading space is important
576                    write!(self.out, " ")?;
577                    self.write_expr(module, return_value, func_ctx)?;
578                }
579                writeln!(self.out, ";")?;
580            }
581            // TODO: copy-paste from glsl-out
582            Statement::Kill => {
583                write!(self.out, "{level}")?;
584                writeln!(self.out, "discard;")?
585            }
586            Statement::Store { pointer, value } => {
587                write!(self.out, "{level}")?;
588
589                let is_atomic_pointer = func_ctx
590                    .resolve_type(pointer, &module.types)
591                    .is_atomic_pointer(&module.types);
592
593                if is_atomic_pointer {
594                    write!(self.out, "atomicStore(")?;
595                    self.write_expr(module, pointer, func_ctx)?;
596                    write!(self.out, ", ")?;
597                    self.write_expr(module, value, func_ctx)?;
598                    write!(self.out, ")")?;
599                } else {
600                    self.write_expr_with_indirection(
601                        module,
602                        pointer,
603                        func_ctx,
604                        Indirection::Reference,
605                    )?;
606                    write!(self.out, " = ")?;
607                    self.write_expr(module, value, func_ctx)?;
608                }
609                writeln!(self.out, ";")?
610            }
611            Statement::Call {
612                function,
613                ref arguments,
614                result,
615            } => {
616                write!(self.out, "{level}")?;
617                if let Some(expr) = result {
618                    let name = Baked(expr).to_string();
619                    self.start_named_expr(module, expr, func_ctx, &name)?;
620                    self.named_expressions.insert(expr, name);
621                }
622                let func_name = &self.names[&NameKey::Function(function)];
623                write!(self.out, "{func_name}(")?;
624                for (index, &argument) in arguments.iter().enumerate() {
625                    if index != 0 {
626                        write!(self.out, ", ")?;
627                    }
628                    self.write_expr(module, argument, func_ctx)?;
629                }
630                writeln!(self.out, ");")?
631            }
632            Statement::Atomic {
633                pointer,
634                ref fun,
635                value,
636                result,
637            } => {
638                write!(self.out, "{level}")?;
639                if let Some(result) = result {
640                    let res_name = Baked(result).to_string();
641                    self.start_named_expr(module, result, func_ctx, &res_name)?;
642                    self.named_expressions.insert(result, res_name);
643                }
644
645                let fun_str = fun.to_wgsl();
646                write!(self.out, "atomic{fun_str}(")?;
647                self.write_expr(module, pointer, func_ctx)?;
648                if let crate::AtomicFunction::Exchange { compare: Some(cmp) } = *fun {
649                    write!(self.out, ", ")?;
650                    self.write_expr(module, cmp, func_ctx)?;
651                }
652                write!(self.out, ", ")?;
653                self.write_expr(module, value, func_ctx)?;
654                writeln!(self.out, ");")?
655            }
656            Statement::ImageAtomic {
657                image,
658                coordinate,
659                array_index,
660                ref fun,
661                value,
662            } => {
663                write!(self.out, "{level}")?;
664                let fun_str = fun.to_wgsl();
665                write!(self.out, "textureAtomic{fun_str}(")?;
666                self.write_expr(module, image, func_ctx)?;
667                write!(self.out, ", ")?;
668                self.write_expr(module, coordinate, func_ctx)?;
669                if let Some(array_index_expr) = array_index {
670                    write!(self.out, ", ")?;
671                    self.write_expr(module, array_index_expr, func_ctx)?;
672                }
673                write!(self.out, ", ")?;
674                self.write_expr(module, value, func_ctx)?;
675                writeln!(self.out, ");")?;
676            }
677            Statement::WorkGroupUniformLoad { pointer, result } => {
678                write!(self.out, "{level}")?;
679                // TODO: Obey named expressions here.
680                let res_name = Baked(result).to_string();
681                self.start_named_expr(module, result, func_ctx, &res_name)?;
682                self.named_expressions.insert(result, res_name);
683                write!(self.out, "workgroupUniformLoad(")?;
684                self.write_expr(module, pointer, func_ctx)?;
685                writeln!(self.out, ");")?;
686            }
687            Statement::ImageStore {
688                image,
689                coordinate,
690                array_index,
691                value,
692            } => {
693                write!(self.out, "{level}")?;
694                write!(self.out, "textureStore(")?;
695                self.write_expr(module, image, func_ctx)?;
696                write!(self.out, ", ")?;
697                self.write_expr(module, coordinate, func_ctx)?;
698                if let Some(array_index_expr) = array_index {
699                    write!(self.out, ", ")?;
700                    self.write_expr(module, array_index_expr, func_ctx)?;
701                }
702                write!(self.out, ", ")?;
703                self.write_expr(module, value, func_ctx)?;
704                writeln!(self.out, ");")?;
705            }
706            // TODO: copy-paste from glsl-out
707            Statement::Block(ref block) => {
708                write!(self.out, "{level}")?;
709                writeln!(self.out, "{{")?;
710                for sta in block.iter() {
711                    // Increase the indentation to help with readability
712                    self.write_stmt(module, sta, func_ctx, level.next())?
713                }
714                writeln!(self.out, "{level}}}")?
715            }
716            Statement::Switch {
717                selector,
718                ref cases,
719            } => {
720                // Start the switch
721                write!(self.out, "{level}")?;
722                write!(self.out, "switch ")?;
723                self.write_expr(module, selector, func_ctx)?;
724                writeln!(self.out, " {{")?;
725
726                let l2 = level.next();
727                let mut new_case = true;
728                for case in cases {
729                    if case.fall_through && !case.body.is_empty() {
730                        // TODO: we could do the same workaround as we did for the HLSL backend
731                        return Err(Error::Unimplemented(
732                            "fall-through switch case block".into(),
733                        ));
734                    }
735
736                    match case.value {
737                        crate::SwitchValue::I32(value) => {
738                            if new_case {
739                                write!(self.out, "{l2}case ")?;
740                            }
741                            write!(self.out, "{value}")?;
742                        }
743                        crate::SwitchValue::U32(value) => {
744                            if new_case {
745                                write!(self.out, "{l2}case ")?;
746                            }
747                            write!(self.out, "{value}u")?;
748                        }
749                        crate::SwitchValue::Default => {
750                            if new_case {
751                                if case.fall_through {
752                                    write!(self.out, "{l2}case ")?;
753                                } else {
754                                    write!(self.out, "{l2}")?;
755                                }
756                            }
757                            write!(self.out, "default")?;
758                        }
759                    }
760
761                    new_case = !case.fall_through;
762
763                    if case.fall_through {
764                        write!(self.out, ", ")?;
765                    } else {
766                        writeln!(self.out, ": {{")?;
767                    }
768
769                    for sta in case.body.iter() {
770                        self.write_stmt(module, sta, func_ctx, l2.next())?;
771                    }
772
773                    if !case.fall_through {
774                        writeln!(self.out, "{l2}}}")?;
775                    }
776                }
777
778                writeln!(self.out, "{level}}}")?
779            }
780            Statement::Loop {
781                ref body,
782                ref continuing,
783                break_if,
784            } => {
785                write!(self.out, "{level}")?;
786                writeln!(self.out, "loop {{")?;
787
788                let l2 = level.next();
789                for sta in body.iter() {
790                    self.write_stmt(module, sta, func_ctx, l2)?;
791                }
792
793                // The continuing is optional so we don't need to write it if
794                // it is empty, but the `break if` counts as a continuing statement
795                // so even if `continuing` is empty we must generate it if a
796                // `break if` exists
797                if !continuing.is_empty() || break_if.is_some() {
798                    writeln!(self.out, "{l2}continuing {{")?;
799                    for sta in continuing.iter() {
800                        self.write_stmt(module, sta, func_ctx, l2.next())?;
801                    }
802
803                    // The `break if` is always the last
804                    // statement of the `continuing` block
805                    if let Some(condition) = break_if {
806                        // The trailing space is important
807                        write!(self.out, "{}break if ", l2.next())?;
808                        self.write_expr(module, condition, func_ctx)?;
809                        // Close the `break if` statement
810                        writeln!(self.out, ";")?;
811                    }
812
813                    writeln!(self.out, "{l2}}}")?;
814                }
815
816                writeln!(self.out, "{level}}}")?
817            }
818            Statement::Break => {
819                writeln!(self.out, "{level}break;")?;
820            }
821            Statement::Continue => {
822                writeln!(self.out, "{level}continue;")?;
823            }
824            Statement::ControlBarrier(barrier) | Statement::MemoryBarrier(barrier) => {
825                if barrier.contains(crate::Barrier::STORAGE) {
826                    writeln!(self.out, "{level}storageBarrier();")?;
827                }
828
829                if barrier.contains(crate::Barrier::WORK_GROUP) {
830                    writeln!(self.out, "{level}workgroupBarrier();")?;
831                }
832
833                if barrier.contains(crate::Barrier::SUB_GROUP) {
834                    writeln!(self.out, "{level}subgroupBarrier();")?;
835                }
836
837                if barrier.contains(crate::Barrier::TEXTURE) {
838                    writeln!(self.out, "{level}textureBarrier();")?;
839                }
840            }
841            Statement::RayQuery { .. } => unreachable!(),
842            Statement::SubgroupBallot { result, predicate } => {
843                write!(self.out, "{level}")?;
844                let res_name = Baked(result).to_string();
845                self.start_named_expr(module, result, func_ctx, &res_name)?;
846                self.named_expressions.insert(result, res_name);
847
848                write!(self.out, "subgroupBallot(")?;
849                if let Some(predicate) = predicate {
850                    self.write_expr(module, predicate, func_ctx)?;
851                }
852                writeln!(self.out, ");")?;
853            }
854            Statement::SubgroupCollectiveOperation {
855                op,
856                collective_op,
857                argument,
858                result,
859            } => {
860                write!(self.out, "{level}")?;
861                let res_name = Baked(result).to_string();
862                self.start_named_expr(module, result, func_ctx, &res_name)?;
863                self.named_expressions.insert(result, res_name);
864
865                match (collective_op, op) {
866                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::All) => {
867                        write!(self.out, "subgroupAll(")?
868                    }
869                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Any) => {
870                        write!(self.out, "subgroupAny(")?
871                    }
872                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Add) => {
873                        write!(self.out, "subgroupAdd(")?
874                    }
875                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Mul) => {
876                        write!(self.out, "subgroupMul(")?
877                    }
878                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Max) => {
879                        write!(self.out, "subgroupMax(")?
880                    }
881                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Min) => {
882                        write!(self.out, "subgroupMin(")?
883                    }
884                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::And) => {
885                        write!(self.out, "subgroupAnd(")?
886                    }
887                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Or) => {
888                        write!(self.out, "subgroupOr(")?
889                    }
890                    (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Xor) => {
891                        write!(self.out, "subgroupXor(")?
892                    }
893                    (crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Add) => {
894                        write!(self.out, "subgroupExclusiveAdd(")?
895                    }
896                    (crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Mul) => {
897                        write!(self.out, "subgroupExclusiveMul(")?
898                    }
899                    (crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Add) => {
900                        write!(self.out, "subgroupInclusiveAdd(")?
901                    }
902                    (crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Mul) => {
903                        write!(self.out, "subgroupInclusiveMul(")?
904                    }
905                    _ => unimplemented!(),
906                }
907                self.write_expr(module, argument, func_ctx)?;
908                writeln!(self.out, ");")?;
909            }
910            Statement::SubgroupGather {
911                mode,
912                argument,
913                result,
914            } => {
915                write!(self.out, "{level}")?;
916                let res_name = Baked(result).to_string();
917                self.start_named_expr(module, result, func_ctx, &res_name)?;
918                self.named_expressions.insert(result, res_name);
919
920                match mode {
921                    crate::GatherMode::BroadcastFirst => {
922                        write!(self.out, "subgroupBroadcastFirst(")?;
923                    }
924                    crate::GatherMode::Broadcast(_) => {
925                        write!(self.out, "subgroupBroadcast(")?;
926                    }
927                    crate::GatherMode::Shuffle(_) => {
928                        write!(self.out, "subgroupShuffle(")?;
929                    }
930                    crate::GatherMode::ShuffleDown(_) => {
931                        write!(self.out, "subgroupShuffleDown(")?;
932                    }
933                    crate::GatherMode::ShuffleUp(_) => {
934                        write!(self.out, "subgroupShuffleUp(")?;
935                    }
936                    crate::GatherMode::ShuffleXor(_) => {
937                        write!(self.out, "subgroupShuffleXor(")?;
938                    }
939                    crate::GatherMode::QuadBroadcast(_) => {
940                        write!(self.out, "quadBroadcast(")?;
941                    }
942                    crate::GatherMode::QuadSwap(direction) => match direction {
943                        crate::Direction::X => {
944                            write!(self.out, "quadSwapX(")?;
945                        }
946                        crate::Direction::Y => {
947                            write!(self.out, "quadSwapY(")?;
948                        }
949                        crate::Direction::Diagonal => {
950                            write!(self.out, "quadSwapDiagonal(")?;
951                        }
952                    },
953                }
954                self.write_expr(module, argument, func_ctx)?;
955                match mode {
956                    crate::GatherMode::BroadcastFirst => {}
957                    crate::GatherMode::Broadcast(index)
958                    | crate::GatherMode::Shuffle(index)
959                    | crate::GatherMode::ShuffleDown(index)
960                    | crate::GatherMode::ShuffleUp(index)
961                    | crate::GatherMode::ShuffleXor(index)
962                    | crate::GatherMode::QuadBroadcast(index) => {
963                        write!(self.out, ", ")?;
964                        self.write_expr(module, index, func_ctx)?;
965                    }
966                    crate::GatherMode::QuadSwap(_) => {}
967                }
968                writeln!(self.out, ");")?;
969            }
970        }
971
972        Ok(())
973    }
974
975    /// Return the sort of indirection that `expr`'s plain form evaluates to.
976    ///
977    /// An expression's 'plain form' is the most general rendition of that
978    /// expression into WGSL, lacking `&` or `*` operators:
979    ///
980    /// - The plain form of `LocalVariable(x)` is simply `x`, which is a reference
981    ///   to the local variable's storage.
982    ///
983    /// - The plain form of `GlobalVariable(g)` is simply `g`, which is usually a
984    ///   reference to the global variable's storage. However, globals in the
985    ///   `Handle` address space are immutable, and `GlobalVariable` expressions for
986    ///   those produce the value directly, not a pointer to it. Such
987    ///   `GlobalVariable` expressions are `Ordinary`.
988    ///
989    /// - `Access` and `AccessIndex` are `Reference` when their `base` operand is a
990    ///   pointer. If they are applied directly to a composite value, they are
991    ///   `Ordinary`.
992    ///
993    /// Note that `FunctionArgument` expressions are never `Reference`, even when
994    /// the argument's type is `Pointer`. `FunctionArgument` always evaluates to the
995    /// argument's value directly, so any pointer it produces is merely the value
996    /// passed by the caller.
997    fn plain_form_indirection(
998        &self,
999        expr: Handle<crate::Expression>,
1000        module: &Module,
1001        func_ctx: &back::FunctionCtx<'_>,
1002    ) -> Indirection {
1003        use crate::Expression as Ex;
1004
1005        // Named expressions are `let` expressions, which apply the Load Rule,
1006        // so if their type is a Naga pointer, then that must be a WGSL pointer
1007        // as well.
1008        if self.named_expressions.contains_key(&expr) {
1009            return Indirection::Ordinary;
1010        }
1011
1012        match func_ctx.expressions[expr] {
1013            Ex::LocalVariable(_) => Indirection::Reference,
1014            Ex::GlobalVariable(handle) => {
1015                let global = &module.global_variables[handle];
1016                match global.space {
1017                    crate::AddressSpace::Handle => Indirection::Ordinary,
1018                    _ => Indirection::Reference,
1019                }
1020            }
1021            Ex::Access { base, .. } | Ex::AccessIndex { base, .. } => {
1022                let base_ty = func_ctx.resolve_type(base, &module.types);
1023                match *base_ty {
1024                    TypeInner::Pointer { .. } | TypeInner::ValuePointer { .. } => {
1025                        Indirection::Reference
1026                    }
1027                    _ => Indirection::Ordinary,
1028                }
1029            }
1030            _ => Indirection::Ordinary,
1031        }
1032    }
1033
1034    fn start_named_expr(
1035        &mut self,
1036        module: &Module,
1037        handle: Handle<crate::Expression>,
1038        func_ctx: &back::FunctionCtx,
1039        name: &str,
1040    ) -> BackendResult {
1041        // Write variable name
1042        write!(self.out, "let {name}")?;
1043        if self.flags.contains(WriterFlags::EXPLICIT_TYPES) {
1044            write!(self.out, ": ")?;
1045            // Write variable type
1046            self.write_type_resolution(module, &func_ctx.info[handle].ty)?;
1047        }
1048
1049        write!(self.out, " = ")?;
1050        Ok(())
1051    }
1052
1053    /// Write the ordinary WGSL form of `expr`.
1054    ///
1055    /// See `write_expr_with_indirection` for details.
1056    fn write_expr(
1057        &mut self,
1058        module: &Module,
1059        expr: Handle<crate::Expression>,
1060        func_ctx: &back::FunctionCtx<'_>,
1061    ) -> BackendResult {
1062        self.write_expr_with_indirection(module, expr, func_ctx, Indirection::Ordinary)
1063    }
1064
1065    /// Write `expr` as a WGSL expression with the requested indirection.
1066    ///
1067    /// In terms of the WGSL grammar, the resulting expression is a
1068    /// `singular_expression`. It may be parenthesized. This makes it suitable
1069    /// for use as the operand of a unary or binary operator without worrying
1070    /// about precedence.
1071    ///
1072    /// This does not produce newlines or indentation.
1073    ///
1074    /// The `requested` argument indicates (roughly) whether Naga
1075    /// `Pointer`-valued expressions represent WGSL references or pointers. See
1076    /// `Indirection` for details.
1077    fn write_expr_with_indirection(
1078        &mut self,
1079        module: &Module,
1080        expr: Handle<crate::Expression>,
1081        func_ctx: &back::FunctionCtx<'_>,
1082        requested: Indirection,
1083    ) -> BackendResult {
1084        // If the plain form of the expression is not what we need, emit the
1085        // operator necessary to correct that.
1086        let plain = self.plain_form_indirection(expr, module, func_ctx);
1087        match (requested, plain) {
1088            (Indirection::Ordinary, Indirection::Reference) => {
1089                write!(self.out, "(&")?;
1090                self.write_expr_plain_form(module, expr, func_ctx, plain)?;
1091                write!(self.out, ")")?;
1092            }
1093            (Indirection::Reference, Indirection::Ordinary) => {
1094                write!(self.out, "(*")?;
1095                self.write_expr_plain_form(module, expr, func_ctx, plain)?;
1096                write!(self.out, ")")?;
1097            }
1098            (_, _) => self.write_expr_plain_form(module, expr, func_ctx, plain)?,
1099        }
1100
1101        Ok(())
1102    }
1103
1104    fn write_const_expression(
1105        &mut self,
1106        module: &Module,
1107        expr: Handle<crate::Expression>,
1108        arena: &crate::Arena<crate::Expression>,
1109    ) -> BackendResult {
1110        self.write_possibly_const_expression(module, expr, arena, |writer, expr| {
1111            writer.write_const_expression(module, expr, arena)
1112        })
1113    }
1114
1115    fn write_possibly_const_expression<E>(
1116        &mut self,
1117        module: &Module,
1118        expr: Handle<crate::Expression>,
1119        expressions: &crate::Arena<crate::Expression>,
1120        write_expression: E,
1121    ) -> BackendResult
1122    where
1123        E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
1124    {
1125        use crate::Expression;
1126
1127        match expressions[expr] {
1128            Expression::Literal(literal) => match literal {
1129                crate::Literal::F16(value) => write!(self.out, "{value}h")?,
1130                crate::Literal::F32(value) => write!(self.out, "{value}f")?,
1131                crate::Literal::U32(value) => write!(self.out, "{value}u")?,
1132                crate::Literal::I32(value) => {
1133                    // `-2147483648i` is not valid WGSL. The most negative `i32`
1134                    // value can only be expressed in WGSL using AbstractInt and
1135                    // a unary negation operator.
1136                    if value == i32::MIN {
1137                        write!(self.out, "i32({value})")?;
1138                    } else {
1139                        write!(self.out, "{value}i")?;
1140                    }
1141                }
1142                crate::Literal::Bool(value) => write!(self.out, "{value}")?,
1143                crate::Literal::F64(value) => write!(self.out, "{value:?}lf")?,
1144                crate::Literal::I64(value) => {
1145                    // `-9223372036854775808li` is not valid WGSL. Nor can we simply use the
1146                    // AbstractInt trick above, as AbstractInt also cannot represent
1147                    // `9223372036854775808`. Instead construct the second most negative
1148                    // AbstractInt, subtract one from it, then cast to i64.
1149                    if value == i64::MIN {
1150                        write!(self.out, "i64({} - 1)", value + 1)?;
1151                    } else {
1152                        write!(self.out, "{value}li")?;
1153                    }
1154                }
1155                crate::Literal::U64(value) => write!(self.out, "{value:?}lu")?,
1156                crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
1157                    return Err(Error::Custom(
1158                        "Abstract types should not appear in IR presented to backends".into(),
1159                    ));
1160                }
1161            },
1162            Expression::Constant(handle) => {
1163                let constant = &module.constants[handle];
1164                if constant.name.is_some() {
1165                    write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
1166                } else {
1167                    self.write_const_expression(module, constant.init, &module.global_expressions)?;
1168                }
1169            }
1170            Expression::ZeroValue(ty) => {
1171                self.write_type(module, ty)?;
1172                write!(self.out, "()")?;
1173            }
1174            Expression::Compose { ty, ref components } => {
1175                self.write_type(module, ty)?;
1176                write!(self.out, "(")?;
1177                for (index, component) in components.iter().enumerate() {
1178                    if index != 0 {
1179                        write!(self.out, ", ")?;
1180                    }
1181                    write_expression(self, *component)?;
1182                }
1183                write!(self.out, ")")?
1184            }
1185            Expression::Splat { size, value } => {
1186                let size = common::vector_size_str(size);
1187                write!(self.out, "vec{size}(")?;
1188                write_expression(self, value)?;
1189                write!(self.out, ")")?;
1190            }
1191            _ => unreachable!(),
1192        }
1193
1194        Ok(())
1195    }
1196
1197    /// Write the 'plain form' of `expr`.
1198    ///
1199    /// An expression's 'plain form' is the most general rendition of that
1200    /// expression into WGSL, lacking `&` or `*` operators. The plain forms of
1201    /// `LocalVariable(x)` and `GlobalVariable(g)` are simply `x` and `g`. Such
1202    /// Naga expressions represent both WGSL pointers and references; it's the
1203    /// caller's responsibility to distinguish those cases appropriately.
1204    fn write_expr_plain_form(
1205        &mut self,
1206        module: &Module,
1207        expr: Handle<crate::Expression>,
1208        func_ctx: &back::FunctionCtx<'_>,
1209        indirection: Indirection,
1210    ) -> BackendResult {
1211        use crate::Expression;
1212
1213        if let Some(name) = self.named_expressions.get(&expr) {
1214            write!(self.out, "{name}")?;
1215            return Ok(());
1216        }
1217
1218        let expression = &func_ctx.expressions[expr];
1219
1220        // Write the plain WGSL form of a Naga expression.
1221        //
1222        // The plain form of `LocalVariable` and `GlobalVariable` expressions is
1223        // simply the variable name; `*` and `&` operators are never emitted.
1224        //
1225        // The plain form of `Access` and `AccessIndex` expressions are WGSL
1226        // `postfix_expression` forms for member/component access and
1227        // subscripting.
1228        match *expression {
1229            Expression::Literal(_)
1230            | Expression::Constant(_)
1231            | Expression::ZeroValue(_)
1232            | Expression::Compose { .. }
1233            | Expression::Splat { .. } => {
1234                self.write_possibly_const_expression(
1235                    module,
1236                    expr,
1237                    func_ctx.expressions,
1238                    |writer, expr| writer.write_expr(module, expr, func_ctx),
1239                )?;
1240            }
1241            Expression::Override(_) => unreachable!(),
1242            Expression::FunctionArgument(pos) => {
1243                let name_key = func_ctx.argument_key(pos);
1244                let name = &self.names[&name_key];
1245                write!(self.out, "{name}")?;
1246            }
1247            Expression::Binary { op, left, right } => {
1248                write!(self.out, "(")?;
1249                self.write_expr(module, left, func_ctx)?;
1250                write!(self.out, " {} ", back::binary_operation_str(op))?;
1251                self.write_expr(module, right, func_ctx)?;
1252                write!(self.out, ")")?;
1253            }
1254            Expression::Access { base, index } => {
1255                self.write_expr_with_indirection(module, base, func_ctx, indirection)?;
1256                write!(self.out, "[")?;
1257                self.write_expr(module, index, func_ctx)?;
1258                write!(self.out, "]")?
1259            }
1260            Expression::AccessIndex { base, index } => {
1261                let base_ty_res = &func_ctx.info[base].ty;
1262                let mut resolved = base_ty_res.inner_with(&module.types);
1263
1264                self.write_expr_with_indirection(module, base, func_ctx, indirection)?;
1265
1266                let base_ty_handle = match *resolved {
1267                    TypeInner::Pointer { base, space: _ } => {
1268                        resolved = &module.types[base].inner;
1269                        Some(base)
1270                    }
1271                    _ => base_ty_res.handle(),
1272                };
1273
1274                match *resolved {
1275                    TypeInner::Vector { .. } => {
1276                        // Write vector access as a swizzle
1277                        write!(self.out, ".{}", back::COMPONENTS[index as usize])?
1278                    }
1279                    TypeInner::Matrix { .. }
1280                    | TypeInner::Array { .. }
1281                    | TypeInner::BindingArray { .. }
1282                    | TypeInner::ValuePointer { .. } => write!(self.out, "[{index}]")?,
1283                    TypeInner::Struct { .. } => {
1284                        // This will never panic in case the type is a `Struct`, this is not true
1285                        // for other types so we can only check while inside this match arm
1286                        let ty = base_ty_handle.unwrap();
1287
1288                        write!(
1289                            self.out,
1290                            ".{}",
1291                            &self.names[&NameKey::StructMember(ty, index)]
1292                        )?
1293                    }
1294                    ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))),
1295                }
1296            }
1297            Expression::ImageSample {
1298                image,
1299                sampler,
1300                gather: None,
1301                coordinate,
1302                array_index,
1303                offset,
1304                level,
1305                depth_ref,
1306                clamp_to_edge,
1307            } => {
1308                use crate::SampleLevel as Sl;
1309
1310                let suffix_cmp = match depth_ref {
1311                    Some(_) => "Compare",
1312                    None => "",
1313                };
1314                let suffix_level = match level {
1315                    Sl::Auto => "",
1316                    Sl::Zero if clamp_to_edge => "BaseClampToEdge",
1317                    Sl::Zero | Sl::Exact(_) => "Level",
1318                    Sl::Bias(_) => "Bias",
1319                    Sl::Gradient { .. } => "Grad",
1320                };
1321
1322                write!(self.out, "textureSample{suffix_cmp}{suffix_level}(")?;
1323                self.write_expr(module, image, func_ctx)?;
1324                write!(self.out, ", ")?;
1325                self.write_expr(module, sampler, func_ctx)?;
1326                write!(self.out, ", ")?;
1327                self.write_expr(module, coordinate, func_ctx)?;
1328
1329                if let Some(array_index) = array_index {
1330                    write!(self.out, ", ")?;
1331                    self.write_expr(module, array_index, func_ctx)?;
1332                }
1333
1334                if let Some(depth_ref) = depth_ref {
1335                    write!(self.out, ", ")?;
1336                    self.write_expr(module, depth_ref, func_ctx)?;
1337                }
1338
1339                match level {
1340                    Sl::Auto => {}
1341                    Sl::Zero => {
1342                        // Level 0 is implied for depth comparison and BaseClampToEdge
1343                        if depth_ref.is_none() && !clamp_to_edge {
1344                            write!(self.out, ", 0.0")?;
1345                        }
1346                    }
1347                    Sl::Exact(expr) => {
1348                        write!(self.out, ", ")?;
1349                        self.write_expr(module, expr, func_ctx)?;
1350                    }
1351                    Sl::Bias(expr) => {
1352                        write!(self.out, ", ")?;
1353                        self.write_expr(module, expr, func_ctx)?;
1354                    }
1355                    Sl::Gradient { x, y } => {
1356                        write!(self.out, ", ")?;
1357                        self.write_expr(module, x, func_ctx)?;
1358                        write!(self.out, ", ")?;
1359                        self.write_expr(module, y, func_ctx)?;
1360                    }
1361                }
1362
1363                if let Some(offset) = offset {
1364                    write!(self.out, ", ")?;
1365                    self.write_const_expression(module, offset, func_ctx.expressions)?;
1366                }
1367
1368                write!(self.out, ")")?;
1369            }
1370
1371            Expression::ImageSample {
1372                image,
1373                sampler,
1374                gather: Some(component),
1375                coordinate,
1376                array_index,
1377                offset,
1378                level: _,
1379                depth_ref,
1380                clamp_to_edge: _,
1381            } => {
1382                let suffix_cmp = match depth_ref {
1383                    Some(_) => "Compare",
1384                    None => "",
1385                };
1386
1387                write!(self.out, "textureGather{suffix_cmp}(")?;
1388                match *func_ctx.resolve_type(image, &module.types) {
1389                    TypeInner::Image {
1390                        class: crate::ImageClass::Depth { multi: _ },
1391                        ..
1392                    } => {}
1393                    _ => {
1394                        write!(self.out, "{}, ", component as u8)?;
1395                    }
1396                }
1397                self.write_expr(module, image, func_ctx)?;
1398                write!(self.out, ", ")?;
1399                self.write_expr(module, sampler, func_ctx)?;
1400                write!(self.out, ", ")?;
1401                self.write_expr(module, coordinate, func_ctx)?;
1402
1403                if let Some(array_index) = array_index {
1404                    write!(self.out, ", ")?;
1405                    self.write_expr(module, array_index, func_ctx)?;
1406                }
1407
1408                if let Some(depth_ref) = depth_ref {
1409                    write!(self.out, ", ")?;
1410                    self.write_expr(module, depth_ref, func_ctx)?;
1411                }
1412
1413                if let Some(offset) = offset {
1414                    write!(self.out, ", ")?;
1415                    self.write_const_expression(module, offset, func_ctx.expressions)?;
1416                }
1417
1418                write!(self.out, ")")?;
1419            }
1420            Expression::ImageQuery { image, query } => {
1421                use crate::ImageQuery as Iq;
1422
1423                let texture_function = match query {
1424                    Iq::Size { .. } => "textureDimensions",
1425                    Iq::NumLevels => "textureNumLevels",
1426                    Iq::NumLayers => "textureNumLayers",
1427                    Iq::NumSamples => "textureNumSamples",
1428                };
1429
1430                write!(self.out, "{texture_function}(")?;
1431                self.write_expr(module, image, func_ctx)?;
1432                if let Iq::Size { level: Some(level) } = query {
1433                    write!(self.out, ", ")?;
1434                    self.write_expr(module, level, func_ctx)?;
1435                };
1436                write!(self.out, ")")?;
1437            }
1438
1439            Expression::ImageLoad {
1440                image,
1441                coordinate,
1442                array_index,
1443                sample,
1444                level,
1445            } => {
1446                write!(self.out, "textureLoad(")?;
1447                self.write_expr(module, image, func_ctx)?;
1448                write!(self.out, ", ")?;
1449                self.write_expr(module, coordinate, func_ctx)?;
1450                if let Some(array_index) = array_index {
1451                    write!(self.out, ", ")?;
1452                    self.write_expr(module, array_index, func_ctx)?;
1453                }
1454                if let Some(index) = sample.or(level) {
1455                    write!(self.out, ", ")?;
1456                    self.write_expr(module, index, func_ctx)?;
1457                }
1458                write!(self.out, ")")?;
1459            }
1460            Expression::GlobalVariable(handle) => {
1461                let name = &self.names[&NameKey::GlobalVariable(handle)];
1462                write!(self.out, "{name}")?;
1463            }
1464
1465            Expression::As {
1466                expr,
1467                kind,
1468                convert,
1469            } => {
1470                let inner = func_ctx.resolve_type(expr, &module.types);
1471                match *inner {
1472                    TypeInner::Matrix {
1473                        columns,
1474                        rows,
1475                        scalar,
1476                    } => {
1477                        let scalar = crate::Scalar {
1478                            kind,
1479                            width: convert.unwrap_or(scalar.width),
1480                        };
1481                        let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
1482                        write!(
1483                            self.out,
1484                            "mat{}x{}<{}>",
1485                            common::vector_size_str(columns),
1486                            common::vector_size_str(rows),
1487                            scalar_kind_str
1488                        )?;
1489                    }
1490                    TypeInner::Vector {
1491                        size,
1492                        scalar: crate::Scalar { width, .. },
1493                    } => {
1494                        let scalar = crate::Scalar {
1495                            kind,
1496                            width: convert.unwrap_or(width),
1497                        };
1498                        let vector_size_str = common::vector_size_str(size);
1499                        let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
1500                        if convert.is_some() {
1501                            write!(self.out, "vec{vector_size_str}<{scalar_kind_str}>")?;
1502                        } else {
1503                            write!(self.out, "bitcast<vec{vector_size_str}<{scalar_kind_str}>>")?;
1504                        }
1505                    }
1506                    TypeInner::Scalar(crate::Scalar { width, .. }) => {
1507                        let scalar = crate::Scalar {
1508                            kind,
1509                            width: convert.unwrap_or(width),
1510                        };
1511                        let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
1512                        if convert.is_some() {
1513                            write!(self.out, "{scalar_kind_str}")?
1514                        } else {
1515                            write!(self.out, "bitcast<{scalar_kind_str}>")?
1516                        }
1517                    }
1518                    _ => {
1519                        return Err(Error::Unimplemented(format!(
1520                            "write_expr expression::as {inner:?}"
1521                        )));
1522                    }
1523                };
1524                write!(self.out, "(")?;
1525                self.write_expr(module, expr, func_ctx)?;
1526                write!(self.out, ")")?;
1527            }
1528            Expression::Load { pointer } => {
1529                let is_atomic_pointer = func_ctx
1530                    .resolve_type(pointer, &module.types)
1531                    .is_atomic_pointer(&module.types);
1532
1533                if is_atomic_pointer {
1534                    write!(self.out, "atomicLoad(")?;
1535                    self.write_expr(module, pointer, func_ctx)?;
1536                    write!(self.out, ")")?;
1537                } else {
1538                    self.write_expr_with_indirection(
1539                        module,
1540                        pointer,
1541                        func_ctx,
1542                        Indirection::Reference,
1543                    )?;
1544                }
1545            }
1546            Expression::LocalVariable(handle) => {
1547                write!(self.out, "{}", self.names[&func_ctx.name_key(handle)])?
1548            }
1549            Expression::ArrayLength(expr) => {
1550                write!(self.out, "arrayLength(")?;
1551                self.write_expr(module, expr, func_ctx)?;
1552                write!(self.out, ")")?;
1553            }
1554
1555            Expression::Math {
1556                fun,
1557                arg,
1558                arg1,
1559                arg2,
1560                arg3,
1561            } => {
1562                use crate::MathFunction as Mf;
1563
1564                enum Function {
1565                    Regular(&'static str),
1566                    InversePolyfill(InversePolyfill),
1567                }
1568
1569                let function = match fun.try_to_wgsl() {
1570                    Some(name) => Function::Regular(name),
1571                    None => match fun {
1572                        Mf::Inverse => {
1573                            let ty = func_ctx.resolve_type(arg, &module.types);
1574                            let Some(overload) = InversePolyfill::find_overload(ty) else {
1575                                return Err(Error::unsupported("math function", fun));
1576                            };
1577
1578                            Function::InversePolyfill(overload)
1579                        }
1580                        _ => return Err(Error::unsupported("math function", fun)),
1581                    },
1582                };
1583
1584                match function {
1585                    Function::Regular(fun_name) => {
1586                        write!(self.out, "{fun_name}(")?;
1587                        self.write_expr(module, arg, func_ctx)?;
1588                        for arg in IntoIterator::into_iter([arg1, arg2, arg3]).flatten() {
1589                            write!(self.out, ", ")?;
1590                            self.write_expr(module, arg, func_ctx)?;
1591                        }
1592                        write!(self.out, ")")?
1593                    }
1594                    Function::InversePolyfill(inverse) => {
1595                        write!(self.out, "{}(", inverse.fun_name)?;
1596                        self.write_expr(module, arg, func_ctx)?;
1597                        write!(self.out, ")")?;
1598                        self.required_polyfills.insert(inverse);
1599                    }
1600                }
1601            }
1602
1603            Expression::Swizzle {
1604                size,
1605                vector,
1606                pattern,
1607            } => {
1608                self.write_expr(module, vector, func_ctx)?;
1609                write!(self.out, ".")?;
1610                for &sc in pattern[..size as usize].iter() {
1611                    self.out.write_char(back::COMPONENTS[sc as usize])?;
1612                }
1613            }
1614            Expression::Unary { op, expr } => {
1615                let unary = match op {
1616                    crate::UnaryOperator::Negate => "-",
1617                    crate::UnaryOperator::LogicalNot => "!",
1618                    crate::UnaryOperator::BitwiseNot => "~",
1619                };
1620
1621                write!(self.out, "{unary}(")?;
1622                self.write_expr(module, expr, func_ctx)?;
1623
1624                write!(self.out, ")")?
1625            }
1626
1627            Expression::Select {
1628                condition,
1629                accept,
1630                reject,
1631            } => {
1632                write!(self.out, "select(")?;
1633                self.write_expr(module, reject, func_ctx)?;
1634                write!(self.out, ", ")?;
1635                self.write_expr(module, accept, func_ctx)?;
1636                write!(self.out, ", ")?;
1637                self.write_expr(module, condition, func_ctx)?;
1638                write!(self.out, ")")?
1639            }
1640            Expression::Derivative { axis, ctrl, expr } => {
1641                use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl};
1642                let op = match (axis, ctrl) {
1643                    (Axis::X, Ctrl::Coarse) => "dpdxCoarse",
1644                    (Axis::X, Ctrl::Fine) => "dpdxFine",
1645                    (Axis::X, Ctrl::None) => "dpdx",
1646                    (Axis::Y, Ctrl::Coarse) => "dpdyCoarse",
1647                    (Axis::Y, Ctrl::Fine) => "dpdyFine",
1648                    (Axis::Y, Ctrl::None) => "dpdy",
1649                    (Axis::Width, Ctrl::Coarse) => "fwidthCoarse",
1650                    (Axis::Width, Ctrl::Fine) => "fwidthFine",
1651                    (Axis::Width, Ctrl::None) => "fwidth",
1652                };
1653                write!(self.out, "{op}(")?;
1654                self.write_expr(module, expr, func_ctx)?;
1655                write!(self.out, ")")?
1656            }
1657            Expression::Relational { fun, argument } => {
1658                use crate::RelationalFunction as Rf;
1659
1660                let fun_name = match fun {
1661                    Rf::All => "all",
1662                    Rf::Any => "any",
1663                    _ => return Err(Error::UnsupportedRelationalFunction(fun)),
1664                };
1665                write!(self.out, "{fun_name}(")?;
1666
1667                self.write_expr(module, argument, func_ctx)?;
1668
1669                write!(self.out, ")")?
1670            }
1671            // Not supported yet
1672            Expression::RayQueryGetIntersection { .. }
1673            | Expression::RayQueryVertexPositions { .. } => unreachable!(),
1674            // Nothing to do here, since call expression already cached
1675            Expression::CallResult(_)
1676            | Expression::AtomicResult { .. }
1677            | Expression::RayQueryProceedResult
1678            | Expression::SubgroupBallotResult
1679            | Expression::SubgroupOperationResult { .. }
1680            | Expression::WorkGroupUniformLoadResult { .. } => {}
1681        }
1682
1683        Ok(())
1684    }
1685
1686    /// Helper method used to write global variables
1687    /// # Notes
1688    /// Always adds a newline
1689    fn write_global(
1690        &mut self,
1691        module: &Module,
1692        global: &crate::GlobalVariable,
1693        handle: Handle<crate::GlobalVariable>,
1694    ) -> BackendResult {
1695        // Write group and binding attributes if present
1696        if let Some(ref binding) = global.binding {
1697            self.write_attributes(&[
1698                Attribute::Group(binding.group),
1699                Attribute::Binding(binding.binding),
1700            ])?;
1701            writeln!(self.out)?;
1702        }
1703
1704        // First write global name and address space if supported
1705        write!(self.out, "var")?;
1706        let (address, maybe_access) = address_space_str(global.space);
1707        if let Some(space) = address {
1708            write!(self.out, "<{space}")?;
1709            if let Some(access) = maybe_access {
1710                write!(self.out, ", {access}")?;
1711            }
1712            write!(self.out, ">")?;
1713        }
1714        write!(
1715            self.out,
1716            " {}: ",
1717            &self.names[&NameKey::GlobalVariable(handle)]
1718        )?;
1719
1720        // Write global type
1721        self.write_type(module, global.ty)?;
1722
1723        // Write initializer
1724        if let Some(init) = global.init {
1725            write!(self.out, " = ")?;
1726            self.write_const_expression(module, init, &module.global_expressions)?;
1727        }
1728
1729        // End with semicolon
1730        writeln!(self.out, ";")?;
1731
1732        Ok(())
1733    }
1734
1735    /// Helper method used to write global constants
1736    ///
1737    /// # Notes
1738    /// Ends in a newline
1739    fn write_global_constant(
1740        &mut self,
1741        module: &Module,
1742        handle: Handle<crate::Constant>,
1743    ) -> BackendResult {
1744        let name = &self.names[&NameKey::Constant(handle)];
1745        // First write only constant name
1746        write!(self.out, "const {name}: ")?;
1747        self.write_type(module, module.constants[handle].ty)?;
1748        write!(self.out, " = ")?;
1749        let init = module.constants[handle].init;
1750        self.write_const_expression(module, init, &module.global_expressions)?;
1751        writeln!(self.out, ";")?;
1752
1753        Ok(())
1754    }
1755
1756    // See https://github.com/rust-lang/rust-clippy/issues/4979.
1757    #[allow(clippy::missing_const_for_fn)]
1758    pub fn finish(self) -> W {
1759        self.out
1760    }
1761}
1762
1763struct WriterTypeContext<'m> {
1764    module: &'m Module,
1765    names: &'m crate::FastHashMap<NameKey, String>,
1766}
1767
1768impl TypeContext for WriterTypeContext<'_> {
1769    fn lookup_type(&self, handle: Handle<crate::Type>) -> &crate::Type {
1770        &self.module.types[handle]
1771    }
1772
1773    fn type_name(&self, handle: Handle<crate::Type>) -> &str {
1774        self.names[&NameKey::Type(handle)].as_str()
1775    }
1776
1777    fn write_unnamed_struct<W: Write>(&self, _: &TypeInner, _: &mut W) -> core::fmt::Result {
1778        unreachable!("the WGSL back end should always provide type handles");
1779    }
1780
1781    fn write_override<W: Write>(&self, _: Handle<crate::Override>, _: &mut W) -> core::fmt::Result {
1782        unreachable!("overrides should be validated out");
1783    }
1784
1785    fn write_non_wgsl_inner<W: Write>(&self, _: &TypeInner, _: &mut W) -> core::fmt::Result {
1786        unreachable!("backends should only be passed validated modules");
1787    }
1788
1789    fn write_non_wgsl_scalar<W: Write>(&self, _: crate::Scalar, _: &mut W) -> core::fmt::Result {
1790        unreachable!("backends should only be passed validated modules");
1791    }
1792}
1793
1794fn map_binding_to_attribute(binding: &crate::Binding) -> Vec<Attribute> {
1795    match *binding {
1796        crate::Binding::BuiltIn(built_in) => {
1797            if let crate::BuiltIn::Position { invariant: true } = built_in {
1798                vec![Attribute::BuiltIn(built_in), Attribute::Invariant]
1799            } else {
1800                vec![Attribute::BuiltIn(built_in)]
1801            }
1802        }
1803        crate::Binding::Location {
1804            location,
1805            interpolation,
1806            sampling,
1807            blend_src: None,
1808        } => vec![
1809            Attribute::Location(location),
1810            Attribute::Interpolate(interpolation, sampling),
1811        ],
1812        crate::Binding::Location {
1813            location,
1814            interpolation,
1815            sampling,
1816            blend_src: Some(blend_src),
1817        } => vec![
1818            Attribute::Location(location),
1819            Attribute::BlendSrc(blend_src),
1820            Attribute::Interpolate(interpolation, sampling),
1821        ],
1822    }
1823}