diff --git a/src/ast/disable_validation_decoration.cc b/src/ast/disable_validation_decoration.cc index 8708cec73f..846530a2e4 100644 --- a/src/ast/disable_validation_decoration.cc +++ b/src/ast/disable_validation_decoration.cc @@ -42,8 +42,6 @@ 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 5e8213fcbc..3ebf0bc112 100644 --- a/src/ast/disable_validation_decoration.h +++ b/src/ast/disable_validation_decoration.h @@ -43,10 +43,6 @@ 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 ba66017ba2..0b8c48d775 100644 --- a/src/resolver/resolver.cc +++ b/src/resolver/resolver.cc @@ -2659,10 +2659,7 @@ bool Resolver::ValidateFunctionCall(const ast::CallExpression* call, } } - if (!is_valid && - IsValidationEnabled( - param->declaration->decorations(), - ast::DisabledValidation::kIgnoreInvalidPointerArgument)) { + if (!is_valid) { 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 d882a6e928..e865cab6b0 100644 --- a/src/transform/module_scope_var_to_entry_point_param.cc +++ b/src/transform/module_scope_var_to_entry_point_param.cc @@ -47,27 +47,6 @@ bool ContainsMatrix(const sem::Type* type) { } return false; } - -// 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. -void CloneStructTypes(const sem::Type* ty, CloneContext& ctx) { - if (auto* str = ty->As()) { - // Recurse into members. - for (auto* member : str->Members()) { - CloneStructTypes(member->Type(), ctx); - } - - // 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(), ctx); - } -} - } // namespace ModuleScopeVarToEntryPointParam::ModuleScopeVarToEntryPointParam() = default; @@ -133,17 +112,6 @@ void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx, // Map module-scope variables onto their function-scope replacement. std::unordered_map var_to_symbol; - // 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; - }; - for (auto* var : func_sem->ReferencedModuleVariables()) { if (var->StorageClass() != ast::StorageClass::kPrivate && var->StorageClass() != ast::StorageClass::kWorkgroup && @@ -154,16 +122,13 @@ void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx, // 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()); - }; + auto* store_type = 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()) { + 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 = @@ -171,7 +136,7 @@ void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx, 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); + auto* param = ctx.dst->Param(new_var_symbol, store_type, decos); ctx.InsertFront(func_ast->params(), param); } else { if (var->StorageClass() == ast::StorageClass::kWorkgroup && @@ -179,24 +144,15 @@ void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx, // 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(), ctx); - - // 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)); + 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}); + ctx.InsertFront(func_ast->params(), param); is_pointer = true; } else { // For any other private or workgroup variable, redeclare it at @@ -208,7 +164,7 @@ 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)); @@ -217,21 +173,13 @@ void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx, } 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()) { + auto* param_type = store_type; + if (!store_type->is_handle()) { param_type = ctx.dst->ty.pointer(param_type, var->StorageClass()); is_pointer = true; - - // 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)); + ctx.dst->Param(new_var_symbol, param_type)); } // Replace all uses of the module-scope variable. @@ -258,22 +206,6 @@ void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx, 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()); 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 91476f479f..cff9c812dc 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, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol : ptr, [[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol_1 : ptr) { +fn bar(a : f32, b : f32, tint_symbol : ptr, tint_symbol_1 : ptr) { *(tint_symbol) = a; *(tint_symbol_1) = b; } -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) { +fn foo(a : f32, tint_symbol_2 : ptr, 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([[internal(disable_validation__ignore_invalid_pointer_argument)]] tint_symbol : ptr) { +fn foo(tint_symbol : ptr) { bar(tint_symbol); } @@ -340,13 +340,8 @@ 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_1 : ptr) { - let tint_symbol : ptr> = &((*(tint_symbol_1)).m); +fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol : ptr>) { let x = *(tint_symbol); } )"; @@ -381,13 +376,8 @@ 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_1 : ptr) { - let tint_symbol : ptr> = &((*(tint_symbol_1)).m); +fn main([[internal(disable_validation__entry_point_parameter)]] tint_symbol : ptr>) { let x = *(tint_symbol); } )"; diff --git a/src/writer/msl/generator_impl_test.cc b/src/writer/msl/generator_impl_test.cc index f2bb2d7247..af43a53eb5 100644 --- a/src/writer/msl/generator_impl_test.cc +++ b/src/writer/msl/generator_impl_test.cc @@ -142,10 +142,6 @@ 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(); @@ -154,8 +150,8 @@ void comp_main_inner(uint local_invocation_index, threadgroup float2x2* const ti float2x2 const x = *(tint_symbol); } -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)); +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); return; } @@ -182,9 +178,6 @@ 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)) { @@ -195,8 +188,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_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)); +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); return; } @@ -234,9 +227,6 @@ 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) { { @@ -247,8 +237,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 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)); +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); return; } @@ -301,22 +291,6 @@ 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(); @@ -329,42 +303,42 @@ void main1_inner(uint local_invocation_index, threadgroup float2x2* const tint_s float2x4 const a3 = *(tint_symbol_2); } -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)); +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); return; } -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) { +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) { { - *(tint_symbol_8) = float3x2(); - *(tint_symbol_9) = float3x3(); - *(tint_symbol_10) = float3x4(); + *(tint_symbol_6) = float3x2(); + *(tint_symbol_7) = float3x3(); + *(tint_symbol_8) = float3x4(); } threadgroup_barrier(mem_flags::mem_threadgroup); - float3x2 const a1 = *(tint_symbol_8); - float3x3 const a2 = *(tint_symbol_9); - float3x4 const a3 = *(tint_symbol_10); + float3x2 const a1 = *(tint_symbol_6); + float3x3 const a2 = *(tint_symbol_7); + float3x4 const a3 = *(tint_symbol_8); } -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)); +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); return; } -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) { +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) { { - *(tint_symbol_16) = float4x2(); - *(tint_symbol_17) = float4x3(); - *(tint_symbol_18) = float4x4(); + *(tint_symbol_12) = float4x2(); + *(tint_symbol_13) = float4x3(); + *(tint_symbol_14) = float4x4(); } threadgroup_barrier(mem_flags::mem_threadgroup); - float4x2 const a1 = *(tint_symbol_16); - float4x3 const a2 = *(tint_symbol_17); - float4x4 const a3 = *(tint_symbol_18); + float4x2 const a1 = *(tint_symbol_12); + float4x3 const a2 = *(tint_symbol_13); + float4x4 const a3 = *(tint_symbol_14); } -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)); +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); return; } @@ -379,12 +353,18 @@ 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(), 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)); + 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)); } } // namespace diff --git a/test/var/initialization/workgroup/matrix.wgsl.expected.msl b/test/var/initialization/workgroup/matrix.wgsl.expected.msl index 776f109888..54bf7bba2d 100644 --- a/test/var/initialization/workgroup/matrix.wgsl.expected.msl +++ b/test/var/initialization/workgroup/matrix.wgsl.expected.msl @@ -1,10 +1,6 @@ #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(); @@ -13,8 +9,8 @@ void tint_symbol_inner(uint local_invocation_index, threadgroup float2x3* const (void) *(tint_symbol_1); } -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)); +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); return; } diff --git a/test/var/uses/many_workgroup_vars.wgsl b/test/var/uses/many_workgroup_vars.wgsl deleted file mode 100644 index 7857da9643..0000000000 --- a/test/var/uses/many_workgroup_vars.wgsl +++ /dev/null @@ -1,204 +0,0 @@ -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 deleted file mode 100644 index 738895a4f4..0000000000 --- a/test/var/uses/many_workgroup_vars.wgsl.expected.hlsl +++ /dev/null @@ -1,316 +0,0 @@ -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 deleted file mode 100644 index f5c8304f39..0000000000 --- a/test/var/uses/many_workgroup_vars.wgsl.expected.msl +++ /dev/null @@ -1,317 +0,0 @@ -#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 deleted file mode 100644 index 6c8669beff..0000000000 --- a/test/var/uses/many_workgroup_vars.wgsl.expected.spvasm +++ /dev/null @@ -1,543 +0,0 @@ -; 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 deleted file mode 100644 index 873700f6ff..0000000000 --- a/test/var/uses/many_workgroup_vars.wgsl.expected.wgsl +++ /dev/null @@ -1,303 +0,0 @@ -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; -}