Quellcodebibliothek Statistik Leitseite products/Sources/formale Sprachen/C/Firefox/third_party/rust/naga/src/front/wgsl/lower/   (Browser von der Mozilla Stiftung Version 136.0.1©)  Datei vom 10.2.2025 mit Größe 135 kB image not shown  

Quellverzeichnis  mod.rs   Sprache: unbekannt

 
Spracherkennung für: .rs vermutete Sprache: Unknown {[0] [0] [0]} [Methode: Schwerpunktbildung, einfache Gewichte, sechs Dimensionen]

use std::num::NonZeroU32;

use crate::front::wgsl::error::{Error, ExpectedToken, InvalidAssignmentType};
use crate::front::wgsl::index::Index;
use crate::front::wgsl::parse::number::Number;
use crate::front::wgsl::parse::{ast, conv};
use crate::front::Typifier;
use crate::proc::{
    ensure_block_returns, Alignment, ConstantEvaluator, Emitter, Layouter, ResolveContext,
};
use crate::{Arena, FastHashMap, FastIndexMap, Handle, Span};

mod construction;
mod conversion;

/// Resolves the inner type of a given expression.
///
/// Expects a &mut [`ExpressionContext`] and a [`Handle<Expression>`].
///
/// Returns a &[`crate::TypeInner`].
///
/// Ideally, we would simply have a function that takes a `&mut ExpressionContext`
/// and returns a `&TypeResolution`. Unfortunately, this leads the borrow checker
/// to conclude that the mutable borrow lasts for as long as we are using the
/// `&TypeResolution`, so we can't use the `ExpressionContext` for anything else -
/// like, say, resolving another operand's type. Using a macro that expands to
/// two separate calls, only the first of which needs a `&mut`,
/// lets the borrow checker see that the mutable borrow is over.
macro_rules! resolve_inner {
    ($ctx:ident, $expr:expr) => {{
        $ctx.grow_types($expr)?;
        $ctx.typifier()[$expr].inner_with(&$ctx.module.types)
    }};
}
pub(super) use resolve_inner;

/// Resolves the inner types of two given expressions.
///
/// Expects a &mut [`ExpressionContext`] and two [`Handle<Expression>`]s.
///
/// Returns a tuple containing two &[`crate::TypeInner`].
///
/// See the documentation of [`resolve_inner!`] for why this macro is necessary.
macro_rules! resolve_inner_binary {
    ($ctx:ident, $left:expr, $right:expr) => {{
        $ctx.grow_types($left)?;
        $ctx.grow_types($right)?;
        (
            $ctx.typifier()[$left].inner_with(&$ctx.module.types),
            $ctx.typifier()[$right].inner_with(&$ctx.module.types),
        )
    }};
}

/// Resolves the type of a given expression.
///
/// Expects a &mut [`ExpressionContext`] and a [`Handle<Expression>`].
///
/// Returns a &[`TypeResolution`].
///
/// See the documentation of [`resolve_inner!`] for why this macro is necessary.
///
/// [`TypeResolution`]: crate::proc::TypeResolution
macro_rules! resolve {
    ($ctx:ident, $expr:expr) => {{
        $ctx.grow_types($expr)?;
        &$ctx.typifier()[$expr]
    }};
}
pub(super) use resolve;

/// State for constructing a `crate::Module`.
pub struct GlobalContext<'source, 'temp, 'out> {
    /// The `TranslationUnit`'s expressions arena.
    ast_expressions: &'temp Arena<ast::Expression<'source>>,

    /// The `TranslationUnit`'s types arena.
    types: &'temp Arena<ast::Type<'source>>,

    // Naga IR values.
    /// The map from the names of module-scope declarations to the Naga IR
    /// `Handle`s we have built for them, owned by `Lowerer::lower`.
    globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,

    /// The module we're constructing.
    module: &'out mut crate::Module,

    const_typifier: &'temp mut Typifier,

    global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
}

impl<'source> GlobalContext<'source, '_, '_> {
    fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
        ExpressionContext {
            ast_expressions: self.ast_expressions,
            globals: self.globals,
            types: self.types,
            module: self.module,
            const_typifier: self.const_typifier,
            expr_type: ExpressionContextType::Constant(None),
            global_expression_kind_tracker: self.global_expression_kind_tracker,
        }
    }

    fn as_override(&mut self) -> ExpressionContext<'source, '_, '_> {
        ExpressionContext {
            ast_expressions: self.ast_expressions,
            globals: self.globals,
            types: self.types,
            module: self.module,
            const_typifier: self.const_typifier,
            expr_type: ExpressionContextType::Override,
            global_expression_kind_tracker: self.global_expression_kind_tracker,
        }
    }

    fn ensure_type_exists(
        &mut self,
        name: Option<String>,
        inner: crate::TypeInner,
    ) -> Handle<crate::Type> {
        self.module
            .types
            .insert(crate::Type { inner, name }, Span::UNDEFINED)
    }
}

/// State for lowering a statement within a function.
pub struct StatementContext<'source, 'temp, 'out> {
    // WGSL AST values.
    /// A reference to [`TranslationUnit::expressions`] for the translation unit
    /// we're lowering.
    ///
    /// [`TranslationUnit::expressions`]: ast::TranslationUnit::expressions
    ast_expressions: &'temp Arena<ast::Expression<'source>>,

    /// A reference to [`TranslationUnit::types`] for the translation unit
    /// we're lowering.
    ///
    /// [`TranslationUnit::types`]: ast::TranslationUnit::types
    types: &'temp Arena<ast::Type<'source>>,

    // Naga IR values.
    /// The map from the names of module-scope declarations to the Naga IR
    /// `Handle`s we have built for them, owned by `Lowerer::lower`.
    globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,

    /// A map from each `ast::Local` handle to the Naga expression
    /// we've built for it:
    ///
    /// - WGSL function arguments become Naga [`FunctionArgument`] expressions.
    ///
    /// - WGSL `var` declarations become Naga [`LocalVariable`] expressions.
    ///
    /// - WGSL `let` declararations become arbitrary Naga expressions.
    ///
    /// This always borrows the `local_table` local variable in
    /// [`Lowerer::function`].
    ///
    /// [`LocalVariable`]: crate::Expression::LocalVariable
    /// [`FunctionArgument`]: crate::Expression::FunctionArgument
    local_table:
        &'temp mut FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<crate::Expression>>>>,

    const_typifier: &'temp mut Typifier,
    typifier: &'temp mut Typifier,
    function: &'out mut crate::Function,
    /// Stores the names of expressions that are assigned in `let` statement
    /// Also stores the spans of the names, for use in errors.
    named_expressions: &'out mut FastIndexMap<Handle<crate::Expression>, (String, Span)>,
    module: &'out mut crate::Module,

    /// Which `Expression`s in `self.naga_expressions` are const expressions, in
    /// the WGSL sense.
    ///
    /// According to the WGSL spec, a const expression must not refer to any
    /// `let` declarations, even if those declarations' initializers are
    /// themselves const expressions. So this tracker is not simply concerned
    /// with the form of the expressions; it is also tracking whether WGSL says
    /// we should consider them to be const. See the use of `force_non_const` in
    /// the code for lowering `let` bindings.
    local_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
    global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
}

impl<'a, 'temp> StatementContext<'a, 'temp, '_> {
    fn as_const<'t>(
        &'t mut self,
        block: &'t mut crate::Block,
        emitter: &'t mut Emitter,
    ) -> ExpressionContext<'a, 't, 't>
    where
        'temp: 't,
    {
        ExpressionContext {
            globals: self.globals,
            types: self.types,
            ast_expressions: self.ast_expressions,
            const_typifier: self.const_typifier,
            global_expression_kind_tracker: self.global_expression_kind_tracker,
            module: self.module,
            expr_type: ExpressionContextType::Constant(Some(LocalExpressionContext {
                local_table: self.local_table,
                function: self.function,
                block,
                emitter,
                typifier: self.typifier,
                local_expression_kind_tracker: self.local_expression_kind_tracker,
            })),
        }
    }

    fn as_expression<'t>(
        &'t mut self,
        block: &'t mut crate::Block,
        emitter: &'t mut Emitter,
    ) -> ExpressionContext<'a, 't, 't>
    where
        'temp: 't,
    {
        ExpressionContext {
            globals: self.globals,
            types: self.types,
            ast_expressions: self.ast_expressions,
            const_typifier: self.const_typifier,
            global_expression_kind_tracker: self.global_expression_kind_tracker,
            module: self.module,
            expr_type: ExpressionContextType::Runtime(LocalExpressionContext {
                local_table: self.local_table,
                function: self.function,
                block,
                emitter,
                typifier: self.typifier,
                local_expression_kind_tracker: self.local_expression_kind_tracker,
            }),
        }
    }

    fn as_global(&mut self) -> GlobalContext<'a, '_, '_> {
        GlobalContext {
            ast_expressions: self.ast_expressions,
            globals: self.globals,
            types: self.types,
            module: self.module,
            const_typifier: self.const_typifier,
            global_expression_kind_tracker: self.global_expression_kind_tracker,
        }
    }

    fn invalid_assignment_type(&self, expr: Handle<crate::Expression>) -> InvalidAssignmentType {
        if let Some(&(_, span)) = self.named_expressions.get(&expr) {
            InvalidAssignmentType::ImmutableBinding(span)
        } else {
            match self.function.expressions[expr] {
                crate::Expression::Swizzle { .. } => InvalidAssignmentType::Swizzle,
                crate::Expression::Access { base, .. } => self.invalid_assignment_type(base),
                crate::Expression::AccessIndex { base, .. } => self.invalid_assignment_type(base),
                _ => InvalidAssignmentType::Other,
            }
        }
    }
}

pub struct LocalExpressionContext<'temp, 'out> {
    /// A map from [`ast::Local`] handles to the Naga expressions we've built for them.
    ///
    /// This is always [`StatementContext::local_table`] for the
    /// enclosing statement; see that documentation for details.
    local_table: &'temp FastHashMap<Handle<ast::Local>, Declared<Typed<Handle<crate::Expression>>>>,

    function: &'out mut crate::Function,
    block: &'temp mut crate::Block,
    emitter: &'temp mut Emitter,
    typifier: &'temp mut Typifier,

    /// Which `Expression`s in `self.naga_expressions` are const expressions, in
    /// the WGSL sense.
    ///
    /// See [`StatementContext::local_expression_kind_tracker`] for details.
    local_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,
}

/// The type of Naga IR expression we are lowering an [`ast::Expression`] to.
pub enum ExpressionContextType<'temp, 'out> {
    /// We are lowering to an arbitrary runtime expression, to be
    /// included in a function's body.
    ///
    /// The given [`LocalExpressionContext`] holds information about local
    /// variables, arguments, and other definitions available only to runtime
    /// expressions, not constant or override expressions.
    Runtime(LocalExpressionContext<'temp, 'out>),

    /// We are lowering to a constant expression, to be included in the module's
    /// constant expression arena.
    ///
    /// Everything global constant expressions are allowed to refer to is
    /// available in the [`ExpressionContext`], but local constant expressions can
    /// also refer to other
    Constant(Option<LocalExpressionContext<'temp, 'out>>),

    /// We are lowering to an override expression, to be included in the module's
    /// constant expression arena.
    ///
    /// Everything override expressions are allowed to refer to is
    /// available in the [`ExpressionContext`], so this variant
    /// carries no further information.
    Override,
}

/// State for lowering an [`ast::Expression`] to Naga IR.
///
/// [`ExpressionContext`]s come in two kinds, distinguished by
/// the value of the [`expr_type`] field:
///
/// - A [`Runtime`] context contributes [`naga::Expression`]s to a [`naga::Function`]'s
///   runtime expression arena.
///
/// - A [`Constant`] context contributes [`naga::Expression`]s to a [`naga::Module`]'s
///   constant expression arena.
///
/// [`ExpressionContext`]s are constructed in restricted ways:
///
/// - To get a [`Runtime`] [`ExpressionContext`], call
///   [`StatementContext::as_expression`].
///
/// - To get a [`Constant`] [`ExpressionContext`], call
///   [`GlobalContext::as_const`].
///
/// - You can demote a [`Runtime`] context to a [`Constant`] context
///   by calling [`as_const`], but there's no way to go in the other
///   direction, producing a runtime context from a constant one. This
///   is because runtime expressions can refer to constant
///   expressions, via [`Expression::Constant`], but constant
///   expressions can't refer to a function's expressions.
///
/// Not to be confused with `wgsl::parse::ExpressionContext`, which is
/// for parsing the `ast::Expression` in the first place.
///
/// [`expr_type`]: ExpressionContext::expr_type
/// [`Runtime`]: ExpressionContextType::Runtime
/// [`naga::Expression`]: crate::Expression
/// [`naga::Function`]: crate::Function
/// [`Constant`]: ExpressionContextType::Constant
/// [`naga::Module`]: crate::Module
/// [`as_const`]: ExpressionContext::as_const
/// [`Expression::Constant`]: crate::Expression::Constant
pub struct ExpressionContext<'source, 'temp, 'out> {
    // WGSL AST values.
    ast_expressions: &'temp Arena<ast::Expression<'source>>,
    types: &'temp Arena<ast::Type<'source>>,

    // Naga IR values.
    /// The map from the names of module-scope declarations to the Naga IR
    /// `Handle`s we have built for them, owned by `Lowerer::lower`.
    globals: &'temp mut FastHashMap<&'source str, LoweredGlobalDecl>,

    /// The IR [`Module`] we're constructing.
    ///
    /// [`Module`]: crate::Module
    module: &'out mut crate::Module,

    /// Type judgments for [`module::global_expressions`].
    ///
    /// [`module::global_expressions`]: crate::Module::global_expressions
    const_typifier: &'temp mut Typifier,
    global_expression_kind_tracker: &'temp mut crate::proc::ExpressionKindTracker,

    /// Whether we are lowering a constant expression or a general
    /// runtime expression, and the data needed in each case.
    expr_type: ExpressionContextType<'temp, 'out>,
}

impl<'source, 'temp, 'out> ExpressionContext<'source, 'temp, 'out> {
    fn as_const(&mut self) -> ExpressionContext<'source, '_, '_> {
        ExpressionContext {
            globals: self.globals,
            types: self.types,
            ast_expressions: self.ast_expressions,
            const_typifier: self.const_typifier,
            module: self.module,
            expr_type: ExpressionContextType::Constant(None),
            global_expression_kind_tracker: self.global_expression_kind_tracker,
        }
    }

    fn as_global(&mut self) -> GlobalContext<'source, '_, '_> {
        GlobalContext {
            ast_expressions: self.ast_expressions,
            globals: self.globals,
            types: self.types,
            module: self.module,
            const_typifier: self.const_typifier,
            global_expression_kind_tracker: self.global_expression_kind_tracker,
        }
    }

    fn as_const_evaluator(&mut self) -> ConstantEvaluator {
        match self.expr_type {
            ExpressionContextType::Runtime(ref mut rctx) => ConstantEvaluator::for_wgsl_function(
                self.module,
                &mut rctx.function.expressions,
                rctx.local_expression_kind_tracker,
                rctx.emitter,
                rctx.block,
                false,
            ),
            ExpressionContextType::Constant(Some(ref mut rctx)) => {
                ConstantEvaluator::for_wgsl_function(
                    self.module,
                    &mut rctx.function.expressions,
                    rctx.local_expression_kind_tracker,
                    rctx.emitter,
                    rctx.block,
                    true,
                )
            }
            ExpressionContextType::Constant(None) => ConstantEvaluator::for_wgsl_module(
                self.module,
                self.global_expression_kind_tracker,
                false,
            ),
            ExpressionContextType::Override => ConstantEvaluator::for_wgsl_module(
                self.module,
                self.global_expression_kind_tracker,
                true,
            ),
        }
    }

    fn append_expression(
        &mut self,
        expr: crate::Expression,
        span: Span,
    ) -> Result<Handle<crate::Expression>, Error<'source>> {
        let mut eval = self.as_const_evaluator();
        eval.try_eval_and_append(expr, span)
            .map_err(|e| Error::ConstantEvaluatorError(e.into(), span))
    }

    fn const_access(&self, handle: Handle<crate::Expression>) -> Option<u32> {
        match self.expr_type {
            ExpressionContextType::Runtime(ref ctx) => {
                if !ctx.local_expression_kind_tracker.is_const(handle) {
                    return None;
                }

                self.module
                    .to_ctx()
                    .eval_expr_to_u32_from(handle, &ctx.function.expressions)
                    .ok()
            }
            ExpressionContextType::Constant(Some(ref ctx)) => {
                assert!(ctx.local_expression_kind_tracker.is_const(handle));
                self.module
                    .to_ctx()
                    .eval_expr_to_u32_from(handle, &ctx.function.expressions)
                    .ok()
            }
            ExpressionContextType::Constant(None) => {
                self.module.to_ctx().eval_expr_to_u32(handle).ok()
            }
            ExpressionContextType::Override => None,
        }
    }

    fn get_expression_span(&self, handle: Handle<crate::Expression>) -> Span {
        match self.expr_type {
            ExpressionContextType::Runtime(ref ctx)
            | ExpressionContextType::Constant(Some(ref ctx)) => {
                ctx.function.expressions.get_span(handle)
            }
            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
                self.module.global_expressions.get_span(handle)
            }
        }
    }

    fn typifier(&self) -> &Typifier {
        match self.expr_type {
            ExpressionContextType::Runtime(ref ctx)
            | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
                self.const_typifier
            }
        }
    }

    fn local(
        &mut self,
        local: &Handle<ast::Local>,
        span: Span,
    ) -> Result<Typed<Handle<crate::Expression>>, Error<'source>> {
        match self.expr_type {
            ExpressionContextType::Runtime(ref ctx) => Ok(ctx.local_table[local].runtime()),
            ExpressionContextType::Constant(Some(ref ctx)) => ctx.local_table[local]
                .const_time()
                .ok_or(Error::UnexpectedOperationInConstContext(span)),
            _ => Err(Error::UnexpectedOperationInConstContext(span)),
        }
    }

    fn runtime_expression_ctx(
        &mut self,
        span: Span,
    ) -> Result<&mut LocalExpressionContext<'temp, 'out>, Error<'source>> {
        match self.expr_type {
            ExpressionContextType::Runtime(ref mut ctx) => Ok(ctx),
            ExpressionContextType::Constant(_) | ExpressionContextType::Override => {
                Err(Error::UnexpectedOperationInConstContext(span))
            }
        }
    }

    fn gather_component(
        &mut self,
        expr: Handle<crate::Expression>,
        component_span: Span,
        gather_span: Span,
    ) -> Result<crate::SwizzleComponent, Error<'source>> {
        match self.expr_type {
            ExpressionContextType::Runtime(ref rctx) => {
                if !rctx.local_expression_kind_tracker.is_const(expr) {
                    return Err(Error::ExpectedConstExprConcreteIntegerScalar(
                        component_span,
                    ));
                }

                let index = self
                    .module
                    .to_ctx()
                    .eval_expr_to_u32_from(expr, &rctx.function.expressions)
                    .map_err(|err| match err {
                        crate::proc::U32EvalError::NonConst => {
                            Error::ExpectedConstExprConcreteIntegerScalar(component_span)
                        }
                        crate::proc::U32EvalError::Negative => {
                            Error::ExpectedNonNegative(component_span)
                        }
                    })?;
                crate::SwizzleComponent::XYZW
                    .get(index as usize)
                    .copied()
                    .ok_or(Error::InvalidGatherComponent(component_span))
            }
            // This means a `gather` operation appeared in a constant expression.
            // This error refers to the `gather` itself, not its "component" argument.
            ExpressionContextType::Constant(_) | ExpressionContextType::Override => {
                Err(Error::UnexpectedOperationInConstContext(gather_span))
            }
        }
    }

    /// Determine the type of `handle`, and add it to the module's arena.
    ///
    /// If you just need a `TypeInner` for `handle`'s type, use the
    /// [`resolve_inner!`] macro instead. This function
    /// should only be used when the type of `handle` needs to appear
    /// in the module's final `Arena<Type>`, for example, if you're
    /// creating a [`LocalVariable`] whose type is inferred from its
    /// initializer.
    ///
    /// [`LocalVariable`]: crate::LocalVariable
    fn register_type(
        &mut self,
        handle: Handle<crate::Expression>,
    ) -> Result<Handle<crate::Type>, Error<'source>> {
        self.grow_types(handle)?;
        // This is equivalent to calling ExpressionContext::typifier(),
        // except that this lets the borrow checker see that it's okay
        // to also borrow self.module.types mutably below.
        let typifier = match self.expr_type {
            ExpressionContextType::Runtime(ref ctx)
            | ExpressionContextType::Constant(Some(ref ctx)) => ctx.typifier,
            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
                &*self.const_typifier
            }
        };
        Ok(typifier.register_type(handle, &mut self.module.types))
    }

    /// Resolve the types of all expressions up through `handle`.
    ///
    /// Ensure that [`self.typifier`] has a [`TypeResolution`] for
    /// every expression in [`self.function.expressions`].
    ///
    /// This does not add types to any arena. The [`Typifier`]
    /// documentation explains the steps we take to avoid filling
    /// arenas with intermediate types.
    ///
    /// This function takes `&mut self`, so it can't conveniently
    /// return a shared reference to the resulting `TypeResolution`:
    /// the shared reference would extend the mutable borrow, and you
    /// wouldn't be able to use `self` for anything else. Instead, you
    /// should use [`register_type`] or one of [`resolve!`],
    /// [`resolve_inner!`] or [`resolve_inner_binary!`].
    ///
    /// [`self.typifier`]: ExpressionContext::typifier
    /// [`TypeResolution`]: crate::proc::TypeResolution
    /// [`register_type`]: Self::register_type
    /// [`Typifier`]: Typifier
    fn grow_types(
        &mut self,
        handle: Handle<crate::Expression>,
    ) -> Result<&mut Self, Error<'source>> {
        let empty_arena = Arena::new();
        let resolve_ctx;
        let typifier;
        let expressions;
        match self.expr_type {
            ExpressionContextType::Runtime(ref mut ctx)
            | ExpressionContextType::Constant(Some(ref mut ctx)) => {
                resolve_ctx = ResolveContext::with_locals(
                    self.module,
                    &ctx.function.local_variables,
                    &ctx.function.arguments,
                );
                typifier = &mut *ctx.typifier;
                expressions = &ctx.function.expressions;
            }
            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {
                resolve_ctx = ResolveContext::with_locals(self.module, &empty_arena, &[]);
                typifier = self.const_typifier;
                expressions = &self.module.global_expressions;
            }
        };
        typifier
            .grow(handle, expressions, &resolve_ctx)
            .map_err(Error::InvalidResolve)?;

        Ok(self)
    }

    fn image_data(
        &mut self,
        image: Handle<crate::Expression>,
        span: Span,
    ) -> Result<(crate::ImageClass, bool), Error<'source>> {
        match *resolve_inner!(self, image) {
            crate::TypeInner::Image { class, arrayed, .. } => Ok((class, arrayed)),
            _ => Err(Error::BadTexture(span)),
        }
    }

    fn prepare_args<'b>(
        &mut self,
        args: &'b [Handle<ast::Expression<'source>>],
        min_args: u32,
        span: Span,
    ) -> ArgumentContext<'b, 'source> {
        ArgumentContext {
            args: args.iter(),
            min_args,
            args_used: 0,
            total_args: args.len() as u32,
            span,
        }
    }

    /// Insert splats, if needed by the non-'*' operations.
    ///
    /// See the "Binary arithmetic expressions with mixed scalar and vector operands"
    /// table in the WebGPU Shading Language specification for relevant operators.
    ///
    /// Multiply is not handled here as backends are expected to handle vec*scalar
    /// operations, so inserting splats into the IR increases size needlessly.
    fn binary_op_splat(
        &mut self,
        op: crate::BinaryOperator,
        left: &mut Handle<crate::Expression>,
        right: &mut Handle<crate::Expression>,
    ) -> Result<(), Error<'source>> {
        if matches!(
            op,
            crate::BinaryOperator::Add
                | crate::BinaryOperator::Subtract
                | crate::BinaryOperator::Divide
                | crate::BinaryOperator::Modulo
        ) {
            match resolve_inner_binary!(self, *left, *right) {
                (&crate::TypeInner::Vector { size, .. }, &crate::TypeInner::Scalar { .. }) => {
                    *right = self.append_expression(
                        crate::Expression::Splat {
                            size,
                            value: *right,
                        },
                        self.get_expression_span(*right),
                    )?;
                }
                (&crate::TypeInner::Scalar { .. }, &crate::TypeInner::Vector { size, .. }) => {
                    *left = self.append_expression(
                        crate::Expression::Splat { size, value: *left },
                        self.get_expression_span(*left),
                    )?;
                }
                _ => {}
            }
        }

        Ok(())
    }

    /// Add a single expression to the expression table that is not covered by `self.emitter`.
    ///
    /// This is useful for `CallResult` and `AtomicResult` expressions, which should not be covered by
    /// `Emit` statements.
    fn interrupt_emitter(
        &mut self,
        expression: crate::Expression,
        span: Span,
    ) -> Result<Handle<crate::Expression>, Error<'source>> {
        match self.expr_type {
            ExpressionContextType::Runtime(ref mut rctx)
            | ExpressionContextType::Constant(Some(ref mut rctx)) => {
                rctx.block
                    .extend(rctx.emitter.finish(&rctx.function.expressions));
            }
            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
        }
        let result = self.append_expression(expression, span);
        match self.expr_type {
            ExpressionContextType::Runtime(ref mut rctx)
            | ExpressionContextType::Constant(Some(ref mut rctx)) => {
                rctx.emitter.start(&rctx.function.expressions);
            }
            ExpressionContextType::Constant(None) | ExpressionContextType::Override => {}
        }
        result
    }

    /// Apply the WGSL Load Rule to `expr`.
    ///
    /// If `expr` is has type `ref<SC, T, A>`, perform a load to produce a value of type
    /// `T`. Otherwise, return `expr` unchanged.
    fn apply_load_rule(
        &mut self,
        expr: Typed<Handle<crate::Expression>>,
    ) -> Result<Handle<crate::Expression>, Error<'source>> {
        match expr {
            Typed::Reference(pointer) => {
                let load = crate::Expression::Load { pointer };
                let span = self.get_expression_span(pointer);
                self.append_expression(load, span)
            }
            Typed::Plain(handle) => Ok(handle),
        }
    }

    fn ensure_type_exists(&mut self, inner: crate::TypeInner) -> Handle<crate::Type> {
        self.as_global().ensure_type_exists(None, inner)
    }
}

struct ArgumentContext<'ctx, 'source> {
    args: std::slice::Iter<'ctx, Handle<ast::Expression<'source>>>,
    min_args: u32,
    args_used: u32,
    total_args: u32,
    span: Span,
}

impl<'source> ArgumentContext<'_, 'source> {
    pub fn finish(self) -> Result<(), Error<'source>> {
        if self.args.len() == 0 {
            Ok(())
        } else {
            Err(Error::WrongArgumentCount {
                found: self.total_args,
                expected: self.min_args..self.args_used + 1,
                span: self.span,
            })
        }
    }

    pub fn next(&mut self) -> Result<Handle<ast::Expression<'source>>, Error<'source>> {
        match self.args.next().copied() {
            Some(arg) => {
                self.args_used += 1;
                Ok(arg)
            }
            None => Err(Error::WrongArgumentCount {
                found: self.total_args,
                expected: self.min_args..self.args_used + 1,
                span: self.span,
            }),
        }
    }
}

#[derive(Debug, Copy, Clone)]
enum Declared<T> {
    /// Value declared as const
    Const(T),

    /// Value declared as non-const
    Runtime(T),
}

impl<T> Declared<T> {
    fn runtime(self) -> T {
        match self {
            Declared::Const(t) | Declared::Runtime(t) => t,
        }
    }

    fn const_time(self) -> Option<T> {
        match self {
            Declared::Const(t) => Some(t),
            Declared::Runtime(_) => None,
        }
    }
}

/// WGSL type annotations on expressions, types, values, etc.
///
/// Naga and WGSL types are very close, but Naga lacks WGSL's `ref` types, which
/// we need to know to apply the Load Rule. This enum carries some WGSL or Naga
/// datum along with enough information to determine its corresponding WGSL
/// type.
///
/// The `T` type parameter can be any expression-like thing:
///
/// - `Typed<Handle<crate::Type>>` can represent a full WGSL type. For example,
///   given some Naga `Pointer` type `ptr`, a WGSL reference type is a
///   `Typed::Reference(ptr)` whereas a WGSL pointer type is a
///   `Typed::Plain(ptr)`.
///
/// - `Typed<crate::Expression>` or `Typed<Handle<crate::Expression>>` can
///   represent references similarly.
///
/// Use the `map` and `try_map` methods to convert from one expression
/// representation to another.
///
/// [`Expression`]: crate::Expression
#[derive(Debug, Copy, Clone)]
enum Typed<T> {
    /// A WGSL reference.
    Reference(T),

    /// A WGSL plain type.
    Plain(T),
}

impl<T> Typed<T> {
    fn map<U>(self, mut f: impl FnMut(T) -> U) -> Typed<U> {
        match self {
            Self::Reference(v) => Typed::Reference(f(v)),
            Self::Plain(v) => Typed::Plain(f(v)),
        }
    }

    fn try_map<U, E>(self, mut f: impl FnMut(T) -> Result<U, E>) -> Result<Typed<U>, E> {
        Ok(match self {
            Self::Reference(expr) => Typed::Reference(f(expr)?),
            Self::Plain(expr) => Typed::Plain(f(expr)?),
        })
    }
}

/// A single vector component or swizzle.
///
/// This represents the things that can appear after the `.` in a vector access
/// expression: either a single component name, or a series of them,
/// representing a swizzle.
enum Components {
    Single(u32),
    Swizzle {
        size: crate::VectorSize,
        pattern: [crate::SwizzleComponent; 4],
    },
}

impl Components {
    const fn letter_component(letter: char) -> Option<crate::SwizzleComponent> {
        use crate::SwizzleComponent as Sc;
        match letter {
            'x' | 'r' => Some(Sc::X),
            'y' | 'g' => Some(Sc::Y),
            'z' | 'b' => Some(Sc::Z),
            'w' | 'a' => Some(Sc::W),
            _ => None,
        }
    }

    fn single_component(name: &str, name_span: Span) -> Result<u32, Error> {
        let ch = name.chars().next().ok_or(Error::BadAccessor(name_span))?;
        match Self::letter_component(ch) {
            Some(sc) => Ok(sc as u32),
            None => Err(Error::BadAccessor(name_span)),
        }
    }

    /// Construct a `Components` value from a 'member' name, like `"wzy"` or `"x"`.
    ///
    /// Use `name_span` for reporting errors in parsing the component string.
    fn new(name: &str, name_span: Span) -> Result<Self, Error> {
        let size = match name.len() {
            1 => return Ok(Components::Single(Self::single_component(name, name_span)?)),
            2 => crate::VectorSize::Bi,
            3 => crate::VectorSize::Tri,
            4 => crate::VectorSize::Quad,
            _ => return Err(Error::BadAccessor(name_span)),
        };

        let mut pattern = [crate::SwizzleComponent::X; 4];
        for (comp, ch) in pattern.iter_mut().zip(name.chars()) {
            *comp = Self::letter_component(ch).ok_or(Error::BadAccessor(name_span))?;
        }

        if name.chars().all(|c| matches!(c, 'x' | 'y' | 'z' | 'w'))
            || name.chars().all(|c| matches!(c, 'r' | 'g' | 'b' | 'a'))
        {
            Ok(Components::Swizzle { size, pattern })
        } else {
            Err(Error::BadAccessor(name_span))
        }
    }
}

/// An `ast::GlobalDecl` for which we have built the Naga IR equivalent.
enum LoweredGlobalDecl {
    Function(Handle<crate::Function>),
    Var(Handle<crate::GlobalVariable>),
    Const(Handle<crate::Constant>),
    Override(Handle<crate::Override>),
    Type(Handle<crate::Type>),
    EntryPoint,
}

enum Texture {
    Gather,
    GatherCompare,

    Sample,
    SampleBias,
    SampleCompare,
    SampleCompareLevel,
    SampleGrad,
    SampleLevel,
    // SampleBaseClampToEdge,
}

impl Texture {
    pub fn map(word: &str) -> Option<Self> {
        Some(match word {
            "textureGather" => Self::Gather,
            "textureGatherCompare" => Self::GatherCompare,

            "textureSample" => Self::Sample,
            "textureSampleBias" => Self::SampleBias,
            "textureSampleCompare" => Self::SampleCompare,
            "textureSampleCompareLevel" => Self::SampleCompareLevel,
            "textureSampleGrad" => Self::SampleGrad,
            "textureSampleLevel" => Self::SampleLevel,
            // "textureSampleBaseClampToEdge" => Some(Self::SampleBaseClampToEdge),
            _ => return None,
        })
    }

    pub const fn min_argument_count(&self) -> u32 {
        match *self {
            Self::Gather => 3,
            Self::GatherCompare => 4,

            Self::Sample => 3,
            Self::SampleBias => 5,
            Self::SampleCompare => 5,
            Self::SampleCompareLevel => 5,
            Self::SampleGrad => 6,
            Self::SampleLevel => 5,
            // Self::SampleBaseClampToEdge => 3,
        }
    }
}

enum SubgroupGather {
    BroadcastFirst,
    Broadcast,
    Shuffle,
    ShuffleDown,
    ShuffleUp,
    ShuffleXor,
}

impl SubgroupGather {
    pub fn map(word: &str) -> Option<Self> {
        Some(match word {
            "subgroupBroadcastFirst" => Self::BroadcastFirst,
            "subgroupBroadcast" => Self::Broadcast,
            "subgroupShuffle" => Self::Shuffle,
            "subgroupShuffleDown" => Self::ShuffleDown,
            "subgroupShuffleUp" => Self::ShuffleUp,
            "subgroupShuffleXor" => Self::ShuffleXor,
            _ => return None,
        })
    }
}

pub struct Lowerer<'source, 'temp> {
    index: &'temp Index<'source>,
    layouter: Layouter,
}

impl<'source, 'temp> Lowerer<'source, 'temp> {
    pub fn new(index: &'temp Index<'source>) -> Self {
        Self {
            index,
            layouter: Layouter::default(),
        }
    }

    pub fn lower(
        &mut self,
        tu: &'temp ast::TranslationUnit<'source>,
    ) -> Result<crate::Module, Error<'source>> {
        let mut module = crate::Module {
            diagnostic_filters: tu.diagnostic_filters.clone(),
            diagnostic_filter_leaf: tu.diagnostic_filter_leaf,
            ..Default::default()
        };

        let mut ctx = GlobalContext {
            ast_expressions: &tu.expressions,
            globals: &mut FastHashMap::default(),
            types: &tu.types,
            module: &mut module,
            const_typifier: &mut Typifier::new(),
            global_expression_kind_tracker: &mut crate::proc::ExpressionKindTracker::new(),
        };

        for decl_handle in self.index.visit_ordered() {
            let span = tu.decls.get_span(decl_handle);
            let decl = &tu.decls[decl_handle];

            match decl.kind {
                ast::GlobalDeclKind::Fn(ref f) => {
                    let lowered_decl = self.function(f, span, &mut ctx)?;
                    ctx.globals.insert(f.name.name, lowered_decl);
                }
                ast::GlobalDeclKind::Var(ref v) => {
                    let explicit_ty =
                        v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx))
                            .transpose()?;

                    let mut ectx = ctx.as_override();

                    let ty;
                    let initializer;
                    match (v.init, explicit_ty) {
                        (Some(init), Some(explicit_ty)) => {
                            let init = self.expression_for_abstract(init, &mut ectx)?;
                            let ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
                            let init = ectx
                                .try_automatic_conversions(init, &ty_res, v.name.span)
                                .map_err(|error| match error {
                                Error::AutoConversion(e) => Error::InitializationTypeMismatch {
                                    name: v.name.span,
                                    expected: e.dest_type,
                                    got: e.source_type,
                                },
                                other => other,
                            })?;
                            ty = explicit_ty;
                            initializer = Some(init);
                        }
                        (Some(init), None) => {
                            let concretized = self.expression(init, &mut ectx)?;
                            ty = ectx.register_type(concretized)?;
                            initializer = Some(concretized);
                        }
                        (None, Some(explicit_ty)) => {
                            ty = explicit_ty;
                            initializer = None;
                        }
                        (None, None) => return Err(Error::DeclMissingTypeAndInit(v.name.span)),
                    }

                    let binding = if let Some(ref binding) = v.binding {
                        Some(crate::ResourceBinding {
                            group: self.const_u32(binding.group, &mut ctx.as_const())?.0,
                            binding: self.const_u32(binding.binding, &mut ctx.as_const())?.0,
                        })
                    } else {
                        None
                    };

                    let handle = ctx.module.global_variables.append(
                        crate::GlobalVariable {
                            name: Some(v.name.name.to_string()),
                            space: v.space,
                            binding,
                            ty,
                            init: initializer,
                        },
                        span,
                    );

                    ctx.globals
                        .insert(v.name.name, LoweredGlobalDecl::Var(handle));
                }
                ast::GlobalDeclKind::Const(ref c) => {
                    let mut ectx = ctx.as_const();
                    let mut init = self.expression_for_abstract(c.init, &mut ectx)?;

                    let ty;
                    if let Some(explicit_ty) = c.ty {
                        let explicit_ty =
                            self.resolve_ast_type(explicit_ty, &mut ectx.as_global())?;
                        let explicit_ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
                        init = ectx
                            .try_automatic_conversions(init, &explicit_ty_res, c.name.span)
                            .map_err(|error| match error {
                                Error::AutoConversion(e) => Error::InitializationTypeMismatch {
                                    name: c.name.span,
                                    expected: e.dest_type,
                                    got: e.source_type,
                                },
                                other => other,
                            })?;
                        ty = explicit_ty;
                    } else {
                        init = ectx.concretize(init)?;
                        ty = ectx.register_type(init)?;
                    }

                    let handle = ctx.module.constants.append(
                        crate::Constant {
                            name: Some(c.name.name.to_string()),
                            ty,
                            init,
                        },
                        span,
                    );

                    ctx.globals
                        .insert(c.name.name, LoweredGlobalDecl::Const(handle));
                }
                ast::GlobalDeclKind::Override(ref o) => {
                    let init = o
                        .init
                        .map(|init| self.expression(init, &mut ctx.as_override()))
                        .transpose()?;
                    let inferred_type = init
                        .map(|init| ctx.as_const().register_type(init))
                        .transpose()?;

                    let explicit_ty =
                        o.ty.map(|ty| self.resolve_ast_type(ty, &mut ctx))
                            .transpose()?;

                    let id =
                        o.id.map(|id| self.const_u32(id, &mut ctx.as_const()))
                            .transpose()?;

                    let id = if let Some((id, id_span)) = id {
                        Some(
                            u16::try_from(id)
                                .map_err(|_| Error::PipelineConstantIDValue(id_span))?,
                        )
                    } else {
                        None
                    };

                    let ty = match (explicit_ty, inferred_type) {
                        (Some(explicit_ty), Some(inferred_type)) => {
                            if explicit_ty == inferred_type {
                                explicit_ty
                            } else {
                                let gctx = ctx.module.to_ctx();
                                return Err(Error::InitializationTypeMismatch {
                                    name: o.name.span,
                                    expected: explicit_ty.to_wgsl(&gctx).into(),
                                    got: inferred_type.to_wgsl(&gctx).into(),
                                });
                            }
                        }
                        (Some(explicit_ty), None) => explicit_ty,
                        (None, Some(inferred_type)) => inferred_type,
                        (None, None) => {
                            return Err(Error::DeclMissingTypeAndInit(o.name.span));
                        }
                    };

                    let handle = ctx.module.overrides.append(
                        crate::Override {
                            name: Some(o.name.name.to_string()),
                            id,
                            ty,
                            init,
                        },
                        span,
                    );

                    ctx.globals
                        .insert(o.name.name, LoweredGlobalDecl::Override(handle));
                }
                ast::GlobalDeclKind::Struct(ref s) => {
                    let handle = self.r#struct(s, span, &mut ctx)?;
                    ctx.globals
                        .insert(s.name.name, LoweredGlobalDecl::Type(handle));
                }
                ast::GlobalDeclKind::Type(ref alias) => {
                    let ty = self.resolve_named_ast_type(
                        alias.ty,
                        Some(alias.name.name.to_string()),
                        &mut ctx,
                    )?;
                    ctx.globals
                        .insert(alias.name.name, LoweredGlobalDecl::Type(ty));
                }
                ast::GlobalDeclKind::ConstAssert(condition) => {
                    let condition = self.expression(condition, &mut ctx.as_const())?;

                    let span = ctx.module.global_expressions.get_span(condition);
                    match ctx
                        .module
                        .to_ctx()
                        .eval_expr_to_bool_from(condition, &ctx.module.global_expressions)
                    {
                        Some(true) => Ok(()),
                        Some(false) => Err(Error::ConstAssertFailed(span)),
                        _ => Err(Error::NotBool(span)),
                    }?;
                }
            }
        }

        // Constant evaluation may leave abstract-typed literals and
        // compositions in expression arenas, so we need to compact the module
        // to remove unused expressions and types.
        crate::compact::compact(&mut module);

        Ok(module)
    }

    fn function(
        &mut self,
        f: &ast::Function<'source>,
        span: Span,
        ctx: &mut GlobalContext<'source, '_, '_>,
    ) -> Result<LoweredGlobalDecl, Error<'source>> {
        let mut local_table = FastHashMap::default();
        let mut expressions = Arena::new();
        let mut named_expressions = FastIndexMap::default();
        let mut local_expression_kind_tracker = crate::proc::ExpressionKindTracker::new();

        let arguments = f
            .arguments
            .iter()
            .enumerate()
            .map(|(i, arg)| -> Result<_, Error<'_>> {
                let ty = self.resolve_ast_type(arg.ty, ctx)?;
                let expr = expressions
                    .append(crate::Expression::FunctionArgument(i as u32), arg.name.span);
                local_table.insert(arg.handle, Declared::Runtime(Typed::Plain(expr)));
                named_expressions.insert(expr, (arg.name.name.to_string(), arg.name.span));
                local_expression_kind_tracker.insert(expr, crate::proc::ExpressionKind::Runtime);

                Ok(crate::FunctionArgument {
                    name: Some(arg.name.name.to_string()),
                    ty,
                    binding: self.binding(&arg.binding, ty, ctx)?,
                })
            })
            .collect::<Result<Vec<_>, _>>()?;

        let result = f
            .result
            .as_ref()
            .map(|res| -> Result<_, Error<'_>> {
                let ty = self.resolve_ast_type(res.ty, ctx)?;
                Ok(crate::FunctionResult {
                    ty,
                    binding: self.binding(&res.binding, ty, ctx)?,
                })
            })
            .transpose()?;

        let mut function = crate::Function {
            name: Some(f.name.name.to_string()),
            arguments,
            result,
            local_variables: Arena::new(),
            expressions,
            named_expressions: crate::NamedExpressions::default(),
            body: crate::Block::default(),
            diagnostic_filter_leaf: f.diagnostic_filter_leaf,
        };

        let mut typifier = Typifier::default();
        let mut stmt_ctx = StatementContext {
            local_table: &mut local_table,
            globals: ctx.globals,
            ast_expressions: ctx.ast_expressions,
            const_typifier: ctx.const_typifier,
            typifier: &mut typifier,
            function: &mut function,
            named_expressions: &mut named_expressions,
            types: ctx.types,
            module: ctx.module,
            local_expression_kind_tracker: &mut local_expression_kind_tracker,
            global_expression_kind_tracker: ctx.global_expression_kind_tracker,
        };
        let mut body = self.block(&f.body, false, &mut stmt_ctx)?;
        ensure_block_returns(&mut body);

        function.body = body;
        function.named_expressions = named_expressions
            .into_iter()
            .map(|(key, (name, _))| (key, name))
            .collect();

        if let Some(ref entry) = f.entry_point {
            let workgroup_size_info = if let Some(workgroup_size) = entry.workgroup_size {
                // TODO: replace with try_map once stabilized
                let mut workgroup_size_out = [1; 3];
                let mut workgroup_size_overrides_out = [None; 3];
                for (i, size) in workgroup_size.into_iter().enumerate() {
                    if let Some(size_expr) = size {
                        match self.const_u32(size_expr, &mut ctx.as_const()) {
                            Ok(value) => {
                                workgroup_size_out[i] = value.0;
                            }
                            err => {
                                if let Err(Error::ConstantEvaluatorError(ref ty, _)) = err {
                                    match **ty {
                                        crate::proc::ConstantEvaluatorError::OverrideExpr => {
                                            workgroup_size_overrides_out[i] =
                                                Some(self.workgroup_size_override(
                                                    size_expr,
                                                    &mut ctx.as_override(),
                                                )?);
                                        }
                                        _ => {
                                            err?;
                                        }
                                    }
                                } else {
                                    err?;
                                }
                            }
                        }
                    }
                }
                if workgroup_size_overrides_out.iter().all(|x| x.is_none()) {
                    (workgroup_size_out, None)
                } else {
                    (workgroup_size_out, Some(workgroup_size_overrides_out))
                }
            } else {
                ([0; 3], None)
            };

            let (workgroup_size, workgroup_size_overrides) = workgroup_size_info;
            ctx.module.entry_points.push(crate::EntryPoint {
                name: f.name.name.to_string(),
                stage: entry.stage,
                early_depth_test: entry.early_depth_test,
                workgroup_size,
                workgroup_size_overrides,
                function,
            });
            Ok(LoweredGlobalDecl::EntryPoint)
        } else {
            let handle = ctx.module.functions.append(function, span);
            Ok(LoweredGlobalDecl::Function(handle))
        }
    }

    fn workgroup_size_override(
        &mut self,
        size_expr: Handle<ast::Expression<'source>>,
        ctx: &mut ExpressionContext<'source, '_, '_>,
    ) -> Result<Handle<crate::Expression>, Error<'source>> {
        let span = ctx.ast_expressions.get_span(size_expr);
        let expr = self.expression(size_expr, ctx)?;
        match resolve_inner!(ctx, expr).scalar_kind().ok_or(0) {
            Ok(crate::ScalarKind::Sint) | Ok(crate::ScalarKind::Uint) => Ok(expr),
            _ => Err(Error::ExpectedConstExprConcreteIntegerScalar(span)),
        }
    }

    fn block(
        &mut self,
        b: &ast::Block<'source>,
        is_inside_loop: bool,
        ctx: &mut StatementContext<'source, '_, '_>,
    ) -> Result<crate::Block, Error<'source>> {
        let mut block = crate::Block::default();

        for stmt in b.stmts.iter() {
            self.statement(stmt, &mut block, is_inside_loop, ctx)?;
        }

        Ok(block)
    }

    fn statement(
        &mut self,
        stmt: &ast::Statement<'source>,
        block: &mut crate::Block,
        is_inside_loop: bool,
        ctx: &mut StatementContext<'source, '_, '_>,
    ) -> Result<(), Error<'source>> {
        let out = match stmt.kind {
            ast::StatementKind::Block(ref block) => {
                let block = self.block(block, is_inside_loop, ctx)?;
                crate::Statement::Block(block)
            }
            ast::StatementKind::LocalDecl(ref decl) => match *decl {
                ast::LocalDecl::Let(ref l) => {
                    let mut emitter = Emitter::default();
                    emitter.start(&ctx.function.expressions);

                    let value =
                        self.expression(l.init, &mut ctx.as_expression(block, &mut emitter))?;

                    // The WGSL spec says that any expression that refers to a
                    // `let`-bound variable is not a const expression. This
                    // affects when errors must be reported, so we can't even
                    // treat suitable `let` bindings as constant as an
                    // optimization.
                    ctx.local_expression_kind_tracker.force_non_const(value);

                    let explicit_ty =
                        l.ty.map(|ty| self.resolve_ast_type(ty, &mut ctx.as_global()))
                            .transpose()?;

                    if let Some(ty) = explicit_ty {
                        let mut ctx = ctx.as_expression(block, &mut emitter);
                        let init_ty = ctx.register_type(value)?;
                        if !ctx.module.types[ty]
                            .inner
                            .equivalent(&ctx.module.types[init_ty].inner, &ctx.module.types)
                        {
                            let gctx = &ctx.module.to_ctx();
                            return Err(Error::InitializationTypeMismatch {
                                name: l.name.span,
                                expected: ty.to_wgsl(gctx).into(),
                                got: init_ty.to_wgsl(gctx).into(),
                            });
                        }
                    }

                    block.extend(emitter.finish(&ctx.function.expressions));
                    ctx.local_table
                        .insert(l.handle, Declared::Runtime(Typed::Plain(value)));
                    ctx.named_expressions
                        .insert(value, (l.name.name.to_string(), l.name.span));

                    return Ok(());
                }
                ast::LocalDecl::Var(ref v) => {
                    let explicit_ty =
                        v.ty.map(|ast| self.resolve_ast_type(ast, &mut ctx.as_global()))
                            .transpose()?;

                    let mut emitter = Emitter::default();
                    emitter.start(&ctx.function.expressions);
                    let mut ectx = ctx.as_expression(block, &mut emitter);

                    let ty;
                    let initializer;
                    match (v.init, explicit_ty) {
                        (Some(init), Some(explicit_ty)) => {
                            let init = self.expression_for_abstract(init, &mut ectx)?;
                            let ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
                            let init = ectx
                                .try_automatic_conversions(init, &ty_res, v.name.span)
                                .map_err(|error| match error {
                                Error::AutoConversion(e) => Error::InitializationTypeMismatch {
                                    name: v.name.span,
                                    expected: e.dest_type,
                                    got: e.source_type,
                                },
                                other => other,
                            })?;
                            ty = explicit_ty;
                            initializer = Some(init);
                        }
                        (Some(init), None) => {
                            let concretized = self.expression(init, &mut ectx)?;
                            ty = ectx.register_type(concretized)?;
                            initializer = Some(concretized);
                        }
                        (None, Some(explicit_ty)) => {
                            ty = explicit_ty;
                            initializer = None;
                        }
                        (None, None) => return Err(Error::DeclMissingTypeAndInit(v.name.span)),
                    }

                    let (const_initializer, initializer) = {
                        match initializer {
                            Some(init) => {
                                // It's not correct to hoist the initializer up
                                // to the top of the function if:
                                // - the initialization is inside a loop, and should
                                //   take place on every iteration, or
                                // - the initialization is not a constant
                                //   expression, so its value depends on the
                                //   state at the point of initialization.
                                if is_inside_loop
                                    || !ctx.local_expression_kind_tracker.is_const_or_override(init)
                                {
                                    (None, Some(init))
                                } else {
                                    (Some(init), None)
                                }
                            }
                            None => (None, None),
                        }
                    };

                    let var = ctx.function.local_variables.append(
                        crate::LocalVariable {
                            name: Some(v.name.name.to_string()),
                            ty,
                            init: const_initializer,
                        },
                        stmt.span,
                    );

                    let handle = ctx.as_expression(block, &mut emitter).interrupt_emitter(
                        crate::Expression::LocalVariable(var),
                        Span::UNDEFINED,
                    )?;
                    block.extend(emitter.finish(&ctx.function.expressions));
                    ctx.local_table
                        .insert(v.handle, Declared::Runtime(Typed::Reference(handle)));

                    match initializer {
                        Some(initializer) => crate::Statement::Store {
                            pointer: handle,
                            value: initializer,
                        },
                        None => return Ok(()),
                    }
                }
                ast::LocalDecl::Const(ref c) => {
                    let mut emitter = Emitter::default();
                    emitter.start(&ctx.function.expressions);

                    let ectx = &mut ctx.as_const(block, &mut emitter);

                    let mut init = self.expression_for_abstract(c.init, ectx)?;

                    if let Some(explicit_ty) = c.ty {
                        let explicit_ty =
                            self.resolve_ast_type(explicit_ty, &mut ectx.as_global())?;
                        let explicit_ty_res = crate::proc::TypeResolution::Handle(explicit_ty);
                        init = ectx
                            .try_automatic_conversions(init, &explicit_ty_res, c.name.span)
                            .map_err(|error| match error {
                                Error::AutoConversion(error) => Error::InitializationTypeMismatch {
                                    name: c.name.span,
                                    expected: error.dest_type,
                                    got: error.source_type,
                                },
                                other => other,
                            })?;
                    } else {
                        init = ectx.concretize(init)?;
                        ectx.register_type(init)?;
                    }

                    block.extend(emitter.finish(&ctx.function.expressions));
                    ctx.local_table
                        .insert(c.handle, Declared::Const(Typed::Plain(init)));
                    ctx.named_expressions
                        .insert(init, (c.name.name.to_string(), c.name.span));

                    return Ok(());
                }
            },
            ast::StatementKind::If {
                condition,
                ref accept,
                ref reject,
            } => {
                let mut emitter = Emitter::default();
                emitter.start(&ctx.function.expressions);

                let condition =
                    self.expression(condition, &mut ctx.as_expression(block, &mut emitter))?;
                block.extend(emitter.finish(&ctx.function.expressions));

                let accept = self.block(accept, is_inside_loop, ctx)?;
                let reject = self.block(reject, is_inside_loop, ctx)?;

                crate::Statement::If {
                    condition,
                    accept,
                    reject,
                }
            }
            ast::StatementKind::Switch {
                selector,
                ref cases,
            } => {
                let mut emitter = Emitter::default();
                emitter.start(&ctx.function.expressions);

                let mut ectx = ctx.as_expression(block, &mut emitter);
                let selector = self.expression(selector, &mut ectx)?;

                let uint =
                    resolve_inner!(ectx, selector).scalar_kind() == Some(crate::ScalarKind::Uint);
                block.extend(emitter.finish(&ctx.function.expressions));

                let cases = cases
                    .iter()
                    .map(|case| {
                        Ok(crate::SwitchCase {
                            value: match case.value {
                                ast::SwitchValue::Expr(expr) => {
                                    let span = ctx.ast_expressions.get_span(expr);
                                    let expr =
                                        self.expression(expr, &mut ctx.as_global().as_const())?;
                                    match ctx.module.to_ctx().eval_expr_to_literal(expr) {
                                        Some(crate::Literal::I32(value)) if !uint => {
                                            crate::SwitchValue::I32(value)
                                        }
                                        Some(crate::Literal::U32(value)) if uint => {
                                            crate::SwitchValue::U32(value)
                                        }
                                        _ => {
                                            return Err(Error::InvalidSwitchValue { uint, span });
                                        }
                                    }
                                }
                                ast::SwitchValue::Default => crate::SwitchValue::Default,
                            },
                            body: self.block(&case.body, is_inside_loop, ctx)?,
                            fall_through: case.fall_through,
                        })
                    })
                    .collect::<Result<_, _>>()?;

                crate::Statement::Switch { selector, cases }
            }
            ast::StatementKind::Loop {
                ref body,
                ref continuing,
                break_if,
            } => {
                let body = self.block(body, true, ctx)?;
                let mut continuing = self.block(continuing, true, ctx)?;

                let mut emitter = Emitter::default();
                emitter.start(&ctx.function.expressions);
                let break_if = break_if
                    .map(|expr| {
                        self.expression(expr, &mut ctx.as_expression(&mut continuing, &mut emitter))
                    })
                    .transpose()?;
--> --------------------

--> maximum size reached

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

[ Verzeichnis aufwärts0.75unsichere Verbindung  ]