transform/shader_io: Generate a wrapper function

This is a major reworking of this transform. The old transform code
was getting unwieldy, with part of the complication coming from the
handling of multiple return statements. By generating a wrapper
function instead, we can avoid a lot of this complexity.

The original entry point function is stripped of all shader IO
attributes (as well as `stage` and `workgroup_size`), but the body is
left unmodified. A new entry point wrapper function is introduced
which calls the original function, packing/unpacking the shader inputs
as necessary, and propagates the result to the corresponding shader
outputs.

The new code has been refactored to use a state object with the
different parts of the transform split into separate functions, which
makes it much more manageable.

Fixed: tint:1076
Bug: tint:920
Change-Id: I3490a0ea7a3509a4e198ce730e476516649d8d96
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60521
Auto-Submit: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/transform/canonicalize_entry_point_io.cc

@@ -16,14 +16,12 @@
 
 #include <algorithm>
 #include <string>
+#include <unordered_set>
 #include <utility>
+#include <vector>
 
 #include "src/program_builder.h"
-#include "src/sem/block_statement.h"
 #include "src/sem/function.h"
-#include "src/sem/statement.h"
-#include "src/sem/struct.h"
-#include "src/sem/variable.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::transform::CanonicalizeEntryPointIO);
 TINT_INSTANTIATE_TYPEINFO(tint::transform::CanonicalizeEntryPointIO::Config);
@@ -62,8 +60,396 @@ bool StructMemberComparator(const ast::StructMember* a,
   }
 }
 
+// Returns true if `deco` is a shader IO decoration.
+bool IsShaderIODecoration(const ast::Decoration* deco) {
+  return deco->IsAnyOf<ast::BuiltinDecoration, ast::InterpolateDecoration,
+                       ast::InvariantDecoration, ast::LocationDecoration>();
+}
+
 }  // namespace
 
+/// State holds the current transform state for a single entry point.
+struct CanonicalizeEntryPointIO::State {
+  /// OutputValue represents a shader result that the wrapper function produces.
+  struct OutputValue {
+    /// The name of the output value.
+    std::string name;
+    /// The type of the output value.
+    ast::Type* type;
+    /// The shader IO attributes.
+    ast::DecorationList attributes;
+    /// The value itself.
+    ast::Expression* value;
+  };
+
+  /// The clone context.
+  CloneContext& ctx;
+  /// The transform config.
+  CanonicalizeEntryPointIO::Config const cfg;
+  /// The entry point function (AST).
+  ast::Function* func_ast;
+  /// The entry point function (SEM).
+  sem::Function const* func_sem;
+
+  /// The new entry point wrapper function's parameters.
+  ast::VariableList wrapper_ep_parameters;
+  /// The members of the wrapper function's struct parameter.
+  ast::StructMemberList wrapper_struct_param_members;
+  /// The name of the wrapper function's struct parameter.
+  Symbol wrapper_struct_param_name;
+  /// The parameters that will be passed to the original function.
+  ast::ExpressionList inner_call_parameters;
+  /// The members of the wrapper function's struct return type.
+  ast::StructMemberList wrapper_struct_output_members;
+  /// The wrapper function output values.
+  std::vector<OutputValue> wrapper_output_values;
+  /// The body of the wrapper function.
+  ast::StatementList wrapper_body;
+
+  /// Constructor
+  /// @param context the clone context
+  /// @param config the transform config
+  /// @param function the entry point function
+  State(CloneContext& context,
+        const CanonicalizeEntryPointIO::Config& config,
+        ast::Function* function)
+      : ctx(context),
+        cfg(config),
+        func_ast(function),
+        func_sem(ctx.src->Sem().Get(function)) {}
+
+  /// Clones the shader IO decorations from `src`.
+  /// @param src the decorations to clone
+  /// @return the cloned decorations
+  ast::DecorationList CloneShaderIOAttributes(const ast::DecorationList& src) {
+    ast::DecorationList new_decorations;
+    for (auto* deco : src) {
+      if (IsShaderIODecoration(deco)) {
+        new_decorations.push_back(ctx.Clone(deco));
+      }
+    }
+    return new_decorations;
+  }
+
+  /// Create or return a symbol for the wrapper function's struct parameter.
+  /// @returns the symbol for the struct parameter
+  Symbol InputStructSymbol() {
+    if (!wrapper_struct_param_name.IsValid()) {
+      wrapper_struct_param_name = ctx.dst->Sym();
+    }
+    return wrapper_struct_param_name;
+  }
+
+  /// Add a shader input to the entry point.
+  /// @param name the name of the shader input
+  /// @param type the type of the shader input
+  /// @param attributes the attributes to apply to the shader input
+  /// @returns an expression which evaluates to the value of the shader input
+  ast::Expression* AddInput(std::string name,
+                            ast::Type* type,
+                            ast::DecorationList attributes) {
+    if (cfg.builtin_style == BuiltinStyle::kParameter &&
+        ast::HasDecoration<ast::BuiltinDecoration>(attributes)) {
+      // If this input is a builtin and we are emitting those as parameters,
+      // then add it to the parameter list and pass it directly to the inner
+      // function.
+      wrapper_ep_parameters.push_back(
+          ctx.dst->Param(name, type, std::move(attributes)));
+      return ctx.dst->Expr(name);
+    } else {
+      // Otherwise, move it to the new structure member list.
+      wrapper_struct_param_members.push_back(
+          ctx.dst->Member(name, type, std::move(attributes)));
+      return ctx.dst->MemberAccessor(InputStructSymbol(), name);
+    }
+  }
+
+  /// Add a shader output to the entry point.
+  /// @param name the name of the shader output
+  /// @param type the type of the shader output
+  /// @param attributes the attributes to apply to the shader output
+  /// @param value the value of the shader output
+  void AddOutput(std::string name,
+                 ast::Type* type,
+                 ast::DecorationList attributes,
+                 ast::Expression* value) {
+    OutputValue output;
+    output.name = name;
+    output.type = type;
+    output.attributes = std::move(attributes);
+    output.value = value;
+    wrapper_output_values.push_back(output);
+  }
+
+  /// Process a non-struct parameter.
+  /// This creates a new object for the shader input, moving the shader IO
+  /// attributes to it. It also adds an expression to the list of parameters
+  /// that will be passed to the original function.
+  /// @param param the original function parameter
+  void ProcessNonStructParameter(const sem::Parameter* param) {
+    // Remove the shader IO attributes from the inner function parameter, and
+    // attach them to the new object instead.
+    ast::DecorationList attributes;
+    for (auto* deco : param->Declaration()->decorations()) {
+      if (IsShaderIODecoration(deco)) {
+        ctx.Remove(param->Declaration()->decorations(), deco);
+        attributes.push_back(ctx.Clone(deco));
+      }
+    }
+
+    auto name = ctx.src->Symbols().NameFor(param->Declaration()->symbol());
+    auto* type = ctx.Clone(param->Declaration()->type());
+    auto* input_expr = AddInput(name, type, std::move(attributes));
+    inner_call_parameters.push_back(input_expr);
+  }
+
+  /// Process a struct parameter.
+  /// This creates new objects for each struct member, moving the shader IO
+  /// attributes to them. It also creates the structure that will be passed to
+  /// the original function.
+  /// @param param the original function parameter
+  void ProcessStructParameter(const sem::Parameter* param) {
+    auto* str = param->Type()->As<sem::Struct>();
+
+    // Recreate struct members in the outer entry point and build an initializer
+    // list to pass them through to the inner function.
+    ast::ExpressionList inner_struct_values;
+    for (auto* member : str->Members()) {
+      if (member->Type()->Is<sem::Struct>()) {
+        TINT_ICE(Transform, ctx.dst->Diagnostics()) << "nested IO struct";
+        continue;
+      }
+
+      auto* member_ast = member->Declaration();
+      auto name = ctx.src->Symbols().NameFor(member_ast->symbol());
+      auto* type = ctx.Clone(member_ast->type());
+      auto attributes = CloneShaderIOAttributes(member_ast->decorations());
+      auto* input_expr = AddInput(name, type, std::move(attributes));
+      inner_struct_values.push_back(input_expr);
+    }
+
+    // Construct the original structure using the new shader input objects.
+    inner_call_parameters.push_back(ctx.dst->Construct(
+        ctx.Clone(param->Declaration()->type()), inner_struct_values));
+  }
+
+  /// Process the entry point return type.
+  /// This generates a list of output values that are returned by the original
+  /// function.
+  /// @param inner_ret_type the original function return type
+  /// @param original_result the result object produced by the original function
+  void ProcessReturnType(const sem::Type* inner_ret_type,
+                         Symbol original_result) {
+    if (auto* str = inner_ret_type->As<sem::Struct>()) {
+      for (auto* member : str->Members()) {
+        if (member->Type()->Is<sem::Struct>()) {
+          TINT_ICE(Transform, ctx.dst->Diagnostics()) << "nested IO struct";
+          continue;
+        }
+
+        auto* member_ast = member->Declaration();
+        auto name = ctx.src->Symbols().NameFor(member_ast->symbol());
+        auto* type = ctx.Clone(member_ast->type());
+        auto attributes = CloneShaderIOAttributes(member_ast->decorations());
+
+        // Extract the original structure member.
+        AddOutput(name, type, std::move(attributes),
+                  ctx.dst->MemberAccessor(original_result, name));
+      }
+    } else if (!inner_ret_type->Is<sem::Void>()) {
+      auto* type = ctx.Clone(func_ast->return_type());
+      auto attributes =
+          CloneShaderIOAttributes(func_ast->return_type_decorations());
+
+      // Propagate the non-struct return value as is.
+      AddOutput("value", type, std::move(attributes),
+                ctx.dst->Expr(original_result));
+    }
+  }
+
+  /// Add a fixed sample mask to the wrapper function output.
+  /// If there is already a sample mask, bitwise-and it with the fixed mask.
+  /// Otherwise, create a new output value from the fixed mask.
+  void AddFixedSampleMask() {
+    // Check the existing output values for a sample mask builtin.
+    for (auto& outval : wrapper_output_values) {
+      auto* builtin =
+          ast::GetDecoration<ast::BuiltinDecoration>(outval.attributes);
+      if (builtin && builtin->value() == ast::Builtin::kSampleMask) {
+        // Combine the authored sample mask with the fixed mask.
+        outval.value = ctx.dst->And(outval.value, cfg.fixed_sample_mask);
+        return;
+      }
+    }
+
+    // No existing sample mask builtin was found, so create a new output value
+    // using the fixed sample mask.
+    AddOutput("fixed_sample_mask", ctx.dst->ty.u32(),
+              {ctx.dst->Builtin(ast::Builtin::kSampleMask)},
+              ctx.dst->Expr(cfg.fixed_sample_mask));
+  }
+
+  /// Create the wrapper function's struct parameter and type objects.
+  void CreateInputStruct() {
+    // Sort the struct members to satisfy HLSL interfacing matching rules.
+    std::sort(wrapper_struct_param_members.begin(),
+              wrapper_struct_param_members.end(), StructMemberComparator);
+
+    // Create the new struct type.
+    auto struct_name = ctx.dst->Sym();
+    auto* in_struct = ctx.dst->create<ast::Struct>(
+        struct_name, wrapper_struct_param_members, ast::DecorationList{});
+    ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), func_ast, in_struct);
+
+    // Create a new function parameter using this struct type.
+    auto* param =
+        ctx.dst->Param(InputStructSymbol(), ctx.dst->ty.type_name(struct_name));
+    wrapper_ep_parameters.push_back(param);
+  }
+
+  /// Create and return the wrapper function's struct result object.
+  /// @returns the struct type
+  ast::Struct* CreateOutputStruct() {
+    ast::StatementList assignments;
+
+    auto wrapper_result = ctx.dst->Symbols().New("wrapper_result");
+
+    // Create the struct members and their corresponding assignment statements.
+    std::unordered_set<std::string> member_names;
+    for (auto& outval : wrapper_output_values) {
+      // Use the original output name, unless that is already taken.
+      Symbol name;
+      if (member_names.count(outval.name)) {
+        name = ctx.dst->Symbols().New(outval.name);
+      } else {
+        name = ctx.dst->Symbols().Register(outval.name);
+      }
+      member_names.insert(ctx.dst->Symbols().NameFor(name));
+
+      wrapper_struct_output_members.push_back(
+          ctx.dst->Member(name, outval.type, std::move(outval.attributes)));
+      assignments.push_back(ctx.dst->Assign(
+          ctx.dst->MemberAccessor(wrapper_result, name), outval.value));
+    }
+
+    // Sort the struct members to satisfy HLSL interfacing matching rules.
+    std::sort(wrapper_struct_output_members.begin(),
+              wrapper_struct_output_members.end(), StructMemberComparator);
+
+    // Create the new struct type.
+    auto* out_struct = ctx.dst->create<ast::Struct>(
+        ctx.dst->Sym(), wrapper_struct_output_members, ast::DecorationList{});
+    ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), func_ast, out_struct);
+
+    // Create the output struct object, assign its members, and return it.
+    auto* result_object =
+        ctx.dst->Var(wrapper_result, ctx.dst->ty.type_name(out_struct->name()));
+    wrapper_body.push_back(ctx.dst->Decl(result_object));
+    wrapper_body.insert(wrapper_body.end(), assignments.begin(),
+                        assignments.end());
+    wrapper_body.push_back(ctx.dst->Return(wrapper_result));
+
+    return out_struct;
+  }
+
+  // Recreate the original function without entry point attributes and call it.
+  /// @returns the inner function call expression
+  ast::CallExpression* CallInnerFunction() {
+    // Add a suffix to the function name, as the wrapper function will take the
+    // original entry point name.
+    auto ep_name = ctx.src->Symbols().NameFor(func_ast->symbol());
+    auto inner_name = ctx.dst->Symbols().New(ep_name + "_inner");
+
+    // Clone everything, dropping the function and return type attributes.
+    // The parameter attributes will have already been stripped during
+    // processing.
+    auto* inner_function = ctx.dst->create<ast::Function>(
+        inner_name, ctx.Clone(func_ast->params()),
+        ctx.Clone(func_ast->return_type()), ctx.Clone(func_ast->body()),
+        ast::DecorationList{}, ast::DecorationList{});
+    ctx.Replace(func_ast, inner_function);
+
+    // Call the function.
+    return ctx.dst->Call(inner_function->symbol(), inner_call_parameters);
+  }
+
+  /// Process the entry point function.
+  void Process() {
+    bool needs_fixed_sample_mask = false;
+    if (func_ast->pipeline_stage() == ast::PipelineStage::kFragment &&
+        cfg.fixed_sample_mask != 0xFFFFFFFF) {
+      needs_fixed_sample_mask = true;
+    }
+
+    // Exit early if there is no shader IO to handle.
+    if (func_sem->Parameters().size() == 0 &&
+        func_sem->ReturnType()->Is<sem::Void>() && !needs_fixed_sample_mask) {
+      return;
+    }
+
+    // Process the entry point parameters, collecting those that need to be
+    // aggregated into a single structure.
+    if (!func_sem->Parameters().empty()) {
+      for (auto* param : func_sem->Parameters()) {
+        if (param->Type()->Is<sem::Struct>()) {
+          ProcessStructParameter(param);
+        } else {
+          ProcessNonStructParameter(param);
+        }
+      }
+
+      // Create a structure parameter for the outer entry point if necessary.
+      if (!wrapper_struct_param_members.empty()) {
+        CreateInputStruct();
+      }
+    }
+
+    // Recreate the original function and call it.
+    auto* call_inner = CallInnerFunction();
+
+    // Process the return type, and start building the wrapper function body.
+    std::function<ast::Type*()> wrapper_ret_type = [&] {
+      return ctx.dst->ty.void_();
+    };
+    if (func_sem->ReturnType()->Is<sem::Void>()) {
+      // The function call is just a statement with no result.
+      wrapper_body.push_back(ctx.dst->create<ast::CallStatement>(call_inner));
+    } else {
+      // Capture the result of calling the original function.
+      auto* inner_result = ctx.dst->Const(
+          ctx.dst->Symbols().New("inner_result"), nullptr, call_inner);
+      wrapper_body.push_back(ctx.dst->Decl(inner_result));
+
+      // Process the original return type to determine the outputs that the
+      // outer function needs to produce.
+      ProcessReturnType(func_sem->ReturnType(), inner_result->symbol());
+    }
+
+    // Add a fixed sample mask, if necessary.
+    if (needs_fixed_sample_mask) {
+      AddFixedSampleMask();
+    }
+
+    // Produce the entry point outputs, if necessary.
+    if (!wrapper_output_values.empty()) {
+      auto* output_struct = CreateOutputStruct();
+      wrapper_ret_type = [&, output_struct] {
+        return ctx.dst->ty.type_name(output_struct->name());
+      };
+    }
+
+    // Create the wrapper entry point function.
+    // Take the name of the original entry point function.
+    auto name = ctx.Clone(func_ast->symbol());
+    auto* wrapper_func = ctx.dst->create<ast::Function>(
+        name, wrapper_ep_parameters, wrapper_ret_type(),
+        ctx.dst->Block(wrapper_body), ctx.Clone(func_ast->decorations()),
+        ast::DecorationList{});
+    ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), func_ast,
+                    wrapper_func);
+  }
+};
+
 void CanonicalizeEntryPointIO::Run(CloneContext& ctx,
                                    const DataMap& inputs,
                                    DataMap&) {
@@ -75,302 +461,27 @@ void CanonicalizeEntryPointIO::Run(CloneContext& ctx,
     return;
   }
 
-  // Strip entry point IO decorations from struct declarations.
-  // TODO(jrprice): This code is duplicated with the SPIR-V transform.
+  // Remove entry point IO attributes from struct declarations.
+  // New structures will be created for each entry point, as necessary.
   for (auto* ty : ctx.src->AST().TypeDecls()) {
     if (auto* struct_ty = ty->As<ast::Struct>()) {
-      // Build new list of struct members without entry point IO decorations.
-      ast::StructMemberList new_struct_members;
       for (auto* member : struct_ty->members()) {
-        ast::DecorationList new_decorations = RemoveDecorations(
-            ctx, member->decorations(), [](const ast::Decoration* deco) {
-              return deco->IsAnyOf<
-                  ast::BuiltinDecoration, ast::InterpolateDecoration,
-                  ast::InvariantDecoration, ast::LocationDecoration>();
-            });
-        new_struct_members.push_back(
-            ctx.dst->Member(ctx.Clone(member->symbol()),
-                            ctx.Clone(member->type()), new_decorations));
+        for (auto* deco : member->decorations()) {
+          if (IsShaderIODecoration(deco)) {
+            ctx.Remove(member->decorations(), deco);
+          }
+        }
       }
-
-      // Redeclare the struct.
-      auto new_struct_name = ctx.Clone(struct_ty->name());
-      auto* new_struct =
-          ctx.dst->create<ast::Struct>(new_struct_name, new_struct_members,
-                                       ctx.Clone(struct_ty->decorations()));
-      ctx.Replace(struct_ty, new_struct);
     }
   }
 
-  // Returns true if `decos` contains a `sample_mask` builtin.
-  auto has_sample_mask_builtin = [](const ast::DecorationList& decos) {
-    if (auto* builtin = ast::GetDecoration<ast::BuiltinDecoration>(decos)) {
-      return builtin->value() == ast::Builtin::kSampleMask;
-    }
-    return false;
-  };
-
   for (auto* func_ast : ctx.src->AST().Functions()) {
     if (!func_ast->IsEntryPoint()) {
       continue;
     }
 
-    auto* func = ctx.src->Sem().Get(func_ast);
-
-    bool needs_fixed_sample_mask =
-        func_ast->pipeline_stage() == ast::PipelineStage::kFragment &&
-        cfg->fixed_sample_mask != 0xFFFFFFFF;
-
-    ast::VariableList new_parameters;
-
-    if (!func->Parameters().empty()) {
-      // Collect all parameters and build a list of new struct members.
-      auto new_struct_param_symbol = ctx.dst->Sym();
-      ast::StructMemberList new_struct_members;
-      for (auto* param : func->Parameters()) {
-        if (cfg->builtin_style == BuiltinStyle::kParameter &&
-            ast::HasDecoration<ast::BuiltinDecoration>(
-                param->Declaration()->decorations())) {
-          // If this parameter is a builtin and we are emitting those as
-          // parameters, then just clone it as is.
-          new_parameters.push_back(
-              ctx.Clone(const_cast<ast::Variable*>(param->Declaration())));
-          continue;
-        }
-
-        auto param_name = ctx.Clone(param->Declaration()->symbol());
-        auto* param_ty = param->Type();
-        auto* param_declared_ty = param->Declaration()->type();
-
-        ast::Expression* func_const_initializer = nullptr;
-
-        if (auto* str = param_ty->As<sem::Struct>()) {
-          // Pull out all struct members and build initializer list.
-          ast::ExpressionList init_values;
-          for (auto* member : str->Members()) {
-            if (member->Type()->Is<sem::Struct>()) {
-              TINT_ICE(Transform, ctx.dst->Diagnostics())
-                  << "nested pipeline IO struct";
-            }
-
-            ast::DecorationList new_decorations = RemoveDecorations(
-                ctx, member->Declaration()->decorations(),
-                [](const ast::Decoration* deco) {
-                  return !deco->IsAnyOf<
-                      ast::BuiltinDecoration, ast::InterpolateDecoration,
-                      ast::InvariantDecoration, ast::LocationDecoration>();
-                });
-
-            if (cfg->builtin_style == BuiltinStyle::kParameter &&
-                ast::HasDecoration<ast::BuiltinDecoration>(
-                    member->Declaration()->decorations())) {
-              // If this struct member is a builtin and we are emitting those as
-              // parameters, then move it to the parameter list.
-              auto* member_ty = CreateASTTypeFor(ctx, member->Type());
-              auto new_param_name = ctx.dst->Sym();
-              new_parameters.push_back(
-                  ctx.dst->Param(new_param_name, member_ty, new_decorations));
-              init_values.push_back(ctx.dst->Expr(new_param_name));
-              continue;
-            }
-
-            auto member_name = ctx.Clone(member->Declaration()->symbol());
-            auto* member_type = ctx.Clone(member->Declaration()->type());
-            new_struct_members.push_back(
-                ctx.dst->Member(member_name, member_type, new_decorations));
-            init_values.push_back(
-                ctx.dst->MemberAccessor(new_struct_param_symbol, member_name));
-          }
-
-          func_const_initializer =
-              ctx.dst->Construct(ctx.Clone(param_declared_ty), init_values);
-        } else {
-          new_struct_members.push_back(
-              ctx.dst->Member(param_name, ctx.Clone(param_declared_ty),
-                              ctx.Clone(param->Declaration()->decorations())));
-          func_const_initializer =
-              ctx.dst->MemberAccessor(new_struct_param_symbol, param_name);
-        }
-
-        // Create a function-scope const to replace the parameter.
-        // Initialize it with the value extracted from the new struct parameter.
-        auto* func_const = ctx.dst->Const(
-            param_name, ctx.Clone(param_declared_ty), func_const_initializer);
-        ctx.InsertFront(func_ast->body()->statements(),
-                        ctx.dst->WrapInStatement(func_const));
-
-        // Replace all uses of the function parameter with the function const.
-        for (auto* user : param->Users()) {
-          ctx.Replace<ast::Expression>(user->Declaration(),
-                                       ctx.dst->Expr(param_name));
-        }
-      }
-
-      if (!new_struct_members.empty()) {
-        // Sort struct members to satisfy HLSL interfacing matching rules.
-        std::sort(new_struct_members.begin(), new_struct_members.end(),
-                  StructMemberComparator);
-
-        // Create the new struct type.
-        auto in_struct_name = ctx.dst->Sym();
-        auto* in_struct = ctx.dst->create<ast::Struct>(
-            in_struct_name, new_struct_members, ast::DecorationList{});
-        ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), func_ast,
-                         in_struct);
-
-        // Create a new function parameter using this struct type.
-        auto* struct_param = ctx.dst->Param(
-            new_struct_param_symbol, ctx.dst->ty.type_name(in_struct_name));
-        new_parameters.push_back(struct_param);
-      }
-    }
-
-    // Handle return type.
-    auto* ret_type = func->ReturnType();
-    std::function<ast::Type*()> new_ret_type;
-    if (ret_type->Is<sem::Void>() && !needs_fixed_sample_mask) {
-      new_ret_type = [&ctx] { return ctx.dst->ty.void_(); };
-    } else {
-      ast::StructMemberList new_struct_members;
-
-      bool has_authored_sample_mask = false;
-
-      if (auto* str = ret_type->As<sem::Struct>()) {
-        // Rebuild struct with only the entry point IO attributes.
-        for (auto* member : str->Members()) {
-          if (member->Type()->Is<sem::Struct>()) {
-            TINT_ICE(Transform, ctx.dst->Diagnostics())
-                << "nested pipeline IO struct";
-          }
-
-          ast::DecorationList new_decorations = RemoveDecorations(
-              ctx, member->Declaration()->decorations(),
-              [](const ast::Decoration* deco) {
-                return !deco->IsAnyOf<
-                    ast::BuiltinDecoration, ast::InterpolateDecoration,
-                    ast::InvariantDecoration, ast::LocationDecoration>();
-              });
-          auto symbol = ctx.Clone(member->Declaration()->symbol());
-          auto* member_ty = ctx.Clone(member->Declaration()->type());
-          new_struct_members.push_back(
-              ctx.dst->Member(symbol, member_ty, new_decorations));
-
-          if (has_sample_mask_builtin(new_decorations)) {
-            has_authored_sample_mask = true;
-          }
-        }
-      } else if (!ret_type->Is<sem::Void>()) {
-        auto* member_ty = ctx.Clone(func->Declaration()->return_type());
-        auto decos = ctx.Clone(func_ast->return_type_decorations());
-        new_struct_members.push_back(
-            ctx.dst->Member("value", member_ty, std::move(decos)));
-
-        if (has_sample_mask_builtin(func_ast->return_type_decorations())) {
-          has_authored_sample_mask = true;
-        }
-      }
-
-      // If a sample mask builtin is required and the shader source did not
-      // contain one, create one now.
-      if (needs_fixed_sample_mask && !has_authored_sample_mask) {
-        new_struct_members.push_back(
-            ctx.dst->Member(ctx.dst->Sym(), ctx.dst->ty.u32(),
-                            {ctx.dst->Builtin(ast::Builtin::kSampleMask)}));
-      }
-
-      // Sort struct members to satisfy HLSL interfacing matching rules.
-      std::sort(new_struct_members.begin(), new_struct_members.end(),
-                StructMemberComparator);
-
-      // Create the new struct type.
-      auto out_struct_name = ctx.dst->Sym();
-      auto* out_struct = ctx.dst->create<ast::Struct>(
-          out_struct_name, new_struct_members, ast::DecorationList{});
-      ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), func_ast,
-                       out_struct);
-      new_ret_type = [out_struct_name, &ctx] {
-        return ctx.dst->ty.type_name(out_struct_name);
-      };
-
-      // Replace all return statements.
-      for (auto* ret : func->ReturnStatements()) {
-        auto* ret_sem = ctx.src->Sem().Get(ret);
-        // Reconstruct the return value using the newly created struct.
-        std::function<ast::Expression*()> new_ret_value = [&ctx, ret] {
-          return ctx.Clone(ret->value());
-        };
-
-        ast::ExpressionList ret_values;
-        if (ret_type->Is<sem::Struct>()) {
-          if (!ret->value()->Is<ast::IdentifierExpression>()) {
-            // Create a const to hold the return value expression to avoid
-            // re-evaluating it multiple times.
-            auto temp = ctx.dst->Sym();
-            auto* ty = CreateASTTypeFor(ctx, ret_type);
-            auto* temp_var =
-                ctx.dst->Decl(ctx.dst->Const(temp, ty, new_ret_value()));
-            ctx.InsertBefore(ret_sem->Block()->Declaration()->statements(), ret,
-                             temp_var);
-            new_ret_value = [&ctx, temp] { return ctx.dst->Expr(temp); };
-          }
-
-          for (auto* member : new_struct_members) {
-            ast::Expression* expr = nullptr;
-
-            if (needs_fixed_sample_mask &&
-                has_sample_mask_builtin(member->decorations())) {
-              // Use the fixed sample mask, combining it with the authored value
-              // if there is one.
-              expr = ctx.dst->Expr(cfg->fixed_sample_mask);
-              if (has_authored_sample_mask) {
-                expr = ctx.dst->And(
-                    ctx.dst->MemberAccessor(new_ret_value(), member->symbol()),
-                    expr);
-              }
-            } else {
-              expr = ctx.dst->MemberAccessor(new_ret_value(), member->symbol());
-            }
-            ret_values.push_back(expr);
-          }
-        } else {
-          if (!ret_type->Is<sem::Void>()) {
-            ret_values.push_back(new_ret_value());
-          }
-
-          if (needs_fixed_sample_mask) {
-            // If the original return value was a sample mask, `and` it with the
-            // fixed mask and return the result.
-            // Otherwise, append the fixed mask to the list of return values,
-            // since it will be the last element of the output struct.
-            if (has_authored_sample_mask) {
-              ret_values[0] =
-                  ctx.dst->And(ret_values[0], cfg->fixed_sample_mask);
-            } else {
-              ret_values.push_back(ctx.dst->Expr(cfg->fixed_sample_mask));
-            }
-          }
-        }
-
-        auto* new_ret =
-            ctx.dst->Return(ctx.dst->Construct(new_ret_type(), ret_values));
-        ctx.Replace(ret, new_ret);
-      }
-
-      if (needs_fixed_sample_mask && func->ReturnStatements().empty()) {
-        // There we no return statements but we need to return a fixed sample
-        // mask, so add a return statement that does this.
-        ctx.InsertBack(func_ast->body()->statements(),
-                       ctx.dst->Return(ctx.dst->Construct(
-                           new_ret_type(), cfg->fixed_sample_mask)));
-      }
-    }
-
-    // Rewrite the function header with the new parameters.
-    auto* new_func = ctx.dst->create<ast::Function>(
-        func_ast->source(), ctx.Clone(func_ast->symbol()), new_parameters,
-        new_ret_type(), ctx.Clone(func_ast->body()),
-        ctx.Clone(func_ast->decorations()), ast::DecorationList{});
-    ctx.Replace(func_ast, new_func);
+    State state(ctx, *cfg, func_ast);
+    state.Process();
   }
 
   ctx.Clone();

src/transform/canonicalize_entry_point_io.h

@@ -21,10 +21,12 @@ namespace tint {
 namespace transform {
 
 /// CanonicalizeEntryPointIO is a transform used to rewrite shader entry point
-/// interfaces into a form that the generators can handle. After the transform,
-/// an entry point's parameters will be aggregated into a single struct, and its
-/// return type will either be a struct or void. All structs in the module that
-/// have entry point IO decorations will have exactly one pipeline stage usage.
+/// interfaces into a form that the generators can handle. Each entry point
+/// function is stripped of all shader IO attributes and wrapped in a function
+/// that provides the shader interface.
+/// The transform config determines how shader IO parameters will be exposed.
+/// Entry point return values are always produced as a structure, and optionally
+/// include additional builtins as per the transform config.
 ///
 /// Before:
 /// ```
@@ -36,12 +38,15 @@ namespace transform {
 /// [[stage(fragment)]]
 /// fn frag_main([[builtin(position)]] coord : vec4<f32>,
 ///              locations : Locations) -> [[location(0)]] f32 {
+///   if (coord.w > 1.0) {
+///     return 0.0;
+///   }
 ///   var col : f32 = (coord.x * locations.loc1);
 ///   return col;
 /// }
 /// ```
 ///
-/// After:
+/// After (using structures for all parameters):
 /// ```
 /// struct Locations{
 ///   loc1 : f32;
@@ -58,14 +63,21 @@ namespace transform {
 ///   [[location(0)]] loc0 : f32;
 /// };
 ///
+/// fn frag_main_inner(coord : vec4<f32>,
+///                    locations : Locations) -> f32 {
+///   if (coord.w > 1.0) {
+///     return 0.0;
+///   }
+///   var col : f32 = (coord.x * locations.loc1);
+///   return col;
+/// }
+///
 /// [[stage(fragment)]]
 /// fn frag_main(in : frag_main_in) -> frag_main_out {
-///   const coord = in.coord;
-///   const locations = Locations(in.loc1, in.loc2);
-///   var col : f32 = (coord.x * locations.loc1);
-///   var retval : frag_main_out;
-///   retval.loc0 = col;
-///   return retval;
+///   let inner_retval = frag_main_inner(in.coord, Locations(in.loc1, in.loc2));
+///   var wrapper_result : frag_main_out;
+///   wrapper_result.loc0 = inner_retval;
+///   return wrapper_result;
 /// }
 /// ```
 class CanonicalizeEntryPointIO
@@ -111,6 +123,8 @@ class CanonicalizeEntryPointIO
   /// @param inputs optional extra transform-specific input data
   /// @param outputs optional extra transform-specific output data
   void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) override;
+
+  struct State;
 };
 
 }  // namespace transform

src/transform/canonicalize_entry_point_io_test.cc

@@ -34,6 +34,29 @@ TEST_F(CanonicalizeEntryPointIOTest, Error_MissingTransformData) {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CanonicalizeEntryPointIOTest, NoShaderIO) {
+  // Test that we do not introduce wrapper functions when there is no shader IO
+  // to process.
+  auto* src = R"(
+[[stage(fragment)]]
+fn frag_main() {
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn comp_main() {
+}
+)";
+
+  auto* expect = src;
+
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(CanonicalizeEntryPointIOTest, Parameters_BuiltinsAsParameters) {
   auto* src = R"(
 [[stage(fragment)]]
@@ -52,11 +75,13 @@ struct tint_symbol_1 {
   loc2 : vec4<u32>;
 };
 
+fn frag_main_inner(loc1 : f32, loc2 : vec4<u32>, coord : vec4<f32>) {
+  var col : f32 = (coord.x * loc1);
+}
+
 [[stage(fragment)]]
 fn frag_main([[builtin(position)]] coord : vec4<f32>, tint_symbol : tint_symbol_1) {
-  let loc1 : f32 = tint_symbol.loc1;
-  let loc2 : vec4<u32> = tint_symbol.loc2;
-  var col : f32 = (coord.x * loc1);
+  frag_main_inner(tint_symbol.loc1, tint_symbol.loc2, coord);
 }
 )";
 
@@ -88,12 +113,13 @@ struct tint_symbol_1 {
   coord : vec4<f32>;
 };
 
+fn frag_main_inner(loc1 : f32, loc2 : vec4<u32>, coord : vec4<f32>) {
+  var col : f32 = (coord.x * loc1);
+}
+
 [[stage(fragment)]]
 fn frag_main(tint_symbol : tint_symbol_1) {
-  let loc1 : f32 = tint_symbol.loc1;
-  let loc2 : vec4<u32> = tint_symbol.loc2;
-  let coord : vec4<f32> = tint_symbol.coord;
-  var col : f32 = (coord.x * loc1);
+  frag_main_inner(tint_symbol.loc1, tint_symbol.loc2, tint_symbol.coord);
 }
 )";
 
@@ -123,10 +149,13 @@ struct tint_symbol_1 {
   loc1 : myf32;
 };
 
+fn frag_main_inner(loc1 : myf32) {
+  var x : myf32 = loc1;
+}
+
 [[stage(fragment)]]
 fn frag_main(tint_symbol : tint_symbol_1) {
-  let loc1 : myf32 = tint_symbol.loc1;
-  var x : myf32 = loc1;
+  frag_main_inner(tint_symbol.loc1);
 }
 )";
 
@@ -156,10 +185,12 @@ struct tint_symbol_1 {
   loc2 : vec4<u32>;
 };
 
+fn frag_main_inner(loc1 : f32, loc2 : vec4<u32>, coord : vec4<f32>) {
+}
+
 [[stage(fragment)]]
 fn frag_main([[builtin(position)]] coord : vec4<f32>, tint_symbol : tint_symbol_1) {
-  let loc1 : f32 = tint_symbol.loc1;
-  let loc2 : vec4<u32> = tint_symbol.loc2;
+  frag_main_inner(tint_symbol.loc1, tint_symbol.loc2, coord);
 }
 )";
 
@@ -191,11 +222,12 @@ struct tint_symbol_1 {
   coord : vec4<f32>;
 };
 
+fn frag_main_inner(loc1 : f32, loc2 : vec4<u32>, coord : vec4<f32>) {
+}
+
 [[stage(fragment)]]
 fn frag_main(tint_symbol : tint_symbol_1) {
-  let loc1 : f32 = tint_symbol.loc1;
-  let loc2 : vec4<u32> = tint_symbol.loc2;
-  let coord : vec4<f32> = tint_symbol.coord;
+  frag_main_inner(tint_symbol.loc1, tint_symbol.loc2, tint_symbol.coord);
 }
 )";
 
@@ -235,7 +267,7 @@ struct FragLocations {
   loc2 : vec4<u32>;
 };
 
-struct tint_symbol_2 {
+struct tint_symbol_1 {
   [[location(0)]]
   loc0 : f32;
   [[location(1)]]
@@ -244,13 +276,14 @@ struct tint_symbol_2 {
   loc2 : vec4<u32>;
 };
 
-[[stage(fragment)]]
-fn frag_main([[builtin(position)]] tint_symbol_1 : vec4<f32>, tint_symbol : tint_symbol_2) {
-  let loc0 : f32 = tint_symbol.loc0;
-  let locations : FragLocations = FragLocations(tint_symbol.loc1, tint_symbol.loc2);
-  let builtins : FragBuiltins = FragBuiltins(tint_symbol_1);
+fn frag_main_inner(loc0 : f32, locations : FragLocations, builtins : FragBuiltins) {
   var col : f32 = ((builtins.coord.x * locations.loc1) + loc0);
 }
+
+[[stage(fragment)]]
+fn frag_main([[builtin(position)]] coord : vec4<f32>, tint_symbol : tint_symbol_1) {
+  frag_main_inner(tint_symbol.loc0, FragLocations(tint_symbol.loc1, tint_symbol.loc2), FragBuiltins(coord));
+}
 )";
 
   DataMap data;
@@ -300,12 +333,13 @@ struct tint_symbol_1 {
   coord : vec4<f32>;
 };
 
+fn frag_main_inner(loc0 : f32, locations : FragLocations, builtins : FragBuiltins) {
+  var col : f32 = ((builtins.coord.x * locations.loc1) + loc0);
+}
+
 [[stage(fragment)]]
 fn frag_main(tint_symbol : tint_symbol_1) {
-  let loc0 : f32 = tint_symbol.loc0;
-  let locations : FragLocations = FragLocations(tint_symbol.loc1, tint_symbol.loc2);
-  let builtins : FragBuiltins = FragBuiltins(tint_symbol.coord);
-  var col : f32 = ((builtins.coord.x * locations.loc1) + loc0);
+  frag_main_inner(tint_symbol.loc0, FragLocations(tint_symbol.loc1, tint_symbol.loc2), FragBuiltins(tint_symbol.coord));
 }
 )";
 
@@ -331,9 +365,16 @@ struct tint_symbol {
   value : f32;
 };
 
+fn frag_main_inner() -> f32 {
+  return 1.0;
+}
+
 [[stage(fragment)]]
 fn frag_main() -> tint_symbol {
-  return tint_symbol(1.0);
+  let inner_result = frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 )";
 
@@ -379,13 +420,22 @@ struct tint_symbol {
   mask : u32;
 };
 
-[[stage(fragment)]]
-fn frag_main() -> tint_symbol {
+fn frag_main_inner() -> FragOutput {
   var output : FragOutput;
   output.depth = 1.0;
   output.mask = 7u;
   output.color = vec4<f32>(0.5, 0.5, 0.5, 1.0);
-  return tint_symbol(output.color, output.depth, output.mask);
+  return output;
+}
+
+[[stage(fragment)]]
+fn frag_main() -> tint_symbol {
+  let inner_result = frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.color = inner_result.color;
+  wrapper_result.depth = inner_result.depth;
+  wrapper_result.mask = inner_result.mask;
+  return wrapper_result;
 }
 )";
 
@@ -436,10 +486,13 @@ struct tint_symbol_1 {
   mul : f32;
 };
 
+fn frag_main1_inner(inputs : FragmentInput) {
+  var x : f32 = foo(inputs);
+}
+
 [[stage(fragment)]]
 fn frag_main1(tint_symbol : tint_symbol_1) {
-  let inputs : FragmentInput = FragmentInput(tint_symbol.value, tint_symbol.mul);
-  var x : f32 = foo(inputs);
+  frag_main1_inner(FragmentInput(tint_symbol.value, tint_symbol.mul));
 }
 
 struct tint_symbol_3 {
@@ -449,10 +502,13 @@ struct tint_symbol_3 {
   mul : f32;
 };
 
+fn frag_main2_inner(inputs : FragmentInput) {
+  var x : f32 = foo(inputs);
+}
+
 [[stage(fragment)]]
 fn frag_main2(tint_symbol_2 : tint_symbol_3) {
-  let inputs : FragmentInput = FragmentInput(tint_symbol_2.value, tint_symbol_2.mul);
-  var x : f32 = foo(inputs);
+  frag_main2_inner(FragmentInput(tint_symbol_2.value, tint_symbol_2.mul));
 }
 )";
 
@@ -512,13 +568,16 @@ struct tint_symbol_1 {
   col2 : f32;
 };
 
-[[stage(fragment)]]
-fn frag_main1(tint_symbol : tint_symbol_1) {
-  let inputs : FragmentInput = FragmentInput(tint_symbol.col1, tint_symbol.col2);
+fn frag_main1_inner(inputs : FragmentInput) {
   global_inputs = inputs;
   var r : f32 = foo();
   var g : f32 = bar();
 }
+
+[[stage(fragment)]]
+fn frag_main1(tint_symbol : tint_symbol_1) {
+  frag_main1_inner(FragmentInput(tint_symbol.col1, tint_symbol.col2));
+}
 )";
 
   DataMap data;
@@ -593,12 +652,18 @@ struct tint_symbol_2 {
   col2 : myf32;
 };
 
+fn frag_main_inner(inputs : MyFragmentInput) -> MyFragmentOutput {
+  var x : myf32 = foo(inputs);
+  return MyFragmentOutput(x, inputs.col2);
+}
+
 [[stage(fragment)]]
 fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
-  let inputs : MyFragmentInput = MyFragmentInput(tint_symbol.col1, tint_symbol.col2);
-  var x : myf32 = foo(inputs);
-  let tint_symbol_3 : FragmentOutput = MyFragmentOutput(x, inputs.col2);
-  return tint_symbol_2(tint_symbol_3.col1, tint_symbol_3.col2);
+  let inner_result = frag_main_inner(MyFragmentInput(tint_symbol.col1, tint_symbol.col2));
+  var wrapper_result : tint_symbol_2;
+  wrapper_result.col1 = inner_result.col1;
+  wrapper_result.col2 = inner_result.col2;
+  return wrapper_result;
 }
 )";
 
@@ -660,13 +725,22 @@ struct tint_symbol {
   pos : vec4<f32>;
 };
 
-[[stage(vertex)]]
-fn vert_main() -> tint_symbol {
-  let tint_symbol_1 : VertexOut = VertexOut();
-  return tint_symbol(tint_symbol_1.loc1, tint_symbol_1.loc2, tint_symbol_1.loc3, tint_symbol_1.pos);
+fn vert_main_inner() -> VertexOut {
+  return VertexOut();
 }
 
-struct tint_symbol_3 {
+[[stage(vertex)]]
+fn vert_main() -> tint_symbol {
+  let inner_result = vert_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.pos = inner_result.pos;
+  wrapper_result.loc1 = inner_result.loc1;
+  wrapper_result.loc2 = inner_result.loc2;
+  wrapper_result.loc3 = inner_result.loc3;
+  return wrapper_result;
+}
+
+struct tint_symbol_2 {
   [[location(1), interpolate(flat)]]
   loc1 : f32;
   [[location(2), interpolate(linear, sample)]]
@@ -675,12 +749,14 @@ struct tint_symbol_3 {
   loc3 : f32;
 };
 
-[[stage(fragment)]]
-fn frag_main(tint_symbol_2 : tint_symbol_3) {
-  let inputs : FragmentIn = FragmentIn(tint_symbol_2.loc1, tint_symbol_2.loc2);
-  let loc3 : f32 = tint_symbol_2.loc3;
+fn frag_main_inner(inputs : FragmentIn, loc3 : f32) {
   let x = ((inputs.loc1 + inputs.loc2) + loc3);
 }
+
+[[stage(fragment)]]
+fn frag_main(tint_symbol_1 : tint_symbol_2) {
+  frag_main_inner(FragmentIn(tint_symbol_1.loc1, tint_symbol_1.loc2), tint_symbol_1.loc3);
+}
 )";
 
   DataMap data;
@@ -718,20 +794,33 @@ struct tint_symbol {
   pos : vec4<f32>;
 };
 
-[[stage(vertex)]]
-fn main1() -> tint_symbol {
-  let tint_symbol_1 : VertexOut = VertexOut();
-  return tint_symbol(tint_symbol_1.pos);
+fn main1_inner() -> VertexOut {
+  return VertexOut();
 }
 
-struct tint_symbol_2 {
+[[stage(vertex)]]
+fn main1() -> tint_symbol {
+  let inner_result = main1_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.pos = inner_result.pos;
+  return wrapper_result;
+}
+
+struct tint_symbol_1 {
   [[builtin(position), invariant]]
   value : vec4<f32>;
 };
 
+fn main2_inner() -> vec4<f32> {
+  return vec4<f32>();
+}
+
 [[stage(vertex)]]
-fn main2() -> tint_symbol_2 {
-  return tint_symbol_2(vec4<f32>());
+fn main2() -> tint_symbol_1 {
+  let inner_result_1 = main2_inner();
+  var wrapper_result_1 : tint_symbol_1;
+  wrapper_result_1.value = inner_result_1;
+  return wrapper_result_1;
 }
 )";
 
@@ -792,11 +881,16 @@ struct tint_symbol_2 {
   value : f32;
 };
 
+fn frag_main_inner(inputs : FragmentInput) -> FragmentOutput {
+  return FragmentOutput(((inputs.coord.x * inputs.value) + inputs.loc0));
+}
+
 [[stage(fragment)]]
 fn frag_main(tint_symbol : tint_symbol_1) -> tint_symbol_2 {
-  let inputs : FragmentInput = FragmentInput(tint_symbol.value, tint_symbol.coord, tint_symbol.loc0);
-  let tint_symbol_3 : FragmentOutput = FragmentOutput(((inputs.coord.x * inputs.value) + inputs.loc0));
-  return tint_symbol_2(tint_symbol_3.value);
+  let inner_result = frag_main_inner(FragmentInput(tint_symbol.value, tint_symbol.coord, tint_symbol.loc0));
+  var wrapper_result : tint_symbol_2;
+  wrapper_result.value = inner_result.value;
+  return wrapper_result;
 }
 )";
 
@@ -865,13 +959,23 @@ struct tint_symbol {
   pos : vec4<f32>;
 };
 
-[[stage(vertex)]]
-fn vert_main() -> tint_symbol {
-  let tint_symbol_1 : VertexOutput = VertexOutput();
-  return tint_symbol(tint_symbol_1.a, tint_symbol_1.b, tint_symbol_1.c, tint_symbol_1.d, tint_symbol_1.pos);
+fn vert_main_inner() -> VertexOutput {
+  return VertexOutput();
 }
 
-struct tint_symbol_3 {
+[[stage(vertex)]]
+fn vert_main() -> tint_symbol {
+  let inner_result = vert_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.b = inner_result.b;
+  wrapper_result.pos = inner_result.pos;
+  wrapper_result.d = inner_result.d;
+  wrapper_result.a = inner_result.a;
+  wrapper_result.c = inner_result.c;
+  return wrapper_result;
+}
+
+struct tint_symbol_2 {
   [[location(0)]]
   a : f32;
   [[location(1)]]
@@ -886,12 +990,12 @@ struct tint_symbol_3 {
   ff : bool;
 };
 
+fn frag_main_inner(ff : bool, c : i32, inputs : FragmentInputExtra, b : u32) {
+}
+
 [[stage(fragment)]]
-fn frag_main(tint_symbol_2 : tint_symbol_3) {
-  let ff : bool = tint_symbol_2.ff;
-  let c : i32 = tint_symbol_2.c;
-  let inputs : FragmentInputExtra = FragmentInputExtra(tint_symbol_2.d, tint_symbol_2.pos, tint_symbol_2.a);
-  let b : u32 = tint_symbol_2.b;
+fn frag_main(tint_symbol_1 : tint_symbol_2) {
+  frag_main_inner(tint_symbol_1.ff, tint_symbol_1.c, FragmentInputExtra(tint_symbol_1.d, tint_symbol_1.pos, tint_symbol_1.a), tint_symbol_1.b);
 }
 )";
 
@@ -916,9 +1020,12 @@ struct tint_symbol_2 {
   col : f32;
 };
 
+fn tint_symbol_1_inner(col : f32) {
+}
+
 [[stage(fragment)]]
 fn tint_symbol_1(tint_symbol : tint_symbol_2) {
-  let col : f32 = tint_symbol.col;
+  tint_symbol_1_inner(tint_symbol.col);
 }
 )";
 
@@ -938,14 +1045,20 @@ fn frag_main() {
 )";
 
   auto* expect = R"(
-struct tint_symbol_1 {
+struct tint_symbol {
   [[builtin(sample_mask)]]
-  tint_symbol : u32;
+  fixed_sample_mask : u32;
 };
 
+fn frag_main_inner() {
+}
+
 [[stage(fragment)]]
-fn frag_main() -> tint_symbol_1 {
-  return tint_symbol_1(3u);
+fn frag_main() -> tint_symbol {
+  frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.fixed_sample_mask = 3u;
+  return wrapper_result;
 }
 )";
 
@@ -966,14 +1079,21 @@ fn frag_main() {
 )";
 
   auto* expect = R"(
-struct tint_symbol_1 {
+struct tint_symbol {
   [[builtin(sample_mask)]]
-  tint_symbol : u32;
+  fixed_sample_mask : u32;
 };
 
+fn frag_main_inner() {
+  return;
+}
+
 [[stage(fragment)]]
-fn frag_main() -> tint_symbol_1 {
-  return tint_symbol_1(3u);
+fn frag_main() -> tint_symbol {
+  frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.fixed_sample_mask = 3u;
+  return wrapper_result;
 }
 )";
 
@@ -999,9 +1119,16 @@ struct tint_symbol {
   value : u32;
 };
 
+fn frag_main_inner() -> u32 {
+  return 7u;
+}
+
 [[stage(fragment)]]
 fn frag_main() -> tint_symbol {
-  return tint_symbol((7u & 3u));
+  let inner_result = frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.value = (inner_result & 3u);
+  return wrapper_result;
 }
 )";
 
@@ -1022,16 +1149,24 @@ fn frag_main() -> [[location(0)]] f32 {
 )";
 
   auto* expect = R"(
-struct tint_symbol_1 {
+struct tint_symbol {
   [[location(0)]]
   value : f32;
   [[builtin(sample_mask)]]
-  tint_symbol : u32;
+  fixed_sample_mask : u32;
 };
 
+fn frag_main_inner() -> f32 {
+  return 1.0;
+}
+
 [[stage(fragment)]]
-fn frag_main() -> tint_symbol_1 {
-  return tint_symbol_1(1.0, 3u);
+fn frag_main() -> tint_symbol {
+  let inner_result = frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.value = inner_result;
+  wrapper_result.fixed_sample_mask = 3u;
+  return wrapper_result;
 }
 )";
 
@@ -1073,10 +1208,18 @@ struct tint_symbol {
   mask : u32;
 };
 
+fn frag_main_inner() -> Output {
+  return Output(0.5, 7u, 1.0);
+}
+
 [[stage(fragment)]]
 fn frag_main() -> tint_symbol {
-  let tint_symbol_1 : Output = Output(0.5, 7u, 1.0);
-  return tint_symbol(tint_symbol_1.value, tint_symbol_1.depth, (tint_symbol_1.mask & 3u));
+  let inner_result = frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.depth = inner_result.depth;
+  wrapper_result.mask = (inner_result.mask & 3u);
+  wrapper_result.value = inner_result.value;
+  return wrapper_result;
 }
 )";
 
@@ -1108,19 +1251,27 @@ struct Output {
   value : f32;
 };
 
-struct tint_symbol_1 {
+struct tint_symbol {
   [[location(0)]]
   value : f32;
   [[builtin(frag_depth)]]
   depth : f32;
   [[builtin(sample_mask)]]
-  tint_symbol : u32;
+  fixed_sample_mask : u32;
 };
 
+fn frag_main_inner() -> Output {
+  return Output(0.5, 1.0);
+}
+
 [[stage(fragment)]]
-fn frag_main() -> tint_symbol_1 {
-  let tint_symbol_2 : Output = Output(0.5, 1.0);
-  return tint_symbol_1(tint_symbol_2.value, tint_symbol_2.depth, 3u);
+fn frag_main() -> tint_symbol {
+  let inner_result = frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.depth = inner_result.depth;
+  wrapper_result.value = inner_result.value;
+  wrapper_result.fixed_sample_mask = 3u;
+  return wrapper_result;
 }
 )";
 
@@ -1160,31 +1311,53 @@ struct tint_symbol {
   value : u32;
 };
 
-[[stage(fragment)]]
-fn frag_main1() -> tint_symbol {
-  return tint_symbol((7u & 3u));
+fn frag_main1_inner() -> u32 {
+  return 7u;
 }
 
-struct tint_symbol_2 {
+[[stage(fragment)]]
+fn frag_main1() -> tint_symbol {
+  let inner_result = frag_main1_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.value = (inner_result & 3u);
+  return wrapper_result;
+}
+
+struct tint_symbol_1 {
   [[location(0)]]
   value : f32;
   [[builtin(sample_mask)]]
-  tint_symbol_1 : u32;
+  fixed_sample_mask : u32;
 };
 
-[[stage(fragment)]]
-fn frag_main2() -> tint_symbol_2 {
-  return tint_symbol_2(1.0, 3u);
+fn frag_main2_inner() -> f32 {
+  return 1.0;
 }
 
-struct tint_symbol_3 {
+[[stage(fragment)]]
+fn frag_main2() -> tint_symbol_1 {
+  let inner_result_1 = frag_main2_inner();
+  var wrapper_result_1 : tint_symbol_1;
+  wrapper_result_1.value = inner_result_1;
+  wrapper_result_1.fixed_sample_mask = 3u;
+  return wrapper_result_1;
+}
+
+struct tint_symbol_2 {
   [[builtin(position)]]
   value : vec4<f32>;
 };
 
+fn vert_main1_inner() -> vec4<f32> {
+  return vec4<f32>();
+}
+
 [[stage(vertex)]]
-fn vert_main1() -> tint_symbol_3 {
-  return tint_symbol_3(vec4<f32>());
+fn vert_main1() -> tint_symbol_2 {
+  let inner_result_2 = vert_main1_inner();
+  var wrapper_result_2 : tint_symbol_2;
+  wrapper_result_2.value = inner_result_2;
+  return wrapper_result_2;
 }
 
 [[stage(compute), workgroup_size(1)]]
@@ -1200,6 +1373,57 @@ fn comp_main1() {
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(CanonicalizeEntryPointIOTest, FixedSampleMask_AvoidNameClash) {
+  auto* src = R"(
+struct FragOut {
+  [[location(0)]] fixed_sample_mask : vec4<f32>;
+  [[location(1)]] fixed_sample_mask_1 : vec4<f32>;
+};
+
+[[stage(fragment)]]
+fn frag_main() -> FragOut {
+  return FragOut();
+}
+)";
+
+  auto* expect = R"(
+struct FragOut {
+  fixed_sample_mask : vec4<f32>;
+  fixed_sample_mask_1 : vec4<f32>;
+};
+
+struct tint_symbol {
+  [[location(0)]]
+  fixed_sample_mask : vec4<f32>;
+  [[location(1)]]
+  fixed_sample_mask_1 : vec4<f32>;
+  [[builtin(sample_mask)]]
+  fixed_sample_mask_2 : u32;
+};
+
+fn frag_main_inner() -> FragOut {
+  return FragOut();
+}
+
+[[stage(fragment)]]
+fn frag_main() -> tint_symbol {
+  let inner_result = frag_main_inner();
+  var wrapper_result : tint_symbol;
+  wrapper_result.fixed_sample_mask = inner_result.fixed_sample_mask;
+  wrapper_result.fixed_sample_mask_1 = inner_result.fixed_sample_mask_1;
+  wrapper_result.fixed_sample_mask_2 = 3u;
+  return wrapper_result;
+}
+)";
+
+  DataMap data;
+  data.Add<CanonicalizeEntryPointIO::Config>(
+      CanonicalizeEntryPointIO::BuiltinStyle::kParameter, 0x03);
+  auto got = Run<CanonicalizeEntryPointIO>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint

src/transform/msl_test.cc

@@ -34,15 +34,19 @@ fn main() {
 )";
 
   auto* expect = R"(
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
-  [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_1 : f32;
-  [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_2 : f32;
+fn main_inner(local_invocation_index : u32, tint_symbol : ptr<workgroup, f32>, tint_symbol_1 : ptr<private, f32>) {
   {
-    tint_symbol_1 = f32();
+    *(tint_symbol) = f32();
   }
   workgroupBarrier();
-  tint_symbol_1 = tint_symbol_2;
+  *(tint_symbol) = *(tint_symbol_1);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
+  [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_2 : f32;
+  [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_3 : f32;
+  main_inner(local_invocation_index, &(tint_symbol_2), &(tint_symbol_3));
 }
 )";
 
@@ -80,26 +84,30 @@ fn main() {
 fn no_uses() {
 }
 
-fn bar(a : f32, b : f32, tint_symbol_1 : ptr<private, f32>, tint_symbol_2 : ptr<workgroup, f32>) {
-  *(tint_symbol_1) = a;
-  *(tint_symbol_2) = b;
+fn bar(a : f32, b : f32, tint_symbol : ptr<private, f32>, tint_symbol_1 : ptr<workgroup, f32>) {
+  *(tint_symbol) = a;
+  *(tint_symbol_1) = b;
 }
 
-fn foo(a : f32, tint_symbol_3 : ptr<private, f32>, tint_symbol_4 : ptr<workgroup, f32>) {
+fn foo(a : f32, tint_symbol_2 : ptr<private, f32>, tint_symbol_3 : ptr<workgroup, f32>) {
   let b : f32 = 2.0;
-  bar(a, b, tint_symbol_3, tint_symbol_4);
+  bar(a, b, tint_symbol_2, tint_symbol_3);
   no_uses();
 }
 
+fn main_inner(local_invocation_index : u32, tint_symbol_4 : ptr<workgroup, f32>, tint_symbol_5 : ptr<private, f32>) {
+  {
+    *(tint_symbol_4) = f32();
+  }
+  workgroupBarrier();
+  foo(1.0, tint_symbol_5, tint_symbol_4);
+}
+
 [[stage(compute), workgroup_size(1)]]
 fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
-  [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_5 : f32;
-  [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_6 : f32;
-  {
-    tint_symbol_5 = f32();
-  }
-  workgroupBarrier();
-  foo(1.0, &(tint_symbol_6), &(tint_symbol_5));
+  [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_6 : f32;
+  [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_7 : f32;
+  main_inner(local_invocation_index, &(tint_symbol_6), &(tint_symbol_7));
 }
 )";
 
@@ -148,16 +156,20 @@ fn main() {
 )";
 
   auto* expect = R"(
-[[stage(compute), workgroup_size(1)]]
-fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
-  [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_1 : f32;
-  [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_2 : f32;
+fn main_inner(local_invocation_index : u32, tint_symbol : ptr<workgroup, f32>, tint_symbol_1 : ptr<private, f32>) {
   {
-    tint_symbol_1 = f32();
+    *(tint_symbol) = f32();
   }
   workgroupBarrier();
-  let x : f32 = (tint_symbol_2 + tint_symbol_1);
-  tint_symbol_2 = x;
+  let x : f32 = (*(tint_symbol_1) + *(tint_symbol));
+  *(tint_symbol_1) = x;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] local_invocation_index : u32) {
+  [[internal(disable_validation__ignore_storage_class)]] var<workgroup> tint_symbol_2 : f32;
+  [[internal(disable_validation__ignore_storage_class)]] var<private> tint_symbol_3 : f32;
+  main_inner(local_invocation_index, &(tint_symbol_2), &(tint_symbol_3));
 }
 )";
 

src/writer/hlsl/generator_impl_function_test.cc

@@ -130,10 +130,15 @@ struct tint_symbol_2 {
   float value : SV_Target1;
 };
 
+float frag_main_inner(float foo) {
+  return foo;
+}
+
 tint_symbol_2 frag_main(tint_symbol_1 tint_symbol) {
-  const float foo = tint_symbol.foo;
-  const tint_symbol_2 tint_symbol_3 = {foo};
-  return tint_symbol_3;
+  const float inner_result = frag_main_inner(tint_symbol.foo);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 )");
 }
@@ -160,10 +165,15 @@ struct tint_symbol_2 {
   float value : SV_Depth;
 };
 
+float frag_main_inner(float4 coord) {
+  return coord.x;
+}
+
 tint_symbol_2 frag_main(tint_symbol_1 tint_symbol) {
-  const float4 coord = tint_symbol.coord;
-  const tint_symbol_2 tint_symbol_3 = {coord.x};
-  return tint_symbol_3;
+  const float inner_result = frag_main_inner(tint_symbol.coord);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 )");
 }
@@ -218,23 +228,35 @@ struct tint_symbol {
   float4 pos : SV_Position;
 };
 
-tint_symbol vert_main() {
-  const Interface tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f), 0.5f, 0.25f};
-  const tint_symbol tint_symbol_4 = {tint_symbol_1.col1, tint_symbol_1.col2, tint_symbol_1.pos};
-  return tint_symbol_4;
+Interface vert_main_inner() {
+  const Interface tint_symbol_3 = {float4(0.0f, 0.0f, 0.0f, 0.0f), 0.5f, 0.25f};
+  return tint_symbol_3;
 }
 
-struct tint_symbol_3 {
+tint_symbol vert_main() {
+  const Interface inner_result = vert_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.pos = inner_result.pos;
+  wrapper_result.col1 = inner_result.col1;
+  wrapper_result.col2 = inner_result.col2;
+  return wrapper_result;
+}
+
+struct tint_symbol_2 {
   float col1 : TEXCOORD1;
   float col2 : TEXCOORD2;
   float4 pos : SV_Position;
 };
 
-void frag_main(tint_symbol_3 tint_symbol_2) {
-  const Interface inputs = {tint_symbol_2.pos, tint_symbol_2.col1, tint_symbol_2.col2};
+void frag_main_inner(Interface inputs) {
   const float r = inputs.col1;
   const float g = inputs.col2;
   const float4 p = inputs.pos;
+}
+
+void frag_main(tint_symbol_2 tint_symbol_1) {
+  const Interface tint_symbol_4 = {tint_symbol_1.pos, tint_symbol_1.col1, tint_symbol_1.col2};
+  frag_main_inner(tint_symbol_4);
   return;
 }
 )");
@@ -278,28 +300,38 @@ TEST_F(HlslGeneratorImplTest_Function,
 };
 
 VertexOutput foo(float x) {
-  const VertexOutput tint_symbol_4 = {float4(x, x, x, 1.0f)};
-  return tint_symbol_4;
+  const VertexOutput tint_symbol_2 = {float4(x, x, x, 1.0f)};
+  return tint_symbol_2;
 }
 
 struct tint_symbol {
   float4 pos : SV_Position;
 };
 
-tint_symbol vert_main1() {
-  const VertexOutput tint_symbol_1 = foo(0.5f);
-  const tint_symbol tint_symbol_5 = {tint_symbol_1.pos};
-  return tint_symbol_5;
+VertexOutput vert_main1_inner() {
+  return foo(0.5f);
 }
 
-struct tint_symbol_2 {
+tint_symbol vert_main1() {
+  const VertexOutput inner_result = vert_main1_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.pos = inner_result.pos;
+  return wrapper_result;
+}
+
+struct tint_symbol_1 {
   float4 pos : SV_Position;
 };
 
-tint_symbol_2 vert_main2() {
-  const VertexOutput tint_symbol_3 = foo(0.25f);
-  const tint_symbol_2 tint_symbol_6 = {tint_symbol_3.pos};
-  return tint_symbol_6;
+VertexOutput vert_main2_inner() {
+  return foo(0.25f);
+}
+
+tint_symbol_1 vert_main2() {
+  const VertexOutput inner_result_1 = vert_main2_inner();
+  tint_symbol_1 wrapper_result_1 = (tint_symbol_1)0;
+  wrapper_result_1.pos = inner_result_1.pos;
+  return wrapper_result_1;
 }
 )");
 }

src/writer/msl/generator_impl_function_test.cc

@@ -111,10 +111,15 @@ struct tint_symbol_2 {
   float value [[color(1)]];
 };
 
+float frag_main_inner(float foo) {
+  return foo;
+}
+
 fragment tint_symbol_2 frag_main(tint_symbol_1 tint_symbol [[stage_in]]) {
-  float const foo = tint_symbol.foo;
-  tint_symbol_2 const tint_symbol_3 = {.value=foo};
-  return tint_symbol_3;
+  float const inner_result = frag_main_inner(tint_symbol.foo);
+  tint_symbol_2 wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 )");
@@ -137,13 +142,19 @@ TEST_F(MslGeneratorImplTest, Emit_Decoration_EntryPoint_WithInOut_Builtins) {
   EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
 
 using namespace metal;
-struct tint_symbol_1 {
+struct tint_symbol {
   float value [[depth(any)]];
 };
 
-fragment tint_symbol_1 frag_main(float4 coord [[position]]) {
-  tint_symbol_1 const tint_symbol_2 = {.value=coord.x};
-  return tint_symbol_2;
+float frag_main_inner(float4 coord) {
+  return coord.x;
+}
+
+fragment tint_symbol frag_main(float4 coord [[position]]) {
+  float const inner_result = frag_main_inner(coord);
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 )");
@@ -201,21 +212,33 @@ struct tint_symbol {
   float col2 [[user(locn2)]];
   float4 pos [[position]];
 };
-struct tint_symbol_4 {
+struct tint_symbol_2 {
   float col1 [[user(locn1)]];
   float col2 [[user(locn2)]];
 };
 
-vertex tint_symbol vert_main() {
-  Interface const tint_symbol_1 = {.col1=0.5f, .col2=0.25f, .pos=float4()};
-  tint_symbol const tint_symbol_5 = {.col1=tint_symbol_1.col1, .col2=tint_symbol_1.col2, .pos=tint_symbol_1.pos};
-  return tint_symbol_5;
+Interface vert_main_inner() {
+  Interface const tint_symbol_3 = {.col1=0.5f, .col2=0.25f, .pos=float4()};
+  return tint_symbol_3;
 }
 
-fragment void frag_main(float4 tint_symbol_3 [[position]], tint_symbol_4 tint_symbol_2 [[stage_in]]) {
-  Interface const colors = {.col1=tint_symbol_2.col1, .col2=tint_symbol_2.col2, .pos=tint_symbol_3};
+vertex tint_symbol vert_main() {
+  Interface const inner_result = vert_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.col1 = inner_result.col1;
+  wrapper_result.col2 = inner_result.col2;
+  wrapper_result.pos = inner_result.pos;
+  return wrapper_result;
+}
+
+void frag_main_inner(Interface colors) {
   float const r = colors.col1;
   float const g = colors.col2;
+}
+
+fragment void frag_main(float4 pos [[position]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  Interface const tint_symbol_4 = {.col1=tint_symbol_1.col1, .col2=tint_symbol_1.col2, .pos=pos};
+  frag_main_inner(tint_symbol_4);
   return;
 }
 
@@ -265,25 +288,35 @@ struct VertexOutput {
 struct tint_symbol {
   float4 pos [[position]];
 };
-struct tint_symbol_2 {
+struct tint_symbol_1 {
   float4 pos [[position]];
 };
 
 VertexOutput foo(float x) {
-  VertexOutput const tint_symbol_4 = {.pos=float4(x, x, x, 1.0f)};
-  return tint_symbol_4;
+  VertexOutput const tint_symbol_2 = {.pos=float4(x, x, x, 1.0f)};
+  return tint_symbol_2;
+}
+
+VertexOutput vert_main1_inner() {
+  return foo(0.5f);
 }
 
 vertex tint_symbol vert_main1() {
-  VertexOutput const tint_symbol_1 = foo(0.5f);
-  tint_symbol const tint_symbol_5 = {.pos=tint_symbol_1.pos};
-  return tint_symbol_5;
+  VertexOutput const inner_result = vert_main1_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.pos = inner_result.pos;
+  return wrapper_result;
 }
 
-vertex tint_symbol_2 vert_main2() {
-  VertexOutput const tint_symbol_3 = foo(0.25f);
-  tint_symbol_2 const tint_symbol_6 = {.pos=tint_symbol_3.pos};
-  return tint_symbol_6;
+VertexOutput vert_main2_inner() {
+  return foo(0.25f);
+}
+
+vertex tint_symbol_1 vert_main2() {
+  VertexOutput const inner_result_1 = vert_main2_inner();
+  tint_symbol_1 wrapper_result_1 = {};
+  wrapper_result_1.pos = inner_result_1.pos;
+  return wrapper_result_1;
 }
 
 )");

test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl

@@ -23,9 +23,7 @@ tint_symbol_11_ret tint_symbol_11(ByteAddressBuffer buffer, uint offset) {
   return arr_1;
 }
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint idx = tint_symbol.idx;
+void main_inner(uint idx) {
   const int3 a = asint(s.Load3((176u * idx)));
   const int b = asint(s.Load(((176u * idx) + 12u)));
   const uint3 c = s.Load3(((176u * idx) + 16u));
@@ -35,5 +33,10 @@ void main(tint_symbol_1 tint_symbol) {
   const float2x3 g = tint_symbol_8(s, ((176u * idx) + 48u));
   const float3x2 h = tint_symbol_9(s, ((176u * idx) + 80u));
   const int4 i[4] = tint_symbol_11(s, ((176u * idx) + 112u));
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.idx);
   return;
 }

test/buffer/storage/dynamic_index/read.wgsl.expected.msl

@@ -20,7 +20,7 @@ struct S {
   /* 0x0000 */ Inner arr[1];
 };
 
-kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S& s [[buffer(0)]]) {
+void tint_symbol_inner(const device S& s, uint idx) {
   int3 const a = s.arr[idx].a;
   int const b = s.arr[idx].b;
   uint3 const c = s.arr[idx].c;
@@ -30,6 +30,10 @@ kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S
   float2x3 const g = s.arr[idx].g;
   float3x2 const h = s.arr[idx].h;
   tint_array_wrapper const i = s.arr[idx].i;
+}
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S& s [[buffer(0)]]) {
+  tint_symbol_inner(s, idx);
   return;
 }
 

test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl

@@ -24,9 +24,7 @@ void tint_symbol_11(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
   }
 }
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint idx = tint_symbol.idx;
+void main_inner(uint idx) {
   s.Store3((176u * idx), asuint(int3(0, 0, 0)));
   s.Store(((176u * idx) + 12u), asuint(0));
   s.Store3(((176u * idx) + 16u), asuint(uint3(0u, 0u, 0u)));
@@ -37,5 +35,10 @@ void main(tint_symbol_1 tint_symbol) {
   tint_symbol_9(s, ((176u * idx) + 80u), float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
   const int4 tint_symbol_13[4] = (int4[4])0;
   tint_symbol_11(s, ((176u * idx) + 112u), tint_symbol_13);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.idx);
   return;
 }

test/buffer/storage/dynamic_index/write.wgsl.expected.msl

@@ -20,7 +20,7 @@ struct S {
   /* 0x0000 */ Inner arr[1];
 };
 
-kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& s [[buffer(0)]]) {
+void tint_symbol_inner(device S& s, uint idx) {
   s.arr[idx].a = int3();
   s.arr[idx].b = int();
   s.arr[idx].c = uint3();
@@ -29,8 +29,12 @@ kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& s [[
   s.arr[idx].f = float();
   s.arr[idx].g = float2x3();
   s.arr[idx].h = float3x2();
-  tint_array_wrapper const tint_symbol_2 = {.arr={}};
-  s.arr[idx].i = tint_symbol_2;
+  tint_array_wrapper const tint_symbol_1 = {.arr={}};
+  s.arr[idx].i = tint_symbol_1;
+}
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& s [[buffer(0)]]) {
+  tint_symbol_inner(s, idx);
   return;
 }
 

test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl

@@ -12,9 +12,7 @@ float2x3 tint_symbol_9(uint4 buffer[96], uint offset) {
   return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
 }
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint idx = tint_symbol.idx;
+void main_inner(uint idx) {
   const uint scalar_offset_2 = ((192u * idx)) / 4;
   const int3 a = asint(s[scalar_offset_2 / 4].xyz);
   const uint scalar_offset_3 = (((192u * idx) + 12u)) / 4;
@@ -34,5 +32,10 @@ void main(tint_symbol_1 tint_symbol) {
   uint4 ubo_load_1 = s[scalar_offset_9 / 4];
   const int2 h = asint(((scalar_offset_9 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
   const float2x3 i = tint_symbol_9(s, ((192u * idx) + 64u));
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.idx);
   return;
 }

test/buffer/uniform/dynamic_index/read.wgsl.expected.msl

@@ -25,7 +25,7 @@ struct S {
   /* 0x0000 */ tint_array_wrapper_1 arr;
 };
 
-kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], constant S& s [[buffer(0)]]) {
+void tint_symbol_inner(constant S& s, uint idx) {
   int3 const a = s.arr.arr[idx].a;
   int const b = s.arr.arr[idx].b;
   uint3 const c = s.arr.arr[idx].c;
@@ -35,6 +35,10 @@ kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], constant S& s
   int2 const g = s.arr.arr[idx].g;
   int2 const h = s.arr.arr[idx].h;
   float2x3 const i = s.arr.arr[idx].i;
+}
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], constant S& s [[buffer(0)]]) {
+  tint_symbol_inner(s, idx);
   return;
 }
 

test/bug/dawn/947.wgsl.expected.hlsl

@@ -14,8 +14,7 @@ struct tint_symbol_2 {
   float4 position : SV_Position;
 };
 
-tint_symbol_2 vs_main(tint_symbol_1 tint_symbol) {
-  const uint VertexIndex = tint_symbol.VertexIndex;
+VertexOutputs vs_main_inner(uint VertexIndex) {
   float2 texcoord[3] = {float2(-0.5f, 0.0f), float2(1.5f, 0.0f), float2(0.5f, 2.0f)};
   VertexOutputs output = (VertexOutputs)0;
   output.position = float4(((texcoord[VertexIndex] * 2.0f) - float2(1.0f, 1.0f)), 0.0f, 1.0f);
@@ -25,8 +24,15 @@ tint_symbol_2 vs_main(tint_symbol_1 tint_symbol) {
   } else {
     output.texcoords = ((((texcoord[VertexIndex] * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f)) * asfloat(uniforms[0].xy)) + asfloat(uniforms[0].zw));
   }
-  const tint_symbol_2 tint_symbol_8 = {output.texcoords, output.position};
-  return tint_symbol_8;
+  return output;
+}
+
+tint_symbol_2 vs_main(tint_symbol_1 tint_symbol) {
+  const VertexOutputs inner_result = vs_main_inner(tint_symbol.VertexIndex);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.texcoords = inner_result.texcoords;
+  wrapper_result.position = inner_result.position;
+  return wrapper_result;
 }
 
 SamplerState mySampler : register(s1, space0);
@@ -39,13 +45,18 @@ struct tint_symbol_5 {
   float4 value : SV_Target0;
 };
 
-tint_symbol_5 fs_main(tint_symbol_4 tint_symbol_3) {
-  const float2 texcoord = tint_symbol_3.texcoord;
+float4 fs_main_inner(float2 texcoord) {
   float2 clampedTexcoord = clamp(texcoord, float2(0.0f, 0.0f), float2(1.0f, 1.0f));
   if (!(all((clampedTexcoord == texcoord)))) {
     discard;
   }
   float4 srcColor = myTexture.Sample(mySampler, texcoord);
-  const tint_symbol_5 tint_symbol_9 = {srcColor};
-  return tint_symbol_9;
+  return srcColor;
+}
+
+tint_symbol_5 fs_main(tint_symbol_4 tint_symbol_3) {
+  const float4 inner_result_1 = fs_main_inner(tint_symbol_3.texcoord);
+  tint_symbol_5 wrapper_result_1 = (tint_symbol_5)0;
+  wrapper_result_1.value = inner_result_1;
+  return wrapper_result_1;
 }

test/bug/dawn/947.wgsl.expected.msl

@@ -9,21 +9,21 @@ struct VertexOutputs {
   float2 texcoords;
   float4 position;
 };
-struct tint_symbol_1 {
+struct tint_symbol {
   float2 texcoords [[user(locn0)]];
   float4 position [[position]];
 };
 struct tint_array_wrapper {
   float2 arr[3];
 };
-struct tint_symbol_3 {
+struct tint_symbol_2 {
   float2 texcoord [[user(locn0)]];
 };
-struct tint_symbol_4 {
+struct tint_symbol_3 {
   float4 value [[color(0)]];
 };
 
-vertex tint_symbol_1 vs_main(uint VertexIndex [[vertex_id]], constant Uniforms& uniforms [[buffer(0)]]) {
+VertexOutputs vs_main_inner(constant Uniforms& uniforms, uint VertexIndex) {
   tint_array_wrapper texcoord = {.arr={float2(-0.5f, 0.0f), float2(1.5f, 0.0f), float2(0.5f, 2.0f)}};
   VertexOutputs output = {};
   output.position = float4(((texcoord.arr[VertexIndex] * 2.0f) - float2(1.0f, 1.0f)), 0.0f, 1.0f);
@@ -33,18 +33,30 @@ vertex tint_symbol_1 vs_main(uint VertexIndex [[vertex_id]], constant Uniforms&
   } else {
     output.texcoords = ((((texcoord.arr[VertexIndex] * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f)) * uniforms.u_scale) + uniforms.u_offset);
   }
-  tint_symbol_1 const tint_symbol_5 = {.texcoords=output.texcoords, .position=output.position};
-  return tint_symbol_5;
+  return output;
 }
 
-fragment tint_symbol_4 fs_main(texture2d<float, access::sample> tint_symbol_7 [[texture(2)]], sampler tint_symbol_8 [[sampler(1)]], tint_symbol_3 tint_symbol_2 [[stage_in]]) {
-  float2 const texcoord = tint_symbol_2.texcoord;
+vertex tint_symbol vs_main(uint VertexIndex [[vertex_id]], constant Uniforms& uniforms [[buffer(0)]]) {
+  VertexOutputs const inner_result = vs_main_inner(uniforms, VertexIndex);
+  tint_symbol wrapper_result = {};
+  wrapper_result.texcoords = inner_result.texcoords;
+  wrapper_result.position = inner_result.position;
+  return wrapper_result;
+}
+
+float4 fs_main_inner(float2 texcoord, texture2d<float, access::sample> tint_symbol_4, sampler tint_symbol_5) {
   float2 clampedTexcoord = clamp(texcoord, float2(0.0f, 0.0f), float2(1.0f, 1.0f));
   if (!(all((clampedTexcoord == texcoord)))) {
     discard_fragment();
   }
-  float4 srcColor = tint_symbol_7.sample(tint_symbol_8, texcoord);
-  tint_symbol_4 const tint_symbol_6 = {.value=srcColor};
-  return tint_symbol_6;
+  float4 srcColor = tint_symbol_4.sample(tint_symbol_5, texcoord);
+  return srcColor;
+}
+
+fragment tint_symbol_3 fs_main(texture2d<float, access::sample> tint_symbol_6 [[texture(2)]], sampler tint_symbol_7 [[sampler(1)]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  float4 const inner_result_1 = fs_main_inner(tint_symbol_1.texcoord, tint_symbol_6, tint_symbol_7);
+  tint_symbol_3 wrapper_result_1 = {};
+  wrapper_result_1.value = inner_result_1;
+  return wrapper_result_1;
 }
 

test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.hlsl

@@ -13,9 +13,7 @@ struct tint_symbol_2 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void f(tint_symbol_2 tint_symbol_1) {
-  const uint local_invocation_index = tint_symbol_1.local_invocation_index;
+void f_inner(uint local_invocation_index) {
   {
     for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
       const uint i = idx;
@@ -24,5 +22,10 @@ void f(tint_symbol_2 tint_symbol_1) {
   }
   GroupMemoryBarrierWithGroupSync();
   result.Store(0u, asuint(s.data[asint(ubo[0].x)]));
+}
+
+[numthreads(1, 1, 1)]
+void f(tint_symbol_2 tint_symbol_1) {
+  f_inner(tint_symbol_1.local_invocation_index);
   return;
 }

test/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl

@@ -14,14 +14,18 @@ struct Result {
   /* 0x0000 */ int out;
 };
 
-kernel void f(uint local_invocation_index [[thread_index_in_threadgroup]], constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
-  threadgroup S tint_symbol_1;
+void f_inner(constant UBO& ubo, device Result& result, uint local_invocation_index, threadgroup S* const tint_symbol) {
   for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
     uint const i = idx;
-    tint_symbol_1.data.arr[i] = int();
+    (*(tint_symbol)).data.arr[i] = int();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  result.out = tint_symbol_1.data.arr[ubo.dynamic_idx];
+  result.out = (*(tint_symbol)).data.arr[ubo.dynamic_idx];
+}
+
+kernel void f(uint local_invocation_index [[thread_index_in_threadgroup]], constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  threadgroup S tint_symbol_1;
+  f_inner(ubo, result, local_invocation_index, &(tint_symbol_1));
   return;
 }
 

test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.hlsl

@@ -13,9 +13,7 @@ struct tint_symbol_2 {
   uint local_invocation_index : SV_GroupIndex;
 };
 
-[numthreads(1, 1, 1)]
-void f(tint_symbol_2 tint_symbol_1) {
-  const uint local_invocation_index = tint_symbol_1.local_invocation_index;
+void f_inner(uint local_invocation_index) {
   {
     for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
       const uint i = idx;
@@ -25,5 +23,10 @@ void f(tint_symbol_2 tint_symbol_1) {
   GroupMemoryBarrierWithGroupSync();
   s.data[asint(ubo[0].x)] = 1;
   result.Store(0u, asuint(s.data[3]));
+}
+
+[numthreads(1, 1, 1)]
+void f(tint_symbol_2 tint_symbol_1) {
+  f_inner(tint_symbol_1.local_invocation_index);
   return;
 }

test/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl

@@ -14,15 +14,19 @@ struct Result {
   /* 0x0000 */ int out;
 };
 
-kernel void f(uint local_invocation_index [[thread_index_in_threadgroup]], constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
-  threadgroup S tint_symbol_1;
+void f_inner(constant UBO& ubo, device Result& result, uint local_invocation_index, threadgroup S* const tint_symbol) {
   for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
     uint const i = idx;
-    tint_symbol_1.data.arr[i] = int();
+    (*(tint_symbol)).data.arr[i] = int();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  tint_symbol_1.data.arr[ubo.dynamic_idx] = 1;
-  result.out = tint_symbol_1.data.arr[3];
+  (*(tint_symbol)).data.arr[ubo.dynamic_idx] = 1;
+  result.out = (*(tint_symbol)).data.arr[3];
+}
+
+kernel void f(uint local_invocation_index [[thread_index_in_threadgroup]], constant UBO& ubo [[buffer(0)]], device Result& result [[buffer(1)]]) {
+  threadgroup S tint_symbol_1;
+  f_inner(ubo, result, local_invocation_index, &(tint_symbol_1));
   return;
 }
 

test/bug/tint/1046.wgsl.expected.hlsl

@@ -48,14 +48,20 @@ struct tint_symbol_2 {
   float4 color : SV_Target0;
 };
 
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const FragmentInput fragment = {tint_symbol.position, tint_symbol.view_position, tint_symbol.normal, tint_symbol.uv, tint_symbol.color};
+FragmentOutput main_inner(FragmentInput fragment) {
   FragmentOutput output = (FragmentOutput)0;
   output.color = float4(1.0f, 0.0f, 0.0f, 1.0f);
   uniforms;
   mySampler;
   myTexture;
   pointLights;
-  const tint_symbol_2 tint_symbol_5 = {output.color};
-  return tint_symbol_5;
+  return output;
+}
+
+tint_symbol_2 main(tint_symbol_1 tint_symbol) {
+  const FragmentInput tint_symbol_5 = {tint_symbol.position, tint_symbol.view_position, tint_symbol.normal, tint_symbol.uv, tint_symbol.color};
+  const FragmentOutput inner_result = main_inner(tint_symbol_5);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
 }

test/bug/tint/1046.wgsl.expected.msl

@@ -25,17 +25,17 @@ struct FragmentInput {
 struct FragmentOutput {
   float4 color;
 };
-struct tint_symbol_4 {
+struct tint_symbol_3 {
   float4 view_position [[user(locn0)]];
   float4 normal [[user(locn1)]];
   float2 uv [[user(locn2)]];
   float4 color [[user(locn3)]];
 };
-struct tint_symbol_5 {
+struct tint_symbol_4 {
   float4 color [[color(0)]];
 };
 
-float4 getColor(constant Uniforms& uniforms, FragmentInput tint_symbol, texture2d<float, access::sample> tint_symbol_7, sampler tint_symbol_8) {
+float4 getColor(constant Uniforms& uniforms, FragmentInput tint_symbol, texture2d<float, access::sample> tint_symbol_6, sampler tint_symbol_7) {
   float4 color = 0.0f;
   if ((uniforms.color_source == 0u)) {
     color = tint_symbol.color;
@@ -48,7 +48,7 @@ float4 getColor(constant Uniforms& uniforms, FragmentInput tint_symbol, texture2
         color = uniforms.color;
       } else {
         if ((uniforms.color_source == 3u)) {
-          color = tint_symbol_7.sample(tint_symbol_8, tint_symbol.uv);
+          color = tint_symbol_6.sample(tint_symbol_7, tint_symbol.uv);
         }
       }
     }
@@ -56,15 +56,21 @@ float4 getColor(constant Uniforms& uniforms, FragmentInput tint_symbol, texture2
   return color;
 }
 
-fragment tint_symbol_5 tint_symbol_1(sampler tint_symbol_9 [[sampler(2)]], texture2d<float, access::sample> tint_symbol_10 [[texture(3)]], float4 tint_symbol_3 [[position]], tint_symbol_4 tint_symbol_2 [[stage_in]], constant Uniforms& uniforms [[buffer(0)]], const device PointLights& pointLights [[buffer(1)]]) {
-  FragmentInput const tint_symbol = {.position=tint_symbol_3, .view_position=tint_symbol_2.view_position, .normal=tint_symbol_2.normal, .uv=tint_symbol_2.uv, .color=tint_symbol_2.color};
+FragmentOutput tint_symbol_1_inner(constant Uniforms& uniforms, const device PointLights& pointLights, FragmentInput tint_symbol, sampler tint_symbol_8, texture2d<float, access::sample> tint_symbol_9) {
   FragmentOutput output = {};
   output.color = float4(1.0f, 0.0f, 0.0f, 1.0f);
   (void) uniforms;
+  (void) tint_symbol_8;
   (void) tint_symbol_9;
-  (void) tint_symbol_10;
   (void) pointLights;
-  tint_symbol_5 const tint_symbol_6 = {.color=output.color};
-  return tint_symbol_6;
+  return output;
+}
+
+fragment tint_symbol_4 tint_symbol_1(sampler tint_symbol_10 [[sampler(2)]], texture2d<float, access::sample> tint_symbol_11 [[texture(3)]], float4 position [[position]], tint_symbol_3 tint_symbol_2 [[stage_in]], constant Uniforms& uniforms [[buffer(0)]], const device PointLights& pointLights [[buffer(1)]]) {
+  FragmentInput const tint_symbol_5 = {.position=position, .view_position=tint_symbol_2.view_position, .normal=tint_symbol_2.normal, .uv=tint_symbol_2.uv, .color=tint_symbol_2.color};
+  FragmentOutput const inner_result = tint_symbol_1_inner(uniforms, pointLights, tint_symbol_5, tint_symbol_10, tint_symbol_11);
+  tint_symbol_4 wrapper_result = {};
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
 }
 

test/bug/tint/1076.wgsl

@@ -0,0 +1,12 @@
+struct FragIn {
+  [[location(0)]] a : f32;
+  [[builtin(sample_mask)]] mask : u32;
+};
+
+[[stage(fragment)]]
+fn main(in : FragIn, [[location(1)]] b : f32) -> FragIn {
+  if (in.mask == 0u) {
+    return in;
+  }
+  return FragIn(b, 1u);
+}

test/bug/tint/1076.wgsl.expected.hlsl

@@ -0,0 +1,30 @@
+struct FragIn {
+  float a;
+  uint mask;
+};
+struct tint_symbol_2 {
+  float a : TEXCOORD0;
+  float b : TEXCOORD1;
+  uint mask : SV_Coverage;
+};
+struct tint_symbol_3 {
+  float a : SV_Target0;
+  uint mask : SV_Coverage;
+};
+
+FragIn main_inner(FragIn tint_symbol, float b) {
+  if ((tint_symbol.mask == 0u)) {
+    return tint_symbol;
+  }
+  const FragIn tint_symbol_4 = {b, 1u};
+  return tint_symbol_4;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+  const FragIn tint_symbol_5 = {tint_symbol_1.a, tint_symbol_1.mask};
+  const FragIn inner_result = main_inner(tint_symbol_5, tint_symbol_1.b);
+  tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+  wrapper_result.a = inner_result.a;
+  wrapper_result.mask = inner_result.mask;
+  return wrapper_result;
+}

test/bug/tint/1076.wgsl.expected.msl

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct FragIn {
+  float a;
+  uint mask;
+};
+struct tint_symbol_2 {
+  float a [[user(locn0)]];
+  float b [[user(locn1)]];
+};
+struct tint_symbol_3 {
+  float a [[color(0)]];
+  uint mask [[sample_mask]];
+};
+
+FragIn tint_symbol_inner(FragIn in, float b) {
+  if ((in.mask == 0u)) {
+    return in;
+  }
+  FragIn const tint_symbol_4 = {.a=b, .mask=1u};
+  return tint_symbol_4;
+}
+
+fragment tint_symbol_3 tint_symbol(uint mask [[sample_mask]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  FragIn const tint_symbol_5 = {.a=tint_symbol_1.a, .mask=mask};
+  FragIn const inner_result = tint_symbol_inner(tint_symbol_5, tint_symbol_1.b);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.a = inner_result.a;
+  wrapper_result.mask = inner_result.mask;
+  return wrapper_result;
+}
+

test/bug/tint/1076.wgsl.expected.spvasm

@@ -0,0 +1,82 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 46
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main" %tint_symbol %tint_symbol_1 %tint_symbol_5 %tint_symbol_6 %tint_symbol_3
+               OpExecutionMode %main OriginUpperLeft
+               OpName %tint_symbol "tint_symbol"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %tint_symbol_3 "tint_symbol_3"
+               OpName %tint_symbol_5 "tint_symbol_5"
+               OpName %tint_symbol_6 "tint_symbol_6"
+               OpName %FragIn "FragIn"
+               OpMemberName %FragIn 0 "a"
+               OpMemberName %FragIn 1 "mask"
+               OpName %tint_symbol_7 "tint_symbol_7"
+               OpName %tint_symbol_4 "tint_symbol_4"
+               OpName %main "main"
+               OpDecorate %tint_symbol Location 0
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+               OpDecorate %tint_symbol_1 BuiltIn SampleMask
+               OpDecorate %tint_symbol_3 Location 1
+               OpDecorate %tint_symbol_5 Location 0
+               OpDecorate %tint_symbol_6 BuiltIn SampleMask
+               OpMemberDecorate %FragIn 0 Offset 0
+               OpMemberDecorate %FragIn 1 Offset 4
+      %float = OpTypeFloat 32
+%_ptr_Input_float = OpTypePointer Input %float
+%tint_symbol = OpVariable %_ptr_Input_float Input
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+%_ptr_Input__arr_uint_uint_1 = OpTypePointer Input %_arr_uint_uint_1
+%tint_symbol_1 = OpVariable %_ptr_Input__arr_uint_uint_1 Input
+%tint_symbol_3 = OpVariable %_ptr_Input_float Input
+%_ptr_Output_float = OpTypePointer Output %float
+         %12 = OpConstantNull %float
+%tint_symbol_5 = OpVariable %_ptr_Output_float Output %12
+%_ptr_Output__arr_uint_uint_1 = OpTypePointer Output %_arr_uint_uint_1
+         %15 = OpConstantNull %_arr_uint_uint_1
+%tint_symbol_6 = OpVariable %_ptr_Output__arr_uint_uint_1 Output %15
+       %void = OpTypeVoid
+     %FragIn = OpTypeStruct %float %uint
+         %16 = OpTypeFunction %void %FragIn
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Output_uint = OpTypePointer Output %uint
+         %28 = OpTypeFunction %void
+%_ptr_Input_uint = OpTypePointer Input %uint
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+%tint_symbol_7 = OpFunction %void None %16
+%tint_symbol_4 = OpFunctionParameter %FragIn
+         %21 = OpLabel
+         %22 = OpCompositeExtract %float %tint_symbol_4 0
+               OpStore %tint_symbol_5 %22
+         %26 = OpAccessChain %_ptr_Output_uint %tint_symbol_6 %int_0
+         %27 = OpCompositeExtract %uint %tint_symbol_4 1
+               OpStore %26 %27
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %28
+         %30 = OpLabel
+         %31 = OpLoad %float %tint_symbol
+         %33 = OpAccessChain %_ptr_Input_uint %tint_symbol_1 %int_0
+         %34 = OpLoad %uint %33
+         %35 = OpCompositeConstruct %FragIn %31 %34
+         %36 = OpCompositeExtract %uint %35 1
+         %38 = OpIEqual %bool %36 %uint_0
+               OpSelectionMerge %40 None
+               OpBranchConditional %38 %41 %40
+         %41 = OpLabel
+         %42 = OpFunctionCall %void %tint_symbol_7 %35
+               OpReturn
+         %40 = OpLabel
+         %44 = OpLoad %float %tint_symbol_3
+         %45 = OpCompositeConstruct %FragIn %44 %uint_1
+         %43 = OpFunctionCall %void %tint_symbol_7 %45
+               OpReturn
+               OpFunctionEnd

test/bug/tint/1076.wgsl.expected.wgsl

@@ -0,0 +1,14 @@
+struct FragIn {
+  [[location(0)]]
+  a : f32;
+  [[builtin(sample_mask)]]
+  mask : u32;
+};
+
+[[stage(fragment)]]
+fn main(in : FragIn, [[location(1)]] b : f32) -> FragIn {
+  if ((in.mask == 0u)) {
+    return in;
+  }
+  return FragIn(b, 1u);
+}

test/bug/tint/292.wgsl.expected.hlsl

@@ -2,9 +2,15 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol main() {
+float4 main_inner() {
   float3 light = float3(1.200000048f, 1.0f, 2.0f);
   float3 negative_light = -(light);
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol main() {
+  const float4 inner_result = main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }

test/bug/tint/292.wgsl.expected.msl

@@ -5,10 +5,16 @@ struct tint_symbol_1 {
   float4 value [[position]];
 };
 
-vertex tint_symbol_1 tint_symbol() {
+float4 tint_symbol_inner() {
   float3 light = float3(1.200000048f, 1.0f, 2.0f);
   float3 negative_light = -(light);
-  tint_symbol_1 const tint_symbol_2 = {.value=float4()};
-  return tint_symbol_2;
+  return float4();
+}
+
+vertex tint_symbol_1 tint_symbol() {
+  float4 const inner_result = tint_symbol_inner();
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 

test/bug/tint/403.wgsl.expected.hlsl

@@ -28,8 +28,7 @@ float2x2 tint_symbol_5(uint4 buffer[1], uint offset) {
   return float2x2(asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
 }
 
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const uint gl_VertexIndex = tint_symbol.gl_VertexIndex;
+float4 main_inner(uint gl_VertexIndex) {
   float2 indexable[3] = (float2[3])0;
   const float2x2 x_23 = tint_symbol_3(x_20, 0u);
   const float2x2 x_28 = tint_symbol_5(x_26, 0u);
@@ -38,6 +37,12 @@ tint_symbol_2 main(tint_symbol_1 tint_symbol) {
   indexable = tint_symbol_7;
   const float2 x_51 = indexable[x_46];
   const float2 x_52 = mul(x_51, float2x2((x_23[0u] + x_28[0u]), (x_23[1u] + x_28[1u])));
-  const tint_symbol_2 tint_symbol_8 = {float4(x_52.x, x_52.y, 0.0f, 1.0f)};
-  return tint_symbol_8;
+  return float4(x_52.x, x_52.y, 0.0f, 1.0f);
+}
+
+tint_symbol_2 main(tint_symbol_1 tint_symbol) {
+  const float4 inner_result = main_inner(tint_symbol.gl_VertexIndex);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }

test/bug/tint/534.wgsl.expected.hlsl

@@ -17,9 +17,7 @@ struct tint_symbol_2 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
-  const uint3 GlobalInvocationID = tint_symbol_1.GlobalInvocationID;
+void main_inner(uint3 GlobalInvocationID) {
   int2 tint_tmp;
   src.GetDimensions(tint_tmp.x, tint_tmp.y);
   int2 size = tint_tmp;
@@ -49,5 +47,10 @@ void main(tint_symbol_2 tint_symbol_1) {
   } else {
     output.Store((4u * outputIndex), asuint(uint(0)));
   }
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_2 tint_symbol_1) {
+  main_inner(tint_symbol_1.GlobalInvocationID);
   return;
 }

test/bug/tint/534.wgsl.expected.msl

@@ -15,15 +15,15 @@ uint ConvertToFp16FloatValue(float fp32) {
   return 1u;
 }
 
-kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture(0)]], texture2d<float, access::sample> tint_symbol_3 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) {
-  int2 size = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
+void tint_symbol_inner(constant Uniforms& uniforms, device OutputBuf& output, uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_1, texture2d<float, access::sample> tint_symbol_2) {
+  int2 size = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
   int2 dstTexCoord = int2(GlobalInvocationID.xy);
   int2 srcTexCoord = dstTexCoord;
   if ((uniforms.dstTextureFlipY == 1u)) {
     srcTexCoord.y = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(size.y) - as_type<uint>(dstTexCoord.y)))) - as_type<uint>(1)));
   }
-  float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0);
-  float4 dstColor = tint_symbol_3.read(uint2(dstTexCoord), 0);
+  float4 srcColor = tint_symbol_1.read(uint2(srcTexCoord), 0);
+  float4 dstColor = tint_symbol_2.read(uint2(dstTexCoord), 0);
   bool success = true;
   uint4 srcColorBits = 0u;
   uint4 dstColorBits = uint4(dstColor);
@@ -37,6 +37,10 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture
   } else {
     output.result[outputIndex] = uint(0);
   }
+}
+
+kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_3 [[texture(0)]], texture2d<float, access::sample> tint_symbol_4 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) {
+  tint_symbol_inner(uniforms, output, GlobalInvocationID, tint_symbol_3, tint_symbol_4);
   return;
 }
 

test/bug/tint/744.wgsl.expected.hlsl

@@ -9,9 +9,7 @@ struct tint_symbol_1 {
   uint3 global_id : SV_DispatchThreadID;
 };
 
-[numthreads(2, 2, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint3 global_id = tint_symbol.global_id;
+void main_inner(uint3 global_id) {
   const uint2 resultCell = uint2(global_id.y, global_id.x);
   const uint dimInner = uniforms[0].y;
   const uint dimOutter = uniforms[1].y;
@@ -25,5 +23,10 @@ void main(tint_symbol_1 tint_symbol) {
   }
   const uint index = (resultCell.y + (resultCell.x * dimOutter));
   resultMatrix.Store((4u * index), asuint(result));
+}
+
+[numthreads(2, 2, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.global_id);
   return;
 }

test/bug/tint/744.wgsl.expected.msl

@@ -10,7 +10,7 @@ struct Matrix {
   /* 0x0000 */ uint numbers[1];
 };
 
-kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
+void tint_symbol_inner(constant Uniforms& uniforms, const device Matrix& firstMatrix, const device Matrix& secondMatrix, device Matrix& resultMatrix, uint3 global_id) {
   uint2 const resultCell = uint2(global_id.y, global_id.x);
   uint const dimInner = uniforms.aShape.y;
   uint const dimOutter = uniforms.outShape.y;
@@ -22,6 +22,10 @@ kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Un
   }
   uint const index = (resultCell.y + (resultCell.x * dimOutter));
   resultMatrix.numbers[index] = result;
+}
+
+kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
+  tint_symbol_inner(uniforms, firstMatrix, secondMatrix, resultMatrix, global_id);
   return;
 }
 

test/bug/tint/749.spvasm.expected.hlsl

@@ -57,18 +57,18 @@ void swap_i1_i1_(inout int i, inout int j) {
   const int x_34_save = x_33;
   const int x_35 = obj.numbers[x_34_save];
   const QuicksortObject x_943 = obj;
-  const int tint_symbol_5[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_6 = {tint_symbol_5};
-  obj = tint_symbol_6;
+  const int tint_symbol_4[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_5 = {tint_symbol_4};
+  obj = tint_symbol_5;
   obj = x_943;
   const float2 x_527 = float2(x_526.x, x_526.x);
   const int x_36_save = x_32;
   const float3 x_528 = float3(x_524.x, x_524.z, x_524.x);
   obj.numbers[x_36_save] = x_35;
   const QuicksortObject x_944 = obj;
-  const int tint_symbol_7[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_8 = {tint_symbol_7};
-  obj = tint_symbol_8;
+  const int tint_symbol_6[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_7 = {tint_symbol_6};
+  obj = tint_symbol_7;
   obj = x_944;
   const float3 x_529 = float3(x_526.y, x_526.z, x_526.x);
   const int x_945 = i;
@@ -91,9 +91,9 @@ void swap_i1_i1_(inout int i, inout int j) {
   obj.numbers[x_36_save] = 0;
   obj.numbers[x_36_save] = x_949;
   const QuicksortObject x_950 = obj;
-  const int tint_symbol_9[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_10 = {tint_symbol_9};
-  obj = tint_symbol_10;
+  const int tint_symbol_8[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_9 = {tint_symbol_8};
+  obj = tint_symbol_9;
   obj = x_950;
   const float3 x_532 = float3(x_528.x, x_528.y, x_528.x);
   const int x_951 = obj.numbers[x_34_save];
@@ -149,9 +149,9 @@ int performPartition_i1_i1_(inout int l, inout int h) {
   const float3 x_536 = float3(x_534.x, x_534.z, x_535.x);
   j_1 = 10;
   const QuicksortObject x_960 = obj;
-  const int tint_symbol_11[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_12 = {tint_symbol_11};
-  obj = tint_symbol_12;
+  const int tint_symbol_10[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_11 = {tint_symbol_10};
+  obj = tint_symbol_11;
   obj = x_960;
   while (true) {
     const int x_961 = pivot;
@@ -166,9 +166,9 @@ int performPartition_i1_i1_(inout int l, inout int h) {
     pivot = x_963;
     x_537 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z);
     const QuicksortObject x_964 = obj;
-    const int tint_symbol_13[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_14 = {tint_symbol_13};
-    obj = tint_symbol_14;
+    const int tint_symbol_12[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_13 = {tint_symbol_12};
+    obj = tint_symbol_13;
     obj = x_964;
     const int x_56 = h;
     const int x_965 = h;
@@ -202,9 +202,9 @@ int performPartition_i1_i1_(inout int l, inout int h) {
     param_1 = x_971;
     const int x_62 = obj.numbers[x_61_save];
     const QuicksortObject x_972 = obj;
-    const int tint_symbol_15[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_16 = {tint_symbol_15};
-    obj = tint_symbol_16;
+    const int tint_symbol_14[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_15 = {tint_symbol_14};
+    obj = tint_symbol_15;
     obj = x_972;
     const int x_63 = pivot;
     const float2 x_540 = float2(float3(1.0f, 2.0f, 3.0f).y, x_534.z);
@@ -262,9 +262,9 @@ int performPartition_i1_i1_(inout int l, inout int h) {
       param_1 = x_985;
     }
     const QuicksortObject x_986 = obj;
-    const int tint_symbol_17[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_18 = {tint_symbol_17};
-    obj = tint_symbol_18;
+    const int tint_symbol_16[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_17 = {tint_symbol_16};
+    obj = tint_symbol_17;
     obj = x_986;
     {
       const int x_987 = h;
@@ -297,9 +297,9 @@ int performPartition_i1_i1_(inout int l, inout int h) {
   obj.numbers[x_42_save] = x_993;
   const float2 x_549 = float2(x_534.x, x_534.y);
   const QuicksortObject x_994 = obj;
-  const int tint_symbol_19[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_20 = {tint_symbol_19};
-  obj = tint_symbol_20;
+  const int tint_symbol_18[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_19 = {tint_symbol_18};
+  obj = tint_symbol_19;
   obj = x_994;
   const int x_995 = h;
   h = 0;
@@ -367,8 +367,8 @@ void quicksort_() {
   param_5 = x_1007;
   h_1 = 9;
   const int x_1008[10] = stack;
-  const int tint_symbol_21[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  stack = tint_symbol_21;
+  const int tint_symbol_20[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  stack = tint_symbol_20;
   stack = x_1008;
   const float2 x_556 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).y);
   const int x_1009 = param_5;
@@ -401,15 +401,15 @@ void quicksort_() {
   param_4 = x_1015;
   const int x_95 = l_1;
   const QuicksortObject x_1016 = obj;
-  const int tint_symbol_22[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_23 = {tint_symbol_22};
-  obj = tint_symbol_23;
+  const int tint_symbol_21[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_22 = {tint_symbol_21};
+  obj = tint_symbol_22;
   obj = x_1016;
   const float3 x_560 = float3(x_559.y, x_559.x, x_557.x);
   const int x_96_save = x_94;
   const int x_1017[10] = stack;
-  const int tint_symbol_24[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  stack = tint_symbol_24;
+  const int tint_symbol_23[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  stack = tint_symbol_23;
   stack = x_1017;
   const float3 x_561 = float3(x_556.y, x_556.y, x_556.y);
   const int x_1018 = l_1;
@@ -459,13 +459,13 @@ void quicksort_() {
     h_1 = 0;
     h_1 = x_1028;
     const int x_1029[10] = stack;
-    const int tint_symbol_25[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    stack = tint_symbol_25;
+    const int tint_symbol_24[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    stack = tint_symbol_24;
     stack = x_1029;
     const int x_106 = top;
     const int x_1030[10] = stack;
-    const int tint_symbol_26[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    stack = tint_symbol_26;
+    const int tint_symbol_25[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    stack = tint_symbol_25;
     stack = x_1030;
     const float2 x_567 = float2(x_558.x, x_564.z);
     const int x_1031 = param_4;
@@ -476,9 +476,9 @@ void quicksort_() {
       break;
     }
     const QuicksortObject x_1032 = obj;
-    const int tint_symbol_27[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_28 = {tint_symbol_27};
-    obj = tint_symbol_28;
+    const int tint_symbol_26[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_27 = {tint_symbol_26};
+    obj = tint_symbol_27;
     obj = x_1032;
     const float3 x_568 = float3(x_559.y, x_559.x, x_563.y);
     const int x_1033 = param_4;
@@ -503,8 +503,8 @@ void quicksort_() {
     stack[x_96_save] = x_1037;
     const int x_111 = stack[x_110_save];
     const int x_1038[10] = stack;
-    const int tint_symbol_29[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    stack = tint_symbol_29;
+    const int tint_symbol_28[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    stack = tint_symbol_28;
     stack = x_1038;
     const float3 x_571 = float3(x_559.y, x_559.x, x_564.y);
     const int x_1039 = l_1;
@@ -512,8 +512,8 @@ void quicksort_() {
     l_1 = x_1039;
     h_1 = x_111;
     const int x_1040[10] = stack;
-    const int tint_symbol_30[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    stack = tint_symbol_30;
+    const int tint_symbol_29[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    stack = tint_symbol_29;
     stack = x_1040;
     const float2 x_572 = float2(x_562.y, x_561.y);
     const int x_1041 = p;
@@ -604,8 +604,8 @@ void quicksort_() {
       stack[x_100_save] = 0;
       stack[x_100_save] = x_1061;
       const int x_1062[10] = stack;
-      const int tint_symbol_31[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-      stack = tint_symbol_31;
+      const int tint_symbol_30[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+      stack = tint_symbol_30;
       stack = x_1062;
       const float2 x_584 = float2(x_569.z, x_569.y);
       const float3 x_585 = float3(x_580.y, x_577.x, x_577.x);
@@ -644,8 +644,8 @@ void quicksort_() {
       h_1 = x_1070;
       top = x_133;
       const int x_1071[10] = stack;
-      const int tint_symbol_32[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-      stack = tint_symbol_32;
+      const int tint_symbol_31[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+      stack = tint_symbol_31;
       stack = x_1071;
       const int x_134 = p;
       const float2 x_590 = float2(x_576.x, x_573.y);
@@ -670,9 +670,9 @@ void quicksort_() {
     stack[x_96_save] = x_1076;
     const float2 x_592 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).y);
     const QuicksortObject x_1077 = obj;
-    const int tint_symbol_33[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_34 = {tint_symbol_33};
-    obj = tint_symbol_34;
+    const int tint_symbol_32[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_33 = {tint_symbol_32};
+    obj = tint_symbol_33;
     obj = x_1077;
     const int x_137 = p;
     const int x_1078 = stack[x_114_save];
@@ -737,8 +737,8 @@ void quicksort_() {
       const float2 x_601 = float2(x_563.x, x_563.y);
       stack[x_147_save] = asint((1u + asuint(x_145)));
       const int x_1093[10] = stack;
-      const int tint_symbol_35[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-      stack = tint_symbol_35;
+      const int tint_symbol_34[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+      stack = tint_symbol_34;
       stack = x_1093;
       const int x_148 = top;
       const int x_1094 = stack[x_114_save];
@@ -746,8 +746,8 @@ void quicksort_() {
       stack[x_114_save] = x_1094;
       const float2 x_602 = float2(x_565.y, x_599.y);
       const int x_1095[10] = stack;
-      const int tint_symbol_36[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-      stack = tint_symbol_36;
+      const int tint_symbol_35[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+      stack = tint_symbol_35;
       stack = x_1095;
       const int x_149 = (x_148 + asint(1u));
       const int x_1096 = stack[x_147_save];
@@ -782,9 +782,9 @@ void quicksort_() {
       l_1 = x_1103;
       const float2 x_604 = float2(x_563.z, x_564.x);
       const QuicksortObject x_1104 = obj;
-      const int tint_symbol_37[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-      const QuicksortObject tint_symbol_38 = {tint_symbol_37};
-      obj = tint_symbol_38;
+      const int tint_symbol_36[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+      const QuicksortObject tint_symbol_37 = {tint_symbol_36};
+      obj = tint_symbol_37;
       obj = x_1104;
     }
   }
@@ -803,15 +803,15 @@ void main_1() {
   uv = x_717;
   i_2 = 0;
   const QuicksortObject x_721 = obj;
-  const int tint_symbol_39[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_40 = {tint_symbol_39};
-  obj = tint_symbol_40;
+  const int tint_symbol_38[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_39 = {tint_symbol_38};
+  obj = tint_symbol_39;
   obj = x_721;
   if (true) {
     const QuicksortObject x_722 = obj;
-    const int tint_symbol_41[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_42 = {tint_symbol_41};
-    obj = tint_symbol_42;
+    const int tint_symbol_40[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_41 = {tint_symbol_40};
+    obj = tint_symbol_41;
     obj = x_722;
     const float2 x_431 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).x);
     const int x_158 = i_2;
@@ -823,15 +823,15 @@ void main_1() {
     color = x_725;
     const float2 x_432 = float2(x_431.y, x_431.y);
     const QuicksortObject x_726 = obj;
-    const int tint_symbol_43[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_44 = {tint_symbol_43};
-    obj = tint_symbol_44;
+    const int tint_symbol_42[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_43 = {tint_symbol_42};
+    obj = tint_symbol_43;
     obj = x_726;
   }
   const QuicksortObject x_756 = obj;
-  const int tint_symbol_45[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_46 = {tint_symbol_45};
-  obj = tint_symbol_46;
+  const int tint_symbol_44[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_45 = {tint_symbol_44};
+  obj = tint_symbol_45;
   obj = x_756;
   const float2 x_446 = float2(float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).x);
   const int x_757 = i_2;
@@ -839,9 +839,9 @@ void main_1() {
   i_2 = x_757;
   quicksort_();
   const QuicksortObject x_758 = obj;
-  const int tint_symbol_47[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_48 = {tint_symbol_47};
-  obj = tint_symbol_48;
+  const int tint_symbol_46[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_47 = {tint_symbol_46};
+  obj = tint_symbol_47;
   obj = x_758;
   const float4 x_184 = gl_FragCoord;
   const float2 x_759 = uv;
@@ -854,18 +854,18 @@ void main_1() {
   const float2 x_185 = float2(x_184.x, x_184.y);
   const float3 x_448 = float3(x_185.y, x_446.y, x_446.y);
   const QuicksortObject x_761 = obj;
-  const int tint_symbol_49[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_50 = {tint_symbol_49};
-  obj = tint_symbol_50;
+  const int tint_symbol_48[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_49 = {tint_symbol_48};
+  obj = tint_symbol_49;
   obj = x_761;
   const float2 x_762 = uv;
   uv = float2(0.0f, 0.0f);
   uv = x_762;
   const float2 x_191 = asfloat(x_188[0].xy);
   const QuicksortObject x_763 = obj;
-  const int tint_symbol_51[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_52 = {tint_symbol_51};
-  obj = tint_symbol_52;
+  const int tint_symbol_50[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_51 = {tint_symbol_50};
+  obj = tint_symbol_51;
   obj = x_763;
   const float3 x_449 = float3(x_184.y, float3(1.0f, 2.0f, 3.0f).z, x_184.w);
   const float3 x_764 = color;
@@ -873,9 +873,9 @@ void main_1() {
   color = x_764;
   const float2 x_192 = (x_185 / x_191);
   const QuicksortObject x_765 = obj;
-  const int tint_symbol_53[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_54 = {tint_symbol_53};
-  obj = tint_symbol_54;
+  const int tint_symbol_52[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_53 = {tint_symbol_52};
+  obj = tint_symbol_53;
   obj = x_765;
   const float2 x_450 = float2(x_447.x, x_185.y);
   const float3 x_766 = color;
@@ -891,18 +891,18 @@ void main_1() {
   color = x_768;
   const float3 x_451 = float3(x_185.x, x_185.y, x_446.y);
   const QuicksortObject x_769 = obj;
-  const int tint_symbol_55[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_56 = {tint_symbol_55};
-  obj = tint_symbol_56;
+  const int tint_symbol_54[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_55 = {tint_symbol_54};
+  obj = tint_symbol_55;
   obj = x_769;
   const int x_770 = obj.numbers[0u];
   obj.numbers[0u] = 0;
   obj.numbers[0u] = x_770;
   const int x_201 = obj.numbers[0u];
   const QuicksortObject x_771 = obj;
-  const int tint_symbol_57[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_58 = {tint_symbol_57};
-  obj = tint_symbol_58;
+  const int tint_symbol_56[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_57 = {tint_symbol_56};
+  obj = tint_symbol_57;
   obj = x_771;
   const int x_772 = obj.numbers[0u];
   obj.numbers[0u] = 0;
@@ -916,9 +916,9 @@ void main_1() {
   i_2 = 0;
   i_2 = x_774;
   const QuicksortObject x_775 = obj;
-  const int tint_symbol_59[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_60 = {tint_symbol_59};
-  obj = tint_symbol_60;
+  const int tint_symbol_58[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_59 = {tint_symbol_58};
+  obj = tint_symbol_59;
   obj = x_775;
   const float3 x_453 = float3(x_451.x, x_450.x, x_450.y);
   color.x = (x_206 + float(x_201));
@@ -935,9 +935,9 @@ void main_1() {
   uv.x = 0.0f;
   uv.x = x_778;
   const QuicksortObject x_779 = obj;
-  const int tint_symbol_61[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_62 = {tint_symbol_61};
-  obj = tint_symbol_62;
+  const int tint_symbol_60[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_61 = {tint_symbol_60};
+  obj = tint_symbol_61;
   obj = x_779;
   if ((x_210 > 0.25f)) {
     const int x_780 = i_2;
@@ -952,18 +952,18 @@ void main_1() {
     uv.x = x_782;
     const int x_216 = obj.numbers[1];
     const QuicksortObject x_783 = obj;
-    const int tint_symbol_63[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_64 = {tint_symbol_63};
-    obj = tint_symbol_64;
+    const int tint_symbol_62[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_63 = {tint_symbol_62};
+    obj = tint_symbol_63;
     obj = x_783;
     const float2 x_457 = float2(x_454.x, x_454.x);
     const float2 x_784 = uv;
     uv = float2(0.0f, 0.0f);
     uv = x_784;
     const QuicksortObject x_785 = obj;
-    const int tint_symbol_65[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_66 = {tint_symbol_65};
-    obj = tint_symbol_66;
+    const int tint_symbol_64[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_65 = {tint_symbol_64};
+    obj = tint_symbol_65;
     obj = x_785;
     const float2 x_458 = float2(float3(1.0f, 2.0f, 3.0f).z, float2(0.0f, 0.0f).y);
     const int x_786 = i_2;
@@ -1081,9 +1081,9 @@ void main_1() {
     color.x = 0.0f;
     color.x = x_816;
     const QuicksortObject x_817 = obj;
-    const int tint_symbol_67[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_68 = {tint_symbol_67};
-    obj = tint_symbol_68;
+    const int tint_symbol_66[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_67 = {tint_symbol_66};
+    obj = tint_symbol_67;
     obj = x_817;
     const float3 x_468 = float3(x_467.x, x_467.x, x_467.x);
     const float x_818 = uv[0];
@@ -1192,9 +1192,9 @@ void main_1() {
     uv[0] = x_844;
     const float3 x_482 = float3(x_455.x, x_475.y, x_455.y);
     const QuicksortObject x_845 = obj;
-    const int tint_symbol_69[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_70 = {tint_symbol_69};
-    obj = tint_symbol_70;
+    const int tint_symbol_68[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_69 = {tint_symbol_68};
+    obj = tint_symbol_69;
     obj = x_845;
     const float x_846 = uv.y;
     uv.y = 0.0f;
@@ -1265,9 +1265,9 @@ void main_1() {
     obj.numbers[6u] = x_863;
     const float2 x_490 = float2(x_480.z, x_480.z);
     const QuicksortObject x_864 = obj;
-    const int tint_symbol_71[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_72 = {tint_symbol_71};
-    obj = tint_symbol_72;
+    const int tint_symbol_70[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_71 = {tint_symbol_70};
+    obj = tint_symbol_71;
     obj = x_864;
     color.y = (float(x_280) + x_283);
     const float x_865 = color.x;
@@ -1284,9 +1284,9 @@ void main_1() {
   color.x = x_867;
   const float x_287 = uv.y;
   const QuicksortObject x_868 = obj;
-  const int tint_symbol_73[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_74 = {tint_symbol_73};
-  obj = tint_symbol_74;
+  const int tint_symbol_72[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_73 = {tint_symbol_72};
+  obj = tint_symbol_73;
   obj = x_868;
   const float2 x_493 = float2(x_475.x, x_475.y);
   const float x_869 = uv[0];
@@ -1446,9 +1446,9 @@ void main_1() {
     uv.x = 0.0f;
     uv.x = x_910;
     const QuicksortObject x_911 = obj;
-    const int tint_symbol_75[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-    const QuicksortObject tint_symbol_76 = {tint_symbol_75};
-    obj = tint_symbol_76;
+    const int tint_symbol_74[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+    const QuicksortObject tint_symbol_75 = {tint_symbol_74};
+    obj = tint_symbol_75;
     obj = x_911;
     const float3 x_513 = float3(x_505.z, x_505.x, x_448.x);
     const int x_912 = obj.numbers[8];
@@ -1500,14 +1500,14 @@ void main_1() {
   uv.x = 0.0f;
   uv.x = x_923;
   const QuicksortObject x_924 = obj;
-  const int tint_symbol_77[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_78 = {tint_symbol_77};
-  obj = tint_symbol_78;
+  const int tint_symbol_76[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_77 = {tint_symbol_76};
+  obj = tint_symbol_77;
   obj = x_924;
   const QuicksortObject x_925 = obj;
-  const int tint_symbol_79[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_80 = {tint_symbol_79};
-  obj = tint_symbol_80;
+  const int tint_symbol_78[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_79 = {tint_symbol_78};
+  obj = tint_symbol_79;
   obj = x_925;
   const float x_926 = color.y;
   color.y = 0.0f;
@@ -1526,9 +1526,9 @@ void main_1() {
   uv.x = x_929;
   x_GLF_color = x_330;
   const QuicksortObject x_930 = obj;
-  const int tint_symbol_81[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
-  const QuicksortObject tint_symbol_82 = {tint_symbol_81};
-  obj = tint_symbol_82;
+  const int tint_symbol_80[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
+  const QuicksortObject tint_symbol_81 = {tint_symbol_80};
+  obj = tint_symbol_81;
   obj = x_930;
   const float3 x_522 = float3(x_330.w, x_330.y, x_493.x);
   const float x_931 = color.x;
@@ -1547,11 +1547,16 @@ struct tint_symbol_2 {
   float4 x_GLF_color_1 : SV_Target0;
 };
 
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const float4 gl_FragCoord_param = tint_symbol.gl_FragCoord_param;
+main_out main_inner(float4 gl_FragCoord_param) {
   gl_FragCoord = gl_FragCoord_param;
   main_1();
-  const main_out tint_symbol_3 = {x_GLF_color};
-  const tint_symbol_2 tint_symbol_83 = {tint_symbol_3.x_GLF_color_1};
-  return tint_symbol_83;
+  const main_out tint_symbol_82 = {x_GLF_color};
+  return tint_symbol_82;
+}
+
+tint_symbol_2 main(tint_symbol_1 tint_symbol) {
+  const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1;
+  return wrapper_result;
 }

test/bug/tint/757.wgsl.expected.hlsl

@@ -9,9 +9,7 @@ struct tint_symbol_1 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint3 GlobalInvocationID = tint_symbol.GlobalInvocationID;
+void main_inner(uint3 GlobalInvocationID) {
   uint flatIndex = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
   flatIndex = (flatIndex * 1u);
   float4 texel = myTexture.Load(int4(int3(int2(GlobalInvocationID.xy), 0), 0));
@@ -20,5 +18,10 @@ void main(tint_symbol_1 tint_symbol) {
       result.Store((4u * (flatIndex + i)), asuint(texel.r));
     }
   }
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.GlobalInvocationID);
   return;
 }

test/bug/tint/757.wgsl.expected.msl

@@ -8,13 +8,17 @@ struct Result {
   /* 0x0000 */ float values[1];
 };
 
-kernel void tint_symbol(texture2d_array<float, access::sample> tint_symbol_2 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], device Result& result [[buffer(3)]]) {
+void tint_symbol_inner(device Result& result, uint3 GlobalInvocationID, texture2d_array<float, access::sample> tint_symbol_1) {
   uint flatIndex = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
   flatIndex = (flatIndex * 1u);
-  float4 texel = tint_symbol_2.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
+  float4 texel = tint_symbol_1.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
   for(uint i = 0u; (i < 1u); i = (i + 1u)) {
     result.values[(flatIndex + i)] = texel.r;
   }
+}
+
+kernel void tint_symbol(texture2d_array<float, access::sample> tint_symbol_2 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], device Result& result [[buffer(3)]]) {
+  tint_symbol_inner(result, GlobalInvocationID, tint_symbol_2);
   return;
 }
 

test/bug/tint/824.wgsl.expected.hlsl

@@ -11,15 +11,20 @@ struct tint_symbol_2 {
   float4 Position : SV_Position;
 };
 
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const uint VertexIndex = tint_symbol.VertexIndex;
-  const uint InstanceIndex = tint_symbol.InstanceIndex;
+Output main_inner(uint VertexIndex, uint InstanceIndex) {
   float2 zv[4] = {float2(0.200000003f, 0.200000003f), float2(0.300000012f, 0.300000012f), float2(-0.100000001f, -0.100000001f), float2(1.100000024f, 1.100000024f)};
   const float z = zv[InstanceIndex].x;
   Output output = (Output)0;
   output.Position = float4(0.5f, 0.5f, z, 1.0f);
   float4 colors[4] = {float4(1.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(0.0f, 0.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f)};
   output.color = colors[InstanceIndex];
-  const tint_symbol_2 tint_symbol_3 = {output.color, output.Position};
-  return tint_symbol_3;
+  return output;
+}
+
+tint_symbol_2 main(tint_symbol_1 tint_symbol) {
+  const Output inner_result = main_inner(tint_symbol.VertexIndex, tint_symbol.InstanceIndex);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.Position = inner_result.Position;
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
 }

test/bug/tint/824.wgsl.expected.msl

@@ -5,7 +5,7 @@ struct Output {
   float4 Position;
   float4 color;
 };
-struct tint_symbol_2 {
+struct tint_symbol_1 {
   float4 color [[user(locn0)]];
   float4 Position [[position]];
 };
@@ -16,14 +16,21 @@ struct tint_array_wrapper_1 {
   float4 arr[4];
 };
 
-vertex tint_symbol_2 tint_symbol(uint VertexIndex [[vertex_id]], uint InstanceIndex [[instance_id]]) {
+Output tint_symbol_inner(uint VertexIndex, uint InstanceIndex) {
   tint_array_wrapper zv = {.arr={float2(0.200000003f, 0.200000003f), float2(0.300000012f, 0.300000012f), float2(-0.100000001f, -0.100000001f), float2(1.100000024f, 1.100000024f)}};
   float const z = zv.arr[InstanceIndex].x;
   Output output = {};
   output.Position = float4(0.5f, 0.5f, z, 1.0f);
   tint_array_wrapper_1 colors = {.arr={float4(1.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(0.0f, 0.0f, 1.0f, 1.0f), float4(1.0f, 1.0f, 1.0f, 1.0f)}};
   output.color = colors.arr[InstanceIndex];
-  tint_symbol_2 const tint_symbol_3 = {.color=output.color, .Position=output.Position};
-  return tint_symbol_3;
+  return output;
+}
+
+vertex tint_symbol_1 tint_symbol(uint VertexIndex [[vertex_id]], uint InstanceIndex [[instance_id]]) {
+  Output const inner_result = tint_symbol_inner(VertexIndex, InstanceIndex);
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.Position = inner_result.Position;
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
 }
 

test/bug/tint/827.wgsl.expected.hlsl

@@ -6,9 +6,12 @@ struct tint_symbol_1 {
   uint3 GlobalInvocationId : SV_DispatchThreadID;
 };
 
+void main_inner(uint3 GlobalInvocationId) {
+  result.Store((4u * ((GlobalInvocationId.y * width) + GlobalInvocationId.x)), asuint(tex.Load(int3(int(GlobalInvocationId.x), int(GlobalInvocationId.y), 0)).x));
+}
+
 [numthreads(1, 1, 1)]
 void main(tint_symbol_1 tint_symbol) {
-  const uint3 GlobalInvocationId = tint_symbol.GlobalInvocationId;
-  result.Store((4u * ((GlobalInvocationId.y * width) + GlobalInvocationId.x)), asuint(tex.Load(int3(int(GlobalInvocationId.x), int(GlobalInvocationId.y), 0)).x));
+  main_inner(tint_symbol.GlobalInvocationId);
   return;
 }

test/bug/tint/827.wgsl.expected.msl

@@ -6,8 +6,12 @@ struct Result {
 };
 
 constant uint width = 128u;
+void tint_symbol_inner(device Result& result, uint3 GlobalInvocationId, depth2d<float, access::sample> tint_symbol_1) {
+  result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = tint_symbol_1.read(uint2(int2(int(GlobalInvocationId.x), int(GlobalInvocationId.y))), 0);
+}
+
 kernel void tint_symbol(depth2d<float, access::sample> tint_symbol_2 [[texture(0)]], uint3 GlobalInvocationId [[thread_position_in_grid]], device Result& result [[buffer(1)]]) {
-  result.values[((GlobalInvocationId.y * width) + GlobalInvocationId.x)] = tint_symbol_2.read(uint2(int2(int(GlobalInvocationId.x), int(GlobalInvocationId.y))), 0);
+  tint_symbol_inner(result, GlobalInvocationId, tint_symbol_2);
   return;
 }
 

test/bug/tint/913.wgsl.expected.hlsl

@@ -13,9 +13,7 @@ struct tint_symbol_2 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_2 tint_symbol_1) {
-  const uint3 GlobalInvocationID = tint_symbol_1.GlobalInvocationID;
+void main_inner(uint3 GlobalInvocationID) {
   int2 tint_tmp;
   src.GetDimensions(tint_tmp.x, tint_tmp.y);
   const int2 srcSize = tint_tmp;
@@ -86,5 +84,10 @@ void main(tint_symbol_2 tint_symbol_1) {
   } else {
     output.Store((4u * outputIndex), asuint(0u));
   }
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_2 tint_symbol_1) {
+  main_inner(tint_symbol_1.GlobalInvocationID);
   return;
 }

test/bug/tint/913.wgsl.expected.msl

@@ -16,21 +16,21 @@ bool aboutEqual(float value, float expect) {
   return (fabs((value - expect)) < 0.001f);
 }
 
-kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture(0)]], texture2d<float, access::sample> tint_symbol_3 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) {
-  int2 const srcSize = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
-  int2 const dstSize = int2(tint_symbol_3.get_width(), tint_symbol_3.get_height());
+void tint_symbol_inner(constant Uniforms& uniforms, device OutputBuf& output, uint3 GlobalInvocationID, texture2d<float, access::sample> tint_symbol_1, texture2d<float, access::sample> tint_symbol_2) {
+  int2 const srcSize = int2(tint_symbol_1.get_width(), tint_symbol_1.get_height());
+  int2 const dstSize = int2(tint_symbol_2.get_width(), tint_symbol_2.get_height());
   uint2 const dstTexCoord = uint2(GlobalInvocationID.xy);
   float4 const nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
   bool success = true;
   if (((((dstTexCoord.x < uniforms.dstCopyOrigin.x) || (dstTexCoord.y < uniforms.dstCopyOrigin.y)) || (dstTexCoord.x >= (uniforms.dstCopyOrigin.x + uniforms.copySize.x))) || (dstTexCoord.y >= (uniforms.dstCopyOrigin.y + uniforms.copySize.y)))) {
-    success = (success && all((tint_symbol_3.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
+    success = (success && all((tint_symbol_2.read(uint2(int2(dstTexCoord)), 0) == nonCoveredColor)));
   } else {
     uint2 srcTexCoord = ((dstTexCoord - uniforms.dstCopyOrigin) + uniforms.srcCopyOrigin);
     if ((uniforms.dstTextureFlipY == 1u)) {
       srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
     }
-    float4 const srcColor = tint_symbol_2.read(uint2(int2(srcTexCoord)), 0);
-    float4 const dstColor = tint_symbol_3.read(uint2(int2(dstTexCoord)), 0);
+    float4 const srcColor = tint_symbol_1.read(uint2(int2(srcTexCoord)), 0);
+    float4 const dstColor = tint_symbol_2.read(uint2(int2(dstTexCoord)), 0);
     if ((uniforms.channelCount == 2u)) {
       success = ((success && aboutEqual(dstColor.r, srcColor.r)) && aboutEqual(dstColor.g, srcColor.g));
     } else {
@@ -43,6 +43,10 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture
   } else {
     output.result[outputIndex] = 0u;
   }
+}
+
+kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_3 [[texture(0)]], texture2d<float, access::sample> tint_symbol_4 [[texture(1)]], uint3 GlobalInvocationID [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], device OutputBuf& output [[buffer(2)]]) {
+  tint_symbol_inner(uniforms, output, GlobalInvocationID, tint_symbol_3, tint_symbol_4);
   return;
 }
 

test/bug/tint/914.wgsl.expected.hlsl

@@ -54,11 +54,7 @@ struct tint_symbol_1 {
   uint3 global_id : SV_DispatchThreadID;
 };
 
-[numthreads(16, 16, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint3 local_id = tint_symbol.local_id;
-  const uint3 global_id = tint_symbol.global_id;
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint3 local_id, uint3 global_id, uint local_invocation_index) {
   {
     for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
       const uint i = (idx / 64u);
@@ -143,5 +139,10 @@ void main(tint_symbol_1 tint_symbol) {
       }
     }
   }
+}
+
+[numthreads(16, 16, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.local_id, tint_symbol.global_id, tint_symbol.local_invocation_index);
   return;
 }

test/bug/tint/914.wgsl.expected.msl

@@ -50,14 +50,12 @@ void mm_write(constant Uniforms& uniforms, device Matrix& resultMatrix, uint row
   }
 }
 
-kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
-  threadgroup tint_array_wrapper tint_symbol_2;
-  threadgroup tint_array_wrapper tint_symbol_3;
+void tint_symbol_inner(constant Uniforms& uniforms, const device Matrix& firstMatrix, const device Matrix& secondMatrix, device Matrix& resultMatrix, uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array_wrapper* const tint_symbol_1, threadgroup tint_array_wrapper* const tint_symbol_2) {
   for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
     uint const i = (idx / 64u);
     uint const i_1 = (idx % 64u);
-    tint_symbol_2.arr[i].arr[i_1] = float();
-    tint_symbol_3.arr[i].arr[i_1] = float();
+    (*(tint_symbol_1)).arr[i].arr[i_1] = float();
+    (*(tint_symbol_2)).arr[i].arr[i_1] = float();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   uint const tileRow = (local_id.y * RowPerThread);
@@ -80,23 +78,23 @@ kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3
       for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
         uint const inputRow = (tileRow + innerRow);
         uint const inputCol = (tileColA + innerCol);
-        tint_symbol_2.arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
+        (*(tint_symbol_1)).arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
       }
     }
     for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
       for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
         uint const inputRow = (tileRowB + innerRow);
         uint const inputCol = (tileCol + innerCol);
-        tint_symbol_3.arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
+        (*(tint_symbol_2)).arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
       }
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
     for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
       for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
-        BCached.arr[inner] = tint_symbol_3.arr[k].arr[(tileCol + inner)];
+        BCached.arr[inner] = (*(tint_symbol_2)).arr[k].arr[(tileCol + inner)];
       }
       for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
-        ACached = tint_symbol_2.arr[(tileRow + innerRow)].arr[k];
+        ACached = (*(tint_symbol_1)).arr[(tileRow + innerRow)].arr[k];
         for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
           uint const index = ((innerRow * ColPerThread) + innerCol);
           acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
@@ -111,6 +109,12 @@ kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3
       mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
     }
   }
+}
+
+kernel void tint_symbol(uint3 local_id [[thread_position_in_threadgroup]], uint3 global_id [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) {
+  threadgroup tint_array_wrapper tint_symbol_3;
+  threadgroup tint_array_wrapper tint_symbol_4;
+  tint_symbol_inner(uniforms, firstMatrix, secondMatrix, resultMatrix, local_id, global_id, local_invocation_index, &(tint_symbol_3), &(tint_symbol_4));
   return;
 }
 

test/bug/tint/922.wgsl.expected.hlsl

@@ -154,40 +154,40 @@ Mat4x3_ _Mat4x3_1(Mat4x4_ m20) {
   return o4;
 }
 
-Mat4x3_ tint_symbol_4(uint4 buffer[96], uint offset) {
+Mat4x3_ tint_symbol_3(uint4 buffer[96], uint offset) {
   const uint scalar_offset = ((offset + 0u)) / 4;
   const uint scalar_offset_1 = ((offset + 16u)) / 4;
   const uint scalar_offset_2 = ((offset + 32u)) / 4;
-  const Mat4x3_ tint_symbol_10 = {asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4])};
-  return tint_symbol_10;
+  const Mat4x3_ tint_symbol_9 = {asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4])};
+  return tint_symbol_9;
 }
 
-Mat4x4_ tint_symbol_6(uint4 buffer[4], uint offset) {
+Mat4x4_ tint_symbol_5(uint4 buffer[4], uint offset) {
   const uint scalar_offset_3 = ((offset + 0u)) / 4;
   const uint scalar_offset_4 = ((offset + 16u)) / 4;
   const uint scalar_offset_5 = ((offset + 32u)) / 4;
   const uint scalar_offset_6 = ((offset + 48u)) / 4;
-  const Mat4x4_ tint_symbol_11 = {asfloat(buffer[scalar_offset_3 / 4]), asfloat(buffer[scalar_offset_4 / 4]), asfloat(buffer[scalar_offset_5 / 4]), asfloat(buffer[scalar_offset_6 / 4])};
-  return tint_symbol_11;
+  const Mat4x4_ tint_symbol_10 = {asfloat(buffer[scalar_offset_3 / 4]), asfloat(buffer[scalar_offset_4 / 4]), asfloat(buffer[scalar_offset_5 / 4]), asfloat(buffer[scalar_offset_6 / 4])};
+  return tint_symbol_10;
 }
 
-Mat4x2_ tint_symbol_9(uint4 buffer[3], uint offset) {
+Mat4x2_ tint_symbol_8(uint4 buffer[3], uint offset) {
   const uint scalar_offset_7 = ((offset + 0u)) / 4;
   const uint scalar_offset_8 = ((offset + 16u)) / 4;
-  const Mat4x2_ tint_symbol_12 = {asfloat(buffer[scalar_offset_7 / 4]), asfloat(buffer[scalar_offset_8 / 4])};
-  return tint_symbol_12;
+  const Mat4x2_ tint_symbol_11 = {asfloat(buffer[scalar_offset_7 / 4]), asfloat(buffer[scalar_offset_8 / 4])};
+  return tint_symbol_11;
 }
 
 void main1() {
   Mat4x3_ t_PosMtx = (Mat4x3_)0;
   float2 t_TexSpaceCoord = float2(0.0f, 0.0f);
-  const Mat4x3_ _e18 = tint_symbol_4(global2, (48u * uint(int(a_PosMtxIdx1))));
+  const Mat4x3_ _e18 = tint_symbol_3(global2, (48u * uint(int(a_PosMtxIdx1))));
   t_PosMtx = _e18;
   const Mat4x4_ _e24 = _Mat4x4_1(t_PosMtx);
   const float3 _e25 = a_Position1;
   const Mat4x4_ _e30 = _Mat4x4_1(t_PosMtx);
   const float4 _e34 = Mul(_e30, float4(a_Position1, 1.0f));
-  const Mat4x4_ _e35 = tint_symbol_6(global, 0u);
+  const Mat4x4_ _e35 = tint_symbol_5(global, 0u);
   const Mat4x4_ _e38 = _Mat4x4_1(t_PosMtx);
   const float3 _e39 = a_Position1;
   const Mat4x4_ _e44 = _Mat4x4_1(t_PosMtx);
@@ -199,7 +199,7 @@ void main1() {
   if ((_e52.x == 2.0f)) {
     {
       const float3 _e59 = a_Normal1;
-      const Mat4x2_ _e64 = tint_symbol_9(global1, (32u * uint(0)));
+      const Mat4x2_ _e64 = tint_symbol_8(global1, (32u * uint(0)));
       const float2 _e68 = Mul2(_e64, float4(a_Normal1, 1.0f));
       v_TexCoord = _e68.xy;
       return;
@@ -207,7 +207,7 @@ void main1() {
   } else {
     {
       const float2 _e73 = a_UV1;
-      const Mat4x2_ _e79 = tint_symbol_9(global1, (32u * uint(0)));
+      const Mat4x2_ _e79 = tint_symbol_8(global1, (32u * uint(0)));
       const float2 _e84 = Mul2(_e79, float4(a_UV1, 1.0f, 1.0f));
       v_TexCoord = _e84.xy;
       return;
@@ -228,19 +228,22 @@ struct tint_symbol_2 {
   float4 member : SV_Position;
 };
 
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const float3 a_Position = tint_symbol.a_Position;
-  const float2 a_UV = tint_symbol.a_UV;
-  const float4 a_Color = tint_symbol.a_Color;
-  const float3 a_Normal = tint_symbol.a_Normal;
-  const float a_PosMtxIdx = tint_symbol.a_PosMtxIdx;
+VertexOutput main_inner(float3 a_Position, float2 a_UV, float4 a_Color, float3 a_Normal, float a_PosMtxIdx) {
   a_Position1 = a_Position;
   a_UV1 = a_UV;
   a_Color1 = a_Color;
   a_Normal1 = a_Normal;
   a_PosMtxIdx1 = a_PosMtxIdx;
   main1();
-  const VertexOutput tint_symbol_3 = {v_Color, v_TexCoord, gl_Position};
-  const tint_symbol_2 tint_symbol_13 = {tint_symbol_3.v_Color, tint_symbol_3.v_TexCoord, tint_symbol_3.member};
-  return tint_symbol_13;
+  const VertexOutput tint_symbol_12 = {v_Color, v_TexCoord, gl_Position};
+  return tint_symbol_12;
+}
+
+tint_symbol_2 main(tint_symbol_1 tint_symbol) {
+  const VertexOutput inner_result = main_inner(tint_symbol.a_Position, tint_symbol.a_UV, tint_symbol.a_Color, tint_symbol.a_Normal, tint_symbol.a_PosMtxIdx);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.v_Color = inner_result.v_Color;
+  wrapper_result.v_TexCoord = inner_result.v_TexCoord;
+  wrapper_result.member = inner_result.member;
+  return wrapper_result;
 }

test/bug/tint/922.wgsl.expected.msl

@@ -222,78 +222,81 @@ Mat4x3_ _Mat4x3_1(Mat4x4_ m20) {
   return _e12;
 }
 
-void main1(constant ub_PacketParams& global2, constant ub_SceneParams& global, constant ub_MaterialParams& global1, thread float* const tint_symbol_6, thread float3* const tint_symbol_7, thread float4* const tint_symbol_8, thread float4* const tint_symbol_9, thread float4* const tint_symbol_10, thread float3* const tint_symbol_11, thread float2* const tint_symbol_12, thread float2* const tint_symbol_13) {
+void main1(constant ub_PacketParams& global2, constant ub_SceneParams& global, constant ub_MaterialParams& global1, thread float* const tint_symbol_5, thread float3* const tint_symbol_6, thread float4* const tint_symbol_7, thread float4* const tint_symbol_8, thread float4* const tint_symbol_9, thread float3* const tint_symbol_10, thread float2* const tint_symbol_11, thread float2* const tint_symbol_12) {
   Mat4x3_ t_PosMtx = {};
   float2 t_TexSpaceCoord = 0.0f;
-  float const _e15 = *(tint_symbol_6);
+  float const _e15 = *(tint_symbol_5);
   Mat4x3_ const _e18 = global2.u_PosMtx.arr[int(_e15)];
   t_PosMtx = _e18;
   Mat4x3_ const _e23 = t_PosMtx;
   Mat4x4_ const _e24 = _Mat4x4_1(_e23);
-  float3 const _e25 = *(tint_symbol_7);
+  float3 const _e25 = *(tint_symbol_6);
   Mat4x3_ const _e29 = t_PosMtx;
   Mat4x4_ const _e30 = _Mat4x4_1(_e29);
-  float3 const _e31 = *(tint_symbol_7);
+  float3 const _e31 = *(tint_symbol_6);
   float4 const _e34 = Mul(_e30, float4(_e31, 1.0f));
   Mat4x4_ const _e35 = global.u_Projection;
   Mat4x3_ const _e37 = t_PosMtx;
   Mat4x4_ const _e38 = _Mat4x4_1(_e37);
-  float3 const _e39 = *(tint_symbol_7);
+  float3 const _e39 = *(tint_symbol_6);
   Mat4x3_ const _e43 = t_PosMtx;
   Mat4x4_ const _e44 = _Mat4x4_1(_e43);
-  float3 const _e45 = *(tint_symbol_7);
+  float3 const _e45 = *(tint_symbol_6);
   float4 const _e48 = Mul(_e44, float4(_e45, 1.0f));
   float4 const _e49 = Mul(_e35, _e48);
-  *(tint_symbol_8) = _e49;
-  float4 const _e50 = *(tint_symbol_9);
-  *(tint_symbol_10) = _e50;
+  *(tint_symbol_7) = _e49;
+  float4 const _e50 = *(tint_symbol_8);
+  *(tint_symbol_9) = _e50;
   float4 const _e52 = global1.u_Misc0_;
   if ((_e52.x == 2.0f)) {
     {
-      float3 const _e59 = *(tint_symbol_11);
+      float3 const _e59 = *(tint_symbol_10);
       Mat4x2_ const _e64 = global1.u_TexMtx.arr[0];
-      float3 const _e65 = *(tint_symbol_11);
+      float3 const _e65 = *(tint_symbol_10);
       float2 const _e68 = Mul2(_e64, float4(_e65, 1.0f));
-      *(tint_symbol_12) = _e68.xy;
+      *(tint_symbol_11) = _e68.xy;
       return;
     }
   } else {
     {
-      float2 const _e73 = *(tint_symbol_13);
+      float2 const _e73 = *(tint_symbol_12);
       Mat4x2_ const _e79 = global1.u_TexMtx.arr[0];
-      float2 const _e80 = *(tint_symbol_13);
+      float2 const _e80 = *(tint_symbol_12);
       float2 const _e84 = Mul2(_e79, float4(_e80, 1.0f, 1.0f));
-      *(tint_symbol_12) = _e84.xy;
+      *(tint_symbol_11) = _e84.xy;
       return;
     }
   }
 }
 
-vertex tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]], constant ub_PacketParams& global2 [[buffer(2)]], constant ub_SceneParams& global [[buffer(0)]], constant ub_MaterialParams& global1 [[buffer(1)]]) {
-  thread float3 tint_symbol_14 = 0.0f;
-  thread float2 tint_symbol_15 = 0.0f;
-  thread float4 tint_symbol_16 = 0.0f;
-  thread float3 tint_symbol_17 = 0.0f;
-  thread float tint_symbol_18 = 0.0f;
-  thread float4 tint_symbol_19 = 0.0f;
-  thread float4 tint_symbol_20 = 0.0f;
-  thread float2 tint_symbol_21 = 0.0f;
-  float3 const a_Position = tint_symbol_1.a_Position;
-  float2 const a_UV = tint_symbol_1.a_UV;
-  float4 const a_Color = tint_symbol_1.a_Color;
-  float3 const a_Normal = tint_symbol_1.a_Normal;
-  float const a_PosMtxIdx = tint_symbol_1.a_PosMtxIdx;
-  tint_symbol_14 = a_Position;
-  tint_symbol_15 = a_UV;
-  tint_symbol_16 = a_Color;
-  tint_symbol_17 = a_Normal;
-  tint_symbol_18 = a_PosMtxIdx;
-  main1(global2, global, global1, &(tint_symbol_18), &(tint_symbol_14), &(tint_symbol_19), &(tint_symbol_16), &(tint_symbol_20), &(tint_symbol_17), &(tint_symbol_21), &(tint_symbol_15));
-  float4 const _e11 = tint_symbol_20;
-  float2 const _e13 = tint_symbol_21;
-  float4 const _e15 = tint_symbol_19;
+VertexOutput tint_symbol_inner(constant ub_PacketParams& global2, constant ub_SceneParams& global, constant ub_MaterialParams& global1, float3 a_Position, float2 a_UV, float4 a_Color, float3 a_Normal, float a_PosMtxIdx, thread float3* const tint_symbol_13, thread float2* const tint_symbol_14, thread float4* const tint_symbol_15, thread float3* const tint_symbol_16, thread float* const tint_symbol_17, thread float4* const tint_symbol_18, thread float4* const tint_symbol_19, thread float2* const tint_symbol_20) {
+  *(tint_symbol_13) = a_Position;
+  *(tint_symbol_14) = a_UV;
+  *(tint_symbol_15) = a_Color;
+  *(tint_symbol_16) = a_Normal;
+  *(tint_symbol_17) = a_PosMtxIdx;
+  main1(global2, global, global1, tint_symbol_17, tint_symbol_13, tint_symbol_18, tint_symbol_15, tint_symbol_19, tint_symbol_16, tint_symbol_20, tint_symbol_14);
+  float4 const _e11 = *(tint_symbol_19);
+  float2 const _e13 = *(tint_symbol_20);
+  float4 const _e15 = *(tint_symbol_18);
   VertexOutput const tint_symbol_4 = {.v_Color=_e11, .v_TexCoord=_e13, .member=_e15};
-  tint_symbol_3 const tint_symbol_5 = {.v_Color=tint_symbol_4.v_Color, .v_TexCoord=tint_symbol_4.v_TexCoord, .member=tint_symbol_4.member};
-  return tint_symbol_5;
+  return tint_symbol_4;
+}
+
+vertex tint_symbol_3 tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]], constant ub_PacketParams& global2 [[buffer(2)]], constant ub_SceneParams& global [[buffer(0)]], constant ub_MaterialParams& global1 [[buffer(1)]]) {
+  thread float3 tint_symbol_21 = 0.0f;
+  thread float2 tint_symbol_22 = 0.0f;
+  thread float4 tint_symbol_23 = 0.0f;
+  thread float3 tint_symbol_24 = 0.0f;
+  thread float tint_symbol_25 = 0.0f;
+  thread float4 tint_symbol_26 = 0.0f;
+  thread float4 tint_symbol_27 = 0.0f;
+  thread float2 tint_symbol_28 = 0.0f;
+  VertexOutput const inner_result = tint_symbol_inner(global2, global, global1, tint_symbol_1.a_Position, tint_symbol_1.a_UV, tint_symbol_1.a_Color, tint_symbol_1.a_Normal, tint_symbol_1.a_PosMtxIdx, &(tint_symbol_21), &(tint_symbol_22), &(tint_symbol_23), &(tint_symbol_24), &(tint_symbol_25), &(tint_symbol_26), &(tint_symbol_27), &(tint_symbol_28));
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.v_Color = inner_result.v_Color;
+  wrapper_result.v_TexCoord = inner_result.v_TexCoord;
+  wrapper_result.member = inner_result.member;
+  return wrapper_result;
 }
 

test/bug/tint/926.wgsl.expected.hlsl

@@ -11,9 +11,12 @@ struct tint_symbol_1 {
   uint3 global_id : SV_DispatchThreadID;
 };
 
+void computeMain_inner(uint3 global_id) {
+  const uint firstVertex = atomicAdd_1(drawOut, 0u, cubeVerts);
+}
+
 [numthreads(1, 1, 1)]
 void computeMain(tint_symbol_1 tint_symbol) {
-  const uint3 global_id = tint_symbol.global_id;
-  const uint firstVertex = atomicAdd_1(drawOut, 0u, cubeVerts);
+  computeMain_inner(tint_symbol.global_id);
   return;
 }

test/bug/tint/926.wgsl.expected.msl

@@ -5,9 +5,13 @@ struct DrawIndirectArgs {
   /* 0x0000 */ atomic_uint vertexCount;
 };
 
+void computeMain_inner(device DrawIndirectArgs& drawOut, uint3 global_id, thread uint* const tint_symbol) {
+  uint const firstVertex = atomic_fetch_add_explicit(&(drawOut.vertexCount), *(tint_symbol), memory_order_relaxed);
+}
+
 kernel void computeMain(uint3 global_id [[thread_position_in_grid]], device DrawIndirectArgs& drawOut [[buffer(5)]]) {
   thread uint tint_symbol_1 = 0u;
-  uint const firstVertex = atomic_fetch_add_explicit(&(drawOut.vertexCount), tint_symbol_1, memory_order_relaxed);
+  computeMain_inner(drawOut, global_id, &(tint_symbol_1));
   return;
 }
 

test/bug/tint/942.wgsl.expected.hlsl

@@ -16,11 +16,7 @@ struct tint_symbol_1 {
   uint3 WorkGroupID : SV_GroupID;
 };
 
-[numthreads(64, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint3 WorkGroupID = tint_symbol.WorkGroupID;
-  const uint3 LocalInvocationID = tint_symbol.LocalInvocationID;
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocation_index) {
   {
     for(uint idx = local_invocation_index; (idx < 1024u); idx = (idx + 64u)) {
       const uint i_1 = (idx / 256u);
@@ -79,5 +75,10 @@ void main(tint_symbol_1 tint_symbol) {
       }
     }
   }
+}
+
+[numthreads(64, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.WorkGroupID, tint_symbol.LocalInvocationID, tint_symbol.local_invocation_index);
   return;
 }

test/bug/tint/942.wgsl.expected.msl

@@ -15,16 +15,15 @@ struct tint_array_wrapper {
   tint_array_wrapper_1 arr[4];
 };
 
-kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_3 [[texture(1)]], sampler tint_symbol_4 [[sampler(0)]], texture2d<float, access::write> tint_symbol_5 [[texture(2)]], uint3 WorkGroupID [[threadgroup_position_in_grid]], uint3 LocalInvocationID [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Params& params [[buffer(1)]], constant Flip& flip [[buffer(3)]]) {
-  threadgroup tint_array_wrapper tint_symbol_2;
+void tint_symbol_inner(constant Params& params, constant Flip& flip, uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocation_index, threadgroup tint_array_wrapper* const tint_symbol_1, texture2d<float, access::sample> tint_symbol_2, sampler tint_symbol_3, texture2d<float, access::write> tint_symbol_4) {
   for(uint idx = local_invocation_index; (idx < 1024u); idx = (idx + 64u)) {
     uint const i_1 = (idx / 256u);
     uint const i_2 = (idx % 256u);
-    tint_symbol_2.arr[i_1].arr[i_2] = float3();
+    (*(tint_symbol_1)).arr[i_1].arr[i_2] = float3();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   uint const filterOffset = ((params.filterDim - 1u) / 2u);
-  int2 const dims = int2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
+  int2 const dims = int2(tint_symbol_2.get_width(0), tint_symbol_2.get_height(0));
   int2 const baseIndex = as_type<int2>((as_type<uint2>(int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u))))) - as_type<uint2>(int2(int(filterOffset), 0))));
   for(uint r = 0u; (r < 4u); r = (r + 1u)) {
     for(uint c = 0u; (c < 4u); c = (c + 1u)) {
@@ -32,7 +31,7 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_3 [[texture
       if ((flip.value != 0u)) {
         loadIndex = loadIndex.yx;
       }
-      tint_symbol_2.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_3.sample(tint_symbol_4, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
+      (*(tint_symbol_1)).arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_2.sample(tint_symbol_3, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
@@ -47,12 +46,17 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_3 [[texture
         float3 acc = float3(0.0f, 0.0f, 0.0f);
         for(uint f = 0u; (f < params.filterDim); f = (f + 1u)) {
           uint i = ((center + f) - filterOffset);
-          acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_2.arr[r].arr[i]));
+          acc = (acc + ((1.0f / float(params.filterDim)) * (*(tint_symbol_1)).arr[r].arr[i]));
         }
-        tint_symbol_5.write(float4(acc, 1.0f), uint2(writeIndex));
+        tint_symbol_4.write(float4(acc, 1.0f), uint2(writeIndex));
       }
     }
   }
+}
+
+kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_6 [[texture(1)]], sampler tint_symbol_7 [[sampler(0)]], texture2d<float, access::write> tint_symbol_8 [[texture(2)]], uint3 WorkGroupID [[threadgroup_position_in_grid]], uint3 LocalInvocationID [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Params& params [[buffer(1)]], constant Flip& flip [[buffer(3)]]) {
+  threadgroup tint_array_wrapper tint_symbol_5;
+  tint_symbol_inner(params, flip, WorkGroupID, LocalInvocationID, local_invocation_index, &(tint_symbol_5), tint_symbol_6, tint_symbol_7, tint_symbol_8);
   return;
 }
 

test/bug/tint/943.spvasm.expected.hlsl

@@ -345,11 +345,7 @@ struct tint_symbol_1 {
   uint3 gl_GlobalInvocationID_param : SV_DispatchThreadID;
 };
 
-[numthreads(1, 64, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint3 gl_LocalInvocationID_param = tint_symbol.gl_LocalInvocationID_param;
-  const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
-  const uint local_invocation_index = tint_symbol.local_invocation_index;
+void main_inner(uint3 gl_LocalInvocationID_param, uint3 gl_GlobalInvocationID_param, uint local_invocation_index) {
   {
     const uint i_1 = local_invocation_index;
     const uint i_2 = (local_invocation_index % 1u);
@@ -366,5 +362,10 @@ void main(tint_symbol_1 tint_symbol) {
   gl_LocalInvocationID = gl_LocalInvocationID_param;
   gl_GlobalInvocationID = gl_GlobalInvocationID_param;
   main_1();
+}
+
+[numthreads(1, 64, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.gl_LocalInvocationID_param, tint_symbol.gl_GlobalInvocationID_param, tint_symbol.local_invocation_index);
   return;
 }

test/bug/tint/943.spvasm.expected.msl

@@ -53,7 +53,7 @@ bool coordsInBounds_vi2_vi2_(thread int2* const coord, thread int2* const shape)
   return x_88;
 }
 
-float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread int* const row, thread int* const col, thread int* const tint_symbol_3, thread int* const tint_symbol_4, thread int* const tint_symbol_5) {
+float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread int* const row, thread int* const col, thread int* const tint_symbol_2, thread int* const tint_symbol_3, thread int* const tint_symbol_4) {
   int batchASize = 0;
   int2 param_10 = 0;
   int2 param_11 = 0;
@@ -63,16 +63,16 @@ float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread
   batchASize = as_type<int>((as_type<uint>(x_417) * as_type<uint>(x_419)));
   int const x_421 = *(row);
   int const x_422 = *(col);
-  int const x_424 = *(tint_symbol_3);
-  int const x_425 = *(tint_symbol_4);
+  int const x_424 = *(tint_symbol_2);
+  int const x_425 = *(tint_symbol_3);
   param_10 = int2(x_421, x_422);
   param_11 = int2(x_424, x_425);
   bool const x_429 = coordsInBounds_vi2_vi2_(&(param_10), &(param_11));
   if (x_429) {
-    int const x_438 = *(tint_symbol_5);
+    int const x_438 = *(tint_symbol_4);
     int const x_439 = batchASize;
     int const x_441 = *(row);
-    int const x_442 = *(tint_symbol_4);
+    int const x_442 = *(tint_symbol_3);
     int const x_445 = *(col);
     float const x_448 = x_165.A[as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_438) * as_type<uint>(x_439)))) + as_type<uint>(as_type<int>((as_type<uint>(x_441) * as_type<uint>(x_442))))))) + as_type<uint>(x_445)))];
     x_430 = x_448;
@@ -83,7 +83,7 @@ float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread
   return x_450;
 }
 
-float mm_readB_i1_i1_(constant Uniforms& x_48, const device ssbB& x_185, thread int* const row_1, thread int* const col_1, thread int* const tint_symbol_6, thread int* const tint_symbol_7, thread int* const tint_symbol_8) {
+float mm_readB_i1_i1_(constant Uniforms& x_48, const device ssbB& x_185, thread int* const row_1, thread int* const col_1, thread int* const tint_symbol_5, thread int* const tint_symbol_6, thread int* const tint_symbol_7) {
   int batchBSize = 0;
   int2 param_12 = 0;
   int2 param_13 = 0;
@@ -93,16 +93,16 @@ float mm_readB_i1_i1_(constant Uniforms& x_48, const device ssbB& x_185, thread
   batchBSize = as_type<int>((as_type<uint>(x_455) * as_type<uint>(x_457)));
   int const x_459 = *(row_1);
   int const x_460 = *(col_1);
-  int const x_462 = *(tint_symbol_6);
-  int const x_463 = *(tint_symbol_7);
+  int const x_462 = *(tint_symbol_5);
+  int const x_463 = *(tint_symbol_6);
   param_12 = int2(x_459, x_460);
   param_13 = int2(x_462, x_463);
   bool const x_467 = coordsInBounds_vi2_vi2_(&(param_12), &(param_13));
   if (x_467) {
-    int const x_475 = *(tint_symbol_8);
+    int const x_475 = *(tint_symbol_7);
     int const x_476 = batchBSize;
     int const x_478 = *(row_1);
-    int const x_479 = *(tint_symbol_7);
+    int const x_479 = *(tint_symbol_6);
     int const x_482 = *(col_1);
     float const x_485 = x_185.B[as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_475) * as_type<uint>(x_476)))) + as_type<uint>(as_type<int>((as_type<uint>(x_478) * as_type<uint>(x_479))))))) + as_type<uint>(x_482)))];
     x_468 = x_485;
@@ -146,17 +146,17 @@ void setOutput_i1_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread
   return;
 }
 
-void mm_write_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread int* const row_2, thread int* const col_2, thread float* const value_2, thread int* const tint_symbol_9) {
+void mm_write_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread int* const row_2, thread int* const col_2, thread float* const value_2, thread int* const tint_symbol_8) {
   int3 outCoord = 0;
   int param_14 = 0;
   int param_15 = 0;
   int param_16 = 0;
   float param_17 = 0.0f;
-  int const x_491 = *(tint_symbol_9);
+  int const x_491 = *(tint_symbol_8);
   int const x_492 = *(row_2);
   int const x_493 = *(col_2);
   outCoord = int3(x_491, x_492, x_493);
-  int const x_496 = *(tint_symbol_9);
+  int const x_496 = *(tint_symbol_8);
   param_14 = x_496;
   int const x_498 = *(row_2);
   param_15 = x_498;
@@ -168,7 +168,7 @@ void mm_write_i1_i1_f1_(constant Uniforms& x_48, device ssbOut& x_54, thread int
   return;
 }
 
-void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const dimAOuter, thread int* const dimInner, thread int* const dimBOuter, thread uint3* const tint_symbol_10, thread uint3* const tint_symbol_11, thread int* const tint_symbol_12, thread int* const tint_symbol_13, thread int* const tint_symbol_14, threadgroup tint_array_wrapper* const tint_symbol_15, thread int* const tint_symbol_16, threadgroup tint_array_wrapper_2* const tint_symbol_17) {
+void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const dimAOuter, thread int* const dimInner, thread int* const dimBOuter, thread uint3* const tint_symbol_9, thread uint3* const tint_symbol_10, thread int* const tint_symbol_11, thread int* const tint_symbol_12, thread int* const tint_symbol_13, threadgroup tint_array_wrapper* const tint_symbol_14, thread int* const tint_symbol_15, threadgroup tint_array_wrapper_2* const tint_symbol_16) {
   int tileRow = 0;
   int tileCol = 0;
   int globalRow = 0;
@@ -203,13 +203,13 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
   int param_7 = 0;
   int param_8 = 0;
   float param_9 = 0.0f;
-  uint const x_132 = (*(tint_symbol_10)).y;
+  uint const x_132 = (*(tint_symbol_9)).y;
   tileRow = as_type<int>((as_type<uint>(as_type<int>(x_132)) * as_type<uint>(1)));
-  uint const x_137 = (*(tint_symbol_10)).x;
+  uint const x_137 = (*(tint_symbol_9)).x;
   tileCol = as_type<int>((as_type<uint>(as_type<int>(x_137)) * as_type<uint>(1)));
-  uint const x_143 = (*(tint_symbol_11)).y;
+  uint const x_143 = (*(tint_symbol_10)).y;
   globalRow = as_type<int>((as_type<uint>(as_type<int>(x_143)) * as_type<uint>(1)));
-  uint const x_148 = (*(tint_symbol_11)).x;
+  uint const x_148 = (*(tint_symbol_10)).x;
   globalCol = as_type<int>((as_type<uint>(as_type<int>(x_148)) * as_type<uint>(1)));
   int const x_152 = *(dimInner);
   numTiles = as_type<int>((as_type<uint>((as_type<int>((as_type<uint>(x_152) - as_type<uint>(1))) / 64)) + as_type<uint>(1)));
@@ -240,9 +240,9 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
       innerRow = as_type<int>((as_type<uint>(x_183) + as_type<uint>(1)));
     }
   }
-  uint const x_187 = (*(tint_symbol_10)).x;
+  uint const x_187 = (*(tint_symbol_9)).x;
   tileColA = as_type<int>((as_type<uint>(as_type<int>(x_187)) * as_type<uint>(64)));
-  uint const x_192 = (*(tint_symbol_10)).y;
+  uint const x_192 = (*(tint_symbol_9)).y;
   tileRowB = as_type<int>((as_type<uint>(as_type<int>(x_192)) * as_type<uint>(1)));
   t = 0;
   while (true) {
@@ -280,8 +280,8 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         int const x_240 = inputCol;
         param_3 = as_type<int>((as_type<uint>(x_235) + as_type<uint>(x_236)));
         param_4 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_238) * as_type<uint>(64)))) + as_type<uint>(x_240)));
-        float const x_244 = mm_readA_i1_i1_(x_48, x_165, &(param_3), &(param_4), tint_symbol_12, tint_symbol_13, tint_symbol_14);
-        (*(tint_symbol_15)).arr[x_233].arr[x_234] = x_244;
+        float const x_244 = mm_readA_i1_i1_(x_48, x_165, &(param_3), &(param_4), tint_symbol_11, tint_symbol_12, tint_symbol_13);
+        (*(tint_symbol_14)).arr[x_233].arr[x_234] = x_244;
         {
           int const x_247 = innerCol_1;
           innerCol_1 = as_type<int>((as_type<uint>(x_247) + as_type<uint>(1)));
@@ -320,8 +320,8 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         int const x_285 = innerCol_2;
         param_5 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_280) * as_type<uint>(64)))) + as_type<uint>(x_282)));
         param_6 = as_type<int>((as_type<uint>(x_284) + as_type<uint>(x_285)));
-        float const x_289 = mm_readB_i1_i1_(x_48, x_185, &(param_5), &(param_6), tint_symbol_13, tint_symbol_16, tint_symbol_14);
-        (*(tint_symbol_17)).arr[x_278].arr[x_279] = x_289;
+        float const x_289 = mm_readB_i1_i1_(x_48, x_185, &(param_5), &(param_6), tint_symbol_12, tint_symbol_15, tint_symbol_13);
+        (*(tint_symbol_16)).arr[x_278].arr[x_279] = x_289;
         {
           int const x_291 = innerCol_2;
           innerCol_2 = as_type<int>((as_type<uint>(x_291) + as_type<uint>(1)));
@@ -351,7 +351,7 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         int const x_315 = k;
         int const x_316 = tileCol;
         int const x_317 = inner;
-        float const x_320 = (*(tint_symbol_17)).arr[x_315].arr[as_type<int>((as_type<uint>(x_316) + as_type<uint>(x_317)))];
+        float const x_320 = (*(tint_symbol_16)).arr[x_315].arr[as_type<int>((as_type<uint>(x_316) + as_type<uint>(x_317)))];
         BCached.arr[x_314] = x_320;
         {
           int const x_322 = inner;
@@ -368,7 +368,7 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         int const x_333 = tileRow;
         int const x_334 = innerRow_3;
         int const x_336 = k;
-        float const x_338 = (*(tint_symbol_15)).arr[as_type<int>((as_type<uint>(x_333) + as_type<uint>(x_334)))].arr[x_336];
+        float const x_338 = (*(tint_symbol_14)).arr[as_type<int>((as_type<uint>(x_333) + as_type<uint>(x_334)))].arr[x_336];
         ACached = x_338;
         innerCol_3 = 0;
         while (true) {
@@ -445,7 +445,7 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         param_8 = as_type<int>((as_type<uint>(x_400) + as_type<uint>(x_401)));
         float const x_409 = acc.arr[x_403].arr[x_404];
         param_9 = x_409;
-        mm_write_i1_i1_f1_(x_48, x_54, &(param_7), &(param_8), &(param_9), tint_symbol_14);
+        mm_write_i1_i1_f1_(x_48, x_54, &(param_7), &(param_8), &(param_9), tint_symbol_13);
       }
       {
         int const x_411 = innerCol_4;
@@ -460,51 +460,55 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
   return;
 }
 
-void main_1(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const tint_symbol_18, thread int* const tint_symbol_19, thread int* const tint_symbol_20, thread uint3* const tint_symbol_21, thread int* const tint_symbol_22, thread uint3* const tint_symbol_23, threadgroup tint_array_wrapper* const tint_symbol_24, threadgroup tint_array_wrapper_2* const tint_symbol_25) {
+void main_1(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, thread int* const tint_symbol_17, thread int* const tint_symbol_18, thread int* const tint_symbol_19, thread uint3* const tint_symbol_20, thread int* const tint_symbol_21, thread uint3* const tint_symbol_22, threadgroup tint_array_wrapper* const tint_symbol_23, threadgroup tint_array_wrapper_2* const tint_symbol_24) {
   int param_18 = 0;
   int param_19 = 0;
   int param_20 = 0;
   int const x_67 = x_48.aShape.y;
-  *(tint_symbol_18) = x_67;
+  *(tint_symbol_17) = x_67;
   int const x_71 = x_48.aShape.z;
-  *(tint_symbol_19) = x_71;
+  *(tint_symbol_18) = x_71;
   int const x_75 = x_48.bShape.z;
-  *(tint_symbol_20) = x_75;
-  uint const x_505 = (*(tint_symbol_21)).z;
-  *(tint_symbol_22) = as_type<int>(x_505);
-  int const x_508 = *(tint_symbol_18);
+  *(tint_symbol_19) = x_75;
+  uint const x_505 = (*(tint_symbol_20)).z;
+  *(tint_symbol_21) = as_type<int>(x_505);
+  int const x_508 = *(tint_symbol_17);
   param_18 = x_508;
-  int const x_510 = *(tint_symbol_19);
+  int const x_510 = *(tint_symbol_18);
   param_19 = x_510;
-  int const x_512 = *(tint_symbol_20);
+  int const x_512 = *(tint_symbol_19);
   param_20 = x_512;
-  mm_matMul_i1_i1_i1_(x_48, x_165, x_185, x_54, &(param_18), &(param_19), &(param_20), tint_symbol_23, tint_symbol_21, tint_symbol_18, tint_symbol_19, tint_symbol_22, tint_symbol_24, tint_symbol_20, tint_symbol_25);
+  mm_matMul_i1_i1_i1_(x_48, x_165, x_185, x_54, &(param_18), &(param_19), &(param_20), tint_symbol_22, tint_symbol_20, tint_symbol_17, tint_symbol_18, tint_symbol_21, tint_symbol_23, tint_symbol_19, tint_symbol_24);
   return;
 }
 
-kernel void tint_symbol_1(uint3 gl_LocalInvocationID_param [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Uniforms& x_48 [[buffer(3)]], const device ssbA& x_165 [[buffer(1)]], const device ssbB& x_185 [[buffer(2)]], device ssbOut& x_54 [[buffer(0)]]) {
-  threadgroup tint_array_wrapper_2 tint_symbol_26;
-  threadgroup tint_array_wrapper tint_symbol_27;
-  thread uint3 tint_symbol_28 = 0u;
-  thread uint3 tint_symbol_29 = 0u;
-  thread int tint_symbol_30 = 0;
-  thread int tint_symbol_31 = 0;
-  thread int tint_symbol_32 = 0;
-  thread int tint_symbol_33 = 0;
+void tint_symbol_1_inner(constant Uniforms& x_48, const device ssbA& x_165, const device ssbB& x_185, device ssbOut& x_54, uint3 gl_LocalInvocationID_param, uint3 gl_GlobalInvocationID_param, uint local_invocation_index, threadgroup tint_array_wrapper_2* const tint_symbol_25, threadgroup tint_array_wrapper* const tint_symbol_26, thread uint3* const tint_symbol_27, thread uint3* const tint_symbol_28, thread int* const tint_symbol_29, thread int* const tint_symbol_30, thread int* const tint_symbol_31, thread int* const tint_symbol_32) {
   {
     uint const i_1 = local_invocation_index;
     uint const i_2 = (local_invocation_index % 1u);
-    tint_symbol_26.arr[i_1].arr[i_2] = float();
+    (*(tint_symbol_25)).arr[i_1].arr[i_2] = float();
   }
   for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 64u)) {
     uint const i = (idx / 64u);
     uint const i_1 = (idx % 64u);
-    tint_symbol_27.arr[i].arr[i_1] = float();
+    (*(tint_symbol_26)).arr[i].arr[i_1] = float();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  tint_symbol_28 = gl_LocalInvocationID_param;
-  tint_symbol_29 = gl_GlobalInvocationID_param;
-  main_1(x_48, x_165, x_185, x_54, &(tint_symbol_30), &(tint_symbol_31), &(tint_symbol_32), &(tint_symbol_29), &(tint_symbol_33), &(tint_symbol_28), &(tint_symbol_27), &(tint_symbol_26));
+  *(tint_symbol_27) = gl_LocalInvocationID_param;
+  *(tint_symbol_28) = gl_GlobalInvocationID_param;
+  main_1(x_48, x_165, x_185, x_54, tint_symbol_29, tint_symbol_30, tint_symbol_31, tint_symbol_28, tint_symbol_32, tint_symbol_27, tint_symbol_26, tint_symbol_25);
+}
+
+kernel void tint_symbol_1(uint3 gl_LocalInvocationID_param [[thread_position_in_threadgroup]], uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Uniforms& x_48 [[buffer(3)]], const device ssbA& x_165 [[buffer(1)]], const device ssbB& x_185 [[buffer(2)]], device ssbOut& x_54 [[buffer(0)]]) {
+  threadgroup tint_array_wrapper_2 tint_symbol_33;
+  threadgroup tint_array_wrapper tint_symbol_34;
+  thread uint3 tint_symbol_35 = 0u;
+  thread uint3 tint_symbol_36 = 0u;
+  thread int tint_symbol_37 = 0;
+  thread int tint_symbol_38 = 0;
+  thread int tint_symbol_39 = 0;
+  thread int tint_symbol_40 = 0;
+  tint_symbol_1_inner(x_48, x_165, x_185, x_54, gl_LocalInvocationID_param, gl_GlobalInvocationID_param, local_invocation_index, &(tint_symbol_33), &(tint_symbol_34), &(tint_symbol_35), &(tint_symbol_36), &(tint_symbol_37), &(tint_symbol_38), &(tint_symbol_39), &(tint_symbol_40));
   return;
 }
 

test/bug/tint/948.wgsl.expected.hlsl

@@ -164,13 +164,7 @@ struct tint_symbol_2 {
   float4 glFragColor_1 : SV_Target0;
 };
 
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const float2 tUV_param = tint_symbol.tUV_param;
-  const float2 tileID_1_param = tint_symbol.tileID_1_param;
-  const float2 levelUnits_param = tint_symbol.levelUnits_param;
-  const float2 stageUnits_1_param = tint_symbol.stageUnits_1_param;
-  const float3 vPosition_param = tint_symbol.vPosition_param;
-  const float2 vUV_param = tint_symbol.vUV_param;
+main_out main_inner(float2 tUV_param, float2 tileID_1_param, float2 levelUnits_param, float2 stageUnits_1_param, float3 vPosition_param, float2 vUV_param) {
   tUV = tUV_param;
   tileID_1 = tileID_1_param;
   levelUnits = levelUnits_param;
@@ -178,7 +172,13 @@ tint_symbol_2 main(tint_symbol_1 tint_symbol) {
   vPosition = vPosition_param;
   vUV = vUV_param;
   main_1();
-  const main_out tint_symbol_3 = {glFragColor};
-  const tint_symbol_2 tint_symbol_7 = {tint_symbol_3.glFragColor_1};
-  return tint_symbol_7;
+  const main_out tint_symbol_6 = {glFragColor};
+  return tint_symbol_6;
+}
+
+tint_symbol_2 main(tint_symbol_1 tint_symbol) {
+  const main_out inner_result = main_inner(tint_symbol.tUV_param, tint_symbol.tileID_1_param, tint_symbol.levelUnits_param, tint_symbol.stageUnits_1_param, tint_symbol.vPosition_param, tint_symbol.vUV_param);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.glFragColor_1 = inner_result.glFragColor_1;
+  return wrapper_result;
 }

test/bug/tint/948.wgsl.expected.msl

@@ -29,21 +29,21 @@ struct tint_symbol_3 {
   float4 glFragColor_1 [[color(0)]];
 };
 
-float4x4 getFrameData_f1_(constant LeftOver& x_20, thread float* const frameID, texture2d<float, access::sample> tint_symbol_6, sampler tint_symbol_7) {
+float4x4 getFrameData_f1_(constant LeftOver& x_20, thread float* const frameID, texture2d<float, access::sample> tint_symbol_5, sampler tint_symbol_6) {
   float fX = 0.0f;
   float const x_15 = *(frameID);
   float const x_25 = x_20.spriteCount;
   fX = (x_15 / x_25);
   float const x_37 = fX;
-  float4 const x_40 = tint_symbol_6.sample(tint_symbol_7, float2(x_37, 0.0f), bias(0.0f));
+  float4 const x_40 = tint_symbol_5.sample(tint_symbol_6, float2(x_37, 0.0f), bias(0.0f));
   float const x_44 = fX;
-  float4 const x_47 = tint_symbol_6.sample(tint_symbol_7, float2(x_44, 0.25f), bias(0.0f));
+  float4 const x_47 = tint_symbol_5.sample(tint_symbol_6, float2(x_44, 0.25f), bias(0.0f));
   float const x_51 = fX;
-  float4 const x_54 = tint_symbol_6.sample(tint_symbol_7, float2(x_51, 0.5f), bias(0.0f));
+  float4 const x_54 = tint_symbol_5.sample(tint_symbol_6, float2(x_51, 0.5f), bias(0.0f));
   return float4x4(float4(x_40.x, x_40.y, x_40.z, x_40.w), float4(x_47.x, x_47.y, x_47.z, x_47.w), float4(x_54.x, x_54.y, x_54.z, x_54.w), float4(float4(0.0f, 0.0f, 0.0f, 0.0f).x, float4(0.0f, 0.0f, 0.0f, 0.0f).y, float4(0.0f, 0.0f, 0.0f, 0.0f).z, float4(0.0f, 0.0f, 0.0f, 0.0f).w));
 }
 
-void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture2d<float, access::sample> tint_symbol_9, sampler tint_symbol_10, texture2d<float, access::sample> tint_symbol_11, texture2d<float, access::sample> tint_symbol_12, sampler tint_symbol_13, thread float* const tint_symbol_14, texture2d<float, access::sample> tint_symbol_15, sampler tint_symbol_16, texture2d<float, access::sample> tint_symbol_17, sampler tint_symbol_18, thread float4* const tint_symbol_19) {
+void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_7, texture2d<float, access::sample> tint_symbol_8, sampler tint_symbol_9, texture2d<float, access::sample> tint_symbol_10, texture2d<float, access::sample> tint_symbol_11, sampler tint_symbol_12, thread float* const tint_symbol_13, texture2d<float, access::sample> tint_symbol_14, sampler tint_symbol_15, texture2d<float, access::sample> tint_symbol_16, sampler tint_symbol_17, thread float4* const tint_symbol_18) {
   float4 color = 0.0f;
   float2 tileUV = 0.0f;
   float2 tileID = 0.0f;
@@ -63,11 +63,11 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
   float alpha = 0.0f;
   float3 mixed = 0.0f;
   color = float4(0.0f, 0.0f, 0.0f, 0.0f);
-  float2 const x_86 = *(tint_symbol_8);
+  float2 const x_86 = *(tint_symbol_7);
   tileUV = fract(x_86);
   float const x_91 = tileUV.y;
   tileUV.y = (1.0f - x_91);
-  float2 const x_95 = *(tint_symbol_8);
+  float2 const x_95 = *(tint_symbol_7);
   tileID = floor(x_95);
   float2 const x_101 = x_20.spriteMapSize;
   sheetUnits = (float2(1.0f, 1.0f) / x_101);
@@ -87,14 +87,14 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
       case 1: {
         float2 const x_150 = tileID;
         float2 const x_154 = x_20.stageSize;
-        float4 const x_156 = tint_symbol_9.sample(tint_symbol_10, ((x_150 + float2(0.5f, 0.5f)) / x_154), bias(0.0f));
+        float4 const x_156 = tint_symbol_8.sample(tint_symbol_9, ((x_150 + float2(0.5f, 0.5f)) / x_154), bias(0.0f));
         frameID_1 = x_156.x;
         break;
       }
       case 0: {
         float2 const x_136 = tileID;
         float2 const x_140 = x_20.stageSize;
-        float4 const x_142 = tint_symbol_11.sample(tint_symbol_10, ((x_136 + float2(0.5f, 0.5f)) / x_140), bias(0.0f));
+        float4 const x_142 = tint_symbol_10.sample(tint_symbol_9, ((x_136 + float2(0.5f, 0.5f)) / x_140), bias(0.0f));
         frameID_1 = x_142.x;
         break;
       }
@@ -104,13 +104,13 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
     }
     float const x_166 = frameID_1;
     float const x_169 = x_20.spriteCount;
-    float4 const x_172 = tint_symbol_12.sample(tint_symbol_13, float2(((x_166 + 0.5f) / x_169), 0.0f), bias(0.0f));
+    float4 const x_172 = tint_symbol_11.sample(tint_symbol_12, float2(((x_166 + 0.5f) / x_169), 0.0f), bias(0.0f));
     animationData = x_172;
     float const x_174 = animationData.y;
     if ((x_174 > 0.0f)) {
       float const x_181 = x_20.time;
       float const x_184 = animationData.z;
-      *(tint_symbol_14) = fmod((x_181 * x_184), 1.0f);
+      *(tint_symbol_13) = fmod((x_181 * x_184), 1.0f);
       f = 0.0f;
       while (true) {
         float const x_193 = f;
@@ -119,7 +119,7 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
           break;
         }
         float const x_197 = animationData.y;
-        float const x_198 = *(tint_symbol_14);
+        float const x_198 = *(tint_symbol_13);
         if ((x_197 > x_198)) {
           float const x_203 = animationData.x;
           frameID_1 = x_203;
@@ -128,7 +128,7 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
         float const x_208 = frameID_1;
         float const x_211 = x_20.spriteCount;
         float const x_214 = f;
-        float4 const x_217 = tint_symbol_12.sample(tint_symbol_13, float2(((x_208 + 0.5f) / x_211), (0.125f * x_214)), bias(0.0f));
+        float4 const x_217 = tint_symbol_11.sample(tint_symbol_12, float2(((x_208 + 0.5f) / x_211), (0.125f * x_214)), bias(0.0f));
         animationData = x_217;
         {
           float const x_218 = f;
@@ -138,7 +138,7 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
     }
     float const x_222 = frameID_1;
     param = (x_222 + 0.5f);
-    float4x4 const x_225 = getFrameData_f1_(x_20, &(param), tint_symbol_15, tint_symbol_16);
+    float4x4 const x_225 = getFrameData_f1_(x_20, &(param), tint_symbol_14, tint_symbol_15);
     frameData = x_225;
     float4 const x_228 = frameData[0];
     float2 const x_231 = x_20.spriteMapSize;
@@ -159,13 +159,13 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
       float2 const x_263 = tileUV;
       float2 const x_264 = frameSize;
       float2 const x_266 = offset_1;
-      float4 const x_268 = tint_symbol_17.sample(tint_symbol_18, ((x_263 * x_264) + x_266));
+      float4 const x_268 = tint_symbol_16.sample(tint_symbol_17, ((x_263 * x_264) + x_266));
       color = x_268;
     } else {
       float2 const x_274 = tileUV;
       float2 const x_275 = frameSize;
       float2 const x_277 = offset_1;
-      float4 const x_279 = tint_symbol_17.sample(tint_symbol_18, ((x_274 * x_275) + x_277));
+      float4 const x_279 = tint_symbol_16.sample(tint_symbol_17, ((x_274 * x_275) + x_277));
       nc = x_279;
       float const x_283 = color.w;
       float const x_285 = nc.w;
@@ -189,34 +189,34 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
   float4 const x_314 = color;
   color = float4(x_313.x, x_313.y, x_313.z, x_314.w);
   float4 const x_318 = color;
-  *(tint_symbol_19) = x_318;
+  *(tint_symbol_18) = x_318;
   return;
 }
 
-fragment tint_symbol_3 tint_symbol(texture2d<float, access::sample> tint_symbol_26 [[texture(6)]], sampler tint_symbol_27 [[sampler(4)]], texture2d<float, access::sample> tint_symbol_28 [[texture(5)]], texture2d<float, access::sample> tint_symbol_29 [[texture(8)]], sampler tint_symbol_30 [[sampler(7)]], texture2d<float, access::sample> tint_symbol_32 [[texture(3)]], sampler tint_symbol_33 [[sampler(2)]], texture2d<float, access::sample> tint_symbol_34 [[texture(1)]], sampler tint_symbol_35 [[sampler(0)]], tint_symbol_2 tint_symbol_1 [[stage_in]], constant LeftOver& x_20 [[buffer(9)]]) {
-  thread float2 tint_symbol_20 = 0.0f;
-  thread float2 tint_symbol_21 = 0.0f;
-  thread float2 tint_symbol_22 = 0.0f;
-  thread float2 tint_symbol_23 = 0.0f;
-  thread float3 tint_symbol_24 = 0.0f;
-  thread float2 tint_symbol_25 = 0.0f;
-  thread float tint_symbol_31 = 0.0f;
-  thread float4 tint_symbol_36 = 0.0f;
-  float2 const tUV_param = tint_symbol_1.tUV_param;
-  float2 const tileID_1_param = tint_symbol_1.tileID_1_param;
-  float2 const levelUnits_param = tint_symbol_1.levelUnits_param;
-  float2 const stageUnits_1_param = tint_symbol_1.stageUnits_1_param;
-  float3 const vPosition_param = tint_symbol_1.vPosition_param;
-  float2 const vUV_param = tint_symbol_1.vUV_param;
-  tint_symbol_20 = tUV_param;
-  tint_symbol_21 = tileID_1_param;
-  tint_symbol_22 = levelUnits_param;
-  tint_symbol_23 = stageUnits_1_param;
-  tint_symbol_24 = vPosition_param;
-  tint_symbol_25 = vUV_param;
-  main_1(x_20, &(tint_symbol_20), tint_symbol_26, tint_symbol_27, tint_symbol_28, tint_symbol_29, tint_symbol_30, &(tint_symbol_31), tint_symbol_32, tint_symbol_33, tint_symbol_34, tint_symbol_35, &(tint_symbol_36));
-  main_out const tint_symbol_4 = {.glFragColor_1=tint_symbol_36};
-  tint_symbol_3 const tint_symbol_5 = {.glFragColor_1=tint_symbol_4.glFragColor_1};
-  return tint_symbol_5;
+main_out tint_symbol_inner(constant LeftOver& x_20, float2 tUV_param, float2 tileID_1_param, float2 levelUnits_param, float2 stageUnits_1_param, float3 vPosition_param, float2 vUV_param, thread float2* const tint_symbol_19, thread float2* const tint_symbol_20, thread float2* const tint_symbol_21, thread float2* const tint_symbol_22, thread float3* const tint_symbol_23, thread float2* const tint_symbol_24, texture2d<float, access::sample> tint_symbol_25, sampler tint_symbol_26, texture2d<float, access::sample> tint_symbol_27, texture2d<float, access::sample> tint_symbol_28, sampler tint_symbol_29, thread float* const tint_symbol_30, texture2d<float, access::sample> tint_symbol_31, sampler tint_symbol_32, texture2d<float, access::sample> tint_symbol_33, sampler tint_symbol_34, thread float4* const tint_symbol_35) {
+  *(tint_symbol_19) = tUV_param;
+  *(tint_symbol_20) = tileID_1_param;
+  *(tint_symbol_21) = levelUnits_param;
+  *(tint_symbol_22) = stageUnits_1_param;
+  *(tint_symbol_23) = vPosition_param;
+  *(tint_symbol_24) = vUV_param;
+  main_1(x_20, tint_symbol_19, tint_symbol_25, tint_symbol_26, tint_symbol_27, tint_symbol_28, tint_symbol_29, tint_symbol_30, tint_symbol_31, tint_symbol_32, tint_symbol_33, tint_symbol_34, tint_symbol_35);
+  main_out const tint_symbol_4 = {.glFragColor_1=*(tint_symbol_35)};
+  return tint_symbol_4;
+}
+
+fragment tint_symbol_3 tint_symbol(texture2d<float, access::sample> tint_symbol_42 [[texture(6)]], sampler tint_symbol_43 [[sampler(4)]], texture2d<float, access::sample> tint_symbol_44 [[texture(5)]], texture2d<float, access::sample> tint_symbol_45 [[texture(8)]], sampler tint_symbol_46 [[sampler(7)]], texture2d<float, access::sample> tint_symbol_48 [[texture(3)]], sampler tint_symbol_49 [[sampler(2)]], texture2d<float, access::sample> tint_symbol_50 [[texture(1)]], sampler tint_symbol_51 [[sampler(0)]], tint_symbol_2 tint_symbol_1 [[stage_in]], constant LeftOver& x_20 [[buffer(9)]]) {
+  thread float2 tint_symbol_36 = 0.0f;
+  thread float2 tint_symbol_37 = 0.0f;
+  thread float2 tint_symbol_38 = 0.0f;
+  thread float2 tint_symbol_39 = 0.0f;
+  thread float3 tint_symbol_40 = 0.0f;
+  thread float2 tint_symbol_41 = 0.0f;
+  thread float tint_symbol_47 = 0.0f;
+  thread float4 tint_symbol_52 = 0.0f;
+  main_out const inner_result = tint_symbol_inner(x_20, tint_symbol_1.tUV_param, tint_symbol_1.tileID_1_param, tint_symbol_1.levelUnits_param, tint_symbol_1.stageUnits_1_param, tint_symbol_1.vPosition_param, tint_symbol_1.vUV_param, &(tint_symbol_36), &(tint_symbol_37), &(tint_symbol_38), &(tint_symbol_39), &(tint_symbol_40), &(tint_symbol_41), tint_symbol_42, tint_symbol_43, tint_symbol_44, tint_symbol_45, tint_symbol_46, &(tint_symbol_47), tint_symbol_48, tint_symbol_49, tint_symbol_50, tint_symbol_51, &(tint_symbol_52));
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.glFragColor_1 = inner_result.glFragColor_1;
+  return wrapper_result;
 }
 

test/bug/tint/949.wgsl.expected.hlsl

@@ -324,19 +324,20 @@ struct tint_symbol_2 {
   float4 glFragColor_1 : SV_Target0;
 };
 
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
-  const float2 vMainuv_param = tint_symbol.vMainuv_param;
-  const float4 v_output1_param = tint_symbol.v_output1_param;
-  const bool gl_FrontFacing_param = tint_symbol.gl_FrontFacing_param;
-  const float2 v_uv_param = tint_symbol.v_uv_param;
-  const float4 v_output2_param = tint_symbol.v_output2_param;
+main_out main_inner(float2 vMainuv_param, float4 v_output1_param, bool gl_FrontFacing_param, float2 v_uv_param, float4 v_output2_param) {
   vMainuv = vMainuv_param;
   v_output1 = v_output1_param;
   gl_FrontFacing = gl_FrontFacing_param;
   v_uv = v_uv_param;
   v_output2 = v_output2_param;
   main_1();
-  const main_out tint_symbol_3 = {glFragColor};
-  const tint_symbol_2 tint_symbol_9 = {tint_symbol_3.glFragColor_1};
-  return tint_symbol_9;
+  const main_out tint_symbol_8 = {glFragColor};
+  return tint_symbol_8;
+}
+
+tint_symbol_2 main(tint_symbol_1 tint_symbol) {
+  const main_out inner_result = main_inner(tint_symbol.vMainuv_param, tint_symbol.v_output1_param, tint_symbol.gl_FrontFacing_param, tint_symbol.v_uv_param, tint_symbol.v_output2_param);
+  tint_symbol_2 wrapper_result = (tint_symbol_2)0;
+  wrapper_result.glFragColor_1 = inner_result.glFragColor_1;
+  return wrapper_result;
 }

test/bug/tint/949.wgsl.expected.msl

@@ -171,7 +171,7 @@ lightingInfo computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1_(thread float
   return x_245;
 }
 
-void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* const tint_symbol_6, thread float3* const tint_symbol_7, thread float2* const tint_symbol_8, texture2d<float, access::sample> tint_symbol_9, sampler tint_symbol_10, thread float4* const tint_symbol_11, thread bool* const tint_symbol_12, thread float2* const tint_symbol_13, thread float4* const tint_symbol_14, texture2d<float, access::sample> tint_symbol_15, sampler tint_symbol_16, thread float4* const tint_symbol_17) {
+void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* const tint_symbol_5, thread float3* const tint_symbol_6, thread float2* const tint_symbol_7, texture2d<float, access::sample> tint_symbol_8, sampler tint_symbol_9, thread float4* const tint_symbol_10, thread bool* const tint_symbol_11, thread float2* const tint_symbol_12, thread float4* const tint_symbol_13, texture2d<float, access::sample> tint_symbol_14, sampler tint_symbol_15, thread float4* const tint_symbol_16) {
   float4 tempTextureRead = 0.0f;
   float3 rgb = 0.0f;
   float3 output5 = 0.0f;
@@ -225,35 +225,35 @@ void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* con
   float3 diffuseOutput = 0.0f;
   float3 specularOutput = 0.0f;
   float3 output3 = 0.0f;
-  *(tint_symbol_6) = 100.0f;
-  *(tint_symbol_7) = float3(0.5f, 0.5f, 0.5f);
-  float2 const x_261 = *(tint_symbol_8);
-  float4 const x_262 = tint_symbol_9.sample(tint_symbol_10, x_261);
+  *(tint_symbol_5) = 100.0f;
+  *(tint_symbol_6) = float3(0.5f, 0.5f, 0.5f);
+  float2 const x_261 = *(tint_symbol_7);
+  float4 const x_262 = tint_symbol_8.sample(tint_symbol_9, x_261);
   tempTextureRead = x_262;
   float4 const x_264 = tempTextureRead;
   float const x_273 = x_269.textureInfoName;
   rgb = (float3(x_264.x, x_264.y, x_264.z) * x_273);
   float3 const x_279 = x_269.u_cameraPosition;
-  float4 const x_282 = *(tint_symbol_11);
+  float4 const x_282 = *(tint_symbol_10);
   output5 = normalize((x_279 - float3(x_282.x, x_282.y, x_282.z)));
   output4 = float4(0.0f, 0.0f, 0.0f, 0.0f);
   uvOffset = float2(0.0f, 0.0f);
   float const x_292 = x_269.u_bumpStrength;
   normalScale = (1.0f / x_292);
-  bool const x_298 = *(tint_symbol_12);
+  bool const x_298 = *(tint_symbol_11);
   if (x_298) {
-    float2 const x_303 = *(tint_symbol_13);
+    float2 const x_303 = *(tint_symbol_12);
     x_299 = x_303;
   } else {
-    float2 const x_305 = *(tint_symbol_13);
+    float2 const x_305 = *(tint_symbol_12);
     x_299 = -(x_305);
   }
   float2 const x_307 = x_299;
   TBNUV = x_307;
-  float4 const x_310 = *(tint_symbol_14);
+  float4 const x_310 = *(tint_symbol_13);
   float const x_312 = normalScale;
   param_3 = (float3(x_310.x, x_310.y, x_310.z) * x_312);
-  float4 const x_317 = *(tint_symbol_11);
+  float4 const x_317 = *(tint_symbol_10);
   param_4 = float3(x_317.x, x_317.y, x_317.z);
   float2 const x_320 = TBNUV;
   param_5 = x_320;
@@ -284,7 +284,7 @@ void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* con
   float3x3 const x_361 = invTBN;
   float3 const x_362 = output5;
   float3x3 const x_365 = invTBN;
-  float4 const x_366 = *(tint_symbol_14);
+  float4 const x_366 = *(tint_symbol_13);
   numSamples = (15.0f + (dot((x_361 * -(x_362)), (x_365 * float3(x_366.x, x_366.y, x_366.z))) * -11.0f));
   float const x_374 = numSamples;
   stepSize = (1.0f / x_374);
@@ -300,9 +300,9 @@ void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* con
     } else {
       break;
     }
-    float2 const x_394 = *(tint_symbol_13);
+    float2 const x_394 = *(tint_symbol_12);
     float2 const x_395 = vCurrOffset;
-    float4 const x_397 = tint_symbol_9.sample(tint_symbol_10, (x_394 + x_395));
+    float4 const x_397 = tint_symbol_8.sample(tint_symbol_9, (x_394 + x_395));
     currSampledHeight = x_397.w;
     float const x_400 = currSampledHeight;
     float const x_401 = currRayHeight;
@@ -346,9 +346,9 @@ void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* con
   parallaxOcclusion_0 = x_444;
   float2 const x_445 = parallaxOcclusion_0;
   uvOffset = x_445;
-  float2 const x_449 = *(tint_symbol_13);
+  float2 const x_449 = *(tint_symbol_12);
   float2 const x_450 = uvOffset;
-  float4 const x_452 = tint_symbol_9.sample(tint_symbol_10, (x_449 + x_450));
+  float4 const x_452 = tint_symbol_8.sample(tint_symbol_9, (x_449 + x_450));
   float const x_454 = x_269.u_bumpStrength;
   float3x3 const x_457 = TBN;
   param_8 = x_457;
@@ -357,19 +357,19 @@ void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* con
   float3 const x_461 = perturbNormal_mf33_vf3_f1_(&(param_8), &(param_9), &(param_10));
   float4 const x_462 = output4;
   output4 = float4(x_461.x, x_461.y, x_461.z, x_462.w);
-  float2 const x_465 = *(tint_symbol_13);
+  float2 const x_465 = *(tint_symbol_12);
   float2 const x_466 = uvOffset;
   output6 = (x_465 + x_466);
   float2 const x_474 = output6;
-  float4 const x_475 = tint_symbol_15.sample(tint_symbol_16, x_474);
+  float4 const x_475 = tint_symbol_14.sample(tint_symbol_15, x_474);
   tempTextureRead1 = x_475;
   float4 const x_477 = tempTextureRead1;
   rgb1 = float3(x_477.x, x_477.y, x_477.z);
   float3 const x_481 = x_269.u_cameraPosition;
-  float4 const x_482 = *(tint_symbol_11);
+  float4 const x_482 = *(tint_symbol_10);
   viewDirectionW_1 = normalize((x_481 - float3(x_482.x, x_482.y, x_482.z)));
   shadow = 1.0f;
-  float const x_488 = *(tint_symbol_6);
+  float const x_488 = *(tint_symbol_5);
   glossiness_1 = (1.0f * x_488);
   diffuseBase = float3(0.0f, 0.0f, 0.0f);
   specularBase = float3(0.0f, 0.0f, 0.0f);
@@ -404,37 +404,39 @@ void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* con
   float3 const x_536 = rgb1;
   diffuseOutput = (x_535 * x_536);
   float3 const x_539 = specularBase;
-  float3 const x_540 = *(tint_symbol_7);
+  float3 const x_540 = *(tint_symbol_6);
   specularOutput = (x_539 * x_540);
   float3 const x_543 = diffuseOutput;
   float3 const x_544 = specularOutput;
   output3 = (x_543 + x_544);
   float3 const x_548 = output3;
-  *(tint_symbol_17) = float4(x_548.x, x_548.y, x_548.z, 1.0f);
+  *(tint_symbol_16) = float4(x_548.x, x_548.y, x_548.z, 1.0f);
   return;
 }
 
-fragment tint_symbol_3 tint_symbol(texture2d<float, access::sample> tint_symbol_25 [[texture(1)]], sampler tint_symbol_26 [[sampler(0)]], texture2d<float, access::sample> tint_symbol_27 [[texture(3)]], sampler tint_symbol_28 [[sampler(2)]], bool gl_FrontFacing_param [[front_facing]], tint_symbol_2 tint_symbol_1 [[stage_in]], constant LeftOver& x_269 [[buffer(6)]], constant Light0& light0 [[buffer(5)]]) {
-  thread float2 tint_symbol_18 = 0.0f;
-  thread float4 tint_symbol_19 = 0.0f;
-  thread bool tint_symbol_20 = false;
-  thread float2 tint_symbol_21 = 0.0f;
-  thread float4 tint_symbol_22 = 0.0f;
-  thread float tint_symbol_23 = 0.0f;
-  thread float3 tint_symbol_24 = 0.0f;
-  thread float4 tint_symbol_29 = 0.0f;
-  float2 const vMainuv_param = tint_symbol_1.vMainuv_param;
-  float4 const v_output1_param = tint_symbol_1.v_output1_param;
-  float2 const v_uv_param = tint_symbol_1.v_uv_param;
-  float4 const v_output2_param = tint_symbol_1.v_output2_param;
-  tint_symbol_18 = vMainuv_param;
-  tint_symbol_19 = v_output1_param;
-  tint_symbol_20 = gl_FrontFacing_param;
-  tint_symbol_21 = v_uv_param;
-  tint_symbol_22 = v_output2_param;
-  main_1(x_269, light0, &(tint_symbol_23), &(tint_symbol_24), &(tint_symbol_18), tint_symbol_25, tint_symbol_26, &(tint_symbol_19), &(tint_symbol_20), &(tint_symbol_21), &(tint_symbol_22), tint_symbol_27, tint_symbol_28, &(tint_symbol_29));
-  main_out const tint_symbol_4 = {.glFragColor_1=tint_symbol_29};
-  tint_symbol_3 const tint_symbol_5 = {.glFragColor_1=tint_symbol_4.glFragColor_1};
-  return tint_symbol_5;
+main_out tint_symbol_inner(constant LeftOver& x_269, constant Light0& light0, float2 vMainuv_param, float4 v_output1_param, bool gl_FrontFacing_param, float2 v_uv_param, float4 v_output2_param, thread float2* const tint_symbol_17, thread float4* const tint_symbol_18, thread bool* const tint_symbol_19, thread float2* const tint_symbol_20, thread float4* const tint_symbol_21, thread float* const tint_symbol_22, thread float3* const tint_symbol_23, texture2d<float, access::sample> tint_symbol_24, sampler tint_symbol_25, texture2d<float, access::sample> tint_symbol_26, sampler tint_symbol_27, thread float4* const tint_symbol_28) {
+  *(tint_symbol_17) = vMainuv_param;
+  *(tint_symbol_18) = v_output1_param;
+  *(tint_symbol_19) = gl_FrontFacing_param;
+  *(tint_symbol_20) = v_uv_param;
+  *(tint_symbol_21) = v_output2_param;
+  main_1(x_269, light0, tint_symbol_22, tint_symbol_23, tint_symbol_17, tint_symbol_24, tint_symbol_25, tint_symbol_18, tint_symbol_19, tint_symbol_20, tint_symbol_21, tint_symbol_26, tint_symbol_27, tint_symbol_28);
+  main_out const tint_symbol_4 = {.glFragColor_1=*(tint_symbol_28)};
+  return tint_symbol_4;
+}
+
+fragment tint_symbol_3 tint_symbol(texture2d<float, access::sample> tint_symbol_36 [[texture(1)]], sampler tint_symbol_37 [[sampler(0)]], texture2d<float, access::sample> tint_symbol_38 [[texture(3)]], sampler tint_symbol_39 [[sampler(2)]], bool gl_FrontFacing_param [[front_facing]], tint_symbol_2 tint_symbol_1 [[stage_in]], constant LeftOver& x_269 [[buffer(6)]], constant Light0& light0 [[buffer(5)]]) {
+  thread float2 tint_symbol_29 = 0.0f;
+  thread float4 tint_symbol_30 = 0.0f;
+  thread bool tint_symbol_31 = false;
+  thread float2 tint_symbol_32 = 0.0f;
+  thread float4 tint_symbol_33 = 0.0f;
+  thread float tint_symbol_34 = 0.0f;
+  thread float3 tint_symbol_35 = 0.0f;
+  thread float4 tint_symbol_40 = 0.0f;
+  main_out const inner_result = tint_symbol_inner(x_269, light0, tint_symbol_1.vMainuv_param, tint_symbol_1.v_output1_param, gl_FrontFacing_param, tint_symbol_1.v_uv_param, tint_symbol_1.v_output2_param, &(tint_symbol_29), &(tint_symbol_30), &(tint_symbol_31), &(tint_symbol_32), &(tint_symbol_33), &(tint_symbol_34), &(tint_symbol_35), tint_symbol_36, tint_symbol_37, tint_symbol_38, tint_symbol_39, &(tint_symbol_40));
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.glFragColor_1 = inner_result.glFragColor_1;
+  return wrapper_result;
 }
 

test/bug/tint/951.spvasm.expected.hlsl

@@ -53,10 +53,13 @@ struct tint_symbol_1 {
   uint3 gl_GlobalInvocationID_param : SV_DispatchThreadID;
 };
 
-[numthreads(128, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
+void main_inner(uint3 gl_GlobalInvocationID_param) {
   gl_GlobalInvocationID = gl_GlobalInvocationID_param;
   main_1();
+}
+
+[numthreads(128, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.gl_GlobalInvocationID_param);
   return;
 }

test/bug/tint/951.spvasm.expected.msl

@@ -15,8 +15,8 @@ struct Uniforms {
   /* 0x0010 */ int size;
 };
 
-float getAAtOutCoords_(const device ssbA& x_20, thread uint3* const tint_symbol_3) {
-  uint const x_42 = (*(tint_symbol_3)).x;
+float getAAtOutCoords_(const device ssbA& x_20, thread uint3* const tint_symbol_2) {
+  uint const x_42 = (*(tint_symbol_2)).x;
   float const x_44 = x_20.A[x_42];
   return x_44;
 }
@@ -37,18 +37,18 @@ void setOutput_i1_f1_(device ssbOut& x_16, thread int* const flatIndex, thread f
   return;
 }
 
-void main_1(constant Uniforms& x_24, const device ssbA& x_20, device ssbOut& x_16, thread uint3* const tint_symbol_4) {
+void main_1(constant Uniforms& x_24, const device ssbA& x_20, device ssbOut& x_16, thread uint3* const tint_symbol_3) {
   int index = 0;
   float a_1 = 0.0f;
   float param = 0.0f;
   int param_1 = 0;
   float param_2 = 0.0f;
-  uint const x_61 = (*(tint_symbol_4)).x;
+  uint const x_61 = (*(tint_symbol_3)).x;
   index = as_type<int>(x_61);
   int const x_63 = index;
   int const x_70 = x_24.size;
   if ((x_63 < x_70)) {
-    float const x_75 = getAAtOutCoords_(x_20, tint_symbol_4);
+    float const x_75 = getAAtOutCoords_(x_20, tint_symbol_3);
     a_1 = x_75;
     float const x_77 = a_1;
     param = x_77;
@@ -61,10 +61,14 @@ void main_1(constant Uniforms& x_24, const device ssbA& x_20, device ssbOut& x_1
   return;
 }
 
+void tint_symbol_1_inner(constant Uniforms& x_24, const device ssbA& x_20, device ssbOut& x_16, uint3 gl_GlobalInvocationID_param, thread uint3* const tint_symbol_4) {
+  *(tint_symbol_4) = gl_GlobalInvocationID_param;
+  main_1(x_24, x_20, x_16, tint_symbol_4);
+}
+
 kernel void tint_symbol_1(uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], constant Uniforms& x_24 [[buffer(2)]], const device ssbA& x_20 [[buffer(1)]], device ssbOut& x_16 [[buffer(0)]]) {
   thread uint3 tint_symbol_5 = 0u;
-  tint_symbol_5 = gl_GlobalInvocationID_param;
-  main_1(x_24, x_20, x_16, &(tint_symbol_5));
+  tint_symbol_1_inner(x_24, x_20, x_16, gl_GlobalInvocationID_param, &(tint_symbol_5));
   return;
 }
 

test/bug/tint/977.spvasm.expected.hlsl

@@ -46,10 +46,13 @@ struct tint_symbol_1 {
   uint3 gl_GlobalInvocationID_param : SV_DispatchThreadID;
 };
 
-[numthreads(1, 1, 1)]
-void main(tint_symbol_1 tint_symbol) {
-  const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
+void main_inner(uint3 gl_GlobalInvocationID_param) {
   gl_GlobalInvocationID = gl_GlobalInvocationID_param;
   main_1();
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.gl_GlobalInvocationID_param);
   return;
 }

test/bug/tint/977.spvasm.expected.msl

@@ -37,12 +37,12 @@ float binaryOperation_f1_f1_(thread float* const a, thread float* const b) {
   return x_41;
 }
 
-void main_1(device ResultMatrix& resultMatrix, thread uint3* const tint_symbol_3) {
+void main_1(device ResultMatrix& resultMatrix, thread uint3* const tint_symbol_2) {
   int index = 0;
   int a_1 = 0;
   float param = 0.0f;
   float param_1 = 0.0f;
-  uint const x_54 = (*(tint_symbol_3)).x;
+  uint const x_54 = (*(tint_symbol_2)).x;
   index = as_type<int>(x_54);
   a_1 = -10;
   int const x_63 = index;
@@ -53,10 +53,14 @@ void main_1(device ResultMatrix& resultMatrix, thread uint3* const tint_symbol_3
   return;
 }
 
+void tint_symbol_1_inner(device ResultMatrix& resultMatrix, uint3 gl_GlobalInvocationID_param, thread uint3* const tint_symbol_3) {
+  *(tint_symbol_3) = gl_GlobalInvocationID_param;
+  main_1(resultMatrix, tint_symbol_3);
+}
+
 kernel void tint_symbol_1(uint3 gl_GlobalInvocationID_param [[thread_position_in_grid]], device ResultMatrix& resultMatrix [[buffer(2)]]) {
   thread uint3 tint_symbol_4 = 0u;
-  tint_symbol_4 = gl_GlobalInvocationID_param;
-  main_1(resultMatrix, &(tint_symbol_4));
+  tint_symbol_1_inner(resultMatrix, gl_GlobalInvocationID_param, &(tint_symbol_4));
   return;
 }
 

test/bug/tint/978.wgsl.expected.hlsl

@@ -15,12 +15,18 @@ struct tint_symbol_3 {
   float4 color : SV_Target0;
 };
 
-tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
-  const FragmentInput fIn = {tint_symbol_1.vUv};
+FragmentOutput main_inner(FragmentInput fIn) {
   const float tint_symbol = depthMap.Sample(texSampler, fIn.vUv).x;
   const float3 color = float3(tint_symbol, tint_symbol, tint_symbol);
   FragmentOutput fOut = (FragmentOutput)0;
   fOut.color = float4(color, 1.0f);
-  const tint_symbol_3 tint_symbol_4 = {fOut.color};
-  return tint_symbol_4;
+  return fOut;
+}
+
+tint_symbol_3 main(tint_symbol_2 tint_symbol_1) {
+  const FragmentInput tint_symbol_4 = {tint_symbol_1.vUv};
+  const FragmentOutput inner_result = main_inner(tint_symbol_4);
+  tint_symbol_3 wrapper_result = (tint_symbol_3)0;
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
 }

test/bug/tint/978.wgsl.expected.msl

@@ -14,13 +14,19 @@ struct tint_symbol_3 {
   float4 color [[color(0)]];
 };
 
-fragment tint_symbol_3 tint_symbol(depth2d<float, access::sample> tint_symbol_5 [[texture(5)]], sampler tint_symbol_6 [[sampler(3)]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
-  FragmentInput const fIn = {.vUv=tint_symbol_1.vUv};
+FragmentOutput tint_symbol_inner(FragmentInput fIn, depth2d<float, access::sample> tint_symbol_5, sampler tint_symbol_6) {
   float const sample = tint_symbol_5.sample(tint_symbol_6, fIn.vUv);
   float3 const color = float3(sample, sample, sample);
   FragmentOutput fOut = {};
   fOut.color = float4(color, 1.0f);
-  tint_symbol_3 const tint_symbol_4 = {.color=fOut.color};
-  return tint_symbol_4;
+  return fOut;
+}
+
+fragment tint_symbol_3 tint_symbol(depth2d<float, access::sample> tint_symbol_7 [[texture(5)]], sampler tint_symbol_8 [[sampler(3)]], tint_symbol_2 tint_symbol_1 [[stage_in]]) {
+  FragmentInput const tint_symbol_4 = {.vUv=tint_symbol_1.vUv};
+  FragmentOutput const inner_result = tint_symbol_inner(tint_symbol_4, tint_symbol_7, tint_symbol_8);
+  tint_symbol_3 wrapper_result = {};
+  wrapper_result.color = inner_result.color;
+  return wrapper_result;
 }
 

test/bug/tint/980.wgsl.expected.hlsl

@@ -14,9 +14,12 @@ struct tint_symbol_1 {
   uint idx : SV_GroupIndex;
 };
 
+void main_inner(uint idx) {
+  io.Store3(0u, asuint(Bad(io.Load(12u), asfloat(io.Load3(0u)))));
+}
+
 [numthreads(1, 1, 1)]
 void main(tint_symbol_1 tint_symbol) {
-  const uint idx = tint_symbol.idx;
-  io.Store3(0u, asuint(Bad(io.Load(12u), asfloat(io.Load3(0u)))));
+  main_inner(tint_symbol.idx);
   return;
 }

test/bug/tint/980.wgsl.expected.msl

@@ -12,8 +12,12 @@ float3 Bad(uint index, float3 rd) {
   return normalize(normal);
 }
 
-kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& io [[buffer(0)]]) {
+void tint_symbol_inner(device S& io, uint idx) {
   io.v = Bad(io.i, io.v);
+}
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& io [[buffer(0)]]) {
+  tint_symbol_inner(io, idx);
   return;
 }
 

test/bug/tint/992.wgsl.expected.hlsl

@@ -2,9 +2,15 @@ struct tint_symbol {
   float4 value : SV_Target0;
 };
 
-tint_symbol frag_main() {
+float4 frag_main_inner() {
   float b = 0.0f;
   float3 v = float3((b).xxx);
-  const tint_symbol tint_symbol_1 = {float4(v, 1.0f)};
-  return tint_symbol_1;
+  return float4(v, 1.0f);
+}
+
+tint_symbol frag_main() {
+  const float4 inner_result = frag_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }

test/bug/tint/992.wgsl.expected.msl

@@ -5,10 +5,16 @@ struct tint_symbol {
   float4 value [[color(0)]];
 };
 
-fragment tint_symbol frag_main() {
+float4 frag_main_inner() {
   float b = 0.0f;
   float3 v = float3(b);
-  tint_symbol const tint_symbol_1 = {.value=float4(v, 1.0f)};
-  return tint_symbol_1;
+  return float4(v, 1.0f);
+}
+
+fragment tint_symbol frag_main() {
+  float4 const inner_result = frag_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 

test/expressions/literals/-inf.spvasm.expected.hlsl

@@ -12,9 +12,15 @@ struct tint_symbol {
   float4 out_var_SV_TARGET_1 : SV_Target0;
 };
 
-tint_symbol main() {
+main_out main_inner() {
   main_1();
   const main_out tint_symbol_1 = {out_var_SV_TARGET};
-  const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
-  return tint_symbol_2;
+  return tint_symbol_1;
+}
+
+tint_symbol main() {
+  const main_out inner_result = main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.out_var_SV_TARGET_1 = inner_result.out_var_SV_TARGET_1;
+  return wrapper_result;
 }

test/expressions/literals/-inf.spvasm.expected.msl

@@ -8,16 +8,22 @@ struct tint_symbol_1 {
   float4 out_var_SV_TARGET_1 [[color(0)]];
 };
 
-void main_1(thread float4* const tint_symbol_4) {
-  *(tint_symbol_4) = float4(-INFINITY, -INFINITY, -INFINITY, -INFINITY);
+void main_1(thread float4* const tint_symbol_3) {
+  *(tint_symbol_3) = float4(-INFINITY, -INFINITY, -INFINITY, -INFINITY);
   return;
 }
 
+main_out tint_symbol_inner(thread float4* const tint_symbol_4) {
+  main_1(tint_symbol_4);
+  main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=*(tint_symbol_4)};
+  return tint_symbol_2;
+}
+
 fragment tint_symbol_1 tint_symbol() {
   thread float4 tint_symbol_5 = 0.0f;
-  main_1(&(tint_symbol_5));
-  main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
-  tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
-  return tint_symbol_3;
+  main_out const inner_result = tint_symbol_inner(&(tint_symbol_5));
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.out_var_SV_TARGET_1 = inner_result.out_var_SV_TARGET_1;
+  return wrapper_result;
 }
 

test/expressions/literals/inf.spvasm.expected.hlsl

@@ -12,9 +12,15 @@ struct tint_symbol {
   float4 out_var_SV_TARGET_1 : SV_Target0;
 };
 
-tint_symbol main() {
+main_out main_inner() {
   main_1();
   const main_out tint_symbol_1 = {out_var_SV_TARGET};
-  const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
-  return tint_symbol_2;
+  return tint_symbol_1;
+}
+
+tint_symbol main() {
+  const main_out inner_result = main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.out_var_SV_TARGET_1 = inner_result.out_var_SV_TARGET_1;
+  return wrapper_result;
 }

test/expressions/literals/inf.spvasm.expected.msl

@@ -8,16 +8,22 @@ struct tint_symbol_1 {
   float4 out_var_SV_TARGET_1 [[color(0)]];
 };
 
-void main_1(thread float4* const tint_symbol_4) {
-  *(tint_symbol_4) = float4(INFINITY, INFINITY, INFINITY, INFINITY);
+void main_1(thread float4* const tint_symbol_3) {
+  *(tint_symbol_3) = float4(INFINITY, INFINITY, INFINITY, INFINITY);
   return;
 }
 
+main_out tint_symbol_inner(thread float4* const tint_symbol_4) {
+  main_1(tint_symbol_4);
+  main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=*(tint_symbol_4)};
+  return tint_symbol_2;
+}
+
 fragment tint_symbol_1 tint_symbol() {
   thread float4 tint_symbol_5 = 0.0f;
-  main_1(&(tint_symbol_5));
-  main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
-  tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
-  return tint_symbol_3;
+  main_out const inner_result = tint_symbol_inner(&(tint_symbol_5));
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.out_var_SV_TARGET_1 = inner_result.out_var_SV_TARGET_1;
+  return wrapper_result;
 }
 

test/expressions/literals/nan.spvasm.expected.hlsl

@@ -12,9 +12,15 @@ struct tint_symbol {
   float4 out_var_SV_TARGET_1 : SV_Target0;
 };
 
-tint_symbol main() {
+main_out main_inner() {
   main_1();
   const main_out tint_symbol_1 = {out_var_SV_TARGET};
-  const tint_symbol tint_symbol_2 = {tint_symbol_1.out_var_SV_TARGET_1};
-  return tint_symbol_2;
+  return tint_symbol_1;
+}
+
+tint_symbol main() {
+  const main_out inner_result = main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.out_var_SV_TARGET_1 = inner_result.out_var_SV_TARGET_1;
+  return wrapper_result;
 }

test/expressions/literals/nan.spvasm.expected.msl

@@ -8,16 +8,22 @@ struct tint_symbol_1 {
   float4 out_var_SV_TARGET_1 [[color(0)]];
 };
 
-void main_1(thread float4* const tint_symbol_4) {
-  *(tint_symbol_4) = float4(NAN, NAN, NAN, NAN);
+void main_1(thread float4* const tint_symbol_3) {
+  *(tint_symbol_3) = float4(NAN, NAN, NAN, NAN);
   return;
 }
 
+main_out tint_symbol_inner(thread float4* const tint_symbol_4) {
+  main_1(tint_symbol_4);
+  main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=*(tint_symbol_4)};
+  return tint_symbol_2;
+}
+
 fragment tint_symbol_1 tint_symbol() {
   thread float4 tint_symbol_5 = 0.0f;
-  main_1(&(tint_symbol_5));
-  main_out const tint_symbol_2 = {.out_var_SV_TARGET_1=tint_symbol_5};
-  tint_symbol_1 const tint_symbol_3 = {.out_var_SV_TARGET_1=tint_symbol_2.out_var_SV_TARGET_1};
-  return tint_symbol_3;
+  main_out const inner_result = tint_symbol_inner(&(tint_symbol_5));
+  tint_symbol_1 wrapper_result = {};
+  wrapper_result.out_var_SV_TARGET_1 = inner_result.out_var_SV_TARGET_1;
+  return wrapper_result;
 }
 

test/intrinsics/gen/abs/002533.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_002533();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/002533.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_002533() {
   float4 res = fabs(float4());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_002533();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/005174.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_005174();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/005174.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_005174() {
   float3 res = fabs(float3());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_005174();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/1ce782.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_1ce782();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/1ce782.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_1ce782() {
   uint4 res = abs(uint4());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_1ce782();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/1e9d53.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_1e9d53();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/1e9d53.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_1e9d53() {
   float2 res = fabs(float2());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_1e9d53();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/467cd1.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_467cd1();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/467cd1.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_467cd1() {
   uint res = abs(1u);
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_467cd1();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/4ad288.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_4ad288();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/4ad288.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_4ad288() {
   int res = abs(1);
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_4ad288();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/5ad50a.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_5ad50a();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/5ad50a.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_5ad50a() {
   int3 res = abs(int3());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_5ad50a();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/7326de.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_7326de();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/7326de.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_7326de() {
   uint3 res = abs(uint3());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_7326de();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/7f28e6.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_7f28e6();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/7f28e6.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_7f28e6() {
   uint2 res = abs(uint2());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_7f28e6();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/7faa9e.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_7faa9e();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/7faa9e.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_7faa9e() {
   int2 res = abs(int2());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_7faa9e();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/9c80a6.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_9c80a6();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/9c80a6.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_9c80a6() {
   int4 res = abs(int4());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_9c80a6();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/abs/b96037.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_b96037();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/abs/b96037.wgsl.expected.msl

@@ -9,10 +9,16 @@ void abs_b96037() {
   float res = fabs(1.0f);
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   abs_b96037();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/acos/489247.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   acos_489247();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/acos/489247.wgsl.expected.msl

@@ -9,10 +9,16 @@ void acos_489247() {
   float res = acos(1.0f);
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   acos_489247();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {

test/intrinsics/gen/acos/8e2acf.wgsl.expected.hlsl

@@ -6,10 +6,16 @@ struct tint_symbol {
   float4 value : SV_Position;
 };
 
-tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   acos_8e2acf();
-  const tint_symbol tint_symbol_1 = {float4(0.0f, 0.0f, 0.0f, 0.0f)};
-  return tint_symbol_1;
+  return float4(0.0f, 0.0f, 0.0f, 0.0f);
+}
+
+tint_symbol vertex_main() {
+  const float4 inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = (tint_symbol)0;
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 void fragment_main() {

test/intrinsics/gen/acos/8e2acf.wgsl.expected.msl

@@ -9,10 +9,16 @@ void acos_8e2acf() {
   float4 res = acos(float4());
 }
 
-vertex tint_symbol vertex_main() {
+float4 vertex_main_inner() {
   acos_8e2acf();
-  tint_symbol const tint_symbol_1 = {.value=float4()};
-  return tint_symbol_1;
+  return float4();
+}
+
+vertex tint_symbol vertex_main() {
+  float4 const inner_result = vertex_main_inner();
+  tint_symbol wrapper_result = {};
+  wrapper_result.value = inner_result;
+  return wrapper_result;
 }
 
 fragment void fragment_main() {