naga/front/wgsl/lower/
mod.rs

1use alloc::{
2    borrow::ToOwned,
3    boxed::Box,
4    format,
5    string::{String, ToString},
6    vec::Vec,
7};
8use core::num::NonZeroU32;
9
10use crate::front::wgsl::error::{Error, ExpectedToken, InvalidAssignmentType};
11use crate::front::wgsl::index::Index;
12use crate::front::wgsl::parse::number::Number;
13use crate::front::wgsl::parse::{ast, conv};
14use crate::front::wgsl::Result;
15use crate::front::Typifier;
16use crate::{
17    common::wgsl::{TryToWgsl, TypeContext},
18    compact::KeepUnused,
19};
20use crate::{common::ForDebugWithTypes, proc::LayoutErrorInner};
21use crate::{ir, proc};
22use crate::{Arena, FastHashMap, FastIndexMap, Handle, Span};
23
24mod construction;
25mod conversion;
26
27/// Resolves the inner type of a given expression.
28///
29/// Expects a &mut [`ExpressionContext`] and a [`Handle<Expression>`].
30///
31/// Returns a &[`ir::TypeInner`].
32///
33/// Ideally, we would simply have a function that takes a `&mut ExpressionContext`
34/// and returns a `&TypeResolution`. Unfortunately, this leads the borrow checker
35/// to conclude that the mutable borrow lasts for as long as we are using the
36/// `&TypeResolution`, so we can't use the `ExpressionContext` for anything else -
37/// like, say, resolving another operand's type. Using a macro that expands to
38/// two separate calls, only the first of which needs a `&mut`,
39/// lets the borrow checker see that the mutable borrow is over.
40macro_rules! resolve_inner {
41    ($ctx:ident, $expr:expr) => {{
42        $ctx.grow_types($expr)?;
43        $ctx.typifier()[$expr].inner_with(&$ctx.module.types)
44    }};
45}
46pub(super) use resolve_inner;
47
48/// Resolves the inner types of two given expressions.
49///
50/// Expects a &mut [`ExpressionContext`] and two [`Handle<Expression>`]s.
51///
52/// Returns a tuple containing two &[`ir::TypeInner`].
53///
54/// See the documentation of [`resolve_inner!`] for why this macro is necessary.
55macro_rules! resolve_inner_binary {
56    ($ctx:ident, $left:expr, $right:expr) => {{
57        $ctx.grow_types($left)?;
58        $ctx.grow_types($right)?;
59        (
60            $ctx.typifier()[$left].inner_with(&$ctx.module.types),
61            $ctx.typifier()[$right].inner_with(&$ctx.module.types),
62        )
63    }};
64}
65
66/// Resolves the type of a given expression.
67///
68/// Expects a &mut [`ExpressionContext`] and a [`Handle<Expression>`].
69///
70/// Returns a &[`TypeResolution`].
71///
72/// See the documentation of [`resolve_inner!`] for why this macro is necessary.
73///
74/// [`TypeResolution`]: proc::TypeResolution
75macro_rules! resolve {
76    ($ctx:ident, $expr:expr) => {{
77        let expr = $expr;
78        $ctx.grow_types(expr)?;
79        &$ctx.typifier()[expr]
80    }};
81}
82pub(super) use resolve;
83
84/// State for constructing a `ir::Module`.
85pub struct GlobalContext<'source, 'temp, 'out> {
86    /// The `TranslationUnit`'s expressions arena.
87    ast_expressions: &'temp Arena<ast::Expression<'source>>,
88
89    /// The `TranslationUnit`'s types arena.
90    types: &'temp Arena<ast::Type<'source>>,
91
92    // Naga IR values.
93    /// The map from the names of module-scope declarations to the Naga IR
94    /// `Handle`s we have built for them, owned by `Lowerer::lower`.
95    globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
96
97    /// The module we're constructing.
98    module: &'out mut ir::Module,
99
100    const_typifier: &'temp mut Typifier,
101
102    layouter: &'temp mut proc::Layouter,
103
104    global_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
105}
106
107impl<'source> GlobalContext<'source, '_, '_> {
108    fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
109        ExpressionContext {
110            ast_expressions: self.ast_expressions,
111            globals: self.globals,
112            types: self.types,
113            module: self.module,
114            const_typifier: self.const_typifier,
115            layouter: self.layouter,
116            expr_type: ExpressionContextType::Constant(None),
117            global_expression_kind_tracker: self.global_expression_kind_tracker,
118        }
119    }
120
121    fn as_override(&mut self) -> ExpressionContext<'source, '_, '_> {
122        ExpressionContext {
123            ast_expressions: self.ast_expressions,
124            globals: self.globals,
125            types: self.types,
126            module: self.module,
127            const_typifier: self.const_typifier,
128            layouter: self.layouter,
129            expr_type: ExpressionContextType::Override,
130            global_expression_kind_tracker: self.global_expression_kind_tracker,
131        }
132    }
133
134    fn ensure_type_exists(
135        &mut self,
136        name: Option<String>,
137        inner: ir::TypeInner,
138    ) -> Handle<ir::Type> {
139        self.module
140            .types
141            .insert(ir::Type { inner, name }, Span::UNDEFINED)
142    }
143}
144
145/// State for lowering a statement within a function.
146pub struct StatementContext<'source, 'temp, 'out> {
147    // WGSL AST values.
148    /// A reference to [`TranslationUnit::expressions`] for the translation unit
149    /// we're lowering.
150    ///
151    /// [`TranslationUnit::expressions`]: ast::TranslationUnit::expressions
152    ast_expressions: &'temp Arena<ast::Expression<'source>>,
153
154    /// A reference to [`TranslationUnit::types`] for the translation unit
155    /// we're lowering.
156    ///
157    /// [`TranslationUnit::types`]: ast::TranslationUnit::types
158    types: &'temp Arena<ast::Type<'source>>,
159
160    // Naga IR values.
161    /// The map from the names of module-scope declarations to the Naga IR
162    /// `Handle`s we have built for them, owned by `Lowerer::lower`.
163    globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
164
165    /// A map from each `ast::Local` handle to the Naga expression
166    /// we've built for it:
167    ///
168    /// - WGSL function arguments become Naga [`FunctionArgument`] expressions.
169    ///
170    /// - WGSL `var` declarations become Naga [`LocalVariable`] expressions.
171    ///
172    /// - WGSL `let` declararations become arbitrary Naga expressions.
173    ///
174    /// This always borrows the `local_table` local variable in
175    /// [`Lowerer::function`].
176    ///
177    /// [`LocalVariable`]: ir::Expression::LocalVariable
178    /// [`FunctionArgument`]: ir::Expression::FunctionArgument
179    local_table:
180        &'temp mut FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<ir::Expression>>>>,
181
182    const_typifier: &'temp mut Typifier,
183    typifier: &'temp mut Typifier,
184    layouter: &'temp mut proc::Layouter,
185    function: &'out mut ir::Function,
186    /// Stores the names of expressions that are assigned in `let` statement
187    /// Also stores the spans of the names, for use in errors.
188    named_expressions: &'out mut FastIndexMap<Handle<ir::Expression>, (String, Span)>,
189    module: &'out mut ir::Module,
190
191    /// Which `Expression`s in `self.naga_expressions` are const expressions, in
192    /// the WGSL sense.
193    ///
194    /// According to the WGSL spec, a const expression must not refer to any
195    /// `let` declarations, even if those declarations' initializers are
196    /// themselves const expressions. So this tracker is not simply concerned
197    /// with the form of the expressions; it is also tracking whether WGSL says
198    /// we should consider them to be const. See the use of `force_non_const` in
199    /// the code for lowering `let` bindings.
200    local_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
201    global_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
202}
203
204impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
205    fn as_const<'t>(
206        &'t mut self,
207        block: &'t mut ir::Block,
208        emitter: &'t mut proc::Emitter,
209    ) -> ExpressionContext<'a, 't, 't>
210    where
211        'temp: 't,
212    {
213        ExpressionContext {
214            globals: self.globals,
215            types: self.types,
216            ast_expressions: self.ast_expressions,
217            const_typifier: self.const_typifier,
218            layouter: self.layouter,
219            global_expression_kind_tracker: self.global_expression_kind_tracker,
220            module: self.module,
221            expr_type: ExpressionContextType::Constant(Some(LocalExpressionContext {
222                local_table: self.local_table,
223                function: self.function,
224                block,
225                emitter,
226                typifier: self.typifier,
227                local_expression_kind_tracker: self.local_expression_kind_tracker,
228            })),
229        }
230    }
231
232    fn as_expression<'t>(
233        &'t mut self,
234        block: &'t mut ir::Block,
235        emitter: &'t mut proc::Emitter,
236    ) -> ExpressionContext<'a, 't, 't>
237    where
238        'temp: 't,
239    {
240        ExpressionContext {
241            globals: self.globals,
242            types: self.types,
243            ast_expressions: self.ast_expressions,
244            const_typifier: self.const_typifier,
245            layouter: self.layouter,
246            global_expression_kind_tracker: self.global_expression_kind_tracker,
247            module: self.module,
248            expr_type: ExpressionContextType::Runtime(LocalExpressionContext {
249                local_table: self.local_table,
250                function: self.function,
251                block,
252                emitter,
253                typifier: self.typifier,
254                local_expression_kind_tracker: self.local_expression_kind_tracker,
255            }),
256        }
257    }
258
259    #[allow(dead_code)]
260    fn as_global(&mut self) -> GlobalContext<'a, '_, '_> {
261        GlobalContext {
262            ast_expressions: self.ast_expressions,
263            globals: self.globals,
264            types: self.types,
265            module: self.module,
266            const_typifier: self.const_typifier,
267            layouter: self.layouter,
268            global_expression_kind_tracker: self.global_expression_kind_tracker,
269        }
270    }
271
272    fn invalid_assignment_type(&self, expr: Handle<ir::Expression>) -> InvalidAssignmentType {
273        if let Some(&(_, span)) = self.named_expressions.get(&expr) {
274            InvalidAssignmentType::ImmutableBinding(span)
275        } else {
276            match self.function.expressions[expr] {
277                ir::Expression::Swizzle { .. } => InvalidAssignmentType::Swizzle,
278                ir::Expression::Access { base, .. } => self.invalid_assignment_type(base),
279                ir::Expression::AccessIndex { base, .. } => self.invalid_assignment_type(base),
280                _ => InvalidAssignmentType::Other,
281            }
282        }
283    }
284}
285
286pub struct LocalExpressionContext<'temp, 'out> {
287    /// A map from [`ast::Local`] handles to the Naga expressions we've built for them.
288    ///
289    /// This is always [`StatementContext::local_table`] for the
290    /// enclosing statement; see that documentation for details.
291    local_table: &'temp FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<ir::Expression>>>>,
292
293    function: &'out mut ir::Function,
294    block: &'temp mut ir::Block,
295    emitter: &'temp mut proc::Emitter,
296    typifier: &'temp mut Typifier,
297
298    /// Which `Expression`s in `self.naga_expressions` are const expressions, in
299    /// the WGSL sense.
300    ///
301    /// See [`StatementContext::local_expression_kind_tracker`] for details.
302    local_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
303}
304
305/// The type of Naga IR expression we are lowering an [`ast::Expression`] to.
306pub enum ExpressionContextType<'temp, 'out> {
307    /// We are lowering to an arbitrary runtime expression, to be
308    /// included in a function's body.
309    ///
310    /// The given [`LocalExpressionContext`] holds information about local
311    /// variables, arguments, and other definitions available only to runtime
312    /// expressions, not constant or override expressions.
313    Runtime(LocalExpressionContext<'temp, 'out>),
314
315    /// We are lowering to a constant expression, to be included in the module's
316    /// constant expression arena.
317    ///
318    /// Everything global constant expressions are allowed to refer to is
319    /// available in the [`ExpressionContext`], but local constant expressions can
320    /// also refer to other
321    Constant(Option<LocalExpressionContext<'temp, 'out>>),
322
323    /// We are lowering to an override expression, to be included in the module's
324    /// constant expression arena.
325    ///
326    /// Everything override expressions are allowed to refer to is
327    /// available in the [`ExpressionContext`], so this variant
328    /// carries no further information.
329    Override,
330}
331
332/// State for lowering an [`ast::Expression`] to Naga IR.
333///
334/// [`ExpressionContext`]s come in two kinds, distinguished by
335/// the value of the [`expr_type`] field:
336///
337/// - A [`Runtime`] context contributes [`naga::Expression`]s to a [`naga::Function`]'s
338///   runtime expression arena.
339///
340/// - A [`Constant`] context contributes [`naga::Expression`]s to a [`naga::Module`]'s
341///   constant expression arena.
342///
343/// [`ExpressionContext`]s are constructed in restricted ways:
344///
345/// - To get a [`Runtime`] [`ExpressionContext`], call
346///   [`StatementContext::as_expression`].
347///
348/// - To get a [`Constant`] [`ExpressionContext`], call
349///   [`GlobalContext::as_const`].
350///
351/// - You can demote a [`Runtime`] context to a [`Constant`] context
352///   by calling [`as_const`], but there's no way to go in the other
353///   direction, producing a runtime context from a constant one. This
354///   is because runtime expressions can refer to constant
355///   expressions, via [`Expression::Constant`], but constant
356///   expressions can't refer to a function's expressions.
357///
358/// Not to be confused with `wgsl::parse::ExpressionContext`, which is
359/// for parsing the `ast::Expression` in the first place.
360///
361/// [`expr_type`]: ExpressionContext::expr_type
362/// [`Runtime`]: ExpressionContextType::Runtime
363/// [`naga::Expression`]: ir::Expression
364/// [`naga::Function`]: ir::Function
365/// [`Constant`]: ExpressionContextType::Constant
366/// [`naga::Module`]: ir::Module
367/// [`as_const`]: ExpressionContext::as_const
368/// [`Expression::Constant`]: ir::Expression::Constant
369pub struct ExpressionContext<'source, 'temp, 'out> {
370    // WGSL AST values.
371    ast_expressions: &'temp Arena<ast::Expression<'source>>,
372    types: &'temp Arena<ast::Type<'source>>,
373
374    // Naga IR values.
375    /// The map from the names of module-scope declarations to the Naga IR
376    /// `Handle`s we have built for them, owned by `Lowerer::lower`.
377    globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,
378
379    /// The IR [`Module`] we're constructing.
380    ///
381    /// [`Module`]: ir::Module
382    module: &'out mut ir::Module,
383
384    /// Type judgments for [`module::global_expressions`].
385    ///
386    /// [`module::global_expressions`]: ir::Module::global_expressions
387    const_typifier: &'temp mut Typifier,
388    layouter: &'temp mut proc::Layouter,
389    global_expression_kind_tracker: &'temp mut proc::ExpressionKindTracker,
390
391    /// Whether we are lowering a constant expression or a general
392    /// runtime expression, and the data needed in each case.
393    expr_type: ExpressionContextType<'temp, 'out>,
394}
395
396impl TypeContext for ExpressionContext<'_, '_, '_> {
397    fn lookup_type(&self, handle: Handle<ir::Type>) -> &ir::Type {
398        &self.module.types[handle]
399    }
400
401    fn type_name(&self, handle: Handle<ir::Type>) -> &str {
402        self.module.types[handle]
403            .name
404            .as_deref()
405            .unwrap_or("{anonymous type}")
406    }
407
408    fn write_override<W: core::fmt::Write>(
409        &self,
410        handle: Handle<ir::Override>,
411        out: &mut W,
412    ) -> core::fmt::Result {
413        match self.module.overrides[handle].name {
414            Some(ref name) => out.write_str(name),
415            None => write!(out, "{{anonymous override {handle:?}}}"),
416        }
417    }
418
419    fn write_unnamed_struct<W: core::fmt::Write>(
420        &self,
421        _: &ir::TypeInner,
422        _: &mut W,
423    ) -> core::fmt::Result {
424        unreachable!("the WGSL front end should always know the type name");
425    }
426}
427
428impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
429    #[allow(dead_code)]
430    fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
431        ExpressionContext {
432            globals: self.globals,
433            types: self.types,
434            ast_expressions: self.ast_expressions,
435            const_typifier: self.const_typifier,
436            layouter: self.layouter,
437            module: self.module,
438            expr_type: ExpressionContextType::Constant(match self.expr_type {
439                ExpressionContextType::Runtime(ref mut local_expression_context)
440                | ExpressionContextType::Constant(Some(ref mut local_expression_context)) => {
441                    Some(LocalExpressionContext {
442                        local_table: local_expression_context.local_table,
443                        function: local_expression_context.function,
444                        block: local_expression_context.block,
445                        emitter: local_expression_context.emitter,
446                        typifier: local_expression_context.typifier,
447                        local_expression_kind_tracker: local_expression_context
448                            .local_expression_kind_tracker,
449                    })
450                }
451                ExpressionContextType::Constant(None) | ExpressionContextType::Override => None,
452            }),
453            global_expression_kind_tracker: self.global_expression_kind_tracker,
454        }
455    }
456
457    fn as_global(&mut self) -> GlobalContext<'source, '_, '_> {
458        GlobalContext {
459            ast_expressions: self.ast_expressions,
460            globals: self.globals,
461            types: self.types,
462            module: self.module,
463            const_typifier: self.const_typifier,
464            layouter: self.layouter,
465            global_expression_kind_tracker: self.global_expression_kind_tracker,
466        }
467    }
468
469    fn as_const_evaluator(&mut self) -> proc::ConstantEvaluator<'_> {
470        match self.expr_type {
471            ExpressionContextType::Runtime(ref mut rctx) => {
472                proc::ConstantEvaluator::for_wgsl_function(
473                    self.module,
474                    &mut rctx.function.expressions,
475                    rctx.local_expression_kind_tracker,
476                    self.layouter,
477                    rctx.emitter,
478                    rctx.block,
479                    false,
480                )
481            }
482            ExpressionContextType::Constant(Some(ref mut rctx)) => {
483                proc::ConstantEvaluator::for_wgsl_function(
484                    self.module,
485                    &mut rctx.function.expressions,
486                    rctx.local_expression_kind_tracker,
487                    self.layouter,
488                    rctx.emitter,
489                    rctx.block,
490                    true,
491                )
492            }
493            ExpressionContextType::Constant(None) => proc::ConstantEvaluator::for_wgsl_module(
494                self.module,
495                self.global_expression_kind_tracker,
496                self.layouter,
497                false,
498            ),
499            ExpressionContextType::Override => proc::ConstantEvaluator::for_wgsl_module(
500                self.module,
501                self.global_expression_kind_tracker,
502                self.layouter,
503                true,
504            ),
505        }
506    }
507
508    /// Return a wrapper around `value` suitable for formatting.
509    ///
510    /// Return a wrapper around `value` that implements
511    /// [`core::fmt::Display`] in a form suitable for use in
512    /// diagnostic messages.
513    fn as_diagnostic_display<T>(
514        &self,
515        value: T,
516    ) -> crate::common::DiagnosticDisplay<(T, proc::GlobalCtx<'_>)> {
517        let ctx = self.module.to_ctx();
518        crate::common::DiagnosticDisplay((value, ctx))
519    }
520
521    fn append_expression(
522        &mut self,
523        expr: ir::Expression,
524        span: Span,
525    ) -> Result<'source, Handle<ir::Expression>> {
526        let mut eval = self.as_const_evaluator();
527        eval.try_eval_and_append(expr, span)
528            .map_err(|e| Box::new(Error::ConstantEvaluatorError(e.into(), span)))
529    }
530
531    fn const_eval_expr_to_u32(
532        &self,
533        handle: Handle<ir::Expression>,
534    ) -> core::result::Result<u32, proc::U32EvalError> {
535        match self.expr_type {
536            ExpressionContextType::Runtime(ref ctx) => {
537                if !ctx.local_expression_kind_tracker.is_const(handle) {
538                    return Err(proc::U32EvalError::NonConst);
539                }
540
541                self.module
542                    .to_ctx()
543                    .eval_expr_to_u32_from(handle, &ctx.function.expressions)
544            }
545            ExpressionContextType::Constant(Some(ref ctx)) => {
546                assert!(ctx.local_expression_kind_tracker.is_const(handle));
547                self.module
548                    .to_ctx()
549                    .eval_expr_to_u32_from(handle, &ctx.function.expressions)
550            }
551            ExpressionContextType::Constant(None) => self.module.to_ctx().eval_expr_to_u32(handle),
552            ExpressionContextType::Override => Err(proc::U32EvalError::NonConst),
553        }
554    }
555
556    /// Return `true` if `handle` is a constant expression.
557    fn is_const(&self, handle: Handle<ir::Expression>) -> bool {
558        use ExpressionContextType as Ect;
559        match self.expr_type {
560            Ect::Runtime(ref ctx) | Ect::Constant(Some(ref ctx)) => {
561                ctx.local_expression_kind_tracker.is_const(handle)
562            }
563            Ect::Constant(None) | Ect::Override => {
564                self.global_expression_kind_tracker.is_const(handle)
565            }
566        }
567    }
568
569    fn get_expression_span(&self, handle: Handle<ir::Expression>) -> Span {
570        match self.expr_type {
571            ExpressionContextType::Runtime(ref ctx)
572            | ExpressionContextType::Constant(Some(ref ctx)) => {
573                ctx.function.expressions.get_span(handle)
574            }
575            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
576                self.module.global_expressions.get_span(handle)
577            }
578        }
579    }
580
581    fn typifier(&self) -> &Typifier {
582        match self.expr_type {
583            ExpressionContextType::Runtime(ref ctx)
584            | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
585            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
586                self.const_typifier
587            }
588        }
589    }
590
591    fn local(
592        &mut self,
593        local: &Handle<ast::Local>,
594        span: Span,
595    ) -> Result<'source, Typed<Handle<ir::Expression>>> {
596        match self.expr_type {
597            ExpressionContextType::Runtime(ref ctx) => Ok(ctx.local_table[local].runtime()),
598            ExpressionContextType::Constant(Some(ref ctx)) => ctx.local_table[local]
599                .const_time()
600                .ok_or(Box::new(Error::UnexpectedOperationInConstContext(span))),
601            _ => Err(Box::new(Error::UnexpectedOperationInConstContext(span))),
602        }
603    }
604
605    fn runtime_expression_ctx(
606        &mut self,
607        span: Span,
608    ) -> Result<'source, &mut LocalExpressionContext<'temp, 'out>> {
609        match self.expr_type {
610            ExpressionContextType::Runtime(ref mut ctx) => Ok(ctx),
611            ExpressionContextType::Constant(_) | ExpressionContextType::Override => {
612                Err(Box::new(Error::UnexpectedOperationInConstContext(span)))
613            }
614        }
615    }
616
617    fn gather_component(
618        &mut self,
619        expr: Handle<ir::Expression>,
620        component_span: Span,
621        gather_span: Span,
622    ) -> Result<'source, ir::SwizzleComponent> {
623        match self.expr_type {
624            ExpressionContextType::Runtime(ref rctx) => {
625                if !rctx.local_expression_kind_tracker.is_const(expr) {
626                    return Err(Box::new(Error::ExpectedConstExprConcreteIntegerScalar(
627                        component_span,
628                    )));
629                }
630
631                let index = self
632                    .module
633                    .to_ctx()
634                    .eval_expr_to_u32_from(expr, &rctx.function.expressions)
635                    .map_err(|err| match err {
636                        proc::U32EvalError::NonConst => {
637                            Error::ExpectedConstExprConcreteIntegerScalar(component_span)
638                        }
639                        proc::U32EvalError::Negative => Error::ExpectedNonNegative(component_span),
640                    })?;
641                ir::SwizzleComponent::XYZW
642                    .get(index as usize)
643                    .copied()
644                    .ok_or(Box::new(Error::InvalidGatherComponent(component_span)))
645            }
646            // This means a `gather` operation appeared in a constant expression.
647            // This error refers to the `gather` itself, not its "component" argument.
648            ExpressionContextType::Constant(_) | ExpressionContextType::Override => Err(Box::new(
649                Error::UnexpectedOperationInConstContext(gather_span),
650            )),
651        }
652    }
653
654    /// Determine the type of `handle`, and add it to the module's arena.
655    ///
656    /// If you just need a `TypeInner` for `handle`'s type, use the
657    /// [`resolve_inner!`] macro instead. This function
658    /// should only be used when the type of `handle` needs to appear
659    /// in the module's final `Arena<Type>`, for example, if you're
660    /// creating a [`LocalVariable`] whose type is inferred from its
661    /// initializer.
662    ///
663    /// [`LocalVariable`]: ir::LocalVariable
664    fn register_type(
665        &mut self,
666        handle: Handle<ir::Expression>,
667    ) -> Result<'source, Handle<ir::Type>> {
668        self.grow_types(handle)?;
669        // This is equivalent to calling ExpressionContext::typifier(),
670        // except that this lets the borrow checker see that it's okay
671        // to also borrow self.module.types mutably below.
672        let typifier = match self.expr_type {
673            ExpressionContextType::Runtime(ref ctx)
674            | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
675            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
676                &*self.const_typifier
677            }
678        };
679        Ok(typifier.register_type(handle, &mut self.module.types))
680    }
681
682    /// Resolve the types of all expressions up through `handle`.
683    ///
684    /// Ensure that [`self.typifier`] has a [`TypeResolution`] for
685    /// every expression in [`self.function.expressions`].
686    ///
687    /// This does not add types to any arena. The [`Typifier`]
688    /// documentation explains the steps we take to avoid filling
689    /// arenas with intermediate types.
690    ///
691    /// This function takes `&mut self`, so it can't conveniently
692    /// return a shared reference to the resulting `TypeResolution`:
693    /// the shared reference would extend the mutable borrow, and you
694    /// wouldn't be able to use `self` for anything else. Instead, you
695    /// should use [`register_type`] or one of [`resolve!`],
696    /// [`resolve_inner!`] or [`resolve_inner_binary!`].
697    ///
698    /// [`self.typifier`]: ExpressionContext::typifier
699    /// [`TypeResolution`]: proc::TypeResolution
700    /// [`register_type`]: Self::register_type
701    /// [`Typifier`]: Typifier
702    fn grow_types(&mut self, handle: Handle<ir::Expression>) -> Result<'source, &mut Self> {
703        let empty_arena = Arena::new();
704        let resolve_ctx;
705        let typifier;
706        let expressions;
707        match self.expr_type {
708            ExpressionContextType::Runtime(ref mut ctx)
709            | ExpressionContextType::Constant(Some(ref mut ctx)) => {
710                resolve_ctx = proc::ResolveContext::with_locals(
711                    self.module,
712                    &ctx.function.local_variables,
713                    &ctx.function.arguments,
714                );
715                typifier = &mut *ctx.typifier;
716                expressions = &ctx.function.expressions;
717            }
718            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
719                resolve_ctx = proc::ResolveContext::with_locals(self.module, &empty_arena, &[]);
720                typifier = self.const_typifier;
721                expressions = &self.module.global_expressions;
722            }
723        };
724        typifier
725            .grow(handle, expressions, &resolve_ctx)
726            .map_err(Error::InvalidResolve)?;
727
728        Ok(self)
729    }
730
731    fn image_data(
732        &mut self,
733        image: Handle<ir::Expression>,
734        span: Span,
735    ) -> Result<'source, (ir::ImageClass, bool)> {
736        match *resolve_inner!(self, image) {
737            ir::TypeInner::Image { class, arrayed, .. } => Ok((class, arrayed)),
738            _ => Err(Box::new(Error::BadTexture(span))),
739        }
740    }
741
742    fn prepare_args<'b>(
743        &mut self,
744        args: &'b [Handle<ast::Expression<'source>>],
745        min_args: u32,
746        span: Span,
747    ) -> ArgumentContext<'b, 'source> {
748        ArgumentContext {
749            args: args.iter(),
750            min_args,
751            args_used: 0,
752            total_args: args.len() as u32,
753            span,
754        }
755    }
756
757    /// Insert splats, if needed by the non-'*' operations.
758    ///
759    /// See the "Binary arithmetic expressions with mixed scalar and vector operands"
760    /// table in the WebGPU Shading Language specification for relevant operators.
761    ///
762    /// Multiply is not handled here as backends are expected to handle vec*scalar
763    /// operations, so inserting splats into the IR increases size needlessly.
764    fn binary_op_splat(
765        &mut self,
766        op: ir::BinaryOperator,
767        left: &mut Handle<ir::Expression>,
768        right: &mut Handle<ir::Expression>,
769    ) -> Result<'source, ()> {
770        if matches!(
771            op,
772            ir::BinaryOperator::Add
773                | ir::BinaryOperator::Subtract
774                | ir::BinaryOperator::Divide
775                | ir::BinaryOperator::Modulo
776        ) {
777            match resolve_inner_binary!(self, *left, *right) {
778                (&ir::TypeInner::Vector { size, .. }, &ir::TypeInner::Scalar { .. }) => {
779                    *right = self.append_expression(
780                        ir::Expression::Splat {
781                            size,
782                            value: *right,
783                        },
784                        self.get_expression_span(*right),
785                    )?;
786                }
787                (&ir::TypeInner::Scalar { .. }, &ir::TypeInner::Vector { size, .. }) => {
788                    *left = self.append_expression(
789                        ir::Expression::Splat { size, value: *left },
790                        self.get_expression_span(*left),
791                    )?;
792                }
793                _ => {}
794            }
795        }
796
797        Ok(())
798    }
799
800    /// Add a single expression to the expression table that is not covered by `self.emitter`.
801    ///
802    /// This is useful for `CallResult` and `AtomicResult` expressions, which should not be covered by
803    /// `Emit` statements.
804    fn interrupt_emitter(
805        &mut self,
806        expression: ir::Expression,
807        span: Span,
808    ) -> Result<'source, Handle<ir::Expression>> {
809        match self.expr_type {
810            ExpressionContextType::Runtime(ref mut rctx)
811            | ExpressionContextType::Constant(Some(ref mut rctx)) => {
812                rctx.block
813                    .extend(rctx.emitter.finish(&rctx.function.expressions));
814            }
815            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
816        }
817        let result = self.append_expression(expression, span);
818        match self.expr_type {
819            ExpressionContextType::Runtime(ref mut rctx)
820            | ExpressionContextType::Constant(Some(ref mut rctx)) => {
821                rctx.emitter.start(&rctx.function.expressions);
822            }
823            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
824        }
825        result
826    }
827
828    /// Apply the WGSL Load Rule to `expr`.
829    ///
830    /// If `expr` is has type `ref<SC, T, A>`, perform a load to produce a value of type
831    /// `T`. Otherwise, return `expr` unchanged.
832    fn apply_load_rule(
833        &mut self,
834        expr: Typed<Handle<ir::Expression>>,
835    ) -> Result<'source, Handle<ir::Expression>> {
836        match expr {
837            Typed::Reference(pointer) => {
838                let load = ir::Expression::Load { pointer };
839                let span = self.get_expression_span(pointer);
840                self.append_expression(load, span)
841            }
842            Typed::Plain(handle) => Ok(handle),
843        }
844    }
845
846    fn ensure_type_exists(&mut self, inner: ir::TypeInner) -> Handle<ir::Type> {
847        self.as_global().ensure_type_exists(None, inner)
848    }
849}
850
851struct ArgumentContext<'ctx, 'source> {
852    args: core::slice::Iter<'ctx, Handle<ast::Expression<'source>>>,
853    min_args: u32,
854    args_used: u32,
855    total_args: u32,
856    span: Span,
857}
858
859impl<'source> ArgumentContext<'_, 'source> {
860    pub fn finish(self) -> Result<'source, ()> {
861        if self.args.len() == 0 {
862            Ok(())
863        } else {
864            Err(Box::new(Error::WrongArgumentCount {
865                found: self.total_args,
866                expected: self.min_args..self.args_used + 1,
867                span: self.span,
868            }))
869        }
870    }
871
872    pub fn next(&mut self) -> Result<'source, Handle<ast::Expression<'source>>> {
873        match self.args.next().copied() {
874            Some(arg) => {
875                self.args_used += 1;
876                Ok(arg)
877            }
878            None => Err(Box::new(Error::WrongArgumentCount {
879                found: self.total_args,
880                expected: self.min_args..self.args_used + 1,
881                span: self.span,
882            })),
883        }
884    }
885}
886
887#[derive(Debug, Copy, Clone)]
888enum Declared<T> {
889    /// Value declared as const
890    Const(T),
891
892    /// Value declared as non-const
893    Runtime(T),
894}
895
896impl<T> Declared<T> {
897    fn runtime(self) -> T {
898        match self {
899            Declared::Const(t) | Declared::Runtime(t) => t,
900        }
901    }
902
903    fn const_time(self) -> Option<T> {
904        match self {
905            Declared::Const(t) => Some(t),
906            Declared::Runtime(_) => None,
907        }
908    }
909}
910
911/// WGSL type annotations on expressions, types, values, etc.
912///
913/// Naga and WGSL types are very close, but Naga lacks WGSL's `ref` types, which
914/// we need to know to apply the Load Rule. This enum carries some WGSL or Naga
915/// datum along with enough information to determine its corresponding WGSL
916/// type.
917///
918/// The `T` type parameter can be any expression-like thing:
919///
920/// - `Typed<Handle<ir::Type>>` can represent a full WGSL type. For example,
921///   given some Naga `Pointer` type `ptr`, a WGSL reference type is a
922///   `Typed::Reference(ptr)` whereas a WGSL pointer type is a
923///   `Typed::Plain(ptr)`.
924///
925/// - `Typed<ir::Expression>` or `Typed<Handle<ir::Expression>>` can
926///   represent references similarly.
927///
928/// Use the `map` and `try_map` methods to convert from one expression
929/// representation to another.
930///
931/// [`Expression`]: ir::Expression
932#[derive(Debug, Copy, Clone)]
933enum Typed<T> {
934    /// A WGSL reference.
935    Reference(T),
936
937    /// A WGSL plain type.
938    Plain(T),
939}
940
941impl<T> Typed<T> {
942    fn map<U>(self, mut f: impl FnMut(T) -> U) -> Typed<U> {
943        match self {
944            Self::Reference(v) => Typed::Reference(f(v)),
945            Self::Plain(v) => Typed::Plain(f(v)),
946        }
947    }
948
949    fn try_map<U, E>(
950        self,
951        mut f: impl FnMut(T) -> core::result::Result<U, E>,
952    ) -> core::result::Result<Typed<U>, E> {
953        Ok(match self {
954            Self::Reference(expr) => Typed::Reference(f(expr)?),
955            Self::Plain(expr) => Typed::Plain(f(expr)?),
956        })
957    }
958}
959
960/// A single vector component or swizzle.
961///
962/// This represents the things that can appear after the `.` in a vector access
963/// expression: either a single component name, or a series of them,
964/// representing a swizzle.
965enum Components {
966    Single(u32),
967    Swizzle {
968        size: ir::VectorSize,
969        pattern: [ir::SwizzleComponent; 4],
970    },
971}
972
973impl Components {
974    const fn letter_component(letter: char) -> Option<ir::SwizzleComponent> {
975        use ir::SwizzleComponent as Sc;
976        match letter {
977            'x' | 'r' => Some(Sc::X),
978            'y' | 'g' => Some(Sc::Y),
979            'z' | 'b' => Some(Sc::Z),
980            'w' | 'a' => Some(Sc::W),
981            _ => None,
982        }
983    }
984
985    fn single_component(name: &str, name_span: Span) -> Result<'_, u32> {
986        let ch = name.chars().next().ok_or(Error::BadAccessor(name_span))?;
987        match Self::letter_component(ch) {
988            Some(sc) => Ok(sc as u32),
989            None => Err(Box::new(Error::BadAccessor(name_span))),
990        }
991    }
992
993    /// Construct a `Components` value from a 'member' name, like `"wzy"` or `"x"`.
994    ///
995    /// Use `name_span` for reporting errors in parsing the component string.
996    fn new(name: &str, name_span: Span) -> Result<'_, Self> {
997        let size = match name.len() {
998            1 => return Ok(Components::Single(Self::single_component(name, name_span)?)),
999            2 => ir::VectorSize::Bi,
1000            3 => ir::VectorSize::Tri,
1001            4 => ir::VectorSize::Quad,
1002            _ => return Err(Box::new(Error::BadAccessor(name_span))),
1003        };
1004
1005        let mut pattern = [ir::SwizzleComponent::X; 4];
1006        for (comp, ch) in pattern.iter_mut().zip(name.chars()) {
1007            *comp = Self::letter_component(ch).ok_or(Error::BadAccessor(name_span))?;
1008        }
1009
1010        if name.chars().all(|c| matches!(c, 'x' | 'y' | 'z' | 'w'))
1011            || name.chars().all(|c| matches!(c, 'r' | 'g' | 'b' | 'a'))
1012        {
1013            Ok(Components::Swizzle { size, pattern })
1014        } else {
1015            Err(Box::new(Error::BadAccessor(name_span)))
1016        }
1017    }
1018}
1019
1020/// An `ast::GlobalDecl` for which we have built the Naga IR equivalent.
1021enum LoweredGlobalDecl {
1022    Function {
1023        handle: Handle<ir::Function>,
1024        must_use: bool,
1025    },
1026    Var(Handle<ir::GlobalVariable>),
1027    Const(Handle<ir::Constant>),
1028    Override(Handle<ir::Override>),
1029    Type(Handle<ir::Type>),
1030    EntryPoint(usize),
1031}
1032
1033enum Texture {
1034    Gather,
1035    GatherCompare,
1036
1037    Sample,
1038    SampleBias,
1039    SampleCompare,
1040    SampleCompareLevel,
1041    SampleGrad,
1042    SampleLevel,
1043    SampleBaseClampToEdge,
1044}
1045
1046impl Texture {
1047    pub fn map(word: &str) -> Option<Self> {
1048        Some(match word {
1049            "textureGather" => Self::Gather,
1050            "textureGatherCompare" => Self::GatherCompare,
1051
1052            "textureSample" => Self::Sample,
1053            "textureSampleBias" => Self::SampleBias,
1054            "textureSampleCompare" => Self::SampleCompare,
1055            "textureSampleCompareLevel" => Self::SampleCompareLevel,
1056            "textureSampleGrad" => Self::SampleGrad,
1057            "textureSampleLevel" => Self::SampleLevel,
1058            "textureSampleBaseClampToEdge" => Self::SampleBaseClampToEdge,
1059            _ => return None,
1060        })
1061    }
1062
1063    pub const fn min_argument_count(&self) -> u32 {
1064        match *self {
1065            Self::Gather => 3,
1066            Self::GatherCompare => 4,
1067
1068            Self::Sample => 3,
1069            Self::SampleBias => 5,
1070            Self::SampleCompare => 5,
1071            Self::SampleCompareLevel => 5,
1072            Self::SampleGrad => 6,
1073            Self::SampleLevel => 5,
1074            Self::SampleBaseClampToEdge => 3,
1075        }
1076    }
1077}
1078
1079enum SubgroupGather {
1080    BroadcastFirst,
1081    Broadcast,
1082    Shuffle,
1083    ShuffleDown,
1084    ShuffleUp,
1085    ShuffleXor,
1086    QuadBroadcast,
1087}
1088
1089impl SubgroupGather {
1090    pub fn map(word: &str) -> Option<Self> {
1091        Some(match word {
1092            "subgroupBroadcastFirst" => Self::BroadcastFirst,
1093            "subgroupBroadcast" => Self::Broadcast,
1094            "subgroupShuffle" => Self::Shuffle,
1095            "subgroupShuffleDown" => Self::ShuffleDown,
1096            "subgroupShuffleUp" => Self::ShuffleUp,
1097            "subgroupShuffleXor" => Self::ShuffleXor,
1098            "quadBroadcast" => Self::QuadBroadcast,
1099            _ => return None,
1100        })
1101    }
1102}
1103
1104/// Whether a declaration accepts abstract types, or concretizes.
1105enum AbstractRule {
1106    /// This declaration concretizes its initialization expression.
1107    Concretize,
1108
1109    /// This declaration can accept initializers with abstract types.
1110    Allow,
1111}
1112
1113pub struct Lowerer<'source, 'temp> {
1114    index: &'temp Index<'source>,
1115}
1116
1117impl<'source, 'temp> Lowerer<'source, 'temp> {
1118    pub const fn new(index: &'temp Index<'source>) -> Self {
1119        Self { index }
1120    }
1121
1122    pub fn lower(&mut self, tu: ast::TranslationUnit<'source>) -> Result<'source, ir::Module> {
1123        let mut module = ir::Module {
1124            diagnostic_filters: tu.diagnostic_filters,
1125            diagnostic_filter_leaf: tu.diagnostic_filter_leaf,
1126            ..Default::default()
1127        };
1128
1129        let mut ctx = GlobalContext {
1130            ast_expressions: &tu.expressions,
1131            globals: &mut FastHashMap::default(),
1132            types: &tu.types,
1133            module: &mut module,
1134            const_typifier: &mut Typifier::new(),
1135            layouter: &mut proc::Layouter::default(),
1136            global_expression_kind_tracker: &mut proc::ExpressionKindTracker::new(),
1137        };
1138        if !tu.doc_comments.is_empty() {
1139            ctx.module.get_or_insert_default_doc_comments().module =
1140                tu.doc_comments.iter().map(|s| s.to_string()).collect();
1141        }
1142
1143        for decl_handle in self.index.visit_ordered() {
1144            let span = tu.decls.get_span(decl_handle);
1145            let decl = &tu.decls[decl_handle];
1146
1147            match decl.kind {
1148                ast::GlobalDeclKind::Fn(ref f) => {
1149                    let lowered_decl = self.function(f, span, &mut ctx)?;
1150                    if !f.doc_comments.is_empty() {
1151                        match lowered_decl {
1152                            LoweredGlobalDecl::Function { handle, .. } => {
1153                                ctx.module
1154                                    .get_or_insert_default_doc_comments()
1155                                    .functions
1156                                    .insert(
1157                                        handle,
1158                                        f.doc_comments.iter().map(|s| s.to_string()).collect(),
1159                                    );
1160                            }
1161                            LoweredGlobalDecl::EntryPoint(index) => {
1162                                ctx.module
1163                                    .get_or_insert_default_doc_comments()
1164                                    .entry_points
1165                                    .insert(
1166                                        index,
1167                                        f.doc_comments.iter().map(|s| s.to_string()).collect(),
1168                                    );
1169                            }
1170                            _ => {}
1171                        }
1172                    }
1173                    ctx.globals.insert(f.name.name, lowered_decl);
1174                }
1175                ast::GlobalDeclKind::Var(ref v) => {
1176                    let explicit_ty =
1177                        v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const()))
1178                            .transpose()?;
1179
1180                    let (ty, initializer) = self.type_and_init(
1181                        v.name,
1182                        v.init,
1183                        explicit_ty,
1184                        AbstractRule::Concretize,
1185                        &mut ctx.as_override(),
1186                    )?;
1187
1188                    let binding = if let Some(ref binding) = v.binding {
1189                        Some(ir::ResourceBinding {
1190                            group: self.const_u32(binding.group, &mut ctx.as_const())?.0,
1191                            binding: self.const_u32(binding.binding, &mut ctx.as_const())?.0,
1192                        })
1193                    } else {
1194                        None
1195                    };
1196
1197                    let handle = ctx.module.global_variables.append(
1198                        ir::GlobalVariable {
1199                            name: Some(v.name.name.to_string()),
1200                            space: v.space,
1201                            binding,
1202                            ty,
1203                            init: initializer,
1204                        },
1205                        span,
1206                    );
1207
1208                    if !v.doc_comments.is_empty() {
1209                        ctx.module
1210                            .get_or_insert_default_doc_comments()
1211                            .global_variables
1212                            .insert(
1213                                handle,
1214                                v.doc_comments.iter().map(|s| s.to_string()).collect(),
1215                            );
1216                    }
1217                    ctx.globals
1218                        .insert(v.name.name, LoweredGlobalDecl::Var(handle));
1219                }
1220                ast::GlobalDeclKind::Const(ref c) => {
1221                    let mut ectx = ctx.as_const();
1222
1223                    let explicit_ty =
1224                        c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx))
1225                            .transpose()?;
1226
1227                    let (ty, init) = self.type_and_init(
1228                        c.name,
1229                        Some(c.init),
1230                        explicit_ty,
1231                        AbstractRule::Allow,
1232                        &mut ectx,
1233                    )?;
1234                    let init = init.expect("Global const must have init");
1235
1236                    let handle = ctx.module.constants.append(
1237                        ir::Constant {
1238                            name: Some(c.name.name.to_string()),
1239                            ty,
1240                            init,
1241                        },
1242                        span,
1243                    );
1244
1245                    ctx.globals
1246                        .insert(c.name.name, LoweredGlobalDecl::Const(handle));
1247                    if !c.doc_comments.is_empty() {
1248                        ctx.module
1249                            .get_or_insert_default_doc_comments()
1250                            .constants
1251                            .insert(
1252                                handle,
1253                                c.doc_comments.iter().map(|s| s.to_string()).collect(),
1254                            );
1255                    }
1256                }
1257                ast::GlobalDeclKind::Override(ref o) => {
1258                    let explicit_ty =
1259                        o.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_const()))
1260                            .transpose()?;
1261
1262                    let mut ectx = ctx.as_override();
1263
1264                    let (ty, init) = self.type_and_init(
1265                        o.name,
1266                        o.init,
1267                        explicit_ty,
1268                        AbstractRule::Concretize,
1269                        &mut ectx,
1270                    )?;
1271
1272                    let id =
1273                        o.id.map(|id| self.const_u32(id, &mut ctx.as_const()))
1274                            .transpose()?;
1275
1276                    let id = if let Some((id, id_span)) = id {
1277                        Some(
1278                            u16::try_from(id)
1279                                .map_err(|_| Error::PipelineConstantIDValue(id_span))?,
1280                        )
1281                    } else {
1282                        None
1283                    };
1284
1285                    let handle = ctx.module.overrides.append(
1286                        ir::Override {
1287                            name: Some(o.name.name.to_string()),
1288                            id,
1289                            ty,
1290                            init,
1291                        },
1292                        span,
1293                    );
1294
1295                    ctx.globals
1296                        .insert(o.name.name, LoweredGlobalDecl::Override(handle));
1297                }
1298                ast::GlobalDeclKind::Struct(ref s) => {
1299                    let handle = self.r#struct(s, span, &mut ctx)?;
1300                    ctx.globals
1301                        .insert(s.name.name, LoweredGlobalDecl::Type(handle));
1302                    if !s.doc_comments.is_empty() {
1303                        ctx.module
1304                            .get_or_insert_default_doc_comments()
1305                            .types
1306                            .insert(
1307                                handle,
1308                                s.doc_comments.iter().map(|s| s.to_string()).collect(),
1309                            );
1310                    }
1311                }
1312                ast::GlobalDeclKind::Type(ref alias) => {
1313                    let ty = self.resolve_named_ast_type(
1314                        alias.ty,
1315                        Some(alias.name.name.to_string()),
1316                        &mut ctx.as_const(),
1317                    )?;
1318                    ctx.globals
1319                        .insert(alias.name.name, LoweredGlobalDecl::Type(ty));
1320                }
1321                ast::GlobalDeclKind::ConstAssert(condition) => {
1322                    let condition = self.expression(condition, &mut ctx.as_const())?;
1323
1324                    let span = ctx.module.global_expressions.get_span(condition);
1325                    match ctx
1326                        .module
1327                        .to_ctx()
1328                        .eval_expr_to_bool_from(condition, &ctx.module.global_expressions)
1329                    {
1330                        Some(true) => Ok(()),
1331                        Some(false) => Err(Error::ConstAssertFailed(span)),
1332                        _ => Err(Error::NotBool(span)),
1333                    }?;
1334                }
1335            }
1336        }
1337
1338        // Constant evaluation may leave abstract-typed literals and
1339        // compositions in expression arenas, so we need to compact the module
1340        // to remove unused expressions and types.
1341        crate::compact::compact(&mut module, KeepUnused::Yes);
1342
1343        Ok(module)
1344    }
1345
1346    /// Obtain (inferred) type and initializer after automatic conversion
1347    fn type_and_init(
1348        &mut self,
1349        name: ast::Ident<'source>,
1350        init: Option<Handle<ast::Expression<'source>>>,
1351        explicit_ty: Option<Handle<ir::Type>>,
1352        abstract_rule: AbstractRule,
1353        ectx: &mut ExpressionContext<'source, '_, '_>,
1354    ) -> Result<'source, (Handle<ir::Type>, Option<Handle<ir::Expression>>)> {
1355        let ty;
1356        let initializer;
1357        match (init, explicit_ty) {
1358            (Some(init), Some(explicit_ty)) => {
1359                let init = self.expression_for_abstract(init, ectx)?;
1360                let ty_res = proc::TypeResolution::Handle(explicit_ty);
1361                let init = ectx
1362                    .try_automatic_conversions(init, &ty_res, name.span)
1363                    .map_err(|error| match *error {
1364                        Error::AutoConversion(e) => Box::new(Error::InitializationTypeMismatch {
1365                            name: name.span,
1366                            expected: e.dest_type,
1367                            got: e.source_type,
1368                        }),
1369                        _ => error,
1370                    })?;
1371
1372                let init_ty = ectx.register_type(init)?;
1373                if !ectx.module.compare_types(
1374                    &proc::TypeResolution::Handle(explicit_ty),
1375                    &proc::TypeResolution::Handle(init_ty),
1376                ) {
1377                    return Err(Box::new(Error::InitializationTypeMismatch {
1378                        name: name.span,
1379                        expected: ectx.type_to_string(explicit_ty),
1380                        got: ectx.type_to_string(init_ty),
1381                    }));
1382                }
1383                ty = explicit_ty;
1384                initializer = Some(init);
1385            }
1386            (Some(init), None) => {
1387                let mut init = self.expression_for_abstract(init, ectx)?;
1388                if let AbstractRule::Concretize = abstract_rule {
1389                    init = ectx.concretize(init)?;
1390                }
1391                ty = ectx.register_type(init)?;
1392                initializer = Some(init);
1393            }
1394            (None, Some(explicit_ty)) => {
1395                ty = explicit_ty;
1396                initializer = None;
1397            }
1398            (None, None) => return Err(Box::new(Error::DeclMissingTypeAndInit(name.span))),
1399        }
1400        Ok((ty, initializer))
1401    }
1402
1403    fn function(
1404        &mut self,
1405        f: &ast::Function<'source>,
1406        span: Span,
1407        ctx: &mut GlobalContext<'source, '_, '_>,
1408    ) -> Result<'source, LoweredGlobalDecl> {
1409        let mut local_table = FastHashMap::default();
1410        let mut expressions = Arena::new();
1411        let mut named_expressions = FastIndexMap::default();
1412        let mut local_expression_kind_tracker = proc::ExpressionKindTracker::new();
1413
1414        let arguments = f
1415            .arguments
1416            .iter()
1417            .enumerate()
1418            .map(|(i, arg)| -> Result<'_, _> {
1419                let ty = self.resolve_ast_type(arg.ty, &mut ctx.as_const())?;
1420                let expr =
1421                    expressions.append(ir::Expression::FunctionArgument(i as u32), arg.name.span);
1422                local_table.insert(arg.handle, Declared::Runtime(Typed::Plain(expr)));
1423                named_expressions.insert(expr, (arg.name.name.to_string(), arg.name.span));
1424                local_expression_kind_tracker.insert(expr, proc::ExpressionKind::Runtime);
1425
1426                Ok(ir::FunctionArgument {
1427                    name: Some(arg.name.name.to_string()),
1428                    ty,
1429                    binding: self.binding(&arg.binding, ty, ctx)?,
1430                })
1431            })
1432            .collect::<Result<Vec<_>>>()?;
1433
1434        let result = f
1435            .result
1436            .as_ref()
1437            .map(|res| -> Result<'_, _> {
1438                let ty = self.resolve_ast_type(res.ty, &mut ctx.as_const())?;
1439                Ok(ir::FunctionResult {
1440                    ty,
1441                    binding: self.binding(&res.binding, ty, ctx)?,
1442                })
1443            })
1444            .transpose()?;
1445
1446        let mut function = ir::Function {
1447            name: Some(f.name.name.to_string()),
1448            arguments,
1449            result,
1450            local_variables: Arena::new(),
1451            expressions,
1452            named_expressions: crate::NamedExpressions::default(),
1453            body: ir::Block::default(),
1454            diagnostic_filter_leaf: f.diagnostic_filter_leaf,
1455        };
1456
1457        let mut typifier = Typifier::default();
1458        let mut stmt_ctx = StatementContext {
1459            local_table: &mut local_table,
1460            globals: ctx.globals,
1461            ast_expressions: ctx.ast_expressions,
1462            const_typifier: ctx.const_typifier,
1463            typifier: &mut typifier,
1464            layouter: ctx.layouter,
1465            function: &mut function,
1466            named_expressions: &mut named_expressions,
1467            types: ctx.types,
1468            module: ctx.module,
1469            local_expression_kind_tracker: &mut local_expression_kind_tracker,
1470            global_expression_kind_tracker: ctx.global_expression_kind_tracker,
1471        };
1472        let mut body = self.block(&f.body, false, &mut stmt_ctx)?;
1473        proc::ensure_block_returns(&mut body);
1474
1475        function.body = body;
1476        function.named_expressions = named_expressions
1477            .into_iter()
1478            .map(|(key, (name, _))| (key, name))
1479            .collect();
1480
1481        if let Some(ref entry) = f.entry_point {
1482            let workgroup_size_info = if let Some(workgroup_size) = entry.workgroup_size {
1483                // TODO: replace with try_map once stabilized
1484                let mut workgroup_size_out = [1; 3];
1485                let mut workgroup_size_overrides_out = [None; 3];
1486                for (i, size) in workgroup_size.into_iter().enumerate() {
1487                    if let Some(size_expr) = size {
1488                        match self.const_u32(size_expr, &mut ctx.as_const()) {
1489                            Ok(value) => {
1490                                workgroup_size_out[i] = value.0;
1491                            }
1492                            Err(err) => {
1493                                if let Error::ConstantEvaluatorError(ref ty, _) = *err {
1494                                    match **ty {
1495                                        proc::ConstantEvaluatorError::OverrideExpr => {
1496                                            workgroup_size_overrides_out[i] =
1497                                                Some(self.workgroup_size_override(
1498                                                    size_expr,
1499                                                    &mut ctx.as_override(),
1500                                                )?);
1501                                        }
1502                                        _ => {
1503                                            return Err(err);
1504                                        }
1505                                    }
1506                                } else {
1507                                    return Err(err);
1508                                }
1509                            }
1510                        }
1511                    }
1512                }
1513                if workgroup_size_overrides_out.iter().all(|x| x.is_none()) {
1514                    (workgroup_size_out, None)
1515                } else {
1516                    (workgroup_size_out, Some(workgroup_size_overrides_out))
1517                }
1518            } else {
1519                ([0; 3], None)
1520            };
1521
1522            let (workgroup_size, workgroup_size_overrides) = workgroup_size_info;
1523            ctx.module.entry_points.push(ir::EntryPoint {
1524                name: f.name.name.to_string(),
1525                stage: entry.stage,
1526                early_depth_test: entry.early_depth_test,
1527                workgroup_size,
1528                workgroup_size_overrides,
1529                function,
1530            });
1531            Ok(LoweredGlobalDecl::EntryPoint(
1532                ctx.module.entry_points.len() - 1,
1533            ))
1534        } else {
1535            let handle = ctx.module.functions.append(function, span);
1536            Ok(LoweredGlobalDecl::Function {
1537                handle,
1538                must_use: f.result.as_ref().is_some_and(|res| res.must_use),
1539            })
1540        }
1541    }
1542
1543    fn workgroup_size_override(
1544        &mut self,
1545        size_expr: Handle<ast::Expression<'source>>,
1546        ctx: &mut ExpressionContext<'source, '_, '_>,
1547    ) -> Result<'source, Handle<ir::Expression>> {
1548        let span = ctx.ast_expressions.get_span(size_expr);
1549        let expr = self.expression(size_expr, ctx)?;
1550        match resolve_inner!(ctx, expr).scalar_kind().ok_or(0) {
1551            Ok(ir::ScalarKind::Sint) | Ok(ir::ScalarKind::Uint) => Ok(expr),
1552            _ => Err(Box::new(Error::ExpectedConstExprConcreteIntegerScalar(
1553                span,
1554            ))),
1555        }
1556    }
1557
1558    fn block(
1559        &mut self,
1560        b: &ast::Block<'source>,
1561        is_inside_loop: bool,
1562        ctx: &mut StatementContext<'source, '_, '_>,
1563    ) -> Result<'source, ir::Block> {
1564        let mut block = ir::Block::default();
1565
1566        for stmt in b.stmts.iter() {
1567            self.statement(stmt, &mut block, is_inside_loop, ctx)?;
1568        }
1569
1570        Ok(block)
1571    }
1572
1573    fn statement(
1574        &mut self,
1575        stmt: &ast::Statement<'source>,
1576        block: &mut ir::Block,
1577        is_inside_loop: bool,
1578        ctx: &mut StatementContext<'source, '_, '_>,
1579    ) -> Result<'source, ()> {
1580        let out = match stmt.kind {
1581            ast::StatementKind::Block(ref block) => {
1582                let block = self.block(block, is_inside_loop, ctx)?;
1583                ir::Statement::Block(block)
1584            }
1585            ast::StatementKind::LocalDecl(ref decl) => match *decl {
1586                ast::LocalDecl::Let(ref l) => {
1587                    let mut emitter = proc::Emitter::default();
1588                    emitter.start(&ctx.function.expressions);
1589
1590                    let explicit_ty = l
1591                        .ty
1592                        .map(|ty| self.resolve_ast_type(ty, &mut ctx.as_const(block, &mut emitter)))
1593                        .transpose()?;
1594
1595                    let mut ectx = ctx.as_expression(block, &mut emitter);
1596
1597                    let (_ty, initializer) = self.type_and_init(
1598                        l.name,
1599                        Some(l.init),
1600                        explicit_ty,
1601                        AbstractRule::Concretize,
1602                        &mut ectx,
1603                    )?;
1604
1605                    // We passed `Some()` to `type_and_init`, so we
1606                    // will get a lowered initializer expression back.
1607                    let initializer =
1608                        initializer.expect("type_and_init did not return an initializer");
1609
1610                    // The WGSL spec says that any expression that refers to a
1611                    // `let`-bound variable is not a const expression. This
1612                    // affects when errors must be reported, so we can't even
1613                    // treat suitable `let` bindings as constant as an
1614                    // optimization.
1615                    ctx.local_expression_kind_tracker
1616                        .force_non_const(initializer);
1617
1618                    block.extend(emitter.finish(&ctx.function.expressions));
1619                    ctx.local_table
1620                        .insert(l.handle, Declared::Runtime(Typed::Plain(initializer)));
1621                    ctx.named_expressions
1622                        .insert(initializer, (l.name.name.to_string(), l.name.span));
1623
1624                    return Ok(());
1625                }
1626                ast::LocalDecl::Var(ref v) => {
1627                    let mut emitter = proc::Emitter::default();
1628                    emitter.start(&ctx.function.expressions);
1629
1630                    let explicit_ty =
1631                        v.ty.map(|ast| {
1632                            self.resolve_ast_type(ast, &mut ctx.as_const(block, &mut emitter))
1633                        })
1634                        .transpose()?;
1635
1636                    let mut ectx = ctx.as_expression(block, &mut emitter);
1637                    let (ty, initializer) = self.type_and_init(
1638                        v.name,
1639                        v.init,
1640                        explicit_ty,
1641                        AbstractRule::Concretize,
1642                        &mut ectx,
1643                    )?;
1644
1645                    let (const_initializer, initializer) = {
1646                        match initializer {
1647                            Some(init) => {
1648                                // It's not correct to hoist the initializer up
1649                                // to the top of the function if:
1650                                // - the initialization is inside a loop, and should
1651                                //   take place on every iteration, or
1652                                // - the initialization is not a constant
1653                                //   expression, so its value depends on the
1654                                //   state at the point of initialization.
1655                                if is_inside_loop
1656                                    || !ctx.local_expression_kind_tracker.is_const_or_override(init)
1657                                {
1658                                    (None, Some(init))
1659                                } else {
1660                                    (Some(init), None)
1661                                }
1662                            }
1663                            None => (None, None),
1664                        }
1665                    };
1666
1667                    let var = ctx.function.local_variables.append(
1668                        ir::LocalVariable {
1669                            name: Some(v.name.name.to_string()),
1670                            ty,
1671                            init: const_initializer,
1672                        },
1673                        stmt.span,
1674                    );
1675
1676                    let handle = ctx
1677                        .as_expression(block, &mut emitter)
1678                        .interrupt_emitter(ir::Expression::LocalVariable(var), Span::UNDEFINED)?;
1679                    block.extend(emitter.finish(&ctx.function.expressions));
1680                    ctx.local_table
1681                        .insert(v.handle, Declared::Runtime(Typed::Reference(handle)));
1682
1683                    match initializer {
1684                        Some(initializer) => ir::Statement::Store {
1685                            pointer: handle,
1686                            value: initializer,
1687                        },
1688                        None => return Ok(()),
1689                    }
1690                }
1691                ast::LocalDecl::Const(ref c) => {
1692                    let mut emitter = proc::Emitter::default();
1693                    emitter.start(&ctx.function.expressions);
1694
1695                    let ectx = &mut ctx.as_const(block, &mut emitter);
1696
1697                    let explicit_ty =
1698                        c.ty.map(|ast| self.resolve_ast_type(ast, &mut ectx.as_const()))
1699                            .transpose()?;
1700
1701                    let (_ty, init) = self.type_and_init(
1702                        c.name,
1703                        Some(c.init),
1704                        explicit_ty,
1705                        AbstractRule::Allow,
1706                        &mut ectx.as_const(),
1707                    )?;
1708                    let init = init.expect("Local const must have init");
1709
1710                    block.extend(emitter.finish(&ctx.function.expressions));
1711                    ctx.local_table
1712                        .insert(c.handle, Declared::Const(Typed::Plain(init)));
1713                    return Ok(());
1714                }
1715            },
1716            ast::StatementKind::If {
1717                condition,
1718                ref accept,
1719                ref reject,
1720            } => {
1721                let mut emitter = proc::Emitter::default();
1722                emitter.start(&ctx.function.expressions);
1723
1724                let condition =
1725                    self.expression(condition, &mut ctx.as_expression(block, &mut emitter))?;
1726                block.extend(emitter.finish(&ctx.function.expressions));
1727
1728                let accept = self.block(accept, is_inside_loop, ctx)?;
1729                let reject = self.block(reject, is_inside_loop, ctx)?;
1730
1731                ir::Statement::If {
1732                    condition,
1733                    accept,
1734                    reject,
1735                }
1736            }
1737            ast::StatementKind::Switch {
1738                selector,
1739                ref cases,
1740            } => {
1741                let mut emitter = proc::Emitter::default();
1742                emitter.start(&ctx.function.expressions);
1743
1744                let mut ectx = ctx.as_expression(block, &mut emitter);
1745
1746                // Determine the scalar type of the selector and case expressions, find the
1747                // consensus type for automatic conversion, then convert them.
1748                let (mut exprs, spans) = core::iter::once(selector)
1749                    .chain(cases.iter().filter_map(|case| match case.value {
1750                        ast::SwitchValue::Expr(expr) => Some(expr),
1751                        ast::SwitchValue::Default => None,
1752                    }))
1753                    .enumerate()
1754                    .map(|(i, expr)| {
1755                        let span = ectx.ast_expressions.get_span(expr);
1756                        let expr = self.expression_for_abstract(expr, &mut ectx)?;
1757                        let ty = resolve_inner!(ectx, expr);
1758                        match *ty {
1759                            ir::TypeInner::Scalar(
1760                                ir::Scalar::I32 | ir::Scalar::U32 | ir::Scalar::ABSTRACT_INT,
1761                            ) => Ok((expr, span)),
1762                            _ => match i {
1763                                0 => Err(Box::new(Error::InvalidSwitchSelector { span })),
1764                                _ => Err(Box::new(Error::InvalidSwitchCase { span })),
1765                            },
1766                        }
1767                    })
1768                    .collect::<Result<(Vec<_>, Vec<_>)>>()?;
1769
1770                let mut consensus =
1771                    ectx.automatic_conversion_consensus(&exprs)
1772                        .map_err(|span_idx| Error::SwitchCaseTypeMismatch {
1773                            span: spans[span_idx],
1774                        })?;
1775                // Concretize to I32 if the selector and all cases were abstract
1776                if consensus == ir::Scalar::ABSTRACT_INT {
1777                    consensus = ir::Scalar::I32;
1778                }
1779                for expr in &mut exprs {
1780                    ectx.convert_to_leaf_scalar(expr, consensus)?;
1781                }
1782
1783                block.extend(emitter.finish(&ctx.function.expressions));
1784
1785                let mut exprs = exprs.into_iter();
1786                let selector = exprs
1787                    .next()
1788                    .expect("First element should be selector expression");
1789
1790                let cases = cases
1791                    .iter()
1792                    .map(|case| {
1793                        Ok(ir::SwitchCase {
1794                            value: match case.value {
1795                                ast::SwitchValue::Expr(expr) => {
1796                                    let span = ctx.ast_expressions.get_span(expr);
1797                                    let expr = exprs.next().expect(
1798                                        "Should yield expression for each SwitchValue::Expr case",
1799                                    );
1800                                    match ctx
1801                                        .module
1802                                        .to_ctx()
1803                                        .eval_expr_to_literal_from(expr, &ctx.function.expressions)
1804                                    {
1805                                        Some(ir::Literal::I32(value)) => {
1806                                            ir::SwitchValue::I32(value)
1807                                        }
1808                                        Some(ir::Literal::U32(value)) => {
1809                                            ir::SwitchValue::U32(value)
1810                                        }
1811                                        _ => {
1812                                            return Err(Box::new(Error::InvalidSwitchCase {
1813                                                span,
1814                                            }));
1815                                        }
1816                                    }
1817                                }
1818                                ast::SwitchValue::Default => ir::SwitchValue::Default,
1819                            },
1820                            body: self.block(&case.body, is_inside_loop, ctx)?,
1821                            fall_through: case.fall_through,
1822                        })
1823                    })
1824                    .collect::<Result<_>>()?;
1825
1826                ir::Statement::Switch { selector, cases }
1827            }
1828            ast::StatementKind::Loop {
1829                ref body,
1830                ref continuing,
1831                break_if,
1832            } => {
1833                let body = self.block(body, true, ctx)?;
1834                let mut continuing = self.block(continuing, true, ctx)?;
1835
1836                let mut emitter = proc::Emitter::default();
1837                emitter.start(&ctx.function.expressions);
1838                let break_if = break_if
1839                    .map(|expr| {
1840                        self.expression(expr, &mut ctx.as_expression(&mut continuing, &mut emitter))
1841                    })
1842                    .transpose()?;
1843                continuing.extend(emitter.finish(&ctx.function.expressions));
1844
1845                ir::Statement::Loop {
1846                    body,
1847                    continuing,
1848                    break_if,
1849                }
1850            }
1851            ast::StatementKind::Break => ir::Statement::Break,
1852            ast::StatementKind::Continue => ir::Statement::Continue,
1853            ast::StatementKind::Return { value: ast_value } => {
1854                let mut emitter = proc::Emitter::default();
1855                emitter.start(&ctx.function.expressions);
1856
1857                let value;
1858                if let Some(ast_expr) = ast_value {
1859                    let result_ty = ctx.function.result.as_ref().map(|r| r.ty);
1860                    let mut ectx = ctx.as_expression(block, &mut emitter);
1861                    let expr = self.expression_for_abstract(ast_expr, &mut ectx)?;
1862
1863                    if let Some(result_ty) = result_ty {
1864                        let mut ectx = ctx.as_expression(block, &mut emitter);
1865                        let resolution = proc::TypeResolution::Handle(result_ty);
1866                        let converted =
1867                            ectx.try_automatic_conversions(expr, &resolution, Span::default())?;
1868                        value = Some(converted);
1869                    } else {
1870                        value = Some(expr);
1871                    }
1872                } else {
1873                    value = None;
1874                }
1875                block.extend(emitter.finish(&ctx.function.expressions));
1876
1877                ir::Statement::Return { value }
1878            }
1879            ast::StatementKind::Kill => ir::Statement::Kill,
1880            ast::StatementKind::Call {
1881                ref function,
1882                ref arguments,
1883            } => {
1884                let mut emitter = proc::Emitter::default();
1885                emitter.start(&ctx.function.expressions);
1886
1887                let _ = self.call(
1888                    stmt.span,
1889                    function,
1890                    arguments,
1891                    &mut ctx.as_expression(block, &mut emitter),
1892                    true,
1893                )?;
1894                block.extend(emitter.finish(&ctx.function.expressions));
1895                return Ok(());
1896            }
1897            ast::StatementKind::Assign {
1898                target: ast_target,
1899                op,
1900                value,
1901            } => {
1902                let mut emitter = proc::Emitter::default();
1903                emitter.start(&ctx.function.expressions);
1904                let target_span = ctx.ast_expressions.get_span(ast_target);
1905
1906                let mut ectx = ctx.as_expression(block, &mut emitter);
1907                let target = self.expression_for_reference(ast_target, &mut ectx)?;
1908                let target_handle = match target {
1909                    Typed::Reference(handle) => handle,
1910                    Typed::Plain(handle) => {
1911                        let ty = ctx.invalid_assignment_type(handle);
1912                        return Err(Box::new(Error::InvalidAssignment {
1913                            span: target_span,
1914                            ty,
1915                        }));
1916                    }
1917                };
1918
1919                // Usually the value needs to be converted to match the type of
1920                // the memory view you're assigning it to. The bit shift
1921                // operators are exceptions, in that the right operand is always
1922                // a `u32` or `vecN<u32>`.
1923                let target_scalar = match op {
1924                    Some(ir::BinaryOperator::ShiftLeft | ir::BinaryOperator::ShiftRight) => {
1925                        Some(ir::Scalar::U32)
1926                    }
1927                    _ => resolve_inner!(ectx, target_handle)
1928                        .pointer_automatically_convertible_scalar(&ectx.module.types),
1929                };
1930
1931                let value = self.expression_for_abstract(value, &mut ectx)?;
1932                let mut value = match target_scalar {
1933                    Some(target_scalar) => ectx.try_automatic_conversion_for_leaf_scalar(
1934                        value,
1935                        target_scalar,
1936                        target_span,
1937                    )?,
1938                    None => value,
1939                };
1940
1941                let value = match op {
1942                    Some(op) => {
1943                        let mut left = ectx.apply_load_rule(target)?;
1944                        ectx.binary_op_splat(op, &mut left, &mut value)?;
1945                        ectx.append_expression(
1946                            ir::Expression::Binary {
1947                                op,
1948                                left,
1949                                right: value,
1950                            },
1951                            stmt.span,
1952                        )?
1953                    }
1954                    None => value,
1955                };
1956                block.extend(emitter.finish(&ctx.function.expressions));
1957
1958                ir::Statement::Store {
1959                    pointer: target_handle,
1960                    value,
1961                }
1962            }
1963            ast::StatementKind::Increment(value) | ast::StatementKind::Decrement(value) => {
1964                let mut emitter = proc::Emitter::default();
1965                emitter.start(&ctx.function.expressions);
1966
1967                let op = match stmt.kind {
1968                    ast::StatementKind::Increment(_) => ir::BinaryOperator::Add,
1969                    ast::StatementKind::Decrement(_) => ir::BinaryOperator::Subtract,
1970                    _ => unreachable!(),
1971                };
1972
1973                let value_span = ctx.ast_expressions.get_span(value);
1974                let target = self
1975                    .expression_for_reference(value, &mut ctx.as_expression(block, &mut emitter))?;
1976                let target_handle = match target {
1977                    Typed::Reference(handle) => handle,
1978                    Typed::Plain(_) => {
1979                        return Err(Box::new(Error::BadIncrDecrReferenceType(value_span)))
1980                    }
1981                };
1982
1983                let mut ectx = ctx.as_expression(block, &mut emitter);
1984                let scalar = match *resolve_inner!(ectx, target_handle) {
1985                    ir::TypeInner::ValuePointer {
1986                        size: None, scalar, ..
1987                    } => scalar,
1988                    ir::TypeInner::Pointer { base, .. } => match ectx.module.types[base].inner {
1989                        ir::TypeInner::Scalar(scalar) => scalar,
1990                        _ => return Err(Box::new(Error::BadIncrDecrReferenceType(value_span))),
1991                    },
1992                    _ => return Err(Box::new(Error::BadIncrDecrReferenceType(value_span))),
1993                };
1994                let literal = match scalar.kind {
1995                    ir::ScalarKind::Sint | ir::ScalarKind::Uint => ir::Literal::one(scalar)
1996                        .ok_or(Error::BadIncrDecrReferenceType(value_span))?,
1997                    _ => return Err(Box::new(Error::BadIncrDecrReferenceType(value_span))),
1998                };
1999
2000                let right =
2001                    ectx.interrupt_emitter(ir::Expression::Literal(literal), Span::UNDEFINED)?;
2002                let rctx = ectx.runtime_expression_ctx(stmt.span)?;
2003                let left = rctx.function.expressions.append(
2004                    ir::Expression::Load {
2005                        pointer: target_handle,
2006                    },
2007                    value_span,
2008                );
2009                let value = rctx
2010                    .function
2011                    .expressions
2012                    .append(ir::Expression::Binary { op, left, right }, stmt.span);
2013                rctx.local_expression_kind_tracker
2014                    .insert(left, proc::ExpressionKind::Runtime);
2015                rctx.local_expression_kind_tracker
2016                    .insert(value, proc::ExpressionKind::Runtime);
2017
2018                block.extend(emitter.finish(&ctx.function.expressions));
2019                ir::Statement::Store {
2020                    pointer: target_handle,
2021                    value,
2022                }
2023            }
2024            ast::StatementKind::ConstAssert(condition) => {
2025                let mut emitter = proc::Emitter::default();
2026                emitter.start(&ctx.function.expressions);
2027
2028                let condition =
2029                    self.expression(condition, &mut ctx.as_const(block, &mut emitter))?;
2030
2031                let span = ctx.function.expressions.get_span(condition);
2032                match ctx
2033                    .module
2034                    .to_ctx()
2035                    .eval_expr_to_bool_from(condition, &ctx.function.expressions)
2036                {
2037                    Some(true) => Ok(()),
2038                    Some(false) => Err(Error::ConstAssertFailed(span)),
2039                    _ => Err(Error::NotBool(span)),
2040                }?;
2041
2042                block.extend(emitter.finish(&ctx.function.expressions));
2043
2044                return Ok(());
2045            }
2046            ast::StatementKind::Phony(expr) => {
2047                // Remembered the RHS of the phony assignment as a named expression. This
2048                // is important (1) to preserve the RHS for validation, (2) to track any
2049                // referenced globals.
2050                let mut emitter = proc::Emitter::default();
2051                emitter.start(&ctx.function.expressions);
2052
2053                let value = self.expression(expr, &mut ctx.as_expression(block, &mut emitter))?;
2054                block.extend(emitter.finish(&ctx.function.expressions));
2055                ctx.named_expressions
2056                    .insert(value, ("phony".to_string(), stmt.span));
2057                return Ok(());
2058            }
2059        };
2060
2061        block.push(out, stmt.span);
2062
2063        Ok(())
2064    }
2065
2066    /// Lower `expr` and apply the Load Rule if possible.
2067    ///
2068    /// For the time being, this concretizes abstract values, to support
2069    /// consumers that haven't been adapted to consume them yet. Consumers
2070    /// prepared for abstract values can call [`expression_for_abstract`].
2071    ///
2072    /// [`expression_for_abstract`]: Lowerer::expression_for_abstract
2073    fn expression(
2074        &mut self,
2075        expr: Handle<ast::Expression<'source>>,
2076        ctx: &mut ExpressionContext<'source, '_, '_>,
2077    ) -> Result<'source, Handle<ir::Expression>> {
2078        let expr = self.expression_for_abstract(expr, ctx)?;
2079        ctx.concretize(expr)
2080    }
2081
2082    fn expression_for_abstract(
2083        &mut self,
2084        expr: Handle<ast::Expression<'source>>,
2085        ctx: &mut ExpressionContext<'source, '_, '_>,
2086    ) -> Result<'source, Handle<ir::Expression>> {
2087        let expr = self.expression_for_reference(expr, ctx)?;
2088        ctx.apply_load_rule(expr)
2089    }
2090
2091    fn expression_with_leaf_scalar(
2092        &mut self,
2093        expr: Handle<ast::Expression<'source>>,
2094        scalar: ir::Scalar,
2095        ctx: &mut ExpressionContext<'source, '_, '_>,
2096    ) -> Result<'source, Handle<ir::Expression>> {
2097        let unconverted = self.expression_for_abstract(expr, ctx)?;
2098        ctx.try_automatic_conversion_for_leaf_scalar(unconverted, scalar, Span::default())
2099    }
2100
2101    fn expression_for_reference(
2102        &mut self,
2103        expr: Handle<ast::Expression<'source>>,
2104        ctx: &mut ExpressionContext<'source, '_, '_>,
2105    ) -> Result<'source, Typed<Handle<ir::Expression>>> {
2106        let span = ctx.ast_expressions.get_span(expr);
2107        let expr = &ctx.ast_expressions[expr];
2108
2109        let expr: Typed<ir::Expression> = match *expr {
2110            ast::Expression::Literal(literal) => {
2111                let literal = match literal {
2112                    ast::Literal::Number(Number::F16(f)) => ir::Literal::F16(f),
2113                    ast::Literal::Number(Number::F32(f)) => ir::Literal::F32(f),
2114                    ast::Literal::Number(Number::I32(i)) => ir::Literal::I32(i),
2115                    ast::Literal::Number(Number::U32(u)) => ir::Literal::U32(u),
2116                    ast::Literal::Number(Number::I64(i)) => ir::Literal::I64(i),
2117                    ast::Literal::Number(Number::U64(u)) => ir::Literal::U64(u),
2118                    ast::Literal::Number(Number::F64(f)) => ir::Literal::F64(f),
2119                    ast::Literal::Number(Number::AbstractInt(i)) => ir::Literal::AbstractInt(i),
2120                    ast::Literal::Number(Number::AbstractFloat(f)) => ir::Literal::AbstractFloat(f),
2121                    ast::Literal::Bool(b) => ir::Literal::Bool(b),
2122                };
2123                let handle = ctx.interrupt_emitter(ir::Expression::Literal(literal), span)?;
2124                return Ok(Typed::Plain(handle));
2125            }
2126            ast::Expression::Ident(ast::IdentExpr::Local(local)) => {
2127                return ctx.local(&local, span);
2128            }
2129            ast::Expression::Ident(ast::IdentExpr::Unresolved(name)) => {
2130                let global = ctx
2131                    .globals
2132                    .get(name)
2133                    .ok_or(Error::UnknownIdent(span, name))?;
2134                let expr = match *global {
2135                    LoweredGlobalDecl::Var(handle) => {
2136                        let expr = ir::Expression::GlobalVariable(handle);
2137                        match ctx.module.global_variables[handle].space {
2138                            ir::AddressSpace::Handle => Typed::Plain(expr),
2139                            _ => Typed::Reference(expr),
2140                        }
2141                    }
2142                    LoweredGlobalDecl::Const(handle) => {
2143                        Typed::Plain(ir::Expression::Constant(handle))
2144                    }
2145                    LoweredGlobalDecl::Override(handle) => {
2146                        Typed::Plain(ir::Expression::Override(handle))
2147                    }
2148                    LoweredGlobalDecl::Function { .. }
2149                    | LoweredGlobalDecl::Type(_)
2150                    | LoweredGlobalDecl::EntryPoint(_) => {
2151                        return Err(Box::new(Error::Unexpected(span, ExpectedToken::Variable)));
2152                    }
2153                };
2154
2155                return expr.try_map(|handle| ctx.interrupt_emitter(handle, span));
2156            }
2157            ast::Expression::Construct {
2158                ref ty,
2159                ty_span,
2160                ref components,
2161            } => {
2162                let handle = self.construct(span, ty, ty_span, components, ctx)?;
2163                return Ok(Typed::Plain(handle));
2164            }
2165            ast::Expression::Unary { op, expr } => {
2166                let expr = self.expression_for_abstract(expr, ctx)?;
2167                Typed::Plain(ir::Expression::Unary { op, expr })
2168            }
2169            ast::Expression::AddrOf(expr) => {
2170                // The `&` operator simply converts a reference to a pointer. And since a
2171                // reference is required, the Load Rule is not applied.
2172                match self.expression_for_reference(expr, ctx)? {
2173                    Typed::Reference(handle) => {
2174                        let expr = &ctx.runtime_expression_ctx(span)?.function.expressions[handle];
2175                        if let &ir::Expression::Access { base, .. }
2176                        | &ir::Expression::AccessIndex { base, .. } = expr
2177                        {
2178                            if let Some(ty) = resolve_inner!(ctx, base).pointer_base_type() {
2179                                if matches!(
2180                                    *ty.inner_with(&ctx.module.types),
2181                                    ir::TypeInner::Vector { .. },
2182                                ) {
2183                                    return Err(Box::new(Error::InvalidAddrOfOperand(
2184                                        ctx.get_expression_span(handle),
2185                                    )));
2186                                }
2187                            }
2188                        }
2189                        // No code is generated. We just declare the reference a pointer now.
2190                        return Ok(Typed::Plain(handle));
2191                    }
2192                    Typed::Plain(_) => {
2193                        return Err(Box::new(Error::NotReference(
2194                            "the operand of the `&` operator",
2195                            span,
2196                        )));
2197                    }
2198                }
2199            }
2200            ast::Expression::Deref(expr) => {
2201                // The pointer we dereference must be loaded.
2202                let pointer = self.expression(expr, ctx)?;
2203
2204                if resolve_inner!(ctx, pointer).pointer_space().is_none() {
2205                    return Err(Box::new(Error::NotPointer(span)));
2206                }
2207
2208                // No code is generated. We just declare the pointer a reference now.
2209                return Ok(Typed::Reference(pointer));
2210            }
2211            ast::Expression::Binary { op, left, right } => {
2212                self.binary(op, left, right, span, ctx)?
2213            }
2214            ast::Expression::Call {
2215                ref function,
2216                ref arguments,
2217            } => {
2218                let handle = self
2219                    .call(span, function, arguments, ctx, false)?
2220                    .ok_or(Error::FunctionReturnsVoid(function.span))?;
2221                return Ok(Typed::Plain(handle));
2222            }
2223            ast::Expression::Index { base, index } => {
2224                let mut lowered_base = self.expression_for_reference(base, ctx)?;
2225                let index = self.expression(index, ctx)?;
2226
2227                // <https://www.w3.org/TR/WGSL/#language_extension-pointer_composite_access>
2228                // Declare pointer as reference
2229                if let Typed::Plain(handle) = lowered_base {
2230                    if resolve_inner!(ctx, handle).pointer_space().is_some() {
2231                        lowered_base = Typed::Reference(handle);
2232                    }
2233                }
2234
2235                lowered_base.try_map(|base| match ctx.const_eval_expr_to_u32(index).ok() {
2236                    Some(index) => Ok::<_, Box<Error>>(ir::Expression::AccessIndex { base, index }),
2237                    None => {
2238                        // When an abstract array value e is indexed by an expression
2239                        // that is not a const-expression, then the array is concretized
2240                        // before the index is applied.
2241                        // https://www.w3.org/TR/WGSL/#array-access-expr
2242                        // Also applies to vectors and matrices.
2243                        let base = ctx.concretize(base)?;
2244                        Ok(ir::Expression::Access { base, index })
2245                    }
2246                })?
2247            }
2248            ast::Expression::Member { base, ref field } => {
2249                let mut lowered_base = self.expression_for_reference(base, ctx)?;
2250
2251                // <https://www.w3.org/TR/WGSL/#language_extension-pointer_composite_access>
2252                // Declare pointer as reference
2253                if let Typed::Plain(handle) = lowered_base {
2254                    if resolve_inner!(ctx, handle).pointer_space().is_some() {
2255                        lowered_base = Typed::Reference(handle);
2256                    }
2257                }
2258
2259                let temp_ty;
2260                let composite_type: &ir::TypeInner = match lowered_base {
2261                    Typed::Reference(handle) => {
2262                        temp_ty = resolve_inner!(ctx, handle)
2263                            .pointer_base_type()
2264                            .expect("In Typed::Reference(handle), handle must be a Naga pointer");
2265                        temp_ty.inner_with(&ctx.module.types)
2266                    }
2267
2268                    Typed::Plain(handle) => {
2269                        resolve_inner!(ctx, handle)
2270                    }
2271                };
2272
2273                let access = match *composite_type {
2274                    ir::TypeInner::Struct { ref members, .. } => {
2275                        let index = members
2276                            .iter()
2277                            .position(|m| m.name.as_deref() == Some(field.name))
2278                            .ok_or(Error::BadAccessor(field.span))?
2279                            as u32;
2280
2281                        lowered_base.map(|base| ir::Expression::AccessIndex { base, index })
2282                    }
2283                    ir::TypeInner::Vector { .. } => {
2284                        match Components::new(field.name, field.span)? {
2285                            Components::Swizzle { size, pattern } => {
2286                                Typed::Plain(ir::Expression::Swizzle {
2287                                    size,
2288                                    vector: ctx.apply_load_rule(lowered_base)?,
2289                                    pattern,
2290                                })
2291                            }
2292                            Components::Single(index) => {
2293                                lowered_base.map(|base| ir::Expression::AccessIndex { base, index })
2294                            }
2295                        }
2296                    }
2297                    _ => return Err(Box::new(Error::BadAccessor(field.span))),
2298                };
2299
2300                access
2301            }
2302            ast::Expression::Bitcast { expr, to, ty_span } => {
2303                let expr = self.expression(expr, ctx)?;
2304                let to_resolved = self.resolve_ast_type(to, &mut ctx.as_const())?;
2305
2306                let element_scalar = match ctx.module.types[to_resolved].inner {
2307                    ir::TypeInner::Scalar(scalar) => scalar,
2308                    ir::TypeInner::Vector { scalar, .. } => scalar,
2309                    _ => {
2310                        let ty = resolve!(ctx, expr);
2311                        return Err(Box::new(Error::BadTypeCast {
2312                            from_type: ctx.type_resolution_to_string(ty),
2313                            span: ty_span,
2314                            to_type: ctx.type_to_string(to_resolved),
2315                        }));
2316                    }
2317                };
2318
2319                Typed::Plain(ir::Expression::As {
2320                    expr,
2321                    kind: element_scalar.kind,
2322                    convert: None,
2323                })
2324            }
2325        };
2326
2327        expr.try_map(|handle| ctx.append_expression(handle, span))
2328    }
2329
2330    fn binary(
2331        &mut self,
2332        op: ir::BinaryOperator,
2333        left: Handle<ast::Expression<'source>>,
2334        right: Handle<ast::Expression<'source>>,
2335        span: Span,
2336        ctx: &mut ExpressionContext<'source, '_, '_>,
2337    ) -> Result<'source, Typed<ir::Expression>> {
2338        // Load both operands.
2339        let mut left = self.expression_for_abstract(left, ctx)?;
2340        let mut right = self.expression_for_abstract(right, ctx)?;
2341
2342        // Convert `scalar op vector` to `vector op vector` by introducing
2343        // `Splat` expressions.
2344        ctx.binary_op_splat(op, &mut left, &mut right)?;
2345
2346        // Apply automatic conversions.
2347        match op {
2348            ir::BinaryOperator::ShiftLeft | ir::BinaryOperator::ShiftRight => {
2349                // Shift operators require the right operand to be `u32` or
2350                // `vecN<u32>`. We can let the validator sort out vector length
2351                // issues, but the right operand must be, or convert to, a u32 leaf
2352                // scalar.
2353                right =
2354                    ctx.try_automatic_conversion_for_leaf_scalar(right, ir::Scalar::U32, span)?;
2355
2356                // Additionally, we must concretize the left operand if the right operand
2357                // is not a const-expression.
2358                // See https://www.w3.org/TR/WGSL/#overload-resolution-section.
2359                //
2360                // 2. Eliminate any candidate where one of its subexpressions resolves to
2361                // an abstract type after feasible automatic conversions, but another of
2362                // the candidate’s subexpressions is not a const-expression.
2363                //
2364                // We only have to explicitly do so for shifts as their operands may be
2365                // of different types - for other binary ops this is achieved by finding
2366                // the conversion consensus for both operands.
2367                if !ctx.is_const(right) {
2368                    left = ctx.concretize(left)?;
2369                }
2370            }
2371
2372            // All other operators follow the same pattern: reconcile the
2373            // scalar leaf types. If there's no reconciliation possible,
2374            // leave the expressions as they are: validation will report the
2375            // problem.
2376            _ => {
2377                ctx.grow_types(left)?;
2378                ctx.grow_types(right)?;
2379                if let Ok(consensus_scalar) =
2380                    ctx.automatic_conversion_consensus([left, right].iter())
2381                {
2382                    ctx.convert_to_leaf_scalar(&mut left, consensus_scalar)?;
2383                    ctx.convert_to_leaf_scalar(&mut right, consensus_scalar)?;
2384                }
2385            }
2386        }
2387
2388        Ok(Typed::Plain(ir::Expression::Binary { op, left, right }))
2389    }
2390
2391    /// Generate Naga IR for call expressions and statements, and type
2392    /// constructor expressions.
2393    ///
2394    /// The "function" being called is simply an `Ident` that we know refers to
2395    /// some module-scope definition.
2396    ///
2397    /// - If it is the name of a type, then the expression is a type constructor
2398    ///   expression: either constructing a value from components, a conversion
2399    ///   expression, or a zero value expression.
2400    ///
2401    /// - If it is the name of a function, then we're generating a [`Call`]
2402    ///   statement. We may be in the midst of generating code for an
2403    ///   expression, in which case we must generate an `Emit` statement to
2404    ///   force evaluation of the IR expressions we've generated so far, add the
2405    ///   `Call` statement to the current block, and then resume generating
2406    ///   expressions.
2407    ///
2408    /// [`Call`]: ir::Statement::Call
2409    fn call(
2410        &mut self,
2411        span: Span,
2412        function: &ast::Ident<'source>,
2413        arguments: &[Handle<ast::Expression<'source>>],
2414        ctx: &mut ExpressionContext<'source, '_, '_>,
2415        is_statement: bool,
2416    ) -> Result<'source, Option<Handle<ir::Expression>>> {
2417        let function_span = function.span;
2418        match ctx.globals.get(function.name) {
2419            Some(&LoweredGlobalDecl::Type(ty)) => {
2420                let handle = self.construct(
2421                    span,
2422                    &ast::ConstructorType::Type(ty),
2423                    function_span,
2424                    arguments,
2425                    ctx,
2426                )?;
2427                Ok(Some(handle))
2428            }
2429            Some(
2430                &LoweredGlobalDecl::Const(_)
2431                | &LoweredGlobalDecl::Override(_)
2432                | &LoweredGlobalDecl::Var(_),
2433            ) => Err(Box::new(Error::Unexpected(
2434                function_span,
2435                ExpectedToken::Function,
2436            ))),
2437            Some(&LoweredGlobalDecl::EntryPoint(_)) => {
2438                Err(Box::new(Error::CalledEntryPoint(function_span)))
2439            }
2440            Some(&LoweredGlobalDecl::Function {
2441                handle: function,
2442                must_use,
2443            }) => {
2444                let arguments = arguments
2445                    .iter()
2446                    .enumerate()
2447                    .map(|(i, &arg)| {
2448                        // Try to convert abstract values to the known argument types
2449                        let Some(&ir::FunctionArgument {
2450                            ty: parameter_ty, ..
2451                        }) = ctx.module.functions[function].arguments.get(i)
2452                        else {
2453                            // Wrong number of arguments... just concretize the type here
2454                            // and let the validator report the error.
2455                            return self.expression(arg, ctx);
2456                        };
2457
2458                        let expr = self.expression_for_abstract(arg, ctx)?;
2459                        ctx.try_automatic_conversions(
2460                            expr,
2461                            &proc::TypeResolution::Handle(parameter_ty),
2462                            ctx.ast_expressions.get_span(arg),
2463                        )
2464                    })
2465                    .collect::<Result<Vec<_>>>()?;
2466
2467                let has_result = ctx.module.functions[function].result.is_some();
2468
2469                if must_use && is_statement {
2470                    return Err(Box::new(Error::FunctionMustUseUnused(function_span)));
2471                }
2472
2473                let rctx = ctx.runtime_expression_ctx(span)?;
2474                // we need to always do this before a fn call since all arguments need to be emitted before the fn call
2475                rctx.block
2476                    .extend(rctx.emitter.finish(&rctx.function.expressions));
2477                let result = has_result.then(|| {
2478                    let result = rctx
2479                        .function
2480                        .expressions
2481                        .append(ir::Expression::CallResult(function), span);
2482                    rctx.local_expression_kind_tracker
2483                        .insert(result, proc::ExpressionKind::Runtime);
2484                    result
2485                });
2486                rctx.emitter.start(&rctx.function.expressions);
2487                rctx.block.push(
2488                    ir::Statement::Call {
2489                        function,
2490                        arguments,
2491                        result,
2492                    },
2493                    span,
2494                );
2495
2496                Ok(result)
2497            }
2498            None => {
2499                let span = function_span;
2500                let expr = if let Some(fun) = conv::map_relational_fun(function.name) {
2501                    let mut args = ctx.prepare_args(arguments, 1, span);
2502                    let argument = self.expression(args.next()?, ctx)?;
2503                    args.finish()?;
2504
2505                    // Check for no-op all(bool) and any(bool):
2506                    let argument_unmodified = matches!(
2507                        fun,
2508                        ir::RelationalFunction::All | ir::RelationalFunction::Any
2509                    ) && {
2510                        matches!(
2511                            resolve_inner!(ctx, argument),
2512                            &ir::TypeInner::Scalar(ir::Scalar {
2513                                kind: ir::ScalarKind::Bool,
2514                                ..
2515                            })
2516                        )
2517                    };
2518
2519                    if argument_unmodified {
2520                        return Ok(Some(argument));
2521                    } else {
2522                        ir::Expression::Relational { fun, argument }
2523                    }
2524                } else if let Some((axis, ctrl)) = conv::map_derivative(function.name) {
2525                    let mut args = ctx.prepare_args(arguments, 1, span);
2526                    let expr = self.expression(args.next()?, ctx)?;
2527                    args.finish()?;
2528
2529                    ir::Expression::Derivative { axis, ctrl, expr }
2530                } else if let Some(fun) = conv::map_standard_fun(function.name) {
2531                    self.math_function_helper(span, fun, arguments, ctx)?
2532                } else if let Some(fun) = Texture::map(function.name) {
2533                    self.texture_sample_helper(fun, arguments, span, ctx)?
2534                } else if let Some((op, cop)) = conv::map_subgroup_operation(function.name) {
2535                    return Ok(Some(
2536                        self.subgroup_operation_helper(span, op, cop, arguments, ctx)?,
2537                    ));
2538                } else if let Some(mode) = SubgroupGather::map(function.name) {
2539                    return Ok(Some(
2540                        self.subgroup_gather_helper(span, mode, arguments, ctx)?,
2541                    ));
2542                } else if let Some(fun) = ir::AtomicFunction::map(function.name) {
2543                    return self.atomic_helper(span, fun, arguments, is_statement, ctx);
2544                } else {
2545                    match function.name {
2546                        "select" => {
2547                            let mut args = ctx.prepare_args(arguments, 3, span);
2548
2549                            let reject_orig = args.next()?;
2550                            let accept_orig = args.next()?;
2551                            let mut values = [
2552                                self.expression_for_abstract(reject_orig, ctx)?,
2553                                self.expression_for_abstract(accept_orig, ctx)?,
2554                            ];
2555                            let condition = self.expression(args.next()?, ctx)?;
2556
2557                            args.finish()?;
2558
2559                            let diagnostic_details =
2560                                |ctx: &ExpressionContext<'_, '_, '_>,
2561                                 ty_res: &proc::TypeResolution,
2562                                 orig_expr| {
2563                                    (
2564                                        ctx.ast_expressions.get_span(orig_expr),
2565                                        format!("`{}`", ctx.as_diagnostic_display(ty_res)),
2566                                    )
2567                                };
2568                            for (&value, orig_value) in
2569                                values.iter().zip([reject_orig, accept_orig])
2570                            {
2571                                let value_ty_res = resolve!(ctx, value);
2572                                if value_ty_res
2573                                    .inner_with(&ctx.module.types)
2574                                    .vector_size_and_scalar()
2575                                    .is_none()
2576                                {
2577                                    let (arg_span, arg_type) =
2578                                        diagnostic_details(ctx, value_ty_res, orig_value);
2579                                    return Err(Box::new(Error::SelectUnexpectedArgumentType {
2580                                        arg_span,
2581                                        arg_type,
2582                                    }));
2583                                }
2584                            }
2585                            let mut consensus_scalar = ctx
2586                                .automatic_conversion_consensus(&values)
2587                                .map_err(|_idx| {
2588                                    let [reject, accept] = values;
2589                                    let [(reject_span, reject_type), (accept_span, accept_type)] =
2590                                        [(reject_orig, reject), (accept_orig, accept)].map(
2591                                            |(orig_expr, expr)| {
2592                                                let ty_res = &ctx.typifier()[expr];
2593                                                diagnostic_details(ctx, ty_res, orig_expr)
2594                                            },
2595                                        );
2596                                    Error::SelectRejectAndAcceptHaveNoCommonType {
2597                                        reject_span,
2598                                        reject_type,
2599                                        accept_span,
2600                                        accept_type,
2601                                    }
2602                                })?;
2603                            if !ctx.is_const(condition) {
2604                                consensus_scalar = consensus_scalar.concretize();
2605                            }
2606
2607                            ctx.convert_slice_to_common_leaf_scalar(&mut values, consensus_scalar)?;
2608
2609                            let [reject, accept] = values;
2610
2611                            ir::Expression::Select {
2612                                reject,
2613                                accept,
2614                                condition,
2615                            }
2616                        }
2617                        "arrayLength" => {
2618                            let mut args = ctx.prepare_args(arguments, 1, span);
2619                            let expr = self.expression(args.next()?, ctx)?;
2620                            args.finish()?;
2621
2622                            ir::Expression::ArrayLength(expr)
2623                        }
2624                        "atomicLoad" => {
2625                            let mut args = ctx.prepare_args(arguments, 1, span);
2626                            let (pointer, _scalar) = self.atomic_pointer(args.next()?, ctx)?;
2627                            args.finish()?;
2628
2629                            ir::Expression::Load { pointer }
2630                        }
2631                        "atomicStore" => {
2632                            let mut args = ctx.prepare_args(arguments, 2, span);
2633                            let (pointer, scalar) = self.atomic_pointer(args.next()?, ctx)?;
2634                            let value =
2635                                self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
2636                            args.finish()?;
2637
2638                            let rctx = ctx.runtime_expression_ctx(span)?;
2639                            rctx.block
2640                                .extend(rctx.emitter.finish(&rctx.function.expressions));
2641                            rctx.emitter.start(&rctx.function.expressions);
2642                            rctx.block
2643                                .push(ir::Statement::Store { pointer, value }, span);
2644                            return Ok(None);
2645                        }
2646                        "atomicCompareExchangeWeak" => {
2647                            let mut args = ctx.prepare_args(arguments, 3, span);
2648
2649                            let (pointer, scalar) = self.atomic_pointer(args.next()?, ctx)?;
2650
2651                            let compare =
2652                                self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
2653
2654                            let value = args.next()?;
2655                            let value_span = ctx.ast_expressions.get_span(value);
2656                            let value = self.expression_with_leaf_scalar(value, scalar, ctx)?;
2657
2658                            args.finish()?;
2659
2660                            let expression = match *resolve_inner!(ctx, value) {
2661                                ir::TypeInner::Scalar(scalar) => ir::Expression::AtomicResult {
2662                                    ty: ctx.module.generate_predeclared_type(
2663                                        ir::PredeclaredType::AtomicCompareExchangeWeakResult(
2664                                            scalar,
2665                                        ),
2666                                    ),
2667                                    comparison: true,
2668                                },
2669                                _ => {
2670                                    return Err(Box::new(Error::InvalidAtomicOperandType(
2671                                        value_span,
2672                                    )))
2673                                }
2674                            };
2675
2676                            let result = ctx.interrupt_emitter(expression, span)?;
2677                            let rctx = ctx.runtime_expression_ctx(span)?;
2678                            rctx.block.push(
2679                                ir::Statement::Atomic {
2680                                    pointer,
2681                                    fun: ir::AtomicFunction::Exchange {
2682                                        compare: Some(compare),
2683                                    },
2684                                    value,
2685                                    result: Some(result),
2686                                },
2687                                span,
2688                            );
2689                            return Ok(Some(result));
2690                        }
2691                        "textureAtomicMin" | "textureAtomicMax" | "textureAtomicAdd"
2692                        | "textureAtomicAnd" | "textureAtomicOr" | "textureAtomicXor" => {
2693                            let mut args = ctx.prepare_args(arguments, 3, span);
2694
2695                            let image = args.next()?;
2696                            let image_span = ctx.ast_expressions.get_span(image);
2697                            let image = self.expression(image, ctx)?;
2698
2699                            let coordinate = self.expression(args.next()?, ctx)?;
2700
2701                            let (_, arrayed) = ctx.image_data(image, image_span)?;
2702                            let array_index = arrayed
2703                                .then(|| {
2704                                    args.min_args += 1;
2705                                    self.expression(args.next()?, ctx)
2706                                })
2707                                .transpose()?;
2708
2709                            let value = self.expression(args.next()?, ctx)?;
2710
2711                            args.finish()?;
2712
2713                            let rctx = ctx.runtime_expression_ctx(span)?;
2714                            rctx.block
2715                                .extend(rctx.emitter.finish(&rctx.function.expressions));
2716                            rctx.emitter.start(&rctx.function.expressions);
2717                            let stmt = ir::Statement::ImageAtomic {
2718                                image,
2719                                coordinate,
2720                                array_index,
2721                                fun: match function.name {
2722                                    "textureAtomicMin" => ir::AtomicFunction::Min,
2723                                    "textureAtomicMax" => ir::AtomicFunction::Max,
2724                                    "textureAtomicAdd" => ir::AtomicFunction::Add,
2725                                    "textureAtomicAnd" => ir::AtomicFunction::And,
2726                                    "textureAtomicOr" => ir::AtomicFunction::InclusiveOr,
2727                                    "textureAtomicXor" => ir::AtomicFunction::ExclusiveOr,
2728                                    _ => unreachable!(),
2729                                },
2730                                value,
2731                            };
2732                            rctx.block.push(stmt, span);
2733                            return Ok(None);
2734                        }
2735                        "storageBarrier" => {
2736                            ctx.prepare_args(arguments, 0, span).finish()?;
2737
2738                            let rctx = ctx.runtime_expression_ctx(span)?;
2739                            rctx.block
2740                                .push(ir::Statement::ControlBarrier(ir::Barrier::STORAGE), span);
2741                            return Ok(None);
2742                        }
2743                        "workgroupBarrier" => {
2744                            ctx.prepare_args(arguments, 0, span).finish()?;
2745
2746                            let rctx = ctx.runtime_expression_ctx(span)?;
2747                            rctx.block
2748                                .push(ir::Statement::ControlBarrier(ir::Barrier::WORK_GROUP), span);
2749                            return Ok(None);
2750                        }
2751                        "subgroupBarrier" => {
2752                            ctx.prepare_args(arguments, 0, span).finish()?;
2753
2754                            let rctx = ctx.runtime_expression_ctx(span)?;
2755                            rctx.block
2756                                .push(ir::Statement::ControlBarrier(ir::Barrier::SUB_GROUP), span);
2757                            return Ok(None);
2758                        }
2759                        "textureBarrier" => {
2760                            ctx.prepare_args(arguments, 0, span).finish()?;
2761
2762                            let rctx = ctx.runtime_expression_ctx(span)?;
2763                            rctx.block
2764                                .push(ir::Statement::ControlBarrier(ir::Barrier::TEXTURE), span);
2765                            return Ok(None);
2766                        }
2767                        "workgroupUniformLoad" => {
2768                            let mut args = ctx.prepare_args(arguments, 1, span);
2769                            let expr = args.next()?;
2770                            args.finish()?;
2771
2772                            let pointer = self.expression(expr, ctx)?;
2773                            let result_ty = match *resolve_inner!(ctx, pointer) {
2774                                ir::TypeInner::Pointer {
2775                                    base,
2776                                    space: ir::AddressSpace::WorkGroup,
2777                                } => base,
2778                                ref other => {
2779                                    log::error!("Type {other:?} passed to workgroupUniformLoad");
2780                                    let span = ctx.ast_expressions.get_span(expr);
2781                                    return Err(Box::new(Error::InvalidWorkGroupUniformLoad(span)));
2782                                }
2783                            };
2784                            let result = ctx.interrupt_emitter(
2785                                ir::Expression::WorkGroupUniformLoadResult { ty: result_ty },
2786                                span,
2787                            )?;
2788                            let rctx = ctx.runtime_expression_ctx(span)?;
2789                            rctx.block.push(
2790                                ir::Statement::WorkGroupUniformLoad { pointer, result },
2791                                span,
2792                            );
2793
2794                            return Ok(Some(result));
2795                        }
2796                        "textureStore" => {
2797                            let mut args = ctx.prepare_args(arguments, 3, span);
2798
2799                            let image = args.next()?;
2800                            let image_span = ctx.ast_expressions.get_span(image);
2801                            let image = self.expression(image, ctx)?;
2802
2803                            let coordinate = self.expression(args.next()?, ctx)?;
2804
2805                            let (class, arrayed) = ctx.image_data(image, image_span)?;
2806                            let array_index = arrayed
2807                                .then(|| {
2808                                    args.min_args += 1;
2809                                    self.expression(args.next()?, ctx)
2810                                })
2811                                .transpose()?;
2812                            let scalar = if let ir::ImageClass::Storage { format, .. } = class {
2813                                format.into()
2814                            } else {
2815                                return Err(Box::new(Error::NotStorageTexture(image_span)));
2816                            };
2817
2818                            let value =
2819                                self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
2820
2821                            args.finish()?;
2822
2823                            let rctx = ctx.runtime_expression_ctx(span)?;
2824                            rctx.block
2825                                .extend(rctx.emitter.finish(&rctx.function.expressions));
2826                            rctx.emitter.start(&rctx.function.expressions);
2827                            let stmt = ir::Statement::ImageStore {
2828                                image,
2829                                coordinate,
2830                                array_index,
2831                                value,
2832                            };
2833                            rctx.block.push(stmt, span);
2834                            return Ok(None);
2835                        }
2836                        "textureLoad" => {
2837                            let mut args = ctx.prepare_args(arguments, 2, span);
2838
2839                            let image = args.next()?;
2840                            let image_span = ctx.ast_expressions.get_span(image);
2841                            let image = self.expression(image, ctx)?;
2842
2843                            let coordinate = self.expression(args.next()?, ctx)?;
2844
2845                            let (class, arrayed) = ctx.image_data(image, image_span)?;
2846                            let array_index = arrayed
2847                                .then(|| {
2848                                    args.min_args += 1;
2849                                    self.expression(args.next()?, ctx)
2850                                })
2851                                .transpose()?;
2852
2853                            let level = class
2854                                .is_mipmapped()
2855                                .then(|| {
2856                                    args.min_args += 1;
2857                                    self.expression(args.next()?, ctx)
2858                                })
2859                                .transpose()?;
2860
2861                            let sample = class
2862                                .is_multisampled()
2863                                .then(|| self.expression(args.next()?, ctx))
2864                                .transpose()?;
2865
2866                            args.finish()?;
2867
2868                            ir::Expression::ImageLoad {
2869                                image,
2870                                coordinate,
2871                                array_index,
2872                                level,
2873                                sample,
2874                            }
2875                        }
2876                        "textureDimensions" => {
2877                            let mut args = ctx.prepare_args(arguments, 1, span);
2878                            let image = self.expression(args.next()?, ctx)?;
2879                            let level = args
2880                                .next()
2881                                .map(|arg| self.expression(arg, ctx))
2882                                .ok()
2883                                .transpose()?;
2884                            args.finish()?;
2885
2886                            ir::Expression::ImageQuery {
2887                                image,
2888                                query: ir::ImageQuery::Size { level },
2889                            }
2890                        }
2891                        "textureNumLevels" => {
2892                            let mut args = ctx.prepare_args(arguments, 1, span);
2893                            let image = self.expression(args.next()?, ctx)?;
2894                            args.finish()?;
2895
2896                            ir::Expression::ImageQuery {
2897                                image,
2898                                query: ir::ImageQuery::NumLevels,
2899                            }
2900                        }
2901                        "textureNumLayers" => {
2902                            let mut args = ctx.prepare_args(arguments, 1, span);
2903                            let image = self.expression(args.next()?, ctx)?;
2904                            args.finish()?;
2905
2906                            ir::Expression::ImageQuery {
2907                                image,
2908                                query: ir::ImageQuery::NumLayers,
2909                            }
2910                        }
2911                        "textureNumSamples" => {
2912                            let mut args = ctx.prepare_args(arguments, 1, span);
2913                            let image = self.expression(args.next()?, ctx)?;
2914                            args.finish()?;
2915
2916                            ir::Expression::ImageQuery {
2917                                image,
2918                                query: ir::ImageQuery::NumSamples,
2919                            }
2920                        }
2921                        "rayQueryInitialize" => {
2922                            let mut args = ctx.prepare_args(arguments, 3, span);
2923                            let query = self.ray_query_pointer(args.next()?, ctx)?;
2924                            let acceleration_structure = self.expression(args.next()?, ctx)?;
2925                            let descriptor = self.expression(args.next()?, ctx)?;
2926                            args.finish()?;
2927
2928                            let _ = ctx.module.generate_ray_desc_type();
2929                            let fun = ir::RayQueryFunction::Initialize {
2930                                acceleration_structure,
2931                                descriptor,
2932                            };
2933
2934                            let rctx = ctx.runtime_expression_ctx(span)?;
2935                            rctx.block
2936                                .extend(rctx.emitter.finish(&rctx.function.expressions));
2937                            rctx.emitter.start(&rctx.function.expressions);
2938                            rctx.block
2939                                .push(ir::Statement::RayQuery { query, fun }, span);
2940                            return Ok(None);
2941                        }
2942                        "getCommittedHitVertexPositions" => {
2943                            let mut args = ctx.prepare_args(arguments, 1, span);
2944                            let query = self.ray_query_pointer(args.next()?, ctx)?;
2945                            args.finish()?;
2946
2947                            let _ = ctx.module.generate_vertex_return_type();
2948
2949                            ir::Expression::RayQueryVertexPositions {
2950                                query,
2951                                committed: true,
2952                            }
2953                        }
2954                        "getCandidateHitVertexPositions" => {
2955                            let mut args = ctx.prepare_args(arguments, 1, span);
2956                            let query = self.ray_query_pointer(args.next()?, ctx)?;
2957                            args.finish()?;
2958
2959                            let _ = ctx.module.generate_vertex_return_type();
2960
2961                            ir::Expression::RayQueryVertexPositions {
2962                                query,
2963                                committed: false,
2964                            }
2965                        }
2966                        "rayQueryProceed" => {
2967                            let mut args = ctx.prepare_args(arguments, 1, span);
2968                            let query = self.ray_query_pointer(args.next()?, ctx)?;
2969                            args.finish()?;
2970
2971                            let result =
2972                                ctx.interrupt_emitter(ir::Expression::RayQueryProceedResult, span)?;
2973                            let fun = ir::RayQueryFunction::Proceed { result };
2974                            let rctx = ctx.runtime_expression_ctx(span)?;
2975                            rctx.block
2976                                .push(ir::Statement::RayQuery { query, fun }, span);
2977                            return Ok(Some(result));
2978                        }
2979                        "rayQueryGenerateIntersection" => {
2980                            let mut args = ctx.prepare_args(arguments, 2, span);
2981                            let query = self.ray_query_pointer(args.next()?, ctx)?;
2982                            let hit_t = self.expression(args.next()?, ctx)?;
2983                            args.finish()?;
2984
2985                            let fun = ir::RayQueryFunction::GenerateIntersection { hit_t };
2986                            let rctx = ctx.runtime_expression_ctx(span)?;
2987                            rctx.block
2988                                .push(ir::Statement::RayQuery { query, fun }, span);
2989                            return Ok(None);
2990                        }
2991                        "rayQueryConfirmIntersection" => {
2992                            let mut args = ctx.prepare_args(arguments, 1, span);
2993                            let query = self.ray_query_pointer(args.next()?, ctx)?;
2994                            args.finish()?;
2995
2996                            let fun = ir::RayQueryFunction::ConfirmIntersection;
2997                            let rctx = ctx.runtime_expression_ctx(span)?;
2998                            rctx.block
2999                                .push(ir::Statement::RayQuery { query, fun }, span);
3000                            return Ok(None);
3001                        }
3002                        "rayQueryTerminate" => {
3003                            let mut args = ctx.prepare_args(arguments, 1, span);
3004                            let query = self.ray_query_pointer(args.next()?, ctx)?;
3005                            args.finish()?;
3006
3007                            let fun = ir::RayQueryFunction::Terminate;
3008                            let rctx = ctx.runtime_expression_ctx(span)?;
3009                            rctx.block
3010                                .push(ir::Statement::RayQuery { query, fun }, span);
3011                            return Ok(None);
3012                        }
3013                        "rayQueryGetCommittedIntersection" => {
3014                            let mut args = ctx.prepare_args(arguments, 1, span);
3015                            let query = self.ray_query_pointer(args.next()?, ctx)?;
3016                            args.finish()?;
3017
3018                            let _ = ctx.module.generate_ray_intersection_type();
3019                            ir::Expression::RayQueryGetIntersection {
3020                                query,
3021                                committed: true,
3022                            }
3023                        }
3024                        "rayQueryGetCandidateIntersection" => {
3025                            let mut args = ctx.prepare_args(arguments, 1, span);
3026                            let query = self.ray_query_pointer(args.next()?, ctx)?;
3027                            args.finish()?;
3028
3029                            let _ = ctx.module.generate_ray_intersection_type();
3030                            ir::Expression::RayQueryGetIntersection {
3031                                query,
3032                                committed: false,
3033                            }
3034                        }
3035                        "RayDesc" => {
3036                            let ty = ctx.module.generate_ray_desc_type();
3037                            let handle = self.construct(
3038                                span,
3039                                &ast::ConstructorType::Type(ty),
3040                                function.span,
3041                                arguments,
3042                                ctx,
3043                            )?;
3044                            return Ok(Some(handle));
3045                        }
3046                        "subgroupBallot" => {
3047                            let mut args = ctx.prepare_args(arguments, 0, span);
3048                            let predicate = if arguments.len() == 1 {
3049                                Some(self.expression(args.next()?, ctx)?)
3050                            } else {
3051                                None
3052                            };
3053                            args.finish()?;
3054
3055                            let result =
3056                                ctx.interrupt_emitter(ir::Expression::SubgroupBallotResult, span)?;
3057                            let rctx = ctx.runtime_expression_ctx(span)?;
3058                            rctx.block
3059                                .push(ir::Statement::SubgroupBallot { result, predicate }, span);
3060                            return Ok(Some(result));
3061                        }
3062                        "quadSwapX" => {
3063                            let mut args = ctx.prepare_args(arguments, 1, span);
3064
3065                            let argument = self.expression(args.next()?, ctx)?;
3066                            args.finish()?;
3067
3068                            let ty = ctx.register_type(argument)?;
3069
3070                            let result = ctx.interrupt_emitter(
3071                                crate::Expression::SubgroupOperationResult { ty },
3072                                span,
3073                            )?;
3074                            let rctx = ctx.runtime_expression_ctx(span)?;
3075                            rctx.block.push(
3076                                crate::Statement::SubgroupGather {
3077                                    mode: crate::GatherMode::QuadSwap(crate::Direction::X),
3078                                    argument,
3079                                    result,
3080                                },
3081                                span,
3082                            );
3083                            return Ok(Some(result));
3084                        }
3085
3086                        "quadSwapY" => {
3087                            let mut args = ctx.prepare_args(arguments, 1, span);
3088
3089                            let argument = self.expression(args.next()?, ctx)?;
3090                            args.finish()?;
3091
3092                            let ty = ctx.register_type(argument)?;
3093
3094                            let result = ctx.interrupt_emitter(
3095                                crate::Expression::SubgroupOperationResult { ty },
3096                                span,
3097                            )?;
3098                            let rctx = ctx.runtime_expression_ctx(span)?;
3099                            rctx.block.push(
3100                                crate::Statement::SubgroupGather {
3101                                    mode: crate::GatherMode::QuadSwap(crate::Direction::Y),
3102                                    argument,
3103                                    result,
3104                                },
3105                                span,
3106                            );
3107                            return Ok(Some(result));
3108                        }
3109
3110                        "quadSwapDiagonal" => {
3111                            let mut args = ctx.prepare_args(arguments, 1, span);
3112
3113                            let argument = self.expression(args.next()?, ctx)?;
3114                            args.finish()?;
3115
3116                            let ty = ctx.register_type(argument)?;
3117
3118                            let result = ctx.interrupt_emitter(
3119                                crate::Expression::SubgroupOperationResult { ty },
3120                                span,
3121                            )?;
3122                            let rctx = ctx.runtime_expression_ctx(span)?;
3123                            rctx.block.push(
3124                                crate::Statement::SubgroupGather {
3125                                    mode: crate::GatherMode::QuadSwap(crate::Direction::Diagonal),
3126                                    argument,
3127                                    result,
3128                                },
3129                                span,
3130                            );
3131                            return Ok(Some(result));
3132                        }
3133                        _ => {
3134                            return Err(Box::new(Error::UnknownIdent(function.span, function.name)))
3135                        }
3136                    }
3137                };
3138
3139                let expr = ctx.append_expression(expr, span)?;
3140                Ok(Some(expr))
3141            }
3142        }
3143    }
3144
3145    /// Generate a Naga IR [`Math`] expression.
3146    ///
3147    /// Generate Naga IR for a call to the [`MathFunction`] `fun`, whose
3148    /// unlowered arguments are `ast_arguments`.
3149    ///
3150    /// The `span` argument should give the span of the function name in the
3151    /// call expression.
3152    ///
3153    /// [`Math`]: ir::Expression::Math
3154    /// [`MathFunction`]: ir::MathFunction
3155    fn math_function_helper(
3156        &mut self,
3157        span: Span,
3158        fun: ir::MathFunction,
3159        ast_arguments: &[Handle<ast::Expression<'source>>],
3160        ctx: &mut ExpressionContext<'source, '_, '_>,
3161    ) -> Result<'source, ir::Expression> {
3162        let mut lowered_arguments = Vec::with_capacity(ast_arguments.len());
3163        for &arg in ast_arguments {
3164            let lowered = self.expression_for_abstract(arg, ctx)?;
3165            ctx.grow_types(lowered)?;
3166            lowered_arguments.push(lowered);
3167        }
3168
3169        let fun_overloads = fun.overloads();
3170        let rule = self.resolve_overloads(span, fun, fun_overloads, &lowered_arguments, ctx)?;
3171        self.apply_automatic_conversions_for_call(&rule, &mut lowered_arguments, ctx)?;
3172
3173        // If this function returns a predeclared type, register it
3174        // in `Module::special_types`. The typifier will expect to
3175        // be able to find it there.
3176        if let proc::Conclusion::Predeclared(predeclared) = rule.conclusion {
3177            ctx.module.generate_predeclared_type(predeclared);
3178        }
3179
3180        Ok(ir::Expression::Math {
3181            fun,
3182            arg: lowered_arguments[0],
3183            arg1: lowered_arguments.get(1).cloned(),
3184            arg2: lowered_arguments.get(2).cloned(),
3185            arg3: lowered_arguments.get(3).cloned(),
3186        })
3187    }
3188
3189    /// Choose the right overload for a function call.
3190    ///
3191    /// Return a [`Rule`] representing the most preferred overload in
3192    /// `overloads` to apply to `arguments`, or return an error explaining why
3193    /// the call is not valid.
3194    ///
3195    /// Use `fun` to identify the function being called in error messages;
3196    /// `span` should be the span of the function name in the call expression.
3197    ///
3198    /// [`Rule`]: proc::Rule
3199    fn resolve_overloads<O, F>(
3200        &self,
3201        span: Span,
3202        fun: F,
3203        overloads: O,
3204        arguments: &[Handle<ir::Expression>],
3205        ctx: &ExpressionContext<'source, '_, '_>,
3206    ) -> Result<'source, proc::Rule>
3207    where
3208        O: proc::OverloadSet,
3209        F: TryToWgsl + core::fmt::Debug + Copy,
3210    {
3211        let mut remaining_overloads = overloads.clone();
3212        let min_arguments = remaining_overloads.min_arguments();
3213        let max_arguments = remaining_overloads.max_arguments();
3214        if arguments.len() < min_arguments {
3215            return Err(Box::new(Error::WrongArgumentCount {
3216                span,
3217                expected: min_arguments as u32..max_arguments as u32,
3218                found: arguments.len() as u32,
3219            }));
3220        }
3221        if arguments.len() > max_arguments {
3222            return Err(Box::new(Error::TooManyArguments {
3223                function: fun.to_wgsl_for_diagnostics(),
3224                call_span: span,
3225                arg_span: ctx.get_expression_span(arguments[max_arguments]),
3226                max_arguments: max_arguments as _,
3227            }));
3228        }
3229
3230        log::debug!(
3231            "Initial overloads: {:#?}",
3232            remaining_overloads.for_debug(&ctx.module.types)
3233        );
3234
3235        for (arg_index, &arg) in arguments.iter().enumerate() {
3236            let arg_type_resolution = &ctx.typifier()[arg];
3237            let arg_inner = arg_type_resolution.inner_with(&ctx.module.types);
3238            log::debug!(
3239                "Supplying argument {arg_index} of type {:?}",
3240                arg_type_resolution.for_debug(&ctx.module.types)
3241            );
3242            let next_remaining_overloads =
3243                remaining_overloads.arg(arg_index, arg_inner, &ctx.module.types);
3244
3245            // If any argument is not a constant expression, then no overloads
3246            // that accept abstract values should be considered.
3247            // (`OverloadSet::concrete_only` is supposed to help impose this
3248            // restriction.) However, no `MathFunction` accepts a mix of
3249            // abstract and concrete arguments, so we don't need to worry
3250            // about that here.
3251
3252            log::debug!(
3253                "Remaining overloads: {:#?}",
3254                next_remaining_overloads.for_debug(&ctx.module.types)
3255            );
3256
3257            // If the set of remaining overloads is empty, then this argument's type
3258            // was unacceptable. Diagnose the problem and produce an error message.
3259            if next_remaining_overloads.is_empty() {
3260                let function = fun.to_wgsl_for_diagnostics();
3261                let call_span = span;
3262                let arg_span = ctx.get_expression_span(arg);
3263                let arg_ty = ctx.as_diagnostic_display(arg_type_resolution).to_string();
3264
3265                // Is this type *ever* permitted for the arg_index'th argument?
3266                // For example, `bool` is never permitted for `max`.
3267                let only_this_argument = overloads.arg(arg_index, arg_inner, &ctx.module.types);
3268                if only_this_argument.is_empty() {
3269                    // No overload of `fun` accepts this type as the
3270                    // arg_index'th argument. Determine the set of types that
3271                    // would ever be allowed there.
3272                    let allowed: Vec<String> = overloads
3273                        .allowed_args(arg_index, &ctx.module.to_ctx())
3274                        .iter()
3275                        .map(|ty| ctx.type_resolution_to_string(ty))
3276                        .collect();
3277
3278                    if allowed.is_empty() {
3279                        // No overload of `fun` accepts any argument at this
3280                        // index, so it's a simple case of excess arguments.
3281                        // However, since each `MathFunction`'s overloads all
3282                        // have the same arity, we should have detected this
3283                        // earlier.
3284                        unreachable!("expected all overloads to have the same arity");
3285                    }
3286
3287                    // Some overloads of `fun` do accept this many arguments,
3288                    // but none accept one of this type.
3289                    return Err(Box::new(Error::WrongArgumentType {
3290                        function,
3291                        call_span,
3292                        arg_span,
3293                        arg_index: arg_index as u32,
3294                        arg_ty,
3295                        allowed,
3296                    }));
3297                }
3298
3299                // This argument's type is accepted by some overloads---just
3300                // not those overloads that remain, given the prior arguments.
3301                // For example, `max` accepts `f32` as its second argument -
3302                // but not if the first was `i32`.
3303
3304                // Build a list of the types that would have been accepted here,
3305                // given the prior arguments.
3306                let allowed: Vec<String> = remaining_overloads
3307                    .allowed_args(arg_index, &ctx.module.to_ctx())
3308                    .iter()
3309                    .map(|ty| ctx.type_resolution_to_string(ty))
3310                    .collect();
3311
3312                // Re-run the argument list to determine which prior argument
3313                // made this one unacceptable.
3314                let mut remaining_overloads = overloads;
3315                for (prior_index, &prior_expr) in arguments.iter().enumerate() {
3316                    let prior_type_resolution = &ctx.typifier()[prior_expr];
3317                    let prior_ty = prior_type_resolution.inner_with(&ctx.module.types);
3318                    remaining_overloads =
3319                        remaining_overloads.arg(prior_index, prior_ty, &ctx.module.types);
3320                    if remaining_overloads
3321                        .arg(arg_index, arg_inner, &ctx.module.types)
3322                        .is_empty()
3323                    {
3324                        // This is the argument that killed our dreams.
3325                        let inconsistent_span = ctx.get_expression_span(arguments[prior_index]);
3326                        let inconsistent_ty =
3327                            ctx.as_diagnostic_display(prior_type_resolution).to_string();
3328
3329                        if allowed.is_empty() {
3330                            // Some overloads did accept `ty` at `arg_index`, but
3331                            // given the arguments up through `prior_expr`, we see
3332                            // no types acceptable at `arg_index`. This means that some
3333                            // overloads expect fewer arguments than others. However,
3334                            // each `MathFunction`'s overloads have the same arity, so this
3335                            // should be impossible.
3336                            unreachable!("expected all overloads to have the same arity");
3337                        }
3338
3339                        // Report `arg`'s type as inconsistent with `prior_expr`'s
3340                        return Err(Box::new(Error::InconsistentArgumentType {
3341                            function,
3342                            call_span,
3343                            arg_span,
3344                            arg_index: arg_index as u32,
3345                            arg_ty,
3346                            inconsistent_span,
3347                            inconsistent_index: prior_index as u32,
3348                            inconsistent_ty,
3349                            allowed,
3350                        }));
3351                    }
3352                }
3353                unreachable!("Failed to eliminate argument type when re-tried");
3354            }
3355            remaining_overloads = next_remaining_overloads;
3356        }
3357
3358        // Select the most preferred type rule for this call,
3359        // given the argument types supplied above.
3360        Ok(remaining_overloads.most_preferred())
3361    }
3362
3363    /// Apply automatic type conversions for a function call.
3364    ///
3365    /// Apply whatever automatic conversions are needed to pass `arguments` to
3366    /// the function overload described by `rule`. Update `arguments` to refer
3367    /// to the converted arguments.
3368    fn apply_automatic_conversions_for_call(
3369        &self,
3370        rule: &proc::Rule,
3371        arguments: &mut [Handle<ir::Expression>],
3372        ctx: &mut ExpressionContext<'source, '_, '_>,
3373    ) -> Result<'source, ()> {
3374        for (i, argument) in arguments.iter_mut().enumerate() {
3375            let goal_inner = rule.arguments[i].inner_with(&ctx.module.types);
3376            let converted = match goal_inner.scalar_for_conversions(&ctx.module.types) {
3377                Some(goal_scalar) => {
3378                    let arg_span = ctx.get_expression_span(*argument);
3379                    ctx.try_automatic_conversion_for_leaf_scalar(*argument, goal_scalar, arg_span)?
3380                }
3381                // No conversion is necessary.
3382                None => *argument,
3383            };
3384
3385            *argument = converted;
3386        }
3387
3388        Ok(())
3389    }
3390
3391    fn atomic_pointer(
3392        &mut self,
3393        expr: Handle<ast::Expression<'source>>,
3394        ctx: &mut ExpressionContext<'source, '_, '_>,
3395    ) -> Result<'source, (Handle<ir::Expression>, ir::Scalar)> {
3396        let span = ctx.ast_expressions.get_span(expr);
3397        let pointer = self.expression(expr, ctx)?;
3398
3399        match *resolve_inner!(ctx, pointer) {
3400            ir::TypeInner::Pointer { base, .. } => match ctx.module.types[base].inner {
3401                ir::TypeInner::Atomic(scalar) => Ok((pointer, scalar)),
3402                ref other => {
3403                    log::error!("Pointer type to {other:?} passed to atomic op");
3404                    Err(Box::new(Error::InvalidAtomicPointer(span)))
3405                }
3406            },
3407            ref other => {
3408                log::error!("Type {other:?} passed to atomic op");
3409                Err(Box::new(Error::InvalidAtomicPointer(span)))
3410            }
3411        }
3412    }
3413
3414    fn atomic_helper(
3415        &mut self,
3416        span: Span,
3417        fun: ir::AtomicFunction,
3418        args: &[Handle<ast::Expression<'source>>],
3419        is_statement: bool,
3420        ctx: &mut ExpressionContext<'source, '_, '_>,
3421    ) -> Result<'source, Option<Handle<ir::Expression>>> {
3422        let mut args = ctx.prepare_args(args, 2, span);
3423
3424        let (pointer, scalar) = self.atomic_pointer(args.next()?, ctx)?;
3425        let value = self.expression_with_leaf_scalar(args.next()?, scalar, ctx)?;
3426        let value_inner = resolve_inner!(ctx, value);
3427        args.finish()?;
3428
3429        // If we don't use the return value of a 64-bit `min` or `max`
3430        // operation, generate a no-result form of the `Atomic` statement, so
3431        // that we can pass validation with only `SHADER_INT64_ATOMIC_MIN_MAX`
3432        // whenever possible.
3433        let is_64_bit_min_max = matches!(fun, ir::AtomicFunction::Min | ir::AtomicFunction::Max)
3434            && matches!(
3435                *value_inner,
3436                ir::TypeInner::Scalar(ir::Scalar { width: 8, .. })
3437            );
3438        let result = if is_64_bit_min_max && is_statement {
3439            let rctx = ctx.runtime_expression_ctx(span)?;
3440            rctx.block
3441                .extend(rctx.emitter.finish(&rctx.function.expressions));
3442            rctx.emitter.start(&rctx.function.expressions);
3443            None
3444        } else {
3445            let ty = ctx.register_type(value)?;
3446            Some(ctx.interrupt_emitter(
3447                ir::Expression::AtomicResult {
3448                    ty,
3449                    comparison: false,
3450                },
3451                span,
3452            )?)
3453        };
3454        let rctx = ctx.runtime_expression_ctx(span)?;
3455        rctx.block.push(
3456            ir::Statement::Atomic {
3457                pointer,
3458                fun,
3459                value,
3460                result,
3461            },
3462            span,
3463        );
3464        Ok(result)
3465    }
3466
3467    fn texture_sample_helper(
3468        &mut self,
3469        fun: Texture,
3470        args: &[Handle<ast::Expression<'source>>],
3471        span: Span,
3472        ctx: &mut ExpressionContext<'source, '_, '_>,
3473    ) -> Result<'source, ir::Expression> {
3474        let mut args = ctx.prepare_args(args, fun.min_argument_count(), span);
3475
3476        fn get_image_and_span<'source>(
3477            lowerer: &mut Lowerer<'source, '_>,
3478            args: &mut ArgumentContext<'_, 'source>,
3479            ctx: &mut ExpressionContext<'source, '_, '_>,
3480        ) -> Result<'source, (Handle<ir::Expression>, Span)> {
3481            let image = args.next()?;
3482            let image_span = ctx.ast_expressions.get_span(image);
3483            let image = lowerer.expression_for_abstract(image, ctx)?;
3484            Ok((image, image_span))
3485        }
3486
3487        let image;
3488        let image_span;
3489        let gather;
3490        match fun {
3491            Texture::Gather => {
3492                let image_or_component = args.next()?;
3493                let image_or_component_span = ctx.ast_expressions.get_span(image_or_component);
3494                // Gathers from depth textures don't take an initial `component` argument.
3495                let lowered_image_or_component = self.expression(image_or_component, ctx)?;
3496
3497                match *resolve_inner!(ctx, lowered_image_or_component) {
3498                    ir::TypeInner::Image {
3499                        class: ir::ImageClass::Depth { .. },
3500                        ..
3501                    } => {
3502                        image = lowered_image_or_component;
3503                        image_span = image_or_component_span;
3504                        gather = Some(ir::SwizzleComponent::X);
3505                    }
3506                    _ => {
3507                        (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
3508                        gather = Some(ctx.gather_component(
3509                            lowered_image_or_component,
3510                            image_or_component_span,
3511                            span,
3512                        )?);
3513                    }
3514                }
3515            }
3516            Texture::GatherCompare => {
3517                (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
3518                gather = Some(ir::SwizzleComponent::X);
3519            }
3520
3521            _ => {
3522                (image, image_span) = get_image_and_span(self, &mut args, ctx)?;
3523                gather = None;
3524            }
3525        };
3526
3527        let sampler = self.expression_for_abstract(args.next()?, ctx)?;
3528
3529        let coordinate = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3530        let clamp_to_edge = matches!(fun, Texture::SampleBaseClampToEdge);
3531
3532        let (class, arrayed) = ctx.image_data(image, image_span)?;
3533        let array_index = arrayed
3534            .then(|| self.expression(args.next()?, ctx))
3535            .transpose()?;
3536
3537        let level;
3538        let depth_ref;
3539        match fun {
3540            Texture::Gather => {
3541                level = ir::SampleLevel::Zero;
3542                depth_ref = None;
3543            }
3544            Texture::GatherCompare => {
3545                let reference =
3546                    self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3547                level = ir::SampleLevel::Zero;
3548                depth_ref = Some(reference);
3549            }
3550
3551            Texture::Sample => {
3552                level = ir::SampleLevel::Auto;
3553                depth_ref = None;
3554            }
3555            Texture::SampleBias => {
3556                let bias = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3557                level = ir::SampleLevel::Bias(bias);
3558                depth_ref = None;
3559            }
3560            Texture::SampleCompare => {
3561                let reference =
3562                    self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3563                level = ir::SampleLevel::Auto;
3564                depth_ref = Some(reference);
3565            }
3566            Texture::SampleCompareLevel => {
3567                let reference =
3568                    self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3569                level = ir::SampleLevel::Zero;
3570                depth_ref = Some(reference);
3571            }
3572            Texture::SampleGrad => {
3573                let x = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3574                let y = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
3575                level = ir::SampleLevel::Gradient { x, y };
3576                depth_ref = None;
3577            }
3578            Texture::SampleLevel => {
3579                let exact = match class {
3580                    // When applied to depth textures, `textureSampleLevel`'s
3581                    // `level` argument is an `i32` or `u32`.
3582                    ir::ImageClass::Depth { .. } => self.expression(args.next()?, ctx)?,
3583
3584                    // When applied to other sampled types, its `level` argument
3585                    // is an `f32`.
3586                    ir::ImageClass::Sampled { .. } => {
3587                        self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?
3588                    }
3589
3590                    // Sampling `External` textures with a specified level isn't
3591                    // allowed, and sampling `Storage` textures isn't allowed at
3592                    // all. Let the validator report the error.
3593                    ir::ImageClass::Storage { .. } | ir::ImageClass::External => {
3594                        self.expression(args.next()?, ctx)?
3595                    }
3596                };
3597                level = ir::SampleLevel::Exact(exact);
3598                depth_ref = None;
3599            }
3600            Texture::SampleBaseClampToEdge => {
3601                level = crate::SampleLevel::Zero;
3602                depth_ref = None;
3603            }
3604        };
3605
3606        let offset = args
3607            .next()
3608            .map(|arg| self.expression_with_leaf_scalar(arg, ir::Scalar::I32, &mut ctx.as_const()))
3609            .ok()
3610            .transpose()?;
3611
3612        args.finish()?;
3613
3614        Ok(ir::Expression::ImageSample {
3615            image,
3616            sampler,
3617            gather,
3618            coordinate,
3619            array_index,
3620            offset,
3621            level,
3622            depth_ref,
3623            clamp_to_edge,
3624        })
3625    }
3626
3627    fn subgroup_operation_helper(
3628        &mut self,
3629        span: Span,
3630        op: ir::SubgroupOperation,
3631        collective_op: ir::CollectiveOperation,
3632        arguments: &[Handle<ast::Expression<'source>>],
3633        ctx: &mut ExpressionContext<'source, '_, '_>,
3634    ) -> Result<'source, Handle<ir::Expression>> {
3635        let mut args = ctx.prepare_args(arguments, 1, span);
3636
3637        let argument = self.expression(args.next()?, ctx)?;
3638        args.finish()?;
3639
3640        let ty = ctx.register_type(argument)?;
3641
3642        let result = ctx.interrupt_emitter(ir::Expression::SubgroupOperationResult { ty }, span)?;
3643        let rctx = ctx.runtime_expression_ctx(span)?;
3644        rctx.block.push(
3645            ir::Statement::SubgroupCollectiveOperation {
3646                op,
3647                collective_op,
3648                argument,
3649                result,
3650            },
3651            span,
3652        );
3653        Ok(result)
3654    }
3655
3656    fn subgroup_gather_helper(
3657        &mut self,
3658        span: Span,
3659        mode: SubgroupGather,
3660        arguments: &[Handle<ast::Expression<'source>>],
3661        ctx: &mut ExpressionContext<'source, '_, '_>,
3662    ) -> Result<'source, Handle<ir::Expression>> {
3663        let mut args = ctx.prepare_args(arguments, 2, span);
3664
3665        let argument = self.expression(args.next()?, ctx)?;
3666
3667        use SubgroupGather as Sg;
3668        let mode = if let Sg::BroadcastFirst = mode {
3669            ir::GatherMode::BroadcastFirst
3670        } else {
3671            let index = self.expression(args.next()?, ctx)?;
3672            match mode {
3673                Sg::BroadcastFirst => unreachable!(),
3674                Sg::Broadcast => ir::GatherMode::Broadcast(index),
3675                Sg::Shuffle => ir::GatherMode::Shuffle(index),
3676                Sg::ShuffleDown => ir::GatherMode::ShuffleDown(index),
3677                Sg::ShuffleUp => ir::GatherMode::ShuffleUp(index),
3678                Sg::ShuffleXor => ir::GatherMode::ShuffleXor(index),
3679                Sg::QuadBroadcast => ir::GatherMode::QuadBroadcast(index),
3680            }
3681        };
3682
3683        args.finish()?;
3684
3685        let ty = ctx.register_type(argument)?;
3686
3687        let result = ctx.interrupt_emitter(ir::Expression::SubgroupOperationResult { ty }, span)?;
3688        let rctx = ctx.runtime_expression_ctx(span)?;
3689        rctx.block.push(
3690            ir::Statement::SubgroupGather {
3691                mode,
3692                argument,
3693                result,
3694            },
3695            span,
3696        );
3697        Ok(result)
3698    }
3699
3700    fn r#struct(
3701        &mut self,
3702        s: &ast::Struct<'source>,
3703        span: Span,
3704        ctx: &mut GlobalContext<'source, '_, '_>,
3705    ) -> Result<'source, Handle<ir::Type>> {
3706        let mut offset = 0;
3707        let mut struct_alignment = proc::Alignment::ONE;
3708        let mut members = Vec::with_capacity(s.members.len());
3709
3710        let mut doc_comments: Vec<Option<Vec<String>>> = Vec::new();
3711
3712        for member in s.members.iter() {
3713            let ty = self.resolve_ast_type(member.ty, &mut ctx.as_const())?;
3714
3715            ctx.layouter.update(ctx.module.to_ctx()).map_err(|err| {
3716                let LayoutErrorInner::TooLarge = err.inner else {
3717                    unreachable!("unexpected layout error: {err:?}");
3718                };
3719                // Since anonymous types of struct members don't get a span,
3720                // associate the error with the member. The layouter could have
3721                // failed on any type that was pending layout, but if it wasn't
3722                // the current struct member, it wasn't a struct member at all,
3723                // because we resolve struct members one-by-one.
3724                if ty == err.ty {
3725                    Box::new(Error::StructMemberTooLarge {
3726                        member_name_span: member.name.span,
3727                    })
3728                } else {
3729                    // Lots of type definitions don't get spans, so this error
3730                    // message may not be very useful.
3731                    Box::new(Error::TypeTooLarge {
3732                        span: ctx.module.types.get_span(err.ty),
3733                    })
3734                }
3735            })?;
3736
3737            let member_min_size = ctx.layouter[ty].size;
3738            let member_min_alignment = ctx.layouter[ty].alignment;
3739
3740            let member_size = if let Some(size_expr) = member.size {
3741                let (size, span) = self.const_u32(size_expr, &mut ctx.as_const())?;
3742                if size < member_min_size {
3743                    return Err(Box::new(Error::SizeAttributeTooLow(span, member_min_size)));
3744                } else {
3745                    size
3746                }
3747            } else {
3748                member_min_size
3749            };
3750
3751            let member_alignment = if let Some(align_expr) = member.align {
3752                let (align, span) = self.const_u32(align_expr, &mut ctx.as_const())?;
3753                if let Some(alignment) = proc::Alignment::new(align) {
3754                    if alignment < member_min_alignment {
3755                        return Err(Box::new(Error::AlignAttributeTooLow(
3756                            span,
3757                            member_min_alignment,
3758                        )));
3759                    } else {
3760                        alignment
3761                    }
3762                } else {
3763                    return Err(Box::new(Error::NonPowerOfTwoAlignAttribute(span)));
3764                }
3765            } else {
3766                member_min_alignment
3767            };
3768
3769            let binding = self.binding(&member.binding, ty, ctx)?;
3770
3771            offset = member_alignment.round_up(offset);
3772            struct_alignment = struct_alignment.max(member_alignment);
3773
3774            if !member.doc_comments.is_empty() {
3775                doc_comments.push(Some(
3776                    member.doc_comments.iter().map(|s| s.to_string()).collect(),
3777                ));
3778            }
3779            members.push(ir::StructMember {
3780                name: Some(member.name.name.to_owned()),
3781                ty,
3782                binding,
3783                offset,
3784            });
3785
3786            offset += member_size;
3787            if offset > crate::valid::MAX_TYPE_SIZE {
3788                return Err(Box::new(Error::TypeTooLarge { span }));
3789            }
3790        }
3791
3792        let size = struct_alignment.round_up(offset);
3793        let inner = ir::TypeInner::Struct {
3794            members,
3795            span: size,
3796        };
3797
3798        let handle = ctx.module.types.insert(
3799            ir::Type {
3800                name: Some(s.name.name.to_string()),
3801                inner,
3802            },
3803            span,
3804        );
3805        for (i, c) in doc_comments.drain(..).enumerate() {
3806            if let Some(comment) = c {
3807                ctx.module
3808                    .get_or_insert_default_doc_comments()
3809                    .struct_members
3810                    .insert((handle, i), comment);
3811            }
3812        }
3813        Ok(handle)
3814    }
3815
3816    fn const_u32(
3817        &mut self,
3818        expr: Handle<ast::Expression<'source>>,
3819        ctx: &mut ExpressionContext<'source, '_, '_>,
3820    ) -> Result<'source, (u32, Span)> {
3821        let span = ctx.ast_expressions.get_span(expr);
3822        let expr = self.expression(expr, ctx)?;
3823        let value = ctx
3824            .module
3825            .to_ctx()
3826            .eval_expr_to_u32(expr)
3827            .map_err(|err| match err {
3828                proc::U32EvalError::NonConst => Error::ExpectedConstExprConcreteIntegerScalar(span),
3829                proc::U32EvalError::Negative => Error::ExpectedNonNegative(span),
3830            })?;
3831        Ok((value, span))
3832    }
3833
3834    fn array_size(
3835        &mut self,
3836        size: ast::ArraySize<'source>,
3837        ctx: &mut ExpressionContext<'source, '_, '_>,
3838    ) -> Result<'source, ir::ArraySize> {
3839        Ok(match size {
3840            ast::ArraySize::Constant(expr) => {
3841                let span = ctx.ast_expressions.get_span(expr);
3842                let const_expr = self.expression(expr, &mut ctx.as_const());
3843                match const_expr {
3844                    Ok(value) => {
3845                        let len = ctx.const_eval_expr_to_u32(value).map_err(|err| {
3846                            Box::new(match err {
3847                                proc::U32EvalError::NonConst => {
3848                                    Error::ExpectedConstExprConcreteIntegerScalar(span)
3849                                }
3850                                proc::U32EvalError::Negative => {
3851                                    Error::ExpectedPositiveArrayLength(span)
3852                                }
3853                            })
3854                        })?;
3855                        let size =
3856                            NonZeroU32::new(len).ok_or(Error::ExpectedPositiveArrayLength(span))?;
3857                        ir::ArraySize::Constant(size)
3858                    }
3859                    Err(err) => {
3860                        if let Error::ConstantEvaluatorError(ref ty, _) = *err {
3861                            match **ty {
3862                                proc::ConstantEvaluatorError::OverrideExpr => {
3863                                    ir::ArraySize::Pending(self.array_size_override(
3864                                        expr,
3865                                        &mut ctx.as_global().as_override(),
3866                                        span,
3867                                    )?)
3868                                }
3869                                _ => {
3870                                    return Err(err);
3871                                }
3872                            }
3873                        } else {
3874                            return Err(err);
3875                        }
3876                    }
3877                }
3878            }
3879            ast::ArraySize::Dynamic => ir::ArraySize::Dynamic,
3880        })
3881    }
3882
3883    fn array_size_override(
3884        &mut self,
3885        size_expr: Handle<ast::Expression<'source>>,
3886        ctx: &mut ExpressionContext<'source, '_, '_>,
3887        span: Span,
3888    ) -> Result<'source, Handle<ir::Override>> {
3889        let expr = self.expression(size_expr, ctx)?;
3890        match resolve_inner!(ctx, expr).scalar_kind().ok_or(0) {
3891            Ok(ir::ScalarKind::Sint) | Ok(ir::ScalarKind::Uint) => Ok({
3892                if let ir::Expression::Override(handle) = ctx.module.global_expressions[expr] {
3893                    handle
3894                } else {
3895                    let ty = ctx.register_type(expr)?;
3896                    ctx.module.overrides.append(
3897                        ir::Override {
3898                            name: None,
3899                            id: None,
3900                            ty,
3901                            init: Some(expr),
3902                        },
3903                        span,
3904                    )
3905                }
3906            }),
3907            _ => Err(Box::new(Error::ExpectedConstExprConcreteIntegerScalar(
3908                span,
3909            ))),
3910        }
3911    }
3912
3913    /// Build the Naga equivalent of a named AST type.
3914    ///
3915    /// Return a Naga `Handle<Type>` representing the front-end type
3916    /// `handle`, which should be named `name`, if given.
3917    ///
3918    /// If `handle` refers to a type cached in [`SpecialTypes`],
3919    /// `name` may be ignored.
3920    ///
3921    /// [`SpecialTypes`]: ir::SpecialTypes
3922    fn resolve_named_ast_type(
3923        &mut self,
3924        handle: Handle<ast::Type<'source>>,
3925        name: Option<String>,
3926        ctx: &mut ExpressionContext<'source, '_, '_>,
3927    ) -> Result<'source, Handle<ir::Type>> {
3928        let inner = match ctx.types[handle] {
3929            ast::Type::Scalar(scalar) => scalar.to_inner_scalar(),
3930            ast::Type::Vector { size, ty, ty_span } => {
3931                let ty = self.resolve_ast_type(ty, ctx)?;
3932                let scalar = match ctx.module.types[ty].inner {
3933                    ir::TypeInner::Scalar(sc) => sc,
3934                    _ => return Err(Box::new(Error::UnknownScalarType(ty_span))),
3935                };
3936                ir::TypeInner::Vector { size, scalar }
3937            }
3938            ast::Type::Matrix {
3939                rows,
3940                columns,
3941                ty,
3942                ty_span,
3943            } => {
3944                let ty = self.resolve_ast_type(ty, ctx)?;
3945                let scalar = match ctx.module.types[ty].inner {
3946                    ir::TypeInner::Scalar(sc) => sc,
3947                    _ => return Err(Box::new(Error::UnknownScalarType(ty_span))),
3948                };
3949                match scalar.kind {
3950                    ir::ScalarKind::Float => ir::TypeInner::Matrix {
3951                        columns,
3952                        rows,
3953                        scalar,
3954                    },
3955                    _ => return Err(Box::new(Error::BadMatrixScalarKind(ty_span, scalar))),
3956                }
3957            }
3958            ast::Type::Atomic(scalar) => scalar.to_inner_atomic(),
3959            ast::Type::Pointer { base, space } => {
3960                let base = self.resolve_ast_type(base, ctx)?;
3961                ir::TypeInner::Pointer { base, space }
3962            }
3963            ast::Type::Array { base, size } => {
3964                let base = self.resolve_ast_type(base, &mut ctx.as_const())?;
3965                let size = self.array_size(size, ctx)?;
3966
3967                // Determine the size of the base type, if needed.
3968                ctx.layouter.update(ctx.module.to_ctx()).map_err(|err| {
3969                    let LayoutErrorInner::TooLarge = err.inner else {
3970                        unreachable!("unexpected layout error: {err:?}");
3971                    };
3972                    // Lots of type definitions don't get spans, so this error
3973                    // message may not be very useful.
3974                    Box::new(Error::TypeTooLarge {
3975                        span: ctx.module.types.get_span(err.ty),
3976                    })
3977                })?;
3978                let stride = ctx.layouter[base].to_stride();
3979
3980                ir::TypeInner::Array { base, size, stride }
3981            }
3982            ast::Type::Image {
3983                dim,
3984                arrayed,
3985                class,
3986            } => ir::TypeInner::Image {
3987                dim,
3988                arrayed,
3989                class,
3990            },
3991            ast::Type::Sampler { comparison } => ir::TypeInner::Sampler { comparison },
3992            ast::Type::AccelerationStructure { vertex_return } => {
3993                ir::TypeInner::AccelerationStructure { vertex_return }
3994            }
3995            ast::Type::RayQuery { vertex_return } => ir::TypeInner::RayQuery { vertex_return },
3996            ast::Type::BindingArray { base, size } => {
3997                let base = self.resolve_ast_type(base, ctx)?;
3998                let size = self.array_size(size, ctx)?;
3999                ir::TypeInner::BindingArray { base, size }
4000            }
4001            ast::Type::RayDesc => {
4002                return Ok(ctx.module.generate_ray_desc_type());
4003            }
4004            ast::Type::RayIntersection => {
4005                return Ok(ctx.module.generate_ray_intersection_type());
4006            }
4007            ast::Type::User(ref ident) => {
4008                return match ctx.globals.get(ident.name) {
4009                    Some(&LoweredGlobalDecl::Type(handle)) => Ok(handle),
4010                    Some(_) => Err(Box::new(Error::Unexpected(ident.span, ExpectedToken::Type))),
4011                    None => Err(Box::new(Error::UnknownType(ident.span))),
4012                }
4013            }
4014        };
4015
4016        Ok(ctx.as_global().ensure_type_exists(name, inner))
4017    }
4018
4019    /// Return a Naga `Handle<Type>` representing the front-end type `handle`.
4020    fn resolve_ast_type(
4021        &mut self,
4022        handle: Handle<ast::Type<'source>>,
4023        ctx: &mut ExpressionContext<'source, '_, '_>,
4024    ) -> Result<'source, Handle<ir::Type>> {
4025        self.resolve_named_ast_type(handle, None, ctx)
4026    }
4027
4028    fn binding(
4029        &mut self,
4030        binding: &Option<ast::Binding<'source>>,
4031        ty: Handle<ir::Type>,
4032        ctx: &mut GlobalContext<'source, '_, '_>,
4033    ) -> Result<'source, Option<ir::Binding>> {
4034        Ok(match *binding {
4035            Some(ast::Binding::BuiltIn(b)) => Some(ir::Binding::BuiltIn(b)),
4036            Some(ast::Binding::Location {
4037                location,
4038                interpolation,
4039                sampling,
4040                blend_src,
4041            }) => {
4042                let blend_src = if let Some(blend_src) = blend_src {
4043                    Some(self.const_u32(blend_src, &mut ctx.as_const())?.0)
4044                } else {
4045                    None
4046                };
4047
4048                let mut binding = ir::Binding::Location {
4049                    location: self.const_u32(location, &mut ctx.as_const())?.0,
4050                    interpolation,
4051                    sampling,
4052                    blend_src,
4053                };
4054                binding.apply_default_interpolation(&ctx.module.types[ty].inner);
4055                Some(binding)
4056            }
4057            None => None,
4058        })
4059    }
4060
4061    fn ray_query_pointer(
4062        &mut self,
4063        expr: Handle<ast::Expression<'source>>,
4064        ctx: &mut ExpressionContext<'source, '_, '_>,
4065    ) -> Result<'source, Handle<ir::Expression>> {
4066        let span = ctx.ast_expressions.get_span(expr);
4067        let pointer = self.expression(expr, ctx)?;
4068
4069        match *resolve_inner!(ctx, pointer) {
4070            ir::TypeInner::Pointer { base, .. } => match ctx.module.types[base].inner {
4071                ir::TypeInner::RayQuery { .. } => Ok(pointer),
4072                ref other => {
4073                    log::error!("Pointer type to {other:?} passed to ray query op");
4074                    Err(Box::new(Error::InvalidRayQueryPointer(span)))
4075                }
4076            },
4077            ref other => {
4078                log::error!("Type {other:?} passed to ray query op");
4079                Err(Box::new(Error::InvalidRayQueryPointer(span)))
4080            }
4081        }
4082    }
4083}
4084
4085impl ir::AtomicFunction {
4086    pub fn map(word: &str) -> Option<Self> {
4087        Some(match word {
4088            "atomicAdd" => ir::AtomicFunction::Add,
4089            "atomicSub" => ir::AtomicFunction::Subtract,
4090            "atomicAnd" => ir::AtomicFunction::And,
4091            "atomicOr" => ir::AtomicFunction::InclusiveOr,
4092            "atomicXor" => ir::AtomicFunction::ExclusiveOr,
4093            "atomicMin" => ir::AtomicFunction::Min,
4094            "atomicMax" => ir::AtomicFunction::Max,
4095            "atomicExchange" => ir::AtomicFunction::Exchange { compare: None },
4096            _ => return None,
4097        })
4098    }
4099}