naga/front/wgsl/lower/
mod.rs

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