msl: Use a struct for threadgroup memory arguments

MSL has a limit on the number of threadgroup memory arguments, so use
a struct to support an arbitrary number of workgroup variables.

This commit introduces a `State` object to this transform, which is
used to track which structs have been cloned eagerly, in order to
avoid duplicating them.

Bug: tint:938
Change-Id: Ia467db186e176a08f160455eab5fd3b3662f56b8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65360
Auto-Submit: James Price <jrprice@google.com>
Kokoro: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/ast/disable_validation_decoration.cc

@@ -42,6 +42,8 @@ std::string DisableValidationDecoration::InternalName() const {
       return "disable_validation__ignore_constructible_function_parameter";
     case DisabledValidation::kIgnoreStrideDecoration:
       return "disable_validation__ignore_stride";
+    case DisabledValidation::kIgnoreInvalidPointerArgument:
+      return "disable_validation__ignore_invalid_pointer_argument";
   }
   return "<invalid>";
 }

src/ast/disable_validation_decoration.h

@@ -43,6 +43,10 @@ enum class DisabledValidation {
   /// When applied to a member decoration, a stride decoration may be applied to
   /// non-array types.
   kIgnoreStrideDecoration,
+  /// When applied to a pointer function parameter, the validator will not
+  /// require a function call argument passed for that parameter to have a
+  /// certain form.
+  kIgnoreInvalidPointerArgument,
 };
 
 /// An internal decoration used to tell the validator to ignore specific

src/resolver/resolver.cc

@@ -2659,7 +2659,10 @@ bool Resolver::ValidateFunctionCall(const ast::CallExpression* call,
         }
       }
 
-      if (!is_valid) {
+      if (!is_valid &&
+          IsValidationEnabled(
+              param->declaration->decorations(),
+              ast::DisabledValidation::kIgnoreInvalidPointerArgument)) {
         AddError(
             "expected an address-of expression of a variable identifier "
             "expression or a function parameter",

src/transform/module_scope_var_to_entry_point_param.cc

@@ -49,110 +49,164 @@ bool ContainsMatrix(const sem::Type* type) {
 }
 }  // namespace
 
-ModuleScopeVarToEntryPointParam::ModuleScopeVarToEntryPointParam() = default;
+/// State holds the current transform state.
+struct ModuleScopeVarToEntryPointParam::State {
+  /// The clone context.
+  CloneContext& ctx;
 
-ModuleScopeVarToEntryPointParam::~ModuleScopeVarToEntryPointParam() = default;
+  /// Constructor
+  /// @param context the clone context
+  explicit State(CloneContext& context) : ctx(context) {}
 
-void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx,
+  /// Clone any struct types that are contained in `ty` (including `ty` itself),
-                                          const DataMap&,
+  /// and add it to the global declarations now, so that they precede new global
-                                          DataMap&) {
+  /// declarations that need to reference them.
-  // Predetermine the list of function calls that need to be replaced.
+  /// @param ty the type to clone
-  using CallList = std::vector<const ast::CallExpression*>;
+  void CloneStructTypes(const sem::Type* ty) {
-  std::unordered_map<const ast::Function*, CallList> calls_to_replace;
+    if (auto* str = ty->As<sem::Struct>()) {
-
+      if (!cloned_structs_.emplace(str).second) {
-  std::vector<ast::Function*> functions_to_process;
+        // The struct has already been cloned.
-
+        return;
-  // Build a list of functions that transitively reference any private or
-  // workgroup variables, or texture/sampler variables.
-  for (auto* func_ast : ctx.src->AST().Functions()) {
-    auto* func_sem = ctx.src->Sem().Get(func_ast);
-
-    bool needs_processing = false;
-    for (auto* var : func_sem->ReferencedModuleVariables()) {
-      if (var->StorageClass() == ast::StorageClass::kPrivate ||
-          var->StorageClass() == ast::StorageClass::kWorkgroup ||
-          var->StorageClass() == ast::StorageClass::kUniformConstant) {
-        needs_processing = true;
-        break;
       }
-    }
 
-    if (needs_processing) {
+      // Recurse into members.
-      functions_to_process.push_back(func_ast);
+      for (auto* member : str->Members()) {
-
+        CloneStructTypes(member->Type());
-      // Find all of the calls to this function that will need to be replaced.
-      for (auto* call : func_sem->CallSites()) {
-        auto* call_sem = ctx.src->Sem().Get(call);
-        calls_to_replace[call_sem->Stmt()->Function()].push_back(call);
       }
+
+      // Clone the struct and add it to the global declaration list.
+      // Remove the old declaration.
+      auto* ast_str = str->Declaration();
+      ctx.dst->AST().AddTypeDecl(ctx.Clone(const_cast<ast::Struct*>(ast_str)));
+      ctx.Remove(ctx.src->AST().GlobalDeclarations(), ast_str);
+    } else if (auto* arr = ty->As<sem::Array>()) {
+      CloneStructTypes(arr->ElemType());
     }
   }
 
-  // Build a list of `&ident` expressions. We'll use this later to avoid
+  /// Process the module.
-  // generating expressions of the form `&*ident`, which break WGSL validation
+  void Process() {
-  // rules when this expression is passed to a function.
+    // Predetermine the list of function calls that need to be replaced.
-  // TODO(jrprice): We should add support for bidirectional SEM tree traversal
+    using CallList = std::vector<const ast::CallExpression*>;
-  // so that we can do this on the fly instead.
+    std::unordered_map<const ast::Function*, CallList> calls_to_replace;
-  std::unordered_map<ast::IdentifierExpression*, ast::UnaryOpExpression*>
+
-      ident_to_address_of;
+    std::vector<ast::Function*> functions_to_process;
-  for (auto* node : ctx.src->ASTNodes().Objects()) {
+
-    auto* address_of = node->As<ast::UnaryOpExpression>();
+    // Build a list of functions that transitively reference any private or
-    if (!address_of || address_of->op() != ast::UnaryOp::kAddressOf) {
+    // workgroup variables, or texture/sampler variables.
-      continue;
+    for (auto* func_ast : ctx.src->AST().Functions()) {
+      auto* func_sem = ctx.src->Sem().Get(func_ast);
+
+      bool needs_processing = false;
+      for (auto* var : func_sem->ReferencedModuleVariables()) {
+        if (var->StorageClass() == ast::StorageClass::kPrivate ||
+            var->StorageClass() == ast::StorageClass::kWorkgroup ||
+            var->StorageClass() == ast::StorageClass::kUniformConstant) {
+          needs_processing = true;
+          break;
+        }
+      }
+
+      if (needs_processing) {
+        functions_to_process.push_back(func_ast);
+
+        // Find all of the calls to this function that will need to be replaced.
+        for (auto* call : func_sem->CallSites()) {
+          auto* call_sem = ctx.src->Sem().Get(call);
+          calls_to_replace[call_sem->Stmt()->Function()].push_back(call);
+        }
+      }
     }
-    if (auto* ident = address_of->expr()->As<ast::IdentifierExpression>()) {
-      ident_to_address_of[ident] = address_of;
-    }
-  }
 
-  for (auto* func_ast : functions_to_process) {
+    // Build a list of `&ident` expressions. We'll use this later to avoid
-    auto* func_sem = ctx.src->Sem().Get(func_ast);
+    // generating expressions of the form `&*ident`, which break WGSL validation
-    bool is_entry_point = func_ast->IsEntryPoint();
+    // rules when this expression is passed to a function.
-
+    // TODO(jrprice): We should add support for bidirectional SEM tree traversal
-    // Map module-scope variables onto their function-scope replacement.
+    // so that we can do this on the fly instead.
-    std::unordered_map<const sem::Variable*, Symbol> var_to_symbol;
+    std::unordered_map<ast::IdentifierExpression*, ast::UnaryOpExpression*>
-
+        ident_to_address_of;
-    for (auto* var : func_sem->ReferencedModuleVariables()) {
+    for (auto* node : ctx.src->ASTNodes().Objects()) {
-      if (var->StorageClass() != ast::StorageClass::kPrivate &&
+      auto* address_of = node->As<ast::UnaryOpExpression>();
-          var->StorageClass() != ast::StorageClass::kWorkgroup &&
+      if (!address_of || address_of->op() != ast::UnaryOp::kAddressOf) {
-          var->StorageClass() != ast::StorageClass::kUniformConstant) {
         continue;
       }
+      if (auto* ident = address_of->expr()->As<ast::IdentifierExpression>()) {
+        ident_to_address_of[ident] = address_of;
+      }
+    }
 
-      // This is the symbol for the variable that replaces the module-scope var.
+    for (auto* func_ast : functions_to_process) {
-      auto new_var_symbol = ctx.dst->Sym();
+      auto* func_sem = ctx.src->Sem().Get(func_ast);
+      bool is_entry_point = func_ast->IsEntryPoint();
 
-      auto* store_type = CreateASTTypeFor(ctx, var->Type()->UnwrapRef());
+      // Map module-scope variables onto their function-scope replacement.
+      std::unordered_map<const sem::Variable*, Symbol> var_to_symbol;
 
-      // Track whether the new variable is a pointer or not.
+      // We aggregate all workgroup variables into a struct to avoid hitting
-      bool is_pointer = false;
+      // MSL's limit for threadgroup memory arguments.
+      Symbol workgroup_parameter_symbol;
+      ast::StructMemberList workgroup_parameter_members;
+      auto workgroup_param = [&]() {
+        if (!workgroup_parameter_symbol.IsValid()) {
+          workgroup_parameter_symbol = ctx.dst->Sym();
+        }
+        return workgroup_parameter_symbol;
+      };
 
-      if (is_entry_point) {
+      for (auto* var : func_sem->ReferencedModuleVariables()) {
-        if (store_type->is_handle()) {
+        if (var->StorageClass() != ast::StorageClass::kPrivate &&
-          // For a texture or sampler variable, redeclare it as an entry point
+            var->StorageClass() != ast::StorageClass::kWorkgroup &&
-          // parameter. Disable entry point parameter validation.
+            var->StorageClass() != ast::StorageClass::kUniformConstant) {
-          auto* disable_validation =
+          continue;
-              ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
+        }
-                  ctx.dst->ID(), ast::DisabledValidation::kEntryPointParameter);
+
-          auto decos = ctx.Clone(var->Declaration()->decorations());
+        // This is the symbol for the variable that replaces the module-scope
-          decos.push_back(disable_validation);
+        // var.
-          auto* param = ctx.dst->Param(new_var_symbol, store_type, decos);
+        auto new_var_symbol = ctx.dst->Sym();
-          ctx.InsertFront(func_ast->params(), param);
+
-        } else {
+        // Helper to create an AST node for the store type of the variable.
-          if (var->StorageClass() == ast::StorageClass::kWorkgroup &&
+        auto store_type = [&]() {
-              ContainsMatrix(var->Type())) {
+          return CreateASTTypeFor(ctx, var->Type()->UnwrapRef());
-            // Due to a bug in the MSL compiler, we use a threadgroup memory
+        };
-            // argument for any workgroup allocation that contains a matrix.
+
-            // See crbug.com/tint/938.
+        // Track whether the new variable is a pointer or not.
+        bool is_pointer = false;
+
+        if (is_entry_point) {
+          if (var->Type()->UnwrapRef()->is_handle()) {
+            // For a texture or sampler variable, redeclare it as an entry point
+            // parameter. Disable entry point parameter validation.
             auto* disable_validation =
                 ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
                     ctx.dst->ID(),
                     ast::DisabledValidation::kEntryPointParameter);
-            auto* param_type =
+            auto decos = ctx.Clone(var->Declaration()->decorations());
-                ctx.dst->ty.pointer(store_type, var->StorageClass());
+            decos.push_back(disable_validation);
-            auto* param = ctx.dst->Param(new_var_symbol, param_type,
+            auto* param = ctx.dst->Param(new_var_symbol, store_type(), decos);
-                                         {disable_validation});
             ctx.InsertFront(func_ast->params(), param);
+          } else if (var->StorageClass() == ast::StorageClass::kWorkgroup &&
+                     ContainsMatrix(var->Type())) {
+            // Due to a bug in the MSL compiler, we use a threadgroup memory
+            // argument for any workgroup allocation that contains a matrix.
+            // See crbug.com/tint/938.
+            // TODO(jrprice): Do this for all other workgroup variables too.
+
+            // Create a member in the workgroup parameter struct.
+            auto member = ctx.Clone(var->Declaration()->symbol());
+            workgroup_parameter_members.push_back(
+                ctx.dst->Member(member, store_type()));
+            CloneStructTypes(var->Type()->UnwrapRef());
+
+            // Create a function-scope variable that is a pointer to the member.
+            auto* member_ptr = ctx.dst->AddressOf(ctx.dst->MemberAccessor(
+                ctx.dst->Deref(workgroup_param()), member));
+            auto* local_var =
+                ctx.dst->Const(new_var_symbol,
+                               ctx.dst->ty.pointer(
+                                   store_type(), ast::StorageClass::kWorkgroup),
+                               member_ptr);
+            ctx.InsertFront(func_ast->body()->statements(),
+                            ctx.dst->Decl(local_var));
             is_pointer = true;
           } else {
             // For any other private or workgroup variable, redeclare it at
@@ -164,83 +218,123 @@ void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx,
                     ast::DisabledValidation::kIgnoreStorageClass);
             auto* constructor = ctx.Clone(var->Declaration()->constructor());
             auto* local_var = ctx.dst->Var(
-                new_var_symbol, store_type, var->StorageClass(), constructor,
+                new_var_symbol, store_type(), var->StorageClass(), constructor,
                 ast::DecorationList{disable_validation});
             ctx.InsertFront(func_ast->body()->statements(),
                             ctx.dst->Decl(local_var));
           }
-        }
+        } else {
-      } else {
+          // For a regular function, redeclare the variable as a parameter.
-        // For a regular function, redeclare the variable as a parameter.
+          // Use a pointer for non-handle types.
-        // Use a pointer for non-handle types.
+          auto* param_type = store_type();
-        auto* param_type = store_type;
+          ast::DecorationList attributes;
-        if (!store_type->is_handle()) {
+          if (!param_type->is_handle()) {
-          param_type = ctx.dst->ty.pointer(param_type, var->StorageClass());
+            param_type = ctx.dst->ty.pointer(param_type, var->StorageClass());
-          is_pointer = true;
+            is_pointer = true;
-        }
-        ctx.InsertBack(func_ast->params(),
-                       ctx.dst->Param(new_var_symbol, param_type));
-      }
 
-      // Replace all uses of the module-scope variable.
+            // Disable validation of arguments passed to this pointer parameter,
-      // For non-entry points, dereference non-handle pointer parameters.
+            // as we will sometimes pass pointers to struct members.
-      for (auto* user : var->Users()) {
+            attributes.push_back(
-        if (user->Stmt()->Function() == func_ast) {
+                ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
-          ast::Expression* expr = ctx.dst->Expr(new_var_symbol);
+                    ctx.dst->ID(),
-          if (is_pointer) {
+                    ast::DisabledValidation::kIgnoreInvalidPointerArgument));
-            // If this identifier is used by an address-of operator, just remove
+          }
-            // the address-of instead of adding a deref, since we already have a
+          ctx.InsertBack(
-            // pointer.
+              func_ast->params(),
-            auto* ident = user->Declaration()->As<ast::IdentifierExpression>();
+              ctx.dst->Param(new_var_symbol, param_type, attributes));
-            if (ident_to_address_of.count(ident)) {
+        }
-              ctx.Replace(ident_to_address_of[ident], expr);
+
-              continue;
+        // Replace all uses of the module-scope variable.
+        // For non-entry points, dereference non-handle pointer parameters.
+        for (auto* user : var->Users()) {
+          if (user->Stmt()->Function() == func_ast) {
+            ast::Expression* expr = ctx.dst->Expr(new_var_symbol);
+            if (is_pointer) {
+              // If this identifier is used by an address-of operator, just
+              // remove the address-of instead of adding a deref, since we
+              // already have a pointer.
+              auto* ident =
+                  user->Declaration()->As<ast::IdentifierExpression>();
+              if (ident_to_address_of.count(ident)) {
+                ctx.Replace(ident_to_address_of[ident], expr);
+                continue;
+              }
+
+              expr = ctx.dst->Deref(expr);
             }
-
+            ctx.Replace(user->Declaration(), expr);
-            expr = ctx.dst->Deref(expr);
           }
-          ctx.Replace(user->Declaration(), expr);
         }
+
+        var_to_symbol[var] = new_var_symbol;
       }
 
-      var_to_symbol[var] = new_var_symbol;
+      if (!workgroup_parameter_members.empty()) {
+        // Create the workgroup memory parameter.
+        // The parameter is a struct that contains members for each workgroup
+        // variable.
+        auto* str = ctx.dst->Structure(ctx.dst->Sym(),
+                                       std::move(workgroup_parameter_members));
+        auto* param_type = ctx.dst->ty.pointer(ctx.dst->ty.Of(str),
+                                               ast::StorageClass::kWorkgroup);
+        auto* disable_validation =
+            ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
+                ctx.dst->ID(), ast::DisabledValidation::kEntryPointParameter);
+        auto* param =
+            ctx.dst->Param(workgroup_param(), param_type, {disable_validation});
+        ctx.InsertFront(func_ast->params(), param);
+      }
+
+      // Pass the variables as pointers to any functions that need them.
+      for (auto* call : calls_to_replace[func_ast]) {
+        auto* target = ctx.src->AST().Functions().Find(call->func()->symbol());
+        auto* target_sem = ctx.src->Sem().Get(target);
+
+        // Add new arguments for any variables that are needed by the callee.
+        // For entry points, pass non-handle types as pointers.
+        for (auto* target_var : target_sem->ReferencedModuleVariables()) {
+          bool is_handle = target_var->Type()->UnwrapRef()->is_handle();
+          bool is_workgroup_matrix =
+              target_var->StorageClass() == ast::StorageClass::kWorkgroup &&
+              ContainsMatrix(target_var->Type());
+          if (target_var->StorageClass() == ast::StorageClass::kPrivate ||
+              target_var->StorageClass() == ast::StorageClass::kWorkgroup ||
+              target_var->StorageClass() ==
+                  ast::StorageClass::kUniformConstant) {
+            ast::Expression* arg = ctx.dst->Expr(var_to_symbol[target_var]);
+            if (is_entry_point && !is_handle && !is_workgroup_matrix) {
+              arg = ctx.dst->AddressOf(arg);
+            }
+            ctx.InsertBack(call->params(), arg);
+          }
+        }
+      }
     }
 
-    // Pass the variables as pointers to any functions that need them.
+    // Now remove all module-scope variables with these storage classes.
-    for (auto* call : calls_to_replace[func_ast]) {
+    for (auto* var_ast : ctx.src->AST().GlobalVariables()) {
-      auto* target = ctx.src->AST().Functions().Find(call->func()->symbol());
+      auto* var_sem = ctx.src->Sem().Get(var_ast);
-      auto* target_sem = ctx.src->Sem().Get(target);
+      if (var_sem->StorageClass() == ast::StorageClass::kPrivate ||
-
+          var_sem->StorageClass() == ast::StorageClass::kWorkgroup ||
-      // Add new arguments for any variables that are needed by the callee.
+          var_sem->StorageClass() == ast::StorageClass::kUniformConstant) {
-      // For entry points, pass non-handle types as pointers.
+        ctx.Remove(ctx.src->AST().GlobalDeclarations(), var_ast);
-      for (auto* target_var : target_sem->ReferencedModuleVariables()) {
-        bool is_handle = target_var->Type()->UnwrapRef()->is_handle();
-        bool is_workgroup_matrix =
-            target_var->StorageClass() == ast::StorageClass::kWorkgroup &&
-            ContainsMatrix(target_var->Type());
-        if (target_var->StorageClass() == ast::StorageClass::kPrivate ||
-            target_var->StorageClass() == ast::StorageClass::kWorkgroup ||
-            target_var->StorageClass() == ast::StorageClass::kUniformConstant) {
-          ast::Expression* arg = ctx.dst->Expr(var_to_symbol[target_var]);
-          if (is_entry_point && !is_handle && !is_workgroup_matrix) {
-            arg = ctx.dst->AddressOf(arg);
-          }
-          ctx.InsertBack(call->params(), arg);
-        }
       }
     }
   }
 
-  // Now remove all module-scope variables with these storage classes.
+ private:
-  for (auto* var_ast : ctx.src->AST().GlobalVariables()) {
+  std::unordered_set<const sem::Struct*> cloned_structs_;
-    auto* var_sem = ctx.src->Sem().Get(var_ast);
+};
-    if (var_sem->StorageClass() == ast::StorageClass::kPrivate ||
-        var_sem->StorageClass() == ast::StorageClass::kWorkgroup ||
-        var_sem->StorageClass() == ast::StorageClass::kUniformConstant) {
-      ctx.Remove(ctx.src->AST().GlobalDeclarations(), var_ast);
-    }
-  }
 
+ModuleScopeVarToEntryPointParam::ModuleScopeVarToEntryPointParam() = default;
+
+ModuleScopeVarToEntryPointParam::~ModuleScopeVarToEntryPointParam() = default;
+
+void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx,
+                                          const DataMap&,
+                                          DataMap&) {
+  State state{ctx};
+  state.Process();
   ctx.Clone();
 }
 

src/transform/module_scope_var_to_entry_point_param.h

@@ -74,6 +74,8 @@ class ModuleScopeVarToEntryPointParam
   /// @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/module_scope_var_to_entry_point_param_test.cc

@@ -78,12 +78,12 @@ fn main() {
 fn no_uses() {
 }
 
-fn bar(a : f32, b : f32, tint_symbol : ptr<private, f32>, tint_symbol_1 : ptr<workgroup, f32>) {
+fn bar(a : f32, b : f32, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol : ptr<private, f32>, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol_1 : ptr<workgroup, f32>) {
   *(tint_symbol) = a;
   *(tint_symbol_1) = b;
 }
 
-fn foo(a : f32, tint_symbol_2 : ptr<private, f32>, tint_symbol_3 : ptr<workgroup, f32>) {
+fn foo(a : f32, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol_2 : ptr<private, f32>, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol_3 : ptr<workgroup, f32>) {
   let b : f32 = 2.0;
   bar(a, b, tint_symbol_2, tint_symbol_3);
   no_uses();
@@ -181,7 +181,7 @@ fn bar(p : ptr<private, f32>) {
   *(p) = 0.0;
 }
 
-fn foo(tint_symbol : ptr<private, f32>) {
+fn foo([[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol : ptr<private, f32>) {
   bar(tint_symbol);
 }
 
@@ -340,8 +340,13 @@ fn main() {
 )";
 
   auto* expect = R"(
+struct tint_symbol_2 {
+  m : mat2x2<f32>;
+};
+
 [[stage(compute), workgroup_size(1)]]
-fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol : ptr<workgroup, mat2x2<f32>>) {
+fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
+  let tint_symbol : ptr<workgroup, mat2x2<f32>> = &((*(tint_symbol_1)).m);
   let x = *(tint_symbol);
 }
 )";
@@ -376,8 +381,13 @@ struct S2 {
   s : S1;
 };
 
+struct tint_symbol_2 {
+  m : array<S2, 4u>;
+};
+
 [[stage(compute), workgroup_size(1)]]
-fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol : ptr<workgroup, array<S2, 4u>>) {
+fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol_1 : ptr<workgroup, tint_symbol_2>) {
+  let tint_symbol : ptr<workgroup, array<S2, 4u>> = &((*(tint_symbol_1)).m);
   let x = *(tint_symbol);
 }
 )";
@@ -387,6 +397,49 @@ fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol : pt
   EXPECT_EQ(expect, str(got));
 }
 
+// Test that we do not duplicate a struct type used by multiple workgroup
+// variables that are promoted to threadgroup memory arguments.
+TEST_F(ModuleScopeVarToEntryPointParamTest, DuplicateThreadgroupArgumentTypes) {
+  auto* src = R"(
+struct S {
+  m : mat2x2<f32>;
+};
+
+var<workgroup> a : S;
+
+var<workgroup> b : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = a;
+  let y = b;
+}
+)";
+
+  auto* expect = R"(
+struct S {
+  m : mat2x2<f32>;
+};
+
+struct tint_symbol_3 {
+  a : S;
+  b : S;
+};
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol_1 : ptr<workgroup, tint_symbol_3>) {
+  let tint_symbol : ptr<workgroup, S> = &((*(tint_symbol_1)).a);
+  let tint_symbol_2 : ptr<workgroup, S> = &((*(tint_symbol_1)).b);
+  let x = *(tint_symbol);
+  let y = *(tint_symbol_2);
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, EmtpyModule) {
   auto* src = "";
 

src/writer/msl/generator_impl_test.cc

@@ -142,6 +142,10 @@ TEST_F(MslGeneratorImplTest, WorkgroupMatrix) {
   EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
 
 using namespace metal;
+struct tint_symbol_3 {
+  float2x2 m;
+};
+
 void comp_main_inner(uint local_invocation_index, threadgroup float2x2* const tint_symbol) {
   {
     *(tint_symbol) = float2x2();
@@ -150,8 +154,8 @@ void comp_main_inner(uint local_invocation_index, threadgroup float2x2* const ti
   float2x2 const x = *(tint_symbol);
 }
 
-kernel void comp_main(threadgroup float2x2* tint_symbol_1 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+kernel void comp_main(threadgroup tint_symbol_3* tint_symbol_2 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  comp_main_inner(local_invocation_index, tint_symbol_1);
+  comp_main_inner(local_invocation_index, &((*(tint_symbol_2)).m));
   return;
 }
 
@@ -178,6 +182,9 @@ using namespace metal;
 struct tint_array_wrapper {
   float2x2 arr[4];
 };
+struct tint_symbol_3 {
+  tint_array_wrapper m;
+};
 
 void comp_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper* const tint_symbol) {
   for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
@@ -188,8 +195,8 @@ void comp_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper
   tint_array_wrapper const x = *(tint_symbol);
 }
 
-kernel void comp_main(threadgroup tint_array_wrapper* tint_symbol_1 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+kernel void comp_main(threadgroup tint_symbol_3* tint_symbol_2 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  comp_main_inner(local_invocation_index, tint_symbol_1);
+  comp_main_inner(local_invocation_index, &((*(tint_symbol_2)).m));
   return;
 }
 
@@ -227,6 +234,9 @@ struct S1 {
 struct S2 {
   S1 s;
 };
+struct tint_symbol_4 {
+  S2 s;
+};
 
 void comp_main_inner(uint local_invocation_index, threadgroup S2* const tint_symbol_1) {
   {
@@ -237,8 +247,8 @@ void comp_main_inner(uint local_invocation_index, threadgroup S2* const tint_sym
   S2 const x = *(tint_symbol_1);
 }
 
-kernel void comp_main(threadgroup S2* tint_symbol_2 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+kernel void comp_main(threadgroup tint_symbol_4* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  comp_main_inner(local_invocation_index, tint_symbol_2);
+  comp_main_inner(local_invocation_index, &((*(tint_symbol_3)).s));
   return;
 }
 
@@ -291,6 +301,22 @@ TEST_F(MslGeneratorImplTest, WorkgroupMatrix_Multiples) {
   EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
 
 using namespace metal;
+struct tint_symbol_7 {
+  float2x2 m1;
+  float2x3 m2;
+  float2x4 m3;
+};
+struct tint_symbol_15 {
+  float3x2 m4;
+  float3x3 m5;
+  float3x4 m6;
+};
+struct tint_symbol_23 {
+  float4x2 m7;
+  float4x3 m8;
+  float4x4 m9;
+};
+
 void main1_inner(uint local_invocation_index, threadgroup float2x2* const tint_symbol, threadgroup float2x3* const tint_symbol_1, threadgroup float2x4* const tint_symbol_2) {
   {
     *(tint_symbol) = float2x2();
@@ -303,42 +329,42 @@ void main1_inner(uint local_invocation_index, threadgroup float2x2* const tint_s
   float2x4 const a3 = *(tint_symbol_2);
 }
 
-kernel void main1(threadgroup float2x2* tint_symbol_3 [[threadgroup(0)]], threadgroup float2x3* tint_symbol_4 [[threadgroup(1)]], threadgroup float2x4* tint_symbol_5 [[threadgroup(2)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+kernel void main1(threadgroup tint_symbol_7* tint_symbol_4 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  main1_inner(local_invocation_index, tint_symbol_3, tint_symbol_4, tint_symbol_5);
+  main1_inner(local_invocation_index, &((*(tint_symbol_4)).m1), &((*(tint_symbol_4)).m2), &((*(tint_symbol_4)).m3));
   return;
 }
 
-void main2_inner(uint local_invocation_index_1, threadgroup float3x2* const tint_symbol_6, threadgroup float3x3* const tint_symbol_7, threadgroup float3x4* const tint_symbol_8) {
+void main2_inner(uint local_invocation_index_1, threadgroup float3x2* const tint_symbol_8, threadgroup float3x3* const tint_symbol_9, threadgroup float3x4* const tint_symbol_10) {
   {
-    *(tint_symbol_6) = float3x2();
+    *(tint_symbol_8) = float3x2();
-    *(tint_symbol_7) = float3x3();
+    *(tint_symbol_9) = float3x3();
-    *(tint_symbol_8) = float3x4();
+    *(tint_symbol_10) = float3x4();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  float3x2 const a1 = *(tint_symbol_6);
+  float3x2 const a1 = *(tint_symbol_8);
-  float3x3 const a2 = *(tint_symbol_7);
+  float3x3 const a2 = *(tint_symbol_9);
-  float3x4 const a3 = *(tint_symbol_8);
+  float3x4 const a3 = *(tint_symbol_10);
 }
 
-kernel void main2(threadgroup float3x2* tint_symbol_9 [[threadgroup(0)]], threadgroup float3x3* tint_symbol_10 [[threadgroup(1)]], threadgroup float3x4* tint_symbol_11 [[threadgroup(2)]], uint local_invocation_index_1 [[thread_index_in_threadgroup]]) {
+kernel void main2(threadgroup tint_symbol_15* tint_symbol_12 [[threadgroup(0)]], uint local_invocation_index_1 [[thread_index_in_threadgroup]]) {
-  main2_inner(local_invocation_index_1, tint_symbol_9, tint_symbol_10, tint_symbol_11);
+  main2_inner(local_invocation_index_1, &((*(tint_symbol_12)).m4), &((*(tint_symbol_12)).m5), &((*(tint_symbol_12)).m6));
   return;
 }
 
-void main3_inner(uint local_invocation_index_2, threadgroup float4x2* const tint_symbol_12, threadgroup float4x3* const tint_symbol_13, threadgroup float4x4* const tint_symbol_14) {
+void main3_inner(uint local_invocation_index_2, threadgroup float4x2* const tint_symbol_16, threadgroup float4x3* const tint_symbol_17, threadgroup float4x4* const tint_symbol_18) {
   {
-    *(tint_symbol_12) = float4x2();
+    *(tint_symbol_16) = float4x2();
-    *(tint_symbol_13) = float4x3();
+    *(tint_symbol_17) = float4x3();
-    *(tint_symbol_14) = float4x4();
+    *(tint_symbol_18) = float4x4();
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  float4x2 const a1 = *(tint_symbol_12);
+  float4x2 const a1 = *(tint_symbol_16);
-  float4x3 const a2 = *(tint_symbol_13);
+  float4x3 const a2 = *(tint_symbol_17);
-  float4x4 const a3 = *(tint_symbol_14);
+  float4x4 const a3 = *(tint_symbol_18);
 }
 
-kernel void main3(threadgroup float4x2* tint_symbol_15 [[threadgroup(0)]], threadgroup float4x3* tint_symbol_16 [[threadgroup(1)]], threadgroup float4x4* tint_symbol_17 [[threadgroup(2)]], uint local_invocation_index_2 [[thread_index_in_threadgroup]]) {
+kernel void main3(threadgroup tint_symbol_23* tint_symbol_20 [[threadgroup(0)]], uint local_invocation_index_2 [[thread_index_in_threadgroup]]) {
-  main3_inner(local_invocation_index_2, tint_symbol_15, tint_symbol_16, tint_symbol_17);
+  main3_inner(local_invocation_index_2, &((*(tint_symbol_20)).m7), &((*(tint_symbol_20)).m8), &((*(tint_symbol_20)).m9));
   return;
 }
 
@@ -353,18 +379,12 @@ kernel void main4_no_usages() {
   ASSERT_TRUE(allocations.count("main2"));
   ASSERT_TRUE(allocations.count("main3"));
   EXPECT_EQ(allocations.count("main4_no_usages"), 0u);
-  ASSERT_EQ(allocations["main1"].size(), 3u);
+  ASSERT_EQ(allocations["main1"].size(), 1u);
-  EXPECT_EQ(allocations["main1"][0], 2u * 2u * sizeof(float));
+  EXPECT_EQ(allocations["main1"][0], 20u * sizeof(float));
-  EXPECT_EQ(allocations["main1"][1], 2u * 4u * sizeof(float));
+  ASSERT_EQ(allocations["main2"].size(), 1u);
-  EXPECT_EQ(allocations["main1"][2], 2u * 4u * sizeof(float));
+  EXPECT_EQ(allocations["main2"][0], 32u * sizeof(float));
-  ASSERT_EQ(allocations["main2"].size(), 3u);
+  ASSERT_EQ(allocations["main3"].size(), 1u);
-  EXPECT_EQ(allocations["main2"][0], 3u * 2u * sizeof(float));
+  EXPECT_EQ(allocations["main3"][0], 40u * sizeof(float));
-  EXPECT_EQ(allocations["main2"][1], 3u * 4u * sizeof(float));
-  EXPECT_EQ(allocations["main2"][2], 3u * 4u * sizeof(float));
-  ASSERT_EQ(allocations["main3"].size(), 3u);
-  EXPECT_EQ(allocations["main3"][0], 4u * 2u * sizeof(float));
-  EXPECT_EQ(allocations["main3"][1], 4u * 4u * sizeof(float));
-  EXPECT_EQ(allocations["main3"][2], 4u * 4u * sizeof(float));
 }
 
 }  // namespace

test/var/initialization/workgroup/matrix.wgsl.expected.msl

@@ -1,6 +1,10 @@
 #include <metal_stdlib>
 
 using namespace metal;
+struct tint_symbol_4 {
+  float2x3 v;
+};
+
 void tint_symbol_inner(uint local_invocation_index, threadgroup float2x3* const tint_symbol_1) {
   {
     *(tint_symbol_1) = float2x3();
@@ -9,8 +13,8 @@ void tint_symbol_inner(uint local_invocation_index, threadgroup float2x3* const
   (void) *(tint_symbol_1);
 }
 
-kernel void tint_symbol(threadgroup float2x3* tint_symbol_2 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(threadgroup tint_symbol_4* tint_symbol_3 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  tint_symbol_inner(local_invocation_index, tint_symbol_2);
+  tint_symbol_inner(local_invocation_index, &((*(tint_symbol_3)).v));
   return;
 }
 

test/var/uses/many_workgroup_vars.wgsl

@@ -0,0 +1,204 @@
+var<workgroup> m00 : mat2x2<f32>;
+var<workgroup> m01 : mat2x2<f32>;
+var<workgroup> m02 : mat2x2<f32>;
+var<workgroup> m03 : mat2x2<f32>;
+var<workgroup> m04 : mat2x2<f32>;
+var<workgroup> m05 : mat2x2<f32>;
+var<workgroup> m06 : mat2x2<f32>;
+var<workgroup> m07 : mat2x2<f32>;
+var<workgroup> m08 : mat2x2<f32>;
+var<workgroup> m09 : mat2x2<f32>;
+var<workgroup> m10 : mat2x2<f32>;
+var<workgroup> m11 : mat2x2<f32>;
+var<workgroup> m12 : mat2x2<f32>;
+var<workgroup> m13 : mat2x2<f32>;
+var<workgroup> m14 : mat2x2<f32>;
+var<workgroup> m15 : mat2x2<f32>;
+var<workgroup> m16 : mat2x2<f32>;
+var<workgroup> m17 : mat2x2<f32>;
+var<workgroup> m18 : mat2x2<f32>;
+var<workgroup> m19 : mat2x2<f32>;
+var<workgroup> m20 : mat2x2<f32>;
+var<workgroup> m21 : mat2x2<f32>;
+var<workgroup> m22 : mat2x2<f32>;
+var<workgroup> m23 : mat2x2<f32>;
+var<workgroup> m24 : mat2x2<f32>;
+var<workgroup> m25 : mat2x2<f32>;
+var<workgroup> m26 : mat2x2<f32>;
+var<workgroup> m27 : mat2x2<f32>;
+var<workgroup> m28 : mat2x2<f32>;
+var<workgroup> m29 : mat2x2<f32>;
+var<workgroup> m30 : mat2x2<f32>;
+var<workgroup> m31 : mat2x2<f32>;
+var<workgroup> m32 : mat2x2<f32>;
+var<workgroup> m33 : mat2x2<f32>;
+var<workgroup> m34 : mat2x2<f32>;
+var<workgroup> m35 : mat2x2<f32>;
+var<workgroup> m36 : mat2x2<f32>;
+var<workgroup> m37 : mat2x2<f32>;
+var<workgroup> m38 : mat2x2<f32>;
+var<workgroup> m39 : mat2x2<f32>;
+var<workgroup> m40 : mat2x2<f32>;
+var<workgroup> m41 : mat2x2<f32>;
+var<workgroup> m42 : mat2x2<f32>;
+var<workgroup> m43 : mat2x2<f32>;
+var<workgroup> m44 : mat2x2<f32>;
+var<workgroup> m45 : mat2x2<f32>;
+var<workgroup> m46 : mat2x2<f32>;
+var<workgroup> m47 : mat2x2<f32>;
+var<workgroup> m48 : mat2x2<f32>;
+var<workgroup> m49 : mat2x2<f32>;
+var<workgroup> m50 : mat2x2<f32>;
+var<workgroup> m51 : mat2x2<f32>;
+var<workgroup> m52 : mat2x2<f32>;
+var<workgroup> m53 : mat2x2<f32>;
+var<workgroup> m54 : mat2x2<f32>;
+var<workgroup> m55 : mat2x2<f32>;
+var<workgroup> m56 : mat2x2<f32>;
+var<workgroup> m57 : mat2x2<f32>;
+var<workgroup> m58 : mat2x2<f32>;
+var<workgroup> m59 : mat2x2<f32>;
+var<workgroup> m60 : mat2x2<f32>;
+var<workgroup> m61 : mat2x2<f32>;
+var<workgroup> m62 : mat2x2<f32>;
+var<workgroup> m63 : mat2x2<f32>;
+var<workgroup> m64 : mat2x2<f32>;
+var<workgroup> m65 : mat2x2<f32>;
+var<workgroup> m66 : mat2x2<f32>;
+var<workgroup> m67 : mat2x2<f32>;
+var<workgroup> m68 : mat2x2<f32>;
+var<workgroup> m69 : mat2x2<f32>;
+var<workgroup> m70 : mat2x2<f32>;
+var<workgroup> m71 : mat2x2<f32>;
+var<workgroup> m72 : mat2x2<f32>;
+var<workgroup> m73 : mat2x2<f32>;
+var<workgroup> m74 : mat2x2<f32>;
+var<workgroup> m75 : mat2x2<f32>;
+var<workgroup> m76 : mat2x2<f32>;
+var<workgroup> m77 : mat2x2<f32>;
+var<workgroup> m78 : mat2x2<f32>;
+var<workgroup> m79 : mat2x2<f32>;
+var<workgroup> m80 : mat2x2<f32>;
+var<workgroup> m81 : mat2x2<f32>;
+var<workgroup> m82 : mat2x2<f32>;
+var<workgroup> m83 : mat2x2<f32>;
+var<workgroup> m84 : mat2x2<f32>;
+var<workgroup> m85 : mat2x2<f32>;
+var<workgroup> m86 : mat2x2<f32>;
+var<workgroup> m87 : mat2x2<f32>;
+var<workgroup> m88 : mat2x2<f32>;
+var<workgroup> m89 : mat2x2<f32>;
+var<workgroup> m90 : mat2x2<f32>;
+var<workgroup> m91 : mat2x2<f32>;
+var<workgroup> m92 : mat2x2<f32>;
+var<workgroup> m93 : mat2x2<f32>;
+var<workgroup> m94 : mat2x2<f32>;
+var<workgroup> m95 : mat2x2<f32>;
+var<workgroup> m96 : mat2x2<f32>;
+var<workgroup> m97 : mat2x2<f32>;
+var<workgroup> m98 : mat2x2<f32>;
+var<workgroup> m99 : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn tint_symbol([[builtin(local_invocation_index)]] idx : u32) {
+  m00[0][0] = 1.0;
+  m01[0][0] = 1.0;
+  m02[0][0] = 1.0;
+  m03[0][0] = 1.0;
+  m04[0][0] = 1.0;
+  m05[0][0] = 1.0;
+  m06[0][0] = 1.0;
+  m07[0][0] = 1.0;
+  m08[0][0] = 1.0;
+  m09[0][0] = 1.0;
+  m10[0][0] = 1.0;
+  m11[0][0] = 1.0;
+  m12[0][0] = 1.0;
+  m13[0][0] = 1.0;
+  m14[0][0] = 1.0;
+  m15[0][0] = 1.0;
+  m16[0][0] = 1.0;
+  m17[0][0] = 1.0;
+  m18[0][0] = 1.0;
+  m19[0][0] = 1.0;
+  m20[0][0] = 1.0;
+  m21[0][0] = 1.0;
+  m22[0][0] = 1.0;
+  m23[0][0] = 1.0;
+  m24[0][0] = 1.0;
+  m25[0][0] = 1.0;
+  m26[0][0] = 1.0;
+  m27[0][0] = 1.0;
+  m28[0][0] = 1.0;
+  m29[0][0] = 1.0;
+  m30[0][0] = 1.0;
+  m31[0][0] = 1.0;
+  m32[0][0] = 1.0;
+  m33[0][0] = 1.0;
+  m34[0][0] = 1.0;
+  m35[0][0] = 1.0;
+  m36[0][0] = 1.0;
+  m37[0][0] = 1.0;
+  m38[0][0] = 1.0;
+  m39[0][0] = 1.0;
+  m40[0][0] = 1.0;
+  m41[0][0] = 1.0;
+  m42[0][0] = 1.0;
+  m43[0][0] = 1.0;
+  m44[0][0] = 1.0;
+  m45[0][0] = 1.0;
+  m46[0][0] = 1.0;
+  m47[0][0] = 1.0;
+  m48[0][0] = 1.0;
+  m49[0][0] = 1.0;
+  m50[0][0] = 1.0;
+  m51[0][0] = 1.0;
+  m52[0][0] = 1.0;
+  m53[0][0] = 1.0;
+  m54[0][0] = 1.0;
+  m55[0][0] = 1.0;
+  m56[0][0] = 1.0;
+  m57[0][0] = 1.0;
+  m58[0][0] = 1.0;
+  m59[0][0] = 1.0;
+  m60[0][0] = 1.0;
+  m61[0][0] = 1.0;
+  m62[0][0] = 1.0;
+  m63[0][0] = 1.0;
+  m64[0][0] = 1.0;
+  m65[0][0] = 1.0;
+  m66[0][0] = 1.0;
+  m67[0][0] = 1.0;
+  m68[0][0] = 1.0;
+  m69[0][0] = 1.0;
+  m70[0][0] = 1.0;
+  m71[0][0] = 1.0;
+  m72[0][0] = 1.0;
+  m73[0][0] = 1.0;
+  m74[0][0] = 1.0;
+  m75[0][0] = 1.0;
+  m76[0][0] = 1.0;
+  m77[0][0] = 1.0;
+  m78[0][0] = 1.0;
+  m79[0][0] = 1.0;
+  m80[0][0] = 1.0;
+  m81[0][0] = 1.0;
+  m82[0][0] = 1.0;
+  m83[0][0] = 1.0;
+  m84[0][0] = 1.0;
+  m85[0][0] = 1.0;
+  m86[0][0] = 1.0;
+  m87[0][0] = 1.0;
+  m88[0][0] = 1.0;
+  m89[0][0] = 1.0;
+  m90[0][0] = 1.0;
+  m91[0][0] = 1.0;
+  m92[0][0] = 1.0;
+  m93[0][0] = 1.0;
+  m94[0][0] = 1.0;
+  m95[0][0] = 1.0;
+  m96[0][0] = 1.0;
+  m97[0][0] = 1.0;
+  m98[0][0] = 1.0;
+  m99[0][0] = 1.0;
+}

test/var/uses/many_workgroup_vars.wgsl.expected.hlsl

@@ -0,0 +1,316 @@
+groupshared float2x2 m00;
+groupshared float2x2 m01;
+groupshared float2x2 m02;
+groupshared float2x2 m03;
+groupshared float2x2 m04;
+groupshared float2x2 m05;
+groupshared float2x2 m06;
+groupshared float2x2 m07;
+groupshared float2x2 m08;
+groupshared float2x2 m09;
+groupshared float2x2 m10;
+groupshared float2x2 m11;
+groupshared float2x2 m12;
+groupshared float2x2 m13;
+groupshared float2x2 m14;
+groupshared float2x2 m15;
+groupshared float2x2 m16;
+groupshared float2x2 m17;
+groupshared float2x2 m18;
+groupshared float2x2 m19;
+groupshared float2x2 m20;
+groupshared float2x2 m21;
+groupshared float2x2 m22;
+groupshared float2x2 m23;
+groupshared float2x2 m24;
+groupshared float2x2 m25;
+groupshared float2x2 m26;
+groupshared float2x2 m27;
+groupshared float2x2 m28;
+groupshared float2x2 m29;
+groupshared float2x2 m30;
+groupshared float2x2 m31;
+groupshared float2x2 m32;
+groupshared float2x2 m33;
+groupshared float2x2 m34;
+groupshared float2x2 m35;
+groupshared float2x2 m36;
+groupshared float2x2 m37;
+groupshared float2x2 m38;
+groupshared float2x2 m39;
+groupshared float2x2 m40;
+groupshared float2x2 m41;
+groupshared float2x2 m42;
+groupshared float2x2 m43;
+groupshared float2x2 m44;
+groupshared float2x2 m45;
+groupshared float2x2 m46;
+groupshared float2x2 m47;
+groupshared float2x2 m48;
+groupshared float2x2 m49;
+groupshared float2x2 m50;
+groupshared float2x2 m51;
+groupshared float2x2 m52;
+groupshared float2x2 m53;
+groupshared float2x2 m54;
+groupshared float2x2 m55;
+groupshared float2x2 m56;
+groupshared float2x2 m57;
+groupshared float2x2 m58;
+groupshared float2x2 m59;
+groupshared float2x2 m60;
+groupshared float2x2 m61;
+groupshared float2x2 m62;
+groupshared float2x2 m63;
+groupshared float2x2 m64;
+groupshared float2x2 m65;
+groupshared float2x2 m66;
+groupshared float2x2 m67;
+groupshared float2x2 m68;
+groupshared float2x2 m69;
+groupshared float2x2 m70;
+groupshared float2x2 m71;
+groupshared float2x2 m72;
+groupshared float2x2 m73;
+groupshared float2x2 m74;
+groupshared float2x2 m75;
+groupshared float2x2 m76;
+groupshared float2x2 m77;
+groupshared float2x2 m78;
+groupshared float2x2 m79;
+groupshared float2x2 m80;
+groupshared float2x2 m81;
+groupshared float2x2 m82;
+groupshared float2x2 m83;
+groupshared float2x2 m84;
+groupshared float2x2 m85;
+groupshared float2x2 m86;
+groupshared float2x2 m87;
+groupshared float2x2 m88;
+groupshared float2x2 m89;
+groupshared float2x2 m90;
+groupshared float2x2 m91;
+groupshared float2x2 m92;
+groupshared float2x2 m93;
+groupshared float2x2 m94;
+groupshared float2x2 m95;
+groupshared float2x2 m96;
+groupshared float2x2 m97;
+groupshared float2x2 m98;
+groupshared float2x2 m99;
+
+struct tint_symbol_2 {
+  uint idx : SV_GroupIndex;
+};
+
+void tint_symbol_inner(uint idx) {
+  {
+    m00 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m01 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m02 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m03 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m04 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m05 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m06 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m07 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m08 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m09 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m10 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m11 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m12 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m13 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m14 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m15 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m16 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m17 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m18 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m19 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m20 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m21 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m22 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m23 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m24 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m25 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m26 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m27 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m28 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m29 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m30 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m31 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m32 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m33 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m34 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m35 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m36 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m37 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m38 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m39 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m40 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m41 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m42 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m43 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m44 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m45 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m46 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m47 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m48 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m49 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m50 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m51 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m52 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m53 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m54 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m55 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m56 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m57 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m58 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m59 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m60 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m61 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m62 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m63 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m64 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m65 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m66 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m67 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m68 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m69 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m70 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m71 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m72 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m73 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m74 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m75 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m76 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m77 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m78 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m79 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m80 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m81 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m82 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m83 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m84 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m85 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m86 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m87 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m88 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m89 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m90 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m91 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m92 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m93 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m94 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m95 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m96 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m97 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m98 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+    m99 = float2x2(0.0f, 0.0f, 0.0f, 0.0f);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  m00[0][0] = 1.0f;
+  m01[0][0] = 1.0f;
+  m02[0][0] = 1.0f;
+  m03[0][0] = 1.0f;
+  m04[0][0] = 1.0f;
+  m05[0][0] = 1.0f;
+  m06[0][0] = 1.0f;
+  m07[0][0] = 1.0f;
+  m08[0][0] = 1.0f;
+  m09[0][0] = 1.0f;
+  m10[0][0] = 1.0f;
+  m11[0][0] = 1.0f;
+  m12[0][0] = 1.0f;
+  m13[0][0] = 1.0f;
+  m14[0][0] = 1.0f;
+  m15[0][0] = 1.0f;
+  m16[0][0] = 1.0f;
+  m17[0][0] = 1.0f;
+  m18[0][0] = 1.0f;
+  m19[0][0] = 1.0f;
+  m20[0][0] = 1.0f;
+  m21[0][0] = 1.0f;
+  m22[0][0] = 1.0f;
+  m23[0][0] = 1.0f;
+  m24[0][0] = 1.0f;
+  m25[0][0] = 1.0f;
+  m26[0][0] = 1.0f;
+  m27[0][0] = 1.0f;
+  m28[0][0] = 1.0f;
+  m29[0][0] = 1.0f;
+  m30[0][0] = 1.0f;
+  m31[0][0] = 1.0f;
+  m32[0][0] = 1.0f;
+  m33[0][0] = 1.0f;
+  m34[0][0] = 1.0f;
+  m35[0][0] = 1.0f;
+  m36[0][0] = 1.0f;
+  m37[0][0] = 1.0f;
+  m38[0][0] = 1.0f;
+  m39[0][0] = 1.0f;
+  m40[0][0] = 1.0f;
+  m41[0][0] = 1.0f;
+  m42[0][0] = 1.0f;
+  m43[0][0] = 1.0f;
+  m44[0][0] = 1.0f;
+  m45[0][0] = 1.0f;
+  m46[0][0] = 1.0f;
+  m47[0][0] = 1.0f;
+  m48[0][0] = 1.0f;
+  m49[0][0] = 1.0f;
+  m50[0][0] = 1.0f;
+  m51[0][0] = 1.0f;
+  m52[0][0] = 1.0f;
+  m53[0][0] = 1.0f;
+  m54[0][0] = 1.0f;
+  m55[0][0] = 1.0f;
+  m56[0][0] = 1.0f;
+  m57[0][0] = 1.0f;
+  m58[0][0] = 1.0f;
+  m59[0][0] = 1.0f;
+  m60[0][0] = 1.0f;
+  m61[0][0] = 1.0f;
+  m62[0][0] = 1.0f;
+  m63[0][0] = 1.0f;
+  m64[0][0] = 1.0f;
+  m65[0][0] = 1.0f;
+  m66[0][0] = 1.0f;
+  m67[0][0] = 1.0f;
+  m68[0][0] = 1.0f;
+  m69[0][0] = 1.0f;
+  m70[0][0] = 1.0f;
+  m71[0][0] = 1.0f;
+  m72[0][0] = 1.0f;
+  m73[0][0] = 1.0f;
+  m74[0][0] = 1.0f;
+  m75[0][0] = 1.0f;
+  m76[0][0] = 1.0f;
+  m77[0][0] = 1.0f;
+  m78[0][0] = 1.0f;
+  m79[0][0] = 1.0f;
+  m80[0][0] = 1.0f;
+  m81[0][0] = 1.0f;
+  m82[0][0] = 1.0f;
+  m83[0][0] = 1.0f;
+  m84[0][0] = 1.0f;
+  m85[0][0] = 1.0f;
+  m86[0][0] = 1.0f;
+  m87[0][0] = 1.0f;
+  m88[0][0] = 1.0f;
+  m89[0][0] = 1.0f;
+  m90[0][0] = 1.0f;
+  m91[0][0] = 1.0f;
+  m92[0][0] = 1.0f;
+  m93[0][0] = 1.0f;
+  m94[0][0] = 1.0f;
+  m95[0][0] = 1.0f;
+  m96[0][0] = 1.0f;
+  m97[0][0] = 1.0f;
+  m98[0][0] = 1.0f;
+  m99[0][0] = 1.0f;
+}
+
+[numthreads(1, 1, 1)]
+void tint_symbol(tint_symbol_2 tint_symbol_1) {
+  tint_symbol_inner(tint_symbol_1.idx);
+  return;
+}

test/var/uses/many_workgroup_vars.wgsl.expected.msl

@@ -0,0 +1,317 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_202 {
+  float2x2 m00;
+  float2x2 m01;
+  float2x2 m02;
+  float2x2 m03;
+  float2x2 m04;
+  float2x2 m05;
+  float2x2 m06;
+  float2x2 m07;
+  float2x2 m08;
+  float2x2 m09;
+  float2x2 m10;
+  float2x2 m11;
+  float2x2 m12;
+  float2x2 m13;
+  float2x2 m14;
+  float2x2 m15;
+  float2x2 m16;
+  float2x2 m17;
+  float2x2 m18;
+  float2x2 m19;
+  float2x2 m20;
+  float2x2 m21;
+  float2x2 m22;
+  float2x2 m23;
+  float2x2 m24;
+  float2x2 m25;
+  float2x2 m26;
+  float2x2 m27;
+  float2x2 m28;
+  float2x2 m29;
+  float2x2 m30;
+  float2x2 m31;
+  float2x2 m32;
+  float2x2 m33;
+  float2x2 m34;
+  float2x2 m35;
+  float2x2 m36;
+  float2x2 m37;
+  float2x2 m38;
+  float2x2 m39;
+  float2x2 m40;
+  float2x2 m41;
+  float2x2 m42;
+  float2x2 m43;
+  float2x2 m44;
+  float2x2 m45;
+  float2x2 m46;
+  float2x2 m47;
+  float2x2 m48;
+  float2x2 m49;
+  float2x2 m50;
+  float2x2 m51;
+  float2x2 m52;
+  float2x2 m53;
+  float2x2 m54;
+  float2x2 m55;
+  float2x2 m56;
+  float2x2 m57;
+  float2x2 m58;
+  float2x2 m59;
+  float2x2 m60;
+  float2x2 m61;
+  float2x2 m62;
+  float2x2 m63;
+  float2x2 m64;
+  float2x2 m65;
+  float2x2 m66;
+  float2x2 m67;
+  float2x2 m68;
+  float2x2 m69;
+  float2x2 m70;
+  float2x2 m71;
+  float2x2 m72;
+  float2x2 m73;
+  float2x2 m74;
+  float2x2 m75;
+  float2x2 m76;
+  float2x2 m77;
+  float2x2 m78;
+  float2x2 m79;
+  float2x2 m80;
+  float2x2 m81;
+  float2x2 m82;
+  float2x2 m83;
+  float2x2 m84;
+  float2x2 m85;
+  float2x2 m86;
+  float2x2 m87;
+  float2x2 m88;
+  float2x2 m89;
+  float2x2 m90;
+  float2x2 m91;
+  float2x2 m92;
+  float2x2 m93;
+  float2x2 m94;
+  float2x2 m95;
+  float2x2 m96;
+  float2x2 m97;
+  float2x2 m98;
+  float2x2 m99;
+};
+
+void tint_symbol_inner(uint idx, threadgroup float2x2* const tint_symbol_1, threadgroup float2x2* const tint_symbol_2, threadgroup float2x2* const tint_symbol_3, threadgroup float2x2* const tint_symbol_4, threadgroup float2x2* const tint_symbol_5, threadgroup float2x2* const tint_symbol_6, threadgroup float2x2* const tint_symbol_7, threadgroup float2x2* const tint_symbol_8, threadgroup float2x2* const tint_symbol_9, threadgroup float2x2* const tint_symbol_10, threadgroup float2x2* const tint_symbol_11, threadgroup float2x2* const tint_symbol_12, threadgroup float2x2* const tint_symbol_13, threadgroup float2x2* const tint_symbol_14, threadgroup float2x2* const tint_symbol_15, threadgroup float2x2* const tint_symbol_16, threadgroup float2x2* const tint_symbol_17, threadgroup float2x2* const tint_symbol_18, threadgroup float2x2* const tint_symbol_19, threadgroup float2x2* const tint_symbol_20, threadgroup float2x2* const tint_symbol_21, threadgroup float2x2* const tint_symbol_22, threadgroup float2x2* const tint_symbol_23, threadgroup float2x2* const tint_symbol_24, threadgroup float2x2* const tint_symbol_25, threadgroup float2x2* const tint_symbol_26, threadgroup float2x2* const tint_symbol_27, threadgroup float2x2* const tint_symbol_28, threadgroup float2x2* const tint_symbol_29, threadgroup float2x2* const tint_symbol_30, threadgroup float2x2* const tint_symbol_31, threadgroup float2x2* const tint_symbol_32, threadgroup float2x2* const tint_symbol_33, threadgroup float2x2* const tint_symbol_34, threadgroup float2x2* const tint_symbol_35, threadgroup float2x2* const tint_symbol_36, threadgroup float2x2* const tint_symbol_37, threadgroup float2x2* const tint_symbol_38, threadgroup float2x2* const tint_symbol_39, threadgroup float2x2* const tint_symbol_40, threadgroup float2x2* const tint_symbol_41, threadgroup float2x2* const tint_symbol_42, threadgroup float2x2* const tint_symbol_43, threadgroup float2x2* const tint_symbol_44, threadgroup float2x2* const tint_symbol_45, threadgroup float2x2* const tint_symbol_46, threadgroup float2x2* const tint_symbol_47, threadgroup float2x2* const tint_symbol_48, threadgroup float2x2* const tint_symbol_49, threadgroup float2x2* const tint_symbol_50, threadgroup float2x2* const tint_symbol_51, threadgroup float2x2* const tint_symbol_52, threadgroup float2x2* const tint_symbol_53, threadgroup float2x2* const tint_symbol_54, threadgroup float2x2* const tint_symbol_55, threadgroup float2x2* const tint_symbol_56, threadgroup float2x2* const tint_symbol_57, threadgroup float2x2* const tint_symbol_58, threadgroup float2x2* const tint_symbol_59, threadgroup float2x2* const tint_symbol_60, threadgroup float2x2* const tint_symbol_61, threadgroup float2x2* const tint_symbol_62, threadgroup float2x2* const tint_symbol_63, threadgroup float2x2* const tint_symbol_64, threadgroup float2x2* const tint_symbol_65, threadgroup float2x2* const tint_symbol_66, threadgroup float2x2* const tint_symbol_67, threadgroup float2x2* const tint_symbol_68, threadgroup float2x2* const tint_symbol_69, threadgroup float2x2* const tint_symbol_70, threadgroup float2x2* const tint_symbol_71, threadgroup float2x2* const tint_symbol_72, threadgroup float2x2* const tint_symbol_73, threadgroup float2x2* const tint_symbol_74, threadgroup float2x2* const tint_symbol_75, threadgroup float2x2* const tint_symbol_76, threadgroup float2x2* const tint_symbol_77, threadgroup float2x2* const tint_symbol_78, threadgroup float2x2* const tint_symbol_79, threadgroup float2x2* const tint_symbol_80, threadgroup float2x2* const tint_symbol_81, threadgroup float2x2* const tint_symbol_82, threadgroup float2x2* const tint_symbol_83, threadgroup float2x2* const tint_symbol_84, threadgroup float2x2* const tint_symbol_85, threadgroup float2x2* const tint_symbol_86, threadgroup float2x2* const tint_symbol_87, threadgroup float2x2* const tint_symbol_88, threadgroup float2x2* const tint_symbol_89, threadgroup float2x2* const tint_symbol_90, threadgroup float2x2* const tint_symbol_91, threadgroup float2x2* const tint_symbol_92, threadgroup float2x2* const tint_symbol_93, threadgroup float2x2* const tint_symbol_94, threadgroup float2x2* const tint_symbol_95, threadgroup float2x2* const tint_symbol_96, threadgroup float2x2* const tint_symbol_97, threadgroup float2x2* const tint_symbol_98, threadgroup float2x2* const tint_symbol_99, threadgroup float2x2* const tint_symbol_100) {
+  {
+    *(tint_symbol_1) = float2x2();
+    *(tint_symbol_2) = float2x2();
+    *(tint_symbol_3) = float2x2();
+    *(tint_symbol_4) = float2x2();
+    *(tint_symbol_5) = float2x2();
+    *(tint_symbol_6) = float2x2();
+    *(tint_symbol_7) = float2x2();
+    *(tint_symbol_8) = float2x2();
+    *(tint_symbol_9) = float2x2();
+    *(tint_symbol_10) = float2x2();
+    *(tint_symbol_11) = float2x2();
+    *(tint_symbol_12) = float2x2();
+    *(tint_symbol_13) = float2x2();
+    *(tint_symbol_14) = float2x2();
+    *(tint_symbol_15) = float2x2();
+    *(tint_symbol_16) = float2x2();
+    *(tint_symbol_17) = float2x2();
+    *(tint_symbol_18) = float2x2();
+    *(tint_symbol_19) = float2x2();
+    *(tint_symbol_20) = float2x2();
+    *(tint_symbol_21) = float2x2();
+    *(tint_symbol_22) = float2x2();
+    *(tint_symbol_23) = float2x2();
+    *(tint_symbol_24) = float2x2();
+    *(tint_symbol_25) = float2x2();
+    *(tint_symbol_26) = float2x2();
+    *(tint_symbol_27) = float2x2();
+    *(tint_symbol_28) = float2x2();
+    *(tint_symbol_29) = float2x2();
+    *(tint_symbol_30) = float2x2();
+    *(tint_symbol_31) = float2x2();
+    *(tint_symbol_32) = float2x2();
+    *(tint_symbol_33) = float2x2();
+    *(tint_symbol_34) = float2x2();
+    *(tint_symbol_35) = float2x2();
+    *(tint_symbol_36) = float2x2();
+    *(tint_symbol_37) = float2x2();
+    *(tint_symbol_38) = float2x2();
+    *(tint_symbol_39) = float2x2();
+    *(tint_symbol_40) = float2x2();
+    *(tint_symbol_41) = float2x2();
+    *(tint_symbol_42) = float2x2();
+    *(tint_symbol_43) = float2x2();
+    *(tint_symbol_44) = float2x2();
+    *(tint_symbol_45) = float2x2();
+    *(tint_symbol_46) = float2x2();
+    *(tint_symbol_47) = float2x2();
+    *(tint_symbol_48) = float2x2();
+    *(tint_symbol_49) = float2x2();
+    *(tint_symbol_50) = float2x2();
+    *(tint_symbol_51) = float2x2();
+    *(tint_symbol_52) = float2x2();
+    *(tint_symbol_53) = float2x2();
+    *(tint_symbol_54) = float2x2();
+    *(tint_symbol_55) = float2x2();
+    *(tint_symbol_56) = float2x2();
+    *(tint_symbol_57) = float2x2();
+    *(tint_symbol_58) = float2x2();
+    *(tint_symbol_59) = float2x2();
+    *(tint_symbol_60) = float2x2();
+    *(tint_symbol_61) = float2x2();
+    *(tint_symbol_62) = float2x2();
+    *(tint_symbol_63) = float2x2();
+    *(tint_symbol_64) = float2x2();
+    *(tint_symbol_65) = float2x2();
+    *(tint_symbol_66) = float2x2();
+    *(tint_symbol_67) = float2x2();
+    *(tint_symbol_68) = float2x2();
+    *(tint_symbol_69) = float2x2();
+    *(tint_symbol_70) = float2x2();
+    *(tint_symbol_71) = float2x2();
+    *(tint_symbol_72) = float2x2();
+    *(tint_symbol_73) = float2x2();
+    *(tint_symbol_74) = float2x2();
+    *(tint_symbol_75) = float2x2();
+    *(tint_symbol_76) = float2x2();
+    *(tint_symbol_77) = float2x2();
+    *(tint_symbol_78) = float2x2();
+    *(tint_symbol_79) = float2x2();
+    *(tint_symbol_80) = float2x2();
+    *(tint_symbol_81) = float2x2();
+    *(tint_symbol_82) = float2x2();
+    *(tint_symbol_83) = float2x2();
+    *(tint_symbol_84) = float2x2();
+    *(tint_symbol_85) = float2x2();
+    *(tint_symbol_86) = float2x2();
+    *(tint_symbol_87) = float2x2();
+    *(tint_symbol_88) = float2x2();
+    *(tint_symbol_89) = float2x2();
+    *(tint_symbol_90) = float2x2();
+    *(tint_symbol_91) = float2x2();
+    *(tint_symbol_92) = float2x2();
+    *(tint_symbol_93) = float2x2();
+    *(tint_symbol_94) = float2x2();
+    *(tint_symbol_95) = float2x2();
+    *(tint_symbol_96) = float2x2();
+    *(tint_symbol_97) = float2x2();
+    *(tint_symbol_98) = float2x2();
+    *(tint_symbol_99) = float2x2();
+    *(tint_symbol_100) = float2x2();
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*(tint_symbol_1))[0][0] = 1.0f;
+  (*(tint_symbol_2))[0][0] = 1.0f;
+  (*(tint_symbol_3))[0][0] = 1.0f;
+  (*(tint_symbol_4))[0][0] = 1.0f;
+  (*(tint_symbol_5))[0][0] = 1.0f;
+  (*(tint_symbol_6))[0][0] = 1.0f;
+  (*(tint_symbol_7))[0][0] = 1.0f;
+  (*(tint_symbol_8))[0][0] = 1.0f;
+  (*(tint_symbol_9))[0][0] = 1.0f;
+  (*(tint_symbol_10))[0][0] = 1.0f;
+  (*(tint_symbol_11))[0][0] = 1.0f;
+  (*(tint_symbol_12))[0][0] = 1.0f;
+  (*(tint_symbol_13))[0][0] = 1.0f;
+  (*(tint_symbol_14))[0][0] = 1.0f;
+  (*(tint_symbol_15))[0][0] = 1.0f;
+  (*(tint_symbol_16))[0][0] = 1.0f;
+  (*(tint_symbol_17))[0][0] = 1.0f;
+  (*(tint_symbol_18))[0][0] = 1.0f;
+  (*(tint_symbol_19))[0][0] = 1.0f;
+  (*(tint_symbol_20))[0][0] = 1.0f;
+  (*(tint_symbol_21))[0][0] = 1.0f;
+  (*(tint_symbol_22))[0][0] = 1.0f;
+  (*(tint_symbol_23))[0][0] = 1.0f;
+  (*(tint_symbol_24))[0][0] = 1.0f;
+  (*(tint_symbol_25))[0][0] = 1.0f;
+  (*(tint_symbol_26))[0][0] = 1.0f;
+  (*(tint_symbol_27))[0][0] = 1.0f;
+  (*(tint_symbol_28))[0][0] = 1.0f;
+  (*(tint_symbol_29))[0][0] = 1.0f;
+  (*(tint_symbol_30))[0][0] = 1.0f;
+  (*(tint_symbol_31))[0][0] = 1.0f;
+  (*(tint_symbol_32))[0][0] = 1.0f;
+  (*(tint_symbol_33))[0][0] = 1.0f;
+  (*(tint_symbol_34))[0][0] = 1.0f;
+  (*(tint_symbol_35))[0][0] = 1.0f;
+  (*(tint_symbol_36))[0][0] = 1.0f;
+  (*(tint_symbol_37))[0][0] = 1.0f;
+  (*(tint_symbol_38))[0][0] = 1.0f;
+  (*(tint_symbol_39))[0][0] = 1.0f;
+  (*(tint_symbol_40))[0][0] = 1.0f;
+  (*(tint_symbol_41))[0][0] = 1.0f;
+  (*(tint_symbol_42))[0][0] = 1.0f;
+  (*(tint_symbol_43))[0][0] = 1.0f;
+  (*(tint_symbol_44))[0][0] = 1.0f;
+  (*(tint_symbol_45))[0][0] = 1.0f;
+  (*(tint_symbol_46))[0][0] = 1.0f;
+  (*(tint_symbol_47))[0][0] = 1.0f;
+  (*(tint_symbol_48))[0][0] = 1.0f;
+  (*(tint_symbol_49))[0][0] = 1.0f;
+  (*(tint_symbol_50))[0][0] = 1.0f;
+  (*(tint_symbol_51))[0][0] = 1.0f;
+  (*(tint_symbol_52))[0][0] = 1.0f;
+  (*(tint_symbol_53))[0][0] = 1.0f;
+  (*(tint_symbol_54))[0][0] = 1.0f;
+  (*(tint_symbol_55))[0][0] = 1.0f;
+  (*(tint_symbol_56))[0][0] = 1.0f;
+  (*(tint_symbol_57))[0][0] = 1.0f;
+  (*(tint_symbol_58))[0][0] = 1.0f;
+  (*(tint_symbol_59))[0][0] = 1.0f;
+  (*(tint_symbol_60))[0][0] = 1.0f;
+  (*(tint_symbol_61))[0][0] = 1.0f;
+  (*(tint_symbol_62))[0][0] = 1.0f;
+  (*(tint_symbol_63))[0][0] = 1.0f;
+  (*(tint_symbol_64))[0][0] = 1.0f;
+  (*(tint_symbol_65))[0][0] = 1.0f;
+  (*(tint_symbol_66))[0][0] = 1.0f;
+  (*(tint_symbol_67))[0][0] = 1.0f;
+  (*(tint_symbol_68))[0][0] = 1.0f;
+  (*(tint_symbol_69))[0][0] = 1.0f;
+  (*(tint_symbol_70))[0][0] = 1.0f;
+  (*(tint_symbol_71))[0][0] = 1.0f;
+  (*(tint_symbol_72))[0][0] = 1.0f;
+  (*(tint_symbol_73))[0][0] = 1.0f;
+  (*(tint_symbol_74))[0][0] = 1.0f;
+  (*(tint_symbol_75))[0][0] = 1.0f;
+  (*(tint_symbol_76))[0][0] = 1.0f;
+  (*(tint_symbol_77))[0][0] = 1.0f;
+  (*(tint_symbol_78))[0][0] = 1.0f;
+  (*(tint_symbol_79))[0][0] = 1.0f;
+  (*(tint_symbol_80))[0][0] = 1.0f;
+  (*(tint_symbol_81))[0][0] = 1.0f;
+  (*(tint_symbol_82))[0][0] = 1.0f;
+  (*(tint_symbol_83))[0][0] = 1.0f;
+  (*(tint_symbol_84))[0][0] = 1.0f;
+  (*(tint_symbol_85))[0][0] = 1.0f;
+  (*(tint_symbol_86))[0][0] = 1.0f;
+  (*(tint_symbol_87))[0][0] = 1.0f;
+  (*(tint_symbol_88))[0][0] = 1.0f;
+  (*(tint_symbol_89))[0][0] = 1.0f;
+  (*(tint_symbol_90))[0][0] = 1.0f;
+  (*(tint_symbol_91))[0][0] = 1.0f;
+  (*(tint_symbol_92))[0][0] = 1.0f;
+  (*(tint_symbol_93))[0][0] = 1.0f;
+  (*(tint_symbol_94))[0][0] = 1.0f;
+  (*(tint_symbol_95))[0][0] = 1.0f;
+  (*(tint_symbol_96))[0][0] = 1.0f;
+  (*(tint_symbol_97))[0][0] = 1.0f;
+  (*(tint_symbol_98))[0][0] = 1.0f;
+  (*(tint_symbol_99))[0][0] = 1.0f;
+  (*(tint_symbol_100))[0][0] = 1.0f;
+}
+
+kernel void tint_symbol(threadgroup tint_symbol_202* tint_symbol_102 [[threadgroup(0)]], uint idx [[thread_index_in_threadgroup]]) {
+  tint_symbol_inner(idx, &((*(tint_symbol_102)).m00), &((*(tint_symbol_102)).m01), &((*(tint_symbol_102)).m02), &((*(tint_symbol_102)).m03), &((*(tint_symbol_102)).m04), &((*(tint_symbol_102)).m05), &((*(tint_symbol_102)).m06), &((*(tint_symbol_102)).m07), &((*(tint_symbol_102)).m08), &((*(tint_symbol_102)).m09), &((*(tint_symbol_102)).m10), &((*(tint_symbol_102)).m11), &((*(tint_symbol_102)).m12), &((*(tint_symbol_102)).m13), &((*(tint_symbol_102)).m14), &((*(tint_symbol_102)).m15), &((*(tint_symbol_102)).m16), &((*(tint_symbol_102)).m17), &((*(tint_symbol_102)).m18), &((*(tint_symbol_102)).m19), &((*(tint_symbol_102)).m20), &((*(tint_symbol_102)).m21), &((*(tint_symbol_102)).m22), &((*(tint_symbol_102)).m23), &((*(tint_symbol_102)).m24), &((*(tint_symbol_102)).m25), &((*(tint_symbol_102)).m26), &((*(tint_symbol_102)).m27), &((*(tint_symbol_102)).m28), &((*(tint_symbol_102)).m29), &((*(tint_symbol_102)).m30), &((*(tint_symbol_102)).m31), &((*(tint_symbol_102)).m32), &((*(tint_symbol_102)).m33), &((*(tint_symbol_102)).m34), &((*(tint_symbol_102)).m35), &((*(tint_symbol_102)).m36), &((*(tint_symbol_102)).m37), &((*(tint_symbol_102)).m38), &((*(tint_symbol_102)).m39), &((*(tint_symbol_102)).m40), &((*(tint_symbol_102)).m41), &((*(tint_symbol_102)).m42), &((*(tint_symbol_102)).m43), &((*(tint_symbol_102)).m44), &((*(tint_symbol_102)).m45), &((*(tint_symbol_102)).m46), &((*(tint_symbol_102)).m47), &((*(tint_symbol_102)).m48), &((*(tint_symbol_102)).m49), &((*(tint_symbol_102)).m50), &((*(tint_symbol_102)).m51), &((*(tint_symbol_102)).m52), &((*(tint_symbol_102)).m53), &((*(tint_symbol_102)).m54), &((*(tint_symbol_102)).m55), &((*(tint_symbol_102)).m56), &((*(tint_symbol_102)).m57), &((*(tint_symbol_102)).m58), &((*(tint_symbol_102)).m59), &((*(tint_symbol_102)).m60), &((*(tint_symbol_102)).m61), &((*(tint_symbol_102)).m62), &((*(tint_symbol_102)).m63), &((*(tint_symbol_102)).m64), &((*(tint_symbol_102)).m65), &((*(tint_symbol_102)).m66), &((*(tint_symbol_102)).m67), &((*(tint_symbol_102)).m68), &((*(tint_symbol_102)).m69), &((*(tint_symbol_102)).m70), &((*(tint_symbol_102)).m71), &((*(tint_symbol_102)).m72), &((*(tint_symbol_102)).m73), &((*(tint_symbol_102)).m74), &((*(tint_symbol_102)).m75), &((*(tint_symbol_102)).m76), &((*(tint_symbol_102)).m77), &((*(tint_symbol_102)).m78), &((*(tint_symbol_102)).m79), &((*(tint_symbol_102)).m80), &((*(tint_symbol_102)).m81), &((*(tint_symbol_102)).m82), &((*(tint_symbol_102)).m83), &((*(tint_symbol_102)).m84), &((*(tint_symbol_102)).m85), &((*(tint_symbol_102)).m86), &((*(tint_symbol_102)).m87), &((*(tint_symbol_102)).m88), &((*(tint_symbol_102)).m89), &((*(tint_symbol_102)).m90), &((*(tint_symbol_102)).m91), &((*(tint_symbol_102)).m92), &((*(tint_symbol_102)).m93), &((*(tint_symbol_102)).m94), &((*(tint_symbol_102)).m95), &((*(tint_symbol_102)).m96), &((*(tint_symbol_102)).m97), &((*(tint_symbol_102)).m98), &((*(tint_symbol_102)).m99));
+  return;
+}
+

test/var/uses/many_workgroup_vars.wgsl.expected.spvasm

@@ -0,0 +1,543 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 226
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %tint_symbol "tint_symbol" %idx_1
+               OpExecutionMode %tint_symbol LocalSize 1 1 1
+               OpName %idx_1 "idx_1"
+               OpName %m00 "m00"
+               OpName %m01 "m01"
+               OpName %m02 "m02"
+               OpName %m03 "m03"
+               OpName %m04 "m04"
+               OpName %m05 "m05"
+               OpName %m06 "m06"
+               OpName %m07 "m07"
+               OpName %m08 "m08"
+               OpName %m09 "m09"
+               OpName %m10 "m10"
+               OpName %m11 "m11"
+               OpName %m12 "m12"
+               OpName %m13 "m13"
+               OpName %m14 "m14"
+               OpName %m15 "m15"
+               OpName %m16 "m16"
+               OpName %m17 "m17"
+               OpName %m18 "m18"
+               OpName %m19 "m19"
+               OpName %m20 "m20"
+               OpName %m21 "m21"
+               OpName %m22 "m22"
+               OpName %m23 "m23"
+               OpName %m24 "m24"
+               OpName %m25 "m25"
+               OpName %m26 "m26"
+               OpName %m27 "m27"
+               OpName %m28 "m28"
+               OpName %m29 "m29"
+               OpName %m30 "m30"
+               OpName %m31 "m31"
+               OpName %m32 "m32"
+               OpName %m33 "m33"
+               OpName %m34 "m34"
+               OpName %m35 "m35"
+               OpName %m36 "m36"
+               OpName %m37 "m37"
+               OpName %m38 "m38"
+               OpName %m39 "m39"
+               OpName %m40 "m40"
+               OpName %m41 "m41"
+               OpName %m42 "m42"
+               OpName %m43 "m43"
+               OpName %m44 "m44"
+               OpName %m45 "m45"
+               OpName %m46 "m46"
+               OpName %m47 "m47"
+               OpName %m48 "m48"
+               OpName %m49 "m49"
+               OpName %m50 "m50"
+               OpName %m51 "m51"
+               OpName %m52 "m52"
+               OpName %m53 "m53"
+               OpName %m54 "m54"
+               OpName %m55 "m55"
+               OpName %m56 "m56"
+               OpName %m57 "m57"
+               OpName %m58 "m58"
+               OpName %m59 "m59"
+               OpName %m60 "m60"
+               OpName %m61 "m61"
+               OpName %m62 "m62"
+               OpName %m63 "m63"
+               OpName %m64 "m64"
+               OpName %m65 "m65"
+               OpName %m66 "m66"
+               OpName %m67 "m67"
+               OpName %m68 "m68"
+               OpName %m69 "m69"
+               OpName %m70 "m70"
+               OpName %m71 "m71"
+               OpName %m72 "m72"
+               OpName %m73 "m73"
+               OpName %m74 "m74"
+               OpName %m75 "m75"
+               OpName %m76 "m76"
+               OpName %m77 "m77"
+               OpName %m78 "m78"
+               OpName %m79 "m79"
+               OpName %m80 "m80"
+               OpName %m81 "m81"
+               OpName %m82 "m82"
+               OpName %m83 "m83"
+               OpName %m84 "m84"
+               OpName %m85 "m85"
+               OpName %m86 "m86"
+               OpName %m87 "m87"
+               OpName %m88 "m88"
+               OpName %m89 "m89"
+               OpName %m90 "m90"
+               OpName %m91 "m91"
+               OpName %m92 "m92"
+               OpName %m93 "m93"
+               OpName %m94 "m94"
+               OpName %m95 "m95"
+               OpName %m96 "m96"
+               OpName %m97 "m97"
+               OpName %m98 "m98"
+               OpName %m99 "m99"
+               OpName %tint_symbol_inner "tint_symbol_inner"
+               OpName %idx "idx"
+               OpName %tint_symbol "tint_symbol"
+               OpDecorate %idx_1 BuiltIn LocalInvocationIndex
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+      %idx_1 = OpVariable %_ptr_Input_uint Input
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+%_ptr_Workgroup_mat2v2float = OpTypePointer Workgroup %mat2v2float
+        %m00 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m01 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m02 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m03 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m04 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m05 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m06 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m07 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m08 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m09 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m10 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m11 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m12 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m13 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m14 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m15 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m16 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m17 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m18 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m19 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m20 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m21 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m22 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m23 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m24 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m25 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m26 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m27 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m28 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m29 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m30 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m31 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m32 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m33 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m34 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m35 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m36 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m37 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m38 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m39 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m40 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m41 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m42 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m43 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m44 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m45 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m46 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m47 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m48 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m49 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m50 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m51 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m52 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m53 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m54 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m55 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m56 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m57 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m58 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m59 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m60 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m61 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m62 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m63 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m64 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m65 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m66 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m67 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m68 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m69 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m70 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m71 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m72 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m73 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m74 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m75 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m76 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m77 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m78 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m79 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m80 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m81 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m82 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m83 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m84 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m85 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m86 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m87 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m88 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m89 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m90 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m91 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m92 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m93 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m94 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m95 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m96 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m97 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m98 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+        %m99 = OpVariable %_ptr_Workgroup_mat2v2float Workgroup
+       %void = OpTypeVoid
+        %108 = OpTypeFunction %void %uint
+        %113 = OpConstantNull %mat2v2float
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Workgroup_float = OpTypePointer Workgroup %float
+    %float_1 = OpConstant %float 1
+        %221 = OpTypeFunction %void
+%tint_symbol_inner = OpFunction %void None %108
+        %idx = OpFunctionParameter %uint
+        %112 = OpLabel
+               OpStore %m00 %113
+               OpStore %m01 %113
+               OpStore %m02 %113
+               OpStore %m03 %113
+               OpStore %m04 %113
+               OpStore %m05 %113
+               OpStore %m06 %113
+               OpStore %m07 %113
+               OpStore %m08 %113
+               OpStore %m09 %113
+               OpStore %m10 %113
+               OpStore %m11 %113
+               OpStore %m12 %113
+               OpStore %m13 %113
+               OpStore %m14 %113
+               OpStore %m15 %113
+               OpStore %m16 %113
+               OpStore %m17 %113
+               OpStore %m18 %113
+               OpStore %m19 %113
+               OpStore %m20 %113
+               OpStore %m21 %113
+               OpStore %m22 %113
+               OpStore %m23 %113
+               OpStore %m24 %113
+               OpStore %m25 %113
+               OpStore %m26 %113
+               OpStore %m27 %113
+               OpStore %m28 %113
+               OpStore %m29 %113
+               OpStore %m30 %113
+               OpStore %m31 %113
+               OpStore %m32 %113
+               OpStore %m33 %113
+               OpStore %m34 %113
+               OpStore %m35 %113
+               OpStore %m36 %113
+               OpStore %m37 %113
+               OpStore %m38 %113
+               OpStore %m39 %113
+               OpStore %m40 %113
+               OpStore %m41 %113
+               OpStore %m42 %113
+               OpStore %m43 %113
+               OpStore %m44 %113
+               OpStore %m45 %113
+               OpStore %m46 %113
+               OpStore %m47 %113
+               OpStore %m48 %113
+               OpStore %m49 %113
+               OpStore %m50 %113
+               OpStore %m51 %113
+               OpStore %m52 %113
+               OpStore %m53 %113
+               OpStore %m54 %113
+               OpStore %m55 %113
+               OpStore %m56 %113
+               OpStore %m57 %113
+               OpStore %m58 %113
+               OpStore %m59 %113
+               OpStore %m60 %113
+               OpStore %m61 %113
+               OpStore %m62 %113
+               OpStore %m63 %113
+               OpStore %m64 %113
+               OpStore %m65 %113
+               OpStore %m66 %113
+               OpStore %m67 %113
+               OpStore %m68 %113
+               OpStore %m69 %113
+               OpStore %m70 %113
+               OpStore %m71 %113
+               OpStore %m72 %113
+               OpStore %m73 %113
+               OpStore %m74 %113
+               OpStore %m75 %113
+               OpStore %m76 %113
+               OpStore %m77 %113
+               OpStore %m78 %113
+               OpStore %m79 %113
+               OpStore %m80 %113
+               OpStore %m81 %113
+               OpStore %m82 %113
+               OpStore %m83 %113
+               OpStore %m84 %113
+               OpStore %m85 %113
+               OpStore %m86 %113
+               OpStore %m87 %113
+               OpStore %m88 %113
+               OpStore %m89 %113
+               OpStore %m90 %113
+               OpStore %m91 %113
+               OpStore %m92 %113
+               OpStore %m93 %113
+               OpStore %m94 %113
+               OpStore %m95 %113
+               OpStore %m96 %113
+               OpStore %m97 %113
+               OpStore %m98 %113
+               OpStore %m99 %113
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+        %120 = OpAccessChain %_ptr_Workgroup_float %m00 %int_0 %int_0
+               OpStore %120 %float_1
+        %122 = OpAccessChain %_ptr_Workgroup_float %m01 %int_0 %int_0
+               OpStore %122 %float_1
+        %123 = OpAccessChain %_ptr_Workgroup_float %m02 %int_0 %int_0
+               OpStore %123 %float_1
+        %124 = OpAccessChain %_ptr_Workgroup_float %m03 %int_0 %int_0
+               OpStore %124 %float_1
+        %125 = OpAccessChain %_ptr_Workgroup_float %m04 %int_0 %int_0
+               OpStore %125 %float_1
+        %126 = OpAccessChain %_ptr_Workgroup_float %m05 %int_0 %int_0
+               OpStore %126 %float_1
+        %127 = OpAccessChain %_ptr_Workgroup_float %m06 %int_0 %int_0
+               OpStore %127 %float_1
+        %128 = OpAccessChain %_ptr_Workgroup_float %m07 %int_0 %int_0
+               OpStore %128 %float_1
+        %129 = OpAccessChain %_ptr_Workgroup_float %m08 %int_0 %int_0
+               OpStore %129 %float_1
+        %130 = OpAccessChain %_ptr_Workgroup_float %m09 %int_0 %int_0
+               OpStore %130 %float_1
+        %131 = OpAccessChain %_ptr_Workgroup_float %m10 %int_0 %int_0
+               OpStore %131 %float_1
+        %132 = OpAccessChain %_ptr_Workgroup_float %m11 %int_0 %int_0
+               OpStore %132 %float_1
+        %133 = OpAccessChain %_ptr_Workgroup_float %m12 %int_0 %int_0
+               OpStore %133 %float_1
+        %134 = OpAccessChain %_ptr_Workgroup_float %m13 %int_0 %int_0
+               OpStore %134 %float_1
+        %135 = OpAccessChain %_ptr_Workgroup_float %m14 %int_0 %int_0
+               OpStore %135 %float_1
+        %136 = OpAccessChain %_ptr_Workgroup_float %m15 %int_0 %int_0
+               OpStore %136 %float_1
+        %137 = OpAccessChain %_ptr_Workgroup_float %m16 %int_0 %int_0
+               OpStore %137 %float_1
+        %138 = OpAccessChain %_ptr_Workgroup_float %m17 %int_0 %int_0
+               OpStore %138 %float_1
+        %139 = OpAccessChain %_ptr_Workgroup_float %m18 %int_0 %int_0
+               OpStore %139 %float_1
+        %140 = OpAccessChain %_ptr_Workgroup_float %m19 %int_0 %int_0
+               OpStore %140 %float_1
+        %141 = OpAccessChain %_ptr_Workgroup_float %m20 %int_0 %int_0
+               OpStore %141 %float_1
+        %142 = OpAccessChain %_ptr_Workgroup_float %m21 %int_0 %int_0
+               OpStore %142 %float_1
+        %143 = OpAccessChain %_ptr_Workgroup_float %m22 %int_0 %int_0
+               OpStore %143 %float_1
+        %144 = OpAccessChain %_ptr_Workgroup_float %m23 %int_0 %int_0
+               OpStore %144 %float_1
+        %145 = OpAccessChain %_ptr_Workgroup_float %m24 %int_0 %int_0
+               OpStore %145 %float_1
+        %146 = OpAccessChain %_ptr_Workgroup_float %m25 %int_0 %int_0
+               OpStore %146 %float_1
+        %147 = OpAccessChain %_ptr_Workgroup_float %m26 %int_0 %int_0
+               OpStore %147 %float_1
+        %148 = OpAccessChain %_ptr_Workgroup_float %m27 %int_0 %int_0
+               OpStore %148 %float_1
+        %149 = OpAccessChain %_ptr_Workgroup_float %m28 %int_0 %int_0
+               OpStore %149 %float_1
+        %150 = OpAccessChain %_ptr_Workgroup_float %m29 %int_0 %int_0
+               OpStore %150 %float_1
+        %151 = OpAccessChain %_ptr_Workgroup_float %m30 %int_0 %int_0
+               OpStore %151 %float_1
+        %152 = OpAccessChain %_ptr_Workgroup_float %m31 %int_0 %int_0
+               OpStore %152 %float_1
+        %153 = OpAccessChain %_ptr_Workgroup_float %m32 %int_0 %int_0
+               OpStore %153 %float_1
+        %154 = OpAccessChain %_ptr_Workgroup_float %m33 %int_0 %int_0
+               OpStore %154 %float_1
+        %155 = OpAccessChain %_ptr_Workgroup_float %m34 %int_0 %int_0
+               OpStore %155 %float_1
+        %156 = OpAccessChain %_ptr_Workgroup_float %m35 %int_0 %int_0
+               OpStore %156 %float_1
+        %157 = OpAccessChain %_ptr_Workgroup_float %m36 %int_0 %int_0
+               OpStore %157 %float_1
+        %158 = OpAccessChain %_ptr_Workgroup_float %m37 %int_0 %int_0
+               OpStore %158 %float_1
+        %159 = OpAccessChain %_ptr_Workgroup_float %m38 %int_0 %int_0
+               OpStore %159 %float_1
+        %160 = OpAccessChain %_ptr_Workgroup_float %m39 %int_0 %int_0
+               OpStore %160 %float_1
+        %161 = OpAccessChain %_ptr_Workgroup_float %m40 %int_0 %int_0
+               OpStore %161 %float_1
+        %162 = OpAccessChain %_ptr_Workgroup_float %m41 %int_0 %int_0
+               OpStore %162 %float_1
+        %163 = OpAccessChain %_ptr_Workgroup_float %m42 %int_0 %int_0
+               OpStore %163 %float_1
+        %164 = OpAccessChain %_ptr_Workgroup_float %m43 %int_0 %int_0
+               OpStore %164 %float_1
+        %165 = OpAccessChain %_ptr_Workgroup_float %m44 %int_0 %int_0
+               OpStore %165 %float_1
+        %166 = OpAccessChain %_ptr_Workgroup_float %m45 %int_0 %int_0
+               OpStore %166 %float_1
+        %167 = OpAccessChain %_ptr_Workgroup_float %m46 %int_0 %int_0
+               OpStore %167 %float_1
+        %168 = OpAccessChain %_ptr_Workgroup_float %m47 %int_0 %int_0
+               OpStore %168 %float_1
+        %169 = OpAccessChain %_ptr_Workgroup_float %m48 %int_0 %int_0
+               OpStore %169 %float_1
+        %170 = OpAccessChain %_ptr_Workgroup_float %m49 %int_0 %int_0
+               OpStore %170 %float_1
+        %171 = OpAccessChain %_ptr_Workgroup_float %m50 %int_0 %int_0
+               OpStore %171 %float_1
+        %172 = OpAccessChain %_ptr_Workgroup_float %m51 %int_0 %int_0
+               OpStore %172 %float_1
+        %173 = OpAccessChain %_ptr_Workgroup_float %m52 %int_0 %int_0
+               OpStore %173 %float_1
+        %174 = OpAccessChain %_ptr_Workgroup_float %m53 %int_0 %int_0
+               OpStore %174 %float_1
+        %175 = OpAccessChain %_ptr_Workgroup_float %m54 %int_0 %int_0
+               OpStore %175 %float_1
+        %176 = OpAccessChain %_ptr_Workgroup_float %m55 %int_0 %int_0
+               OpStore %176 %float_1
+        %177 = OpAccessChain %_ptr_Workgroup_float %m56 %int_0 %int_0
+               OpStore %177 %float_1
+        %178 = OpAccessChain %_ptr_Workgroup_float %m57 %int_0 %int_0
+               OpStore %178 %float_1
+        %179 = OpAccessChain %_ptr_Workgroup_float %m58 %int_0 %int_0
+               OpStore %179 %float_1
+        %180 = OpAccessChain %_ptr_Workgroup_float %m59 %int_0 %int_0
+               OpStore %180 %float_1
+        %181 = OpAccessChain %_ptr_Workgroup_float %m60 %int_0 %int_0
+               OpStore %181 %float_1
+        %182 = OpAccessChain %_ptr_Workgroup_float %m61 %int_0 %int_0
+               OpStore %182 %float_1
+        %183 = OpAccessChain %_ptr_Workgroup_float %m62 %int_0 %int_0
+               OpStore %183 %float_1
+        %184 = OpAccessChain %_ptr_Workgroup_float %m63 %int_0 %int_0
+               OpStore %184 %float_1
+        %185 = OpAccessChain %_ptr_Workgroup_float %m64 %int_0 %int_0
+               OpStore %185 %float_1
+        %186 = OpAccessChain %_ptr_Workgroup_float %m65 %int_0 %int_0
+               OpStore %186 %float_1
+        %187 = OpAccessChain %_ptr_Workgroup_float %m66 %int_0 %int_0
+               OpStore %187 %float_1
+        %188 = OpAccessChain %_ptr_Workgroup_float %m67 %int_0 %int_0
+               OpStore %188 %float_1
+        %189 = OpAccessChain %_ptr_Workgroup_float %m68 %int_0 %int_0
+               OpStore %189 %float_1
+        %190 = OpAccessChain %_ptr_Workgroup_float %m69 %int_0 %int_0
+               OpStore %190 %float_1
+        %191 = OpAccessChain %_ptr_Workgroup_float %m70 %int_0 %int_0
+               OpStore %191 %float_1
+        %192 = OpAccessChain %_ptr_Workgroup_float %m71 %int_0 %int_0
+               OpStore %192 %float_1
+        %193 = OpAccessChain %_ptr_Workgroup_float %m72 %int_0 %int_0
+               OpStore %193 %float_1
+        %194 = OpAccessChain %_ptr_Workgroup_float %m73 %int_0 %int_0
+               OpStore %194 %float_1
+        %195 = OpAccessChain %_ptr_Workgroup_float %m74 %int_0 %int_0
+               OpStore %195 %float_1
+        %196 = OpAccessChain %_ptr_Workgroup_float %m75 %int_0 %int_0
+               OpStore %196 %float_1
+        %197 = OpAccessChain %_ptr_Workgroup_float %m76 %int_0 %int_0
+               OpStore %197 %float_1
+        %198 = OpAccessChain %_ptr_Workgroup_float %m77 %int_0 %int_0
+               OpStore %198 %float_1
+        %199 = OpAccessChain %_ptr_Workgroup_float %m78 %int_0 %int_0
+               OpStore %199 %float_1
+        %200 = OpAccessChain %_ptr_Workgroup_float %m79 %int_0 %int_0
+               OpStore %200 %float_1
+        %201 = OpAccessChain %_ptr_Workgroup_float %m80 %int_0 %int_0
+               OpStore %201 %float_1
+        %202 = OpAccessChain %_ptr_Workgroup_float %m81 %int_0 %int_0
+               OpStore %202 %float_1
+        %203 = OpAccessChain %_ptr_Workgroup_float %m82 %int_0 %int_0
+               OpStore %203 %float_1
+        %204 = OpAccessChain %_ptr_Workgroup_float %m83 %int_0 %int_0
+               OpStore %204 %float_1
+        %205 = OpAccessChain %_ptr_Workgroup_float %m84 %int_0 %int_0
+               OpStore %205 %float_1
+        %206 = OpAccessChain %_ptr_Workgroup_float %m85 %int_0 %int_0
+               OpStore %206 %float_1
+        %207 = OpAccessChain %_ptr_Workgroup_float %m86 %int_0 %int_0
+               OpStore %207 %float_1
+        %208 = OpAccessChain %_ptr_Workgroup_float %m87 %int_0 %int_0
+               OpStore %208 %float_1
+        %209 = OpAccessChain %_ptr_Workgroup_float %m88 %int_0 %int_0
+               OpStore %209 %float_1
+        %210 = OpAccessChain %_ptr_Workgroup_float %m89 %int_0 %int_0
+               OpStore %210 %float_1
+        %211 = OpAccessChain %_ptr_Workgroup_float %m90 %int_0 %int_0
+               OpStore %211 %float_1
+        %212 = OpAccessChain %_ptr_Workgroup_float %m91 %int_0 %int_0
+               OpStore %212 %float_1
+        %213 = OpAccessChain %_ptr_Workgroup_float %m92 %int_0 %int_0
+               OpStore %213 %float_1
+        %214 = OpAccessChain %_ptr_Workgroup_float %m93 %int_0 %int_0
+               OpStore %214 %float_1
+        %215 = OpAccessChain %_ptr_Workgroup_float %m94 %int_0 %int_0
+               OpStore %215 %float_1
+        %216 = OpAccessChain %_ptr_Workgroup_float %m95 %int_0 %int_0
+               OpStore %216 %float_1
+        %217 = OpAccessChain %_ptr_Workgroup_float %m96 %int_0 %int_0
+               OpStore %217 %float_1
+        %218 = OpAccessChain %_ptr_Workgroup_float %m97 %int_0 %int_0
+               OpStore %218 %float_1
+        %219 = OpAccessChain %_ptr_Workgroup_float %m98 %int_0 %int_0
+               OpStore %219 %float_1
+        %220 = OpAccessChain %_ptr_Workgroup_float %m99 %int_0 %int_0
+               OpStore %220 %float_1
+               OpReturn
+               OpFunctionEnd
+%tint_symbol = OpFunction %void None %221
+        %223 = OpLabel
+        %225 = OpLoad %uint %idx_1
+        %224 = OpFunctionCall %void %tint_symbol_inner %225
+               OpReturn
+               OpFunctionEnd

test/var/uses/many_workgroup_vars.wgsl.expected.wgsl

@@ -0,0 +1,303 @@
+var<workgroup> m00 : mat2x2<f32>;
+
+var<workgroup> m01 : mat2x2<f32>;
+
+var<workgroup> m02 : mat2x2<f32>;
+
+var<workgroup> m03 : mat2x2<f32>;
+
+var<workgroup> m04 : mat2x2<f32>;
+
+var<workgroup> m05 : mat2x2<f32>;
+
+var<workgroup> m06 : mat2x2<f32>;
+
+var<workgroup> m07 : mat2x2<f32>;
+
+var<workgroup> m08 : mat2x2<f32>;
+
+var<workgroup> m09 : mat2x2<f32>;
+
+var<workgroup> m10 : mat2x2<f32>;
+
+var<workgroup> m11 : mat2x2<f32>;
+
+var<workgroup> m12 : mat2x2<f32>;
+
+var<workgroup> m13 : mat2x2<f32>;
+
+var<workgroup> m14 : mat2x2<f32>;
+
+var<workgroup> m15 : mat2x2<f32>;
+
+var<workgroup> m16 : mat2x2<f32>;
+
+var<workgroup> m17 : mat2x2<f32>;
+
+var<workgroup> m18 : mat2x2<f32>;
+
+var<workgroup> m19 : mat2x2<f32>;
+
+var<workgroup> m20 : mat2x2<f32>;
+
+var<workgroup> m21 : mat2x2<f32>;
+
+var<workgroup> m22 : mat2x2<f32>;
+
+var<workgroup> m23 : mat2x2<f32>;
+
+var<workgroup> m24 : mat2x2<f32>;
+
+var<workgroup> m25 : mat2x2<f32>;
+
+var<workgroup> m26 : mat2x2<f32>;
+
+var<workgroup> m27 : mat2x2<f32>;
+
+var<workgroup> m28 : mat2x2<f32>;
+
+var<workgroup> m29 : mat2x2<f32>;
+
+var<workgroup> m30 : mat2x2<f32>;
+
+var<workgroup> m31 : mat2x2<f32>;
+
+var<workgroup> m32 : mat2x2<f32>;
+
+var<workgroup> m33 : mat2x2<f32>;
+
+var<workgroup> m34 : mat2x2<f32>;
+
+var<workgroup> m35 : mat2x2<f32>;
+
+var<workgroup> m36 : mat2x2<f32>;
+
+var<workgroup> m37 : mat2x2<f32>;
+
+var<workgroup> m38 : mat2x2<f32>;
+
+var<workgroup> m39 : mat2x2<f32>;
+
+var<workgroup> m40 : mat2x2<f32>;
+
+var<workgroup> m41 : mat2x2<f32>;
+
+var<workgroup> m42 : mat2x2<f32>;
+
+var<workgroup> m43 : mat2x2<f32>;
+
+var<workgroup> m44 : mat2x2<f32>;
+
+var<workgroup> m45 : mat2x2<f32>;
+
+var<workgroup> m46 : mat2x2<f32>;
+
+var<workgroup> m47 : mat2x2<f32>;
+
+var<workgroup> m48 : mat2x2<f32>;
+
+var<workgroup> m49 : mat2x2<f32>;
+
+var<workgroup> m50 : mat2x2<f32>;
+
+var<workgroup> m51 : mat2x2<f32>;
+
+var<workgroup> m52 : mat2x2<f32>;
+
+var<workgroup> m53 : mat2x2<f32>;
+
+var<workgroup> m54 : mat2x2<f32>;
+
+var<workgroup> m55 : mat2x2<f32>;
+
+var<workgroup> m56 : mat2x2<f32>;
+
+var<workgroup> m57 : mat2x2<f32>;
+
+var<workgroup> m58 : mat2x2<f32>;
+
+var<workgroup> m59 : mat2x2<f32>;
+
+var<workgroup> m60 : mat2x2<f32>;
+
+var<workgroup> m61 : mat2x2<f32>;
+
+var<workgroup> m62 : mat2x2<f32>;
+
+var<workgroup> m63 : mat2x2<f32>;
+
+var<workgroup> m64 : mat2x2<f32>;
+
+var<workgroup> m65 : mat2x2<f32>;
+
+var<workgroup> m66 : mat2x2<f32>;
+
+var<workgroup> m67 : mat2x2<f32>;
+
+var<workgroup> m68 : mat2x2<f32>;
+
+var<workgroup> m69 : mat2x2<f32>;
+
+var<workgroup> m70 : mat2x2<f32>;
+
+var<workgroup> m71 : mat2x2<f32>;
+
+var<workgroup> m72 : mat2x2<f32>;
+
+var<workgroup> m73 : mat2x2<f32>;
+
+var<workgroup> m74 : mat2x2<f32>;
+
+var<workgroup> m75 : mat2x2<f32>;
+
+var<workgroup> m76 : mat2x2<f32>;
+
+var<workgroup> m77 : mat2x2<f32>;
+
+var<workgroup> m78 : mat2x2<f32>;
+
+var<workgroup> m79 : mat2x2<f32>;
+
+var<workgroup> m80 : mat2x2<f32>;
+
+var<workgroup> m81 : mat2x2<f32>;
+
+var<workgroup> m82 : mat2x2<f32>;
+
+var<workgroup> m83 : mat2x2<f32>;
+
+var<workgroup> m84 : mat2x2<f32>;
+
+var<workgroup> m85 : mat2x2<f32>;
+
+var<workgroup> m86 : mat2x2<f32>;
+
+var<workgroup> m87 : mat2x2<f32>;
+
+var<workgroup> m88 : mat2x2<f32>;
+
+var<workgroup> m89 : mat2x2<f32>;
+
+var<workgroup> m90 : mat2x2<f32>;
+
+var<workgroup> m91 : mat2x2<f32>;
+
+var<workgroup> m92 : mat2x2<f32>;
+
+var<workgroup> m93 : mat2x2<f32>;
+
+var<workgroup> m94 : mat2x2<f32>;
+
+var<workgroup> m95 : mat2x2<f32>;
+
+var<workgroup> m96 : mat2x2<f32>;
+
+var<workgroup> m97 : mat2x2<f32>;
+
+var<workgroup> m98 : mat2x2<f32>;
+
+var<workgroup> m99 : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn tint_symbol([[builtin(local_invocation_index)]] idx : u32) {
+  m00[0][0] = 1.0;
+  m01[0][0] = 1.0;
+  m02[0][0] = 1.0;
+  m03[0][0] = 1.0;
+  m04[0][0] = 1.0;
+  m05[0][0] = 1.0;
+  m06[0][0] = 1.0;
+  m07[0][0] = 1.0;
+  m08[0][0] = 1.0;
+  m09[0][0] = 1.0;
+  m10[0][0] = 1.0;
+  m11[0][0] = 1.0;
+  m12[0][0] = 1.0;
+  m13[0][0] = 1.0;
+  m14[0][0] = 1.0;
+  m15[0][0] = 1.0;
+  m16[0][0] = 1.0;
+  m17[0][0] = 1.0;
+  m18[0][0] = 1.0;
+  m19[0][0] = 1.0;
+  m20[0][0] = 1.0;
+  m21[0][0] = 1.0;
+  m22[0][0] = 1.0;
+  m23[0][0] = 1.0;
+  m24[0][0] = 1.0;
+  m25[0][0] = 1.0;
+  m26[0][0] = 1.0;
+  m27[0][0] = 1.0;
+  m28[0][0] = 1.0;
+  m29[0][0] = 1.0;
+  m30[0][0] = 1.0;
+  m31[0][0] = 1.0;
+  m32[0][0] = 1.0;
+  m33[0][0] = 1.0;
+  m34[0][0] = 1.0;
+  m35[0][0] = 1.0;
+  m36[0][0] = 1.0;
+  m37[0][0] = 1.0;
+  m38[0][0] = 1.0;
+  m39[0][0] = 1.0;
+  m40[0][0] = 1.0;
+  m41[0][0] = 1.0;
+  m42[0][0] = 1.0;
+  m43[0][0] = 1.0;
+  m44[0][0] = 1.0;
+  m45[0][0] = 1.0;
+  m46[0][0] = 1.0;
+  m47[0][0] = 1.0;
+  m48[0][0] = 1.0;
+  m49[0][0] = 1.0;
+  m50[0][0] = 1.0;
+  m51[0][0] = 1.0;
+  m52[0][0] = 1.0;
+  m53[0][0] = 1.0;
+  m54[0][0] = 1.0;
+  m55[0][0] = 1.0;
+  m56[0][0] = 1.0;
+  m57[0][0] = 1.0;
+  m58[0][0] = 1.0;
+  m59[0][0] = 1.0;
+  m60[0][0] = 1.0;
+  m61[0][0] = 1.0;
+  m62[0][0] = 1.0;
+  m63[0][0] = 1.0;
+  m64[0][0] = 1.0;
+  m65[0][0] = 1.0;
+  m66[0][0] = 1.0;
+  m67[0][0] = 1.0;
+  m68[0][0] = 1.0;
+  m69[0][0] = 1.0;
+  m70[0][0] = 1.0;
+  m71[0][0] = 1.0;
+  m72[0][0] = 1.0;
+  m73[0][0] = 1.0;
+  m74[0][0] = 1.0;
+  m75[0][0] = 1.0;
+  m76[0][0] = 1.0;
+  m77[0][0] = 1.0;
+  m78[0][0] = 1.0;
+  m79[0][0] = 1.0;
+  m80[0][0] = 1.0;
+  m81[0][0] = 1.0;
+  m82[0][0] = 1.0;
+  m83[0][0] = 1.0;
+  m84[0][0] = 1.0;
+  m85[0][0] = 1.0;
+  m86[0][0] = 1.0;
+  m87[0][0] = 1.0;
+  m88[0][0] = 1.0;
+  m89[0][0] = 1.0;
+  m90[0][0] = 1.0;
+  m91[0][0] = 1.0;
+  m92[0][0] = 1.0;
+  m93[0][0] = 1.0;
+  m94[0][0] = 1.0;
+  m95[0][0] = 1.0;
+  m96[0][0] = 1.0;
+  m97[0][0] = 1.0;
+  m98[0][0] = 1.0;
+  m99[0][0] = 1.0;
+}