From 1ca6fbad8f231ffe3d68eb88e939b5eafba119e2 Mon Sep 17 00:00:00 2001 From: James Price Date: Wed, 29 Sep 2021 18:56:17 +0000 Subject: [PATCH] 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 Kokoro: James Price Commit-Queue: James Price Reviewed-by: Ben Clayton --- src/ast/disable_validation_decoration.cc | 2 + src/ast/disable_validation_decoration.h | 4 + src/resolver/resolver.cc | 5 +- .../module_scope_var_to_entry_point_param.cc | 380 +++++++----- .../module_scope_var_to_entry_point_param.h | 2 + ...ule_scope_var_to_entry_point_param_test.cc | 63 +- src/writer/msl/generator_impl_test.cc | 96 ++-- .../workgroup/matrix.wgsl.expected.msl | 8 +- test/var/uses/many_workgroup_vars.wgsl | 204 +++++++ .../many_workgroup_vars.wgsl.expected.hlsl | 316 ++++++++++ .../many_workgroup_vars.wgsl.expected.msl | 317 ++++++++++ .../many_workgroup_vars.wgsl.expected.spvasm | 543 ++++++++++++++++++ .../many_workgroup_vars.wgsl.expected.wgsl | 303 ++++++++++ 13 files changed, 2054 insertions(+), 189 deletions(-) create mode 100644 test/var/uses/many_workgroup_vars.wgsl create mode 100644 test/var/uses/many_workgroup_vars.wgsl.expected.hlsl create mode 100644 test/var/uses/many_workgroup_vars.wgsl.expected.msl create mode 100644 test/var/uses/many_workgroup_vars.wgsl.expected.spvasm create mode 100644 test/var/uses/many_workgroup_vars.wgsl.expected.wgsl diff --git a/src/ast/disable_validation_decoration.cc b/src/ast/disable_validation_decoration.cc index 846530a2e4..8708cec73f 100644 --- a/src/ast/disable_validation_decoration.cc +++ b/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 ""; } diff --git a/src/ast/disable_validation_decoration.h b/src/ast/disable_validation_decoration.h index 3ebf0bc112..5e8213fcbc 100644 --- a/src/ast/disable_validation_decoration.h +++ b/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 diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc index 0b8c48d775..ba66017ba2 100644 --- a/src/resolver/resolver.cc +++ b/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", diff --git a/src/transform/module_scope_var_to_entry_point_param.cc b/src/transform/module_scope_var_to_entry_point_param.cc index e865cab6b0..787de981b3 100644 --- a/src/transform/module_scope_var_to_entry_point_param.cc +++ b/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, - const DataMap&, - DataMap&) { - // Predetermine the list of function calls that need to be replaced. - using CallList = std::vector; - std::unordered_map calls_to_replace; - - std::vector functions_to_process; - - // 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; + /// Clone any struct types that are contained in `ty` (including `ty` itself), + /// and add it to the global declarations now, so that they precede new global + /// declarations that need to reference them. + /// @param ty the type to clone + void CloneStructTypes(const sem::Type* ty) { + if (auto* str = ty->As()) { + if (!cloned_structs_.emplace(str).second) { + // The struct has already been cloned. + return; } - } - 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); + // Recurse into members. + for (auto* member : str->Members()) { + CloneStructTypes(member->Type()); } + + // 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_str))); + ctx.Remove(ctx.src->AST().GlobalDeclarations(), ast_str); + } else if (auto* arr = ty->As()) { + CloneStructTypes(arr->ElemType()); } } - // Build a list of `&ident` expressions. We'll use this later to avoid - // generating expressions of the form `&*ident`, which break WGSL validation - // rules when this expression is passed to a function. - // TODO(jrprice): We should add support for bidirectional SEM tree traversal - // so that we can do this on the fly instead. - std::unordered_map - ident_to_address_of; - for (auto* node : ctx.src->ASTNodes().Objects()) { - auto* address_of = node->As(); - if (!address_of || address_of->op() != ast::UnaryOp::kAddressOf) { - continue; + /// Process the module. + void Process() { + // Predetermine the list of function calls that need to be replaced. + using CallList = std::vector; + std::unordered_map calls_to_replace; + + std::vector functions_to_process; + + // 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) { + 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()) { - ident_to_address_of[ident] = address_of; - } - } - for (auto* func_ast : functions_to_process) { - auto* func_sem = ctx.src->Sem().Get(func_ast); - bool is_entry_point = func_ast->IsEntryPoint(); - - // Map module-scope variables onto their function-scope replacement. - std::unordered_map var_to_symbol; - - for (auto* var : func_sem->ReferencedModuleVariables()) { - if (var->StorageClass() != ast::StorageClass::kPrivate && - var->StorageClass() != ast::StorageClass::kWorkgroup && - var->StorageClass() != ast::StorageClass::kUniformConstant) { + // Build a list of `&ident` expressions. We'll use this later to avoid + // generating expressions of the form `&*ident`, which break WGSL validation + // rules when this expression is passed to a function. + // TODO(jrprice): We should add support for bidirectional SEM tree traversal + // so that we can do this on the fly instead. + std::unordered_map + ident_to_address_of; + for (auto* node : ctx.src->ASTNodes().Objects()) { + auto* address_of = node->As(); + if (!address_of || address_of->op() != ast::UnaryOp::kAddressOf) { continue; } + if (auto* ident = address_of->expr()->As()) { + ident_to_address_of[ident] = address_of; + } + } - // This is the symbol for the variable that replaces the module-scope var. - auto new_var_symbol = ctx.dst->Sym(); + for (auto* func_ast : functions_to_process) { + 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 var_to_symbol; - // Track whether the new variable is a pointer or not. - bool is_pointer = false; + // We aggregate all workgroup variables into a struct to avoid hitting + // 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) { - if (store_type->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( - ctx.dst->ID(), ast::DisabledValidation::kEntryPointParameter); - auto decos = ctx.Clone(var->Declaration()->decorations()); - decos.push_back(disable_validation); - auto* param = ctx.dst->Param(new_var_symbol, store_type, decos); - 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. + for (auto* var : func_sem->ReferencedModuleVariables()) { + if (var->StorageClass() != ast::StorageClass::kPrivate && + var->StorageClass() != ast::StorageClass::kWorkgroup && + var->StorageClass() != ast::StorageClass::kUniformConstant) { + continue; + } + + // This is the symbol for the variable that replaces the module-scope + // var. + auto new_var_symbol = ctx.dst->Sym(); + + // Helper to create an AST node for the store type of the variable. + auto store_type = [&]() { + return CreateASTTypeFor(ctx, var->Type()->UnwrapRef()); + }; + + // 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( ctx.dst->ID(), ast::DisabledValidation::kEntryPointParameter); - auto* param_type = - ctx.dst->ty.pointer(store_type, var->StorageClass()); - auto* param = ctx.dst->Param(new_var_symbol, param_type, - {disable_validation}); + auto decos = ctx.Clone(var->Declaration()->decorations()); + decos.push_back(disable_validation); + auto* param = ctx.dst->Param(new_var_symbol, store_type(), decos); 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 { - // For a regular function, redeclare the variable as a parameter. - // Use a pointer for non-handle types. - auto* param_type = store_type; - if (!store_type->is_handle()) { - param_type = ctx.dst->ty.pointer(param_type, var->StorageClass()); - is_pointer = true; - } - ctx.InsertBack(func_ast->params(), - ctx.dst->Param(new_var_symbol, param_type)); - } + } else { + // For a regular function, redeclare the variable as a parameter. + // Use a pointer for non-handle types. + auto* param_type = store_type(); + ast::DecorationList attributes; + if (!param_type->is_handle()) { + param_type = ctx.dst->ty.pointer(param_type, var->StorageClass()); + is_pointer = true; - // 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(); - if (ident_to_address_of.count(ident)) { - ctx.Replace(ident_to_address_of[ident], expr); - continue; + // Disable validation of arguments passed to this pointer parameter, + // as we will sometimes pass pointers to struct members. + attributes.push_back( + ctx.dst->ASTNodes().Create( + ctx.dst->ID(), + ast::DisabledValidation::kIgnoreInvalidPointerArgument)); + } + ctx.InsertBack( + func_ast->params(), + ctx.dst->Param(new_var_symbol, param_type, attributes)); + } + + // 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(); + if (ident_to_address_of.count(ident)) { + ctx.Replace(ident_to_address_of[ident], expr); + continue; + } + + expr = ctx.dst->Deref(expr); } - - expr = ctx.dst->Deref(expr); + ctx.Replace(user->Declaration(), 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( + 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. - 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); - } + // Now remove all module-scope variables with these storage classes. + for (auto* var_ast : ctx.src->AST().GlobalVariables()) { + 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); } } } - // Now remove all module-scope variables with these storage classes. - for (auto* var_ast : ctx.src->AST().GlobalVariables()) { - 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); - } - } + private: + std::unordered_set cloned_structs_; +}; +ModuleScopeVarToEntryPointParam::ModuleScopeVarToEntryPointParam() = default; + +ModuleScopeVarToEntryPointParam::~ModuleScopeVarToEntryPointParam() = default; + +void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx, + const DataMap&, + DataMap&) { + State state{ctx}; + state.Process(); ctx.Clone(); } diff --git a/src/transform/module_scope_var_to_entry_point_param.h b/src/transform/module_scope_var_to_entry_point_param.h index 46967d9bb1..9e448b44bc 100644 --- a/src/transform/module_scope_var_to_entry_point_param.h +++ b/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 diff --git a/src/transform/module_scope_var_to_entry_point_param_test.cc b/src/transform/module_scope_var_to_entry_point_param_test.cc index cff9c812dc..b7e13daaf8 100644 --- a/src/transform/module_scope_var_to_entry_point_param_test.cc +++ b/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, tint_symbol_1 : ptr) { +fn bar(a : f32, b : f32, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol : ptr, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol_1 : ptr) { *(tint_symbol) = a; *(tint_symbol_1) = b; } -fn foo(a : f32, tint_symbol_2 : ptr, tint_symbol_3 : ptr) { +fn foo(a : f32, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol_2 : ptr, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol_3 : ptr) { let b : f32 = 2.0; bar(a, b, tint_symbol_2, tint_symbol_3); no_uses(); @@ -181,7 +181,7 @@ fn bar(p : ptr) { *(p) = 0.0; } -fn foo(tint_symbol : ptr) { +fn foo([[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol : ptr) { bar(tint_symbol); } @@ -340,8 +340,13 @@ fn main() { )"; auto* expect = R"( +struct tint_symbol_2 { + m : mat2x2; +}; + [[stage(compute), workgroup_size(1)]] -fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol : ptr>) { +fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol_1 : ptr) { + let tint_symbol : ptr> = &((*(tint_symbol_1)).m); let x = *(tint_symbol); } )"; @@ -376,8 +381,13 @@ struct S2 { s : S1; }; +struct tint_symbol_2 { + m : array; +}; + [[stage(compute), workgroup_size(1)]] -fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol : ptr>) { +fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol_1 : ptr) { + let tint_symbol : ptr> = &((*(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; +}; + +var a : S; + +var b : S; + +[[stage(compute), workgroup_size(1)]] +fn main() { + let x = a; + let y = b; +} +)"; + + auto* expect = R"( +struct S { + m : mat2x2; +}; + +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) { + let tint_symbol : ptr = &((*(tint_symbol_1)).a); + let tint_symbol_2 : ptr = &((*(tint_symbol_1)).b); + let x = *(tint_symbol); + let y = *(tint_symbol_2); +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + TEST_F(ModuleScopeVarToEntryPointParamTest, EmtpyModule) { auto* src = ""; diff --git a/src/writer/msl/generator_impl_test.cc b/src/writer/msl/generator_impl_test.cc index af43a53eb5..f2bb2d7247 100644 --- a/src/writer/msl/generator_impl_test.cc +++ b/src/writer/msl/generator_impl_test.cc @@ -142,6 +142,10 @@ TEST_F(MslGeneratorImplTest, WorkgroupMatrix) { EXPECT_EQ(gen.result(), R"(#include 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]]) { - comp_main_inner(local_invocation_index, tint_symbol_1); +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_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]]) { - comp_main_inner(local_invocation_index, tint_symbol_1); +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_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]]) { - comp_main_inner(local_invocation_index, tint_symbol_2); +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_3)).s)); return; } @@ -291,6 +301,22 @@ TEST_F(MslGeneratorImplTest, WorkgroupMatrix_Multiples) { EXPECT_EQ(gen.result(), R"(#include 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]]) { - main1_inner(local_invocation_index, tint_symbol_3, tint_symbol_4, tint_symbol_5); +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_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_7) = float3x3(); - *(tint_symbol_8) = float3x4(); + *(tint_symbol_8) = float3x2(); + *(tint_symbol_9) = float3x3(); + *(tint_symbol_10) = float3x4(); } threadgroup_barrier(mem_flags::mem_threadgroup); - float3x2 const a1 = *(tint_symbol_6); - float3x3 const a2 = *(tint_symbol_7); - float3x4 const a3 = *(tint_symbol_8); + float3x2 const a1 = *(tint_symbol_8); + float3x3 const a2 = *(tint_symbol_9); + 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]]) { - main2_inner(local_invocation_index_1, tint_symbol_9, tint_symbol_10, tint_symbol_11); +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_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_13) = float4x3(); - *(tint_symbol_14) = float4x4(); + *(tint_symbol_16) = float4x2(); + *(tint_symbol_17) = float4x3(); + *(tint_symbol_18) = float4x4(); } threadgroup_barrier(mem_flags::mem_threadgroup); - float4x2 const a1 = *(tint_symbol_12); - float4x3 const a2 = *(tint_symbol_13); - float4x4 const a3 = *(tint_symbol_14); + float4x2 const a1 = *(tint_symbol_16); + float4x3 const a2 = *(tint_symbol_17); + 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]]) { - main3_inner(local_invocation_index_2, tint_symbol_15, tint_symbol_16, tint_symbol_17); +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_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); - EXPECT_EQ(allocations["main1"][0], 2u * 2u * sizeof(float)); - EXPECT_EQ(allocations["main1"][1], 2u * 4u * sizeof(float)); - EXPECT_EQ(allocations["main1"][2], 2u * 4u * sizeof(float)); - ASSERT_EQ(allocations["main2"].size(), 3u); - EXPECT_EQ(allocations["main2"][0], 3u * 2u * 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)); + ASSERT_EQ(allocations["main1"].size(), 1u); + EXPECT_EQ(allocations["main1"][0], 20u * sizeof(float)); + ASSERT_EQ(allocations["main2"].size(), 1u); + EXPECT_EQ(allocations["main2"][0], 32u * sizeof(float)); + ASSERT_EQ(allocations["main3"].size(), 1u); + EXPECT_EQ(allocations["main3"][0], 40u * sizeof(float)); } } // namespace diff --git a/test/var/initialization/workgroup/matrix.wgsl.expected.msl b/test/var/initialization/workgroup/matrix.wgsl.expected.msl index 54bf7bba2d..776f109888 100644 --- a/test/var/initialization/workgroup/matrix.wgsl.expected.msl +++ b/test/var/initialization/workgroup/matrix.wgsl.expected.msl @@ -1,6 +1,10 @@ #include 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]]) { - tint_symbol_inner(local_invocation_index, tint_symbol_2); +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_3)).v)); return; } diff --git a/test/var/uses/many_workgroup_vars.wgsl b/test/var/uses/many_workgroup_vars.wgsl new file mode 100644 index 0000000000..7857da9643 --- /dev/null +++ b/test/var/uses/many_workgroup_vars.wgsl @@ -0,0 +1,204 @@ +var m00 : mat2x2; +var m01 : mat2x2; +var m02 : mat2x2; +var m03 : mat2x2; +var m04 : mat2x2; +var m05 : mat2x2; +var m06 : mat2x2; +var m07 : mat2x2; +var m08 : mat2x2; +var m09 : mat2x2; +var m10 : mat2x2; +var m11 : mat2x2; +var m12 : mat2x2; +var m13 : mat2x2; +var m14 : mat2x2; +var m15 : mat2x2; +var m16 : mat2x2; +var m17 : mat2x2; +var m18 : mat2x2; +var m19 : mat2x2; +var m20 : mat2x2; +var m21 : mat2x2; +var m22 : mat2x2; +var m23 : mat2x2; +var m24 : mat2x2; +var m25 : mat2x2; +var m26 : mat2x2; +var m27 : mat2x2; +var m28 : mat2x2; +var m29 : mat2x2; +var m30 : mat2x2; +var m31 : mat2x2; +var m32 : mat2x2; +var m33 : mat2x2; +var m34 : mat2x2; +var m35 : mat2x2; +var m36 : mat2x2; +var m37 : mat2x2; +var m38 : mat2x2; +var m39 : mat2x2; +var m40 : mat2x2; +var m41 : mat2x2; +var m42 : mat2x2; +var m43 : mat2x2; +var m44 : mat2x2; +var m45 : mat2x2; +var m46 : mat2x2; +var m47 : mat2x2; +var m48 : mat2x2; +var m49 : mat2x2; +var m50 : mat2x2; +var m51 : mat2x2; +var m52 : mat2x2; +var m53 : mat2x2; +var m54 : mat2x2; +var m55 : mat2x2; +var m56 : mat2x2; +var m57 : mat2x2; +var m58 : mat2x2; +var m59 : mat2x2; +var m60 : mat2x2; +var m61 : mat2x2; +var m62 : mat2x2; +var m63 : mat2x2; +var m64 : mat2x2; +var m65 : mat2x2; +var m66 : mat2x2; +var m67 : mat2x2; +var m68 : mat2x2; +var m69 : mat2x2; +var m70 : mat2x2; +var m71 : mat2x2; +var m72 : mat2x2; +var m73 : mat2x2; +var m74 : mat2x2; +var m75 : mat2x2; +var m76 : mat2x2; +var m77 : mat2x2; +var m78 : mat2x2; +var m79 : mat2x2; +var m80 : mat2x2; +var m81 : mat2x2; +var m82 : mat2x2; +var m83 : mat2x2; +var m84 : mat2x2; +var m85 : mat2x2; +var m86 : mat2x2; +var m87 : mat2x2; +var m88 : mat2x2; +var m89 : mat2x2; +var m90 : mat2x2; +var m91 : mat2x2; +var m92 : mat2x2; +var m93 : mat2x2; +var m94 : mat2x2; +var m95 : mat2x2; +var m96 : mat2x2; +var m97 : mat2x2; +var m98 : mat2x2; +var m99 : mat2x2; + +[[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; +} \ No newline at end of file diff --git a/test/var/uses/many_workgroup_vars.wgsl.expected.hlsl b/test/var/uses/many_workgroup_vars.wgsl.expected.hlsl new file mode 100644 index 0000000000..738895a4f4 --- /dev/null +++ b/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; +} diff --git a/test/var/uses/many_workgroup_vars.wgsl.expected.msl b/test/var/uses/many_workgroup_vars.wgsl.expected.msl new file mode 100644 index 0000000000..f5c8304f39 --- /dev/null +++ b/test/var/uses/many_workgroup_vars.wgsl.expected.msl @@ -0,0 +1,317 @@ +#include + +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; +} + diff --git a/test/var/uses/many_workgroup_vars.wgsl.expected.spvasm b/test/var/uses/many_workgroup_vars.wgsl.expected.spvasm new file mode 100644 index 0000000000..6c8669beff --- /dev/null +++ b/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 diff --git a/test/var/uses/many_workgroup_vars.wgsl.expected.wgsl b/test/var/uses/many_workgroup_vars.wgsl.expected.wgsl new file mode 100644 index 0000000000..873700f6ff --- /dev/null +++ b/test/var/uses/many_workgroup_vars.wgsl.expected.wgsl @@ -0,0 +1,303 @@ +var m00 : mat2x2; + +var m01 : mat2x2; + +var m02 : mat2x2; + +var m03 : mat2x2; + +var m04 : mat2x2; + +var m05 : mat2x2; + +var m06 : mat2x2; + +var m07 : mat2x2; + +var m08 : mat2x2; + +var m09 : mat2x2; + +var m10 : mat2x2; + +var m11 : mat2x2; + +var m12 : mat2x2; + +var m13 : mat2x2; + +var m14 : mat2x2; + +var m15 : mat2x2; + +var m16 : mat2x2; + +var m17 : mat2x2; + +var m18 : mat2x2; + +var m19 : mat2x2; + +var m20 : mat2x2; + +var m21 : mat2x2; + +var m22 : mat2x2; + +var m23 : mat2x2; + +var m24 : mat2x2; + +var m25 : mat2x2; + +var m26 : mat2x2; + +var m27 : mat2x2; + +var m28 : mat2x2; + +var m29 : mat2x2; + +var m30 : mat2x2; + +var m31 : mat2x2; + +var m32 : mat2x2; + +var m33 : mat2x2; + +var m34 : mat2x2; + +var m35 : mat2x2; + +var m36 : mat2x2; + +var m37 : mat2x2; + +var m38 : mat2x2; + +var m39 : mat2x2; + +var m40 : mat2x2; + +var m41 : mat2x2; + +var m42 : mat2x2; + +var m43 : mat2x2; + +var m44 : mat2x2; + +var m45 : mat2x2; + +var m46 : mat2x2; + +var m47 : mat2x2; + +var m48 : mat2x2; + +var m49 : mat2x2; + +var m50 : mat2x2; + +var m51 : mat2x2; + +var m52 : mat2x2; + +var m53 : mat2x2; + +var m54 : mat2x2; + +var m55 : mat2x2; + +var m56 : mat2x2; + +var m57 : mat2x2; + +var m58 : mat2x2; + +var m59 : mat2x2; + +var m60 : mat2x2; + +var m61 : mat2x2; + +var m62 : mat2x2; + +var m63 : mat2x2; + +var m64 : mat2x2; + +var m65 : mat2x2; + +var m66 : mat2x2; + +var m67 : mat2x2; + +var m68 : mat2x2; + +var m69 : mat2x2; + +var m70 : mat2x2; + +var m71 : mat2x2; + +var m72 : mat2x2; + +var m73 : mat2x2; + +var m74 : mat2x2; + +var m75 : mat2x2; + +var m76 : mat2x2; + +var m77 : mat2x2; + +var m78 : mat2x2; + +var m79 : mat2x2; + +var m80 : mat2x2; + +var m81 : mat2x2; + +var m82 : mat2x2; + +var m83 : mat2x2; + +var m84 : mat2x2; + +var m85 : mat2x2; + +var m86 : mat2x2; + +var m87 : mat2x2; + +var m88 : mat2x2; + +var m89 : mat2x2; + +var m90 : mat2x2; + +var m91 : mat2x2; + +var m92 : mat2x2; + +var m93 : mat2x2; + +var m94 : mat2x2; + +var m95 : mat2x2; + +var m96 : mat2x2; + +var m97 : mat2x2; + +var m98 : mat2x2; + +var m99 : mat2x2; + +[[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; +}