encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

msl: Use a struct for threadgroup memory arguments 

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

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

Bug: tint:938
Change-Id: Ia467db186e176a08f160455eab5fd3b3662f56b8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/65360
Auto-Submit: James Price <jrprice@google.com>
Kokoro: James Price <jrprice@google.com>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
This commit is contained in:
James Price 2021-09-29 18:56:17 +00:00 committed by Tint LUCI CQ
parent efe1f14685
commit 1ca6fbad8f
13 changed files with 2054 additions and 189 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/ast/disable_validation_decoration.cc
View File

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

4
src/ast/disable_validation_decoration.h
View File

@ -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

5
src/resolver/resolver.cc
View File

@ -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",

380
src/transform/module_scope_var_to_entry_point_param.cc
View File

@ -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<const ast::CallExpression*>;
std::unordered_map<const ast::Function*, CallList> calls_to_replace;
std::vector<ast::Function*> 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<sem::Struct>()) {
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::Struct*>(ast_str)));
ctx.Remove(ctx.src->AST().GlobalDeclarations(), ast_str);
} else if (auto* arr = ty->As<sem::Array>()) {
CloneStructTypes(arr->ElemType());
}
}
// Build a list of `&ident` expressions. We'll use this later to avoid
// 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<ast::IdentifierExpression*, ast::UnaryOpExpression*>
ident_to_address_of;
for (auto* node : ctx.src->ASTNodes().Objects()) {
auto* address_of = node->As<ast::UnaryOpExpression>();
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<const ast::CallExpression*>;
std::unordered_map<const ast::Function*, CallList> calls_to_replace;
std::vector<ast::Function*> 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<ast::IdentifierExpression>()) {
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<const sem::Variable*, Symbol> 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<ast::IdentifierExpression*, ast::UnaryOpExpression*>
ident_to_address_of;
for (auto* node : ctx.src->ASTNodes().Objects()) {
auto* address_of = node->As<ast::UnaryOpExpression>();
if (!address_of || address_of->op() != ast::UnaryOp::kAddressOf) {
continue;
}
if (auto* ident = address_of->expr()->As<ast::IdentifierExpression>()) {
ident_to_address_of[ident] = address_of;
}
}
// This is the symbol for the variable that replaces the module-scope var.
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<const sem::Variable*, Symbol> 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<ast::DisableValidationDecoration>(
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<ast::DisableValidationDecoration>(
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<ast::IdentifierExpression>();
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<ast::DisableValidationDecoration>(
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<ast::IdentifierExpression>();
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<ast::DisableValidationDecoration>(
ctx.dst->ID(), ast::DisabledValidation::kEntryPointParameter);
auto* param =
ctx.dst->Param(workgroup_param(), param_type, {disable_validation});
ctx.InsertFront(func_ast->params(), param);
}
// Pass the variables as pointers to any functions that need them.
for (auto* call : calls_to_replace[func_ast]) {
auto* target = ctx.src->AST().Functions().Find(call->func()->symbol());
auto* target_sem = ctx.src->Sem().Get(target);
// Add new arguments for any variables that are needed by the callee.
// For entry points, pass non-handle types as pointers.
for (auto* target_var : target_sem->ReferencedModuleVariables()) {
bool is_handle = target_var->Type()->UnwrapRef()->is_handle();
bool is_workgroup_matrix =
target_var->StorageClass() == ast::StorageClass::kWorkgroup &&
ContainsMatrix(target_var->Type());
if (target_var->StorageClass() == ast::StorageClass::kPrivate ||
target_var->StorageClass() == ast::StorageClass::kWorkgroup ||
target_var->StorageClass() ==
ast::StorageClass::kUniformConstant) {
ast::Expression* arg = ctx.dst->Expr(var_to_symbol[target_var]);
if (is_entry_point && !is_handle && !is_workgroup_matrix) {
arg = ctx.dst->AddressOf(arg);
}
ctx.InsertBack(call->params(), arg);
}
}
}
}
// Pass the variables as pointers to any functions that need them.
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<const sem::Struct*> cloned_structs_;
};
ModuleScopeVarToEntryPointParam::ModuleScopeVarToEntryPointParam() = default;
ModuleScopeVarToEntryPointParam::~ModuleScopeVarToEntryPointParam() = default;
void ModuleScopeVarToEntryPointParam::Run(CloneContext& ctx,
const DataMap&,
DataMap&) {
State state{ctx};
state.Process();
ctx.Clone();
}

2
src/transform/module_scope_var_to_entry_point_param.h
View File

@ -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

63
src/transform/module_scope_var_to_entry_point_param_test.cc
View File

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

96
src/writer/msl/generator_impl_test.cc
View File

@ -142,6 +142,10 @@ TEST_F(MslGeneratorImplTest, WorkgroupMatrix) {
EXPECT_EQ(gen.result(), R"(#include <metal_stdlib>
using namespace metal;
struct tint_symbol_3 {
float2x2 m;
};
void comp_main_inner(uint local_invocation_index, threadgroup float2x2* const tint_symbol) {
{
*(tint_symbol) = float2x2();
@ -150,8 +154,8 @@ void comp_main_inner(uint local_invocation_index, threadgroup float2x2* const ti
float2x2 const x = *(tint_symbol);
}
kernel void comp_main(threadgroup float2x2* tint_symbol_1 [[threadgroup(0)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
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 <metal_stdlib>
using namespace metal;
struct tint_symbol_7 {
float2x2 m1;
float2x3 m2;
float2x4 m3;
};
struct tint_symbol_15 {
float3x2 m4;
float3x3 m5;
float3x4 m6;
};
struct tint_symbol_23 {
float4x2 m7;
float4x3 m8;
float4x4 m9;
};
void main1_inner(uint local_invocation_index, threadgroup float2x2* const tint_symbol, threadgroup float2x3* const tint_symbol_1, threadgroup float2x4* const tint_symbol_2) {
{
*(tint_symbol) = float2x2();
@ -303,42 +329,42 @@ void main1_inner(uint local_invocation_index, threadgroup float2x2* const tint_s
float2x4 const a3 = *(tint_symbol_2);
}
kernel void main1(threadgroup float2x2* tint_symbol_3 [[threadgroup(0)]], threadgroup float2x3* tint_symbol_4 [[threadgroup(1)]], threadgroup float2x4* tint_symbol_5 [[threadgroup(2)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
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

8
test/var/initialization/workgroup/matrix.wgsl.expected.msl
View File

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

204
test/var/uses/many_workgroup_vars.wgsl Normal file
View File

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

316
test/var/uses/many_workgroup_vars.wgsl.expected.hlsl Normal file
View File

@ -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;
}

317
test/var/uses/many_workgroup_vars.wgsl.expected.msl Normal file
View File

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

543
test/var/uses/many_workgroup_vars.wgsl.expected.spvasm Normal file
View File

@ -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

303
test/var/uses/many_workgroup_vars.wgsl.expected.wgsl Normal file
View File

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