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

tint/msl: Fix emission of private variables 

In order to avoid declaring too many function parameters, we
previously modified this transform to redeclare private variables that
are only used inside a single function as function-scope
variables. This was broken as it meant that their values did not
persist across multiple calls to the same function.

Instead, wrap all private variables in a structure and pass it around
as a pointer.

Fixed: tint:1875
Change-Id: I83f5eb1071d57b9c6af56d6cf21b3a32c6e94260
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/124800
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: James Price <jrprice@google.com>
This commit is contained in:
James Price 2023-03-20 21:46:01 +00:00 committed by Dawn LUCI CQ
parent 2a7005f416
commit 0e22bdbae7
788 changed files with 9507 additions and 6127 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

198
src/tint/transform/module_scope_var_to_entry_point_param.cc
View File

@ -33,7 +33,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::transform::ModuleScopeVarToEntryPointParam);
namespace tint::transform {
namespace {
using WorkgroupParameterMemberList = utils::Vector<const ast::StructMember*, 8>;
using StructMemberList = utils::Vector<const ast::StructMember*, 8>;
// The name of the struct member for arrays that are wrapped in structures.
const char* kWrappedArrayMemberName = "arr";
@ -114,7 +114,7 @@ struct ModuleScopeVarToEntryPointParam::State {
const sem::Variable* var,
Symbol new_var_symbol,
std::function<Symbol()> workgroup_param,
WorkgroupParameterMemberList& workgroup_parameter_members,
StructMemberList& workgroup_parameter_members,
bool& is_pointer,
bool& is_wrapped) {
auto* ty = var->Type()->UnwrapRef();
@ -188,21 +188,14 @@ struct ModuleScopeVarToEntryPointParam::State {
member_ptr);
ctx.InsertFront(func->body->statements, ctx.dst->Decl(local_var));
is_pointer = true;
break;
} else {
auto* disable_validation =
ctx.dst->Disable(ast::DisabledValidation::kIgnoreAddressSpace);
auto* initializer = ctx.Clone(var->Declaration()->initializer);
auto* local_var = ctx.dst->Var(new_var_symbol, store_type(), sc, initializer,
utils::Vector{disable_validation});
ctx.InsertFront(func->body->statements, ctx.dst->Decl(local_var));
}
[[fallthrough]];
}
case builtin::AddressSpace::kPrivate: {
// Variables in the Private and Workgroup address spaces are redeclared at function
// scope. Disable address space validation on this variable.
auto* disable_validation =
ctx.dst->Disable(ast::DisabledValidation::kIgnoreAddressSpace);
auto* initializer = ctx.Clone(var->Declaration()->initializer);
auto* local_var = ctx.dst->Var(new_var_symbol, store_type(), sc, initializer,
utils::Vector{disable_validation});
ctx.InsertFront(func->body->statements, ctx.dst->Decl(local_var));
break;
}
case builtin::AddressSpace::kPushConstant: {
@ -234,6 +227,8 @@ struct ModuleScopeVarToEntryPointParam::State {
auto sc = var->AddressSpace();
switch (sc) {
case builtin::AddressSpace::kPrivate:
// Private variables are passed all together in a struct.
return;
case builtin::AddressSpace::kStorage:
case builtin::AddressSpace::kUniform:
case builtin::AddressSpace::kHandle:
@ -275,12 +270,12 @@ struct ModuleScopeVarToEntryPointParam::State {
/// @param var the variable to replace
/// @param new_var the symbol to use for replacement
/// @param is_pointer true if `new_var` is a pointer to the new variable
/// @param is_wrapped true if `new_var` is an array wrapped in a structure
/// @param member_name if valid, the name of the struct member that holds this variable
void ReplaceUsesInFunction(const ast::Function* func,
const sem::Variable* var,
Symbol new_var,
bool is_pointer,
bool is_wrapped) {
Symbol member_name) {
for (auto* user : var->Users()) {
if (user->Stmt()->Function()->Declaration() == func) {
const ast::Expression* expr = ctx.dst->Expr(new_var);
@ -288,16 +283,16 @@ struct ModuleScopeVarToEntryPointParam::State {
// 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)) {
if (ident_to_address_of_.count(ident) && !member_name.IsValid()) {
ctx.Replace(ident_to_address_of_[ident], expr);
continue;
}
expr = ctx.dst->Deref(expr);
}
if (is_wrapped) {
// Get the member from the wrapper structure.
expr = ctx.dst->MemberAccessor(expr, kWrappedArrayMemberName);
if (member_name.IsValid()) {
// Get the member from the containing structure.
expr = ctx.dst->MemberAccessor(expr, member_name);
}
ctx.Replace(user->Declaration(), expr);
}
@ -312,8 +307,34 @@ struct ModuleScopeVarToEntryPointParam::State {
utils::Vector<const ast::Function*, 8> functions_to_process;
// Collect private variables into a single structure.
StructMemberList private_struct_members;
utils::Vector<std::function<const ast::AssignmentStatement*()>, 4> private_initializers;
std::unordered_set<const ast::Function*> uses_privates;
// Build a list of functions that transitively reference any module-scope variables.
for (auto* decl : ctx.src->Sem().Module()->DependencyOrderedDeclarations()) {
if (auto* var = decl->As<ast::Var>()) {
auto* sem_var = ctx.src->Sem().Get(var);
if (sem_var->AddressSpace() == builtin::AddressSpace::kPrivate) {
// Create a member in the private variable struct.
auto* ty = sem_var->Type()->UnwrapRef();
auto name = ctx.Clone(var->name->symbol);
private_struct_members.Push(ctx.dst->Member(name, CreateASTTypeFor(ctx, ty)));
CloneStructTypes(ty);
// Create a statement to assign the initializer if present.
if (var->initializer) {
private_initializers.Push([&, name, var]() {
return ctx.dst->Assign(
ctx.dst->MemberAccessor(PrivateStructVariableName(), name),
ctx.Clone(var->initializer));
});
}
}
continue;
}
auto* func_ast = decl->As<ast::Function>();
if (!func_ast) {
continue;
@ -324,8 +345,10 @@ struct ModuleScopeVarToEntryPointParam::State {
bool needs_processing = false;
for (auto* var : func_sem->TransitivelyReferencedGlobals()) {
if (var->AddressSpace() != builtin::AddressSpace::kUndefined) {
if (var->AddressSpace() == builtin::AddressSpace::kPrivate) {
uses_privates.insert(func_ast);
}
needs_processing = true;
break;
}
}
if (needs_processing) {
@ -339,6 +362,14 @@ struct ModuleScopeVarToEntryPointParam::State {
}
}
if (!private_struct_members.IsEmpty()) {
// Create the private variable struct.
ctx.dst->Structure(PrivateStructName(), std::move(private_struct_members));
// Passing a pointer to a private variable will now involve passing a pointer to the
// member of a structure, so enable the extension that allows this.
ctx.dst->Enable(builtin::Extension::kChromiumExperimentalFullPtrParameters);
}
// 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.
@ -370,7 +401,7 @@ struct ModuleScopeVarToEntryPointParam::State {
// We aggregate all workgroup variables into a struct to avoid hitting MSL's limit for
// threadgroup memory arguments.
Symbol workgroup_parameter_symbol;
WorkgroupParameterMemberList workgroup_parameter_members;
StructMemberList workgroup_parameter_members;
auto workgroup_param = [&]() {
if (!workgroup_parameter_symbol.IsValid()) {
workgroup_parameter_symbol = ctx.dst->Sym();
@ -378,12 +409,43 @@ struct ModuleScopeVarToEntryPointParam::State {
return workgroup_parameter_symbol;
};
// If this function references any private variables, it needs to take the private
// variable struct as a parameter (or declare it, if it is an entry point function).
if (uses_privates.count(func_ast)) {
if (is_entry_point) {
// Create a local declaration for the private variable struct.
auto* var = ctx.dst->Var(
PrivateStructVariableName(), ctx.dst->ty(PrivateStructName()),
builtin::AddressSpace::kPrivate,
utils::Vector{
ctx.dst->Disable(ast::DisabledValidation::kIgnoreAddressSpace),
});
ctx.InsertFront(func_ast->body->statements, ctx.dst->Decl(var));
// Initialize the members of that struct with the original initializers.
for (auto init : private_initializers) {
ctx.InsertFront(func_ast->body->statements, init());
}
} else {
// Create a parameter that is a pointer to the private variable struct.
auto ptr = ctx.dst->ty.pointer(ctx.dst->ty(PrivateStructName()),
builtin::AddressSpace::kPrivate);
auto* param = ctx.dst->Param(PrivateStructVariableName(), ptr);
ctx.InsertBack(func_ast->params, param);
}
}
// Process and redeclare all variables referenced by the function.
for (auto* var : func_sem->TransitivelyReferencedGlobals()) {
if (var->AddressSpace() == builtin::AddressSpace::kUndefined) {
continue;
}
if (local_private_vars_.count(var)) {
if (var->AddressSpace() == builtin::AddressSpace::kPrivate) {
// Private variable are collected into a single struct that is passed by
// pointer (handled above), so we just need to replace the uses here.
ReplaceUsesInFunction(func_ast, var, PrivateStructVariableName(),
/* is_pointer */ !is_entry_point,
ctx.Clone(var->Declaration()->name->symbol));
continue;
}
@ -396,49 +458,25 @@ struct ModuleScopeVarToEntryPointParam::State {
// Track whether the new variable was wrapped in a struct or not.
bool is_wrapped = false;
// Check if this is a private variable that is only referenced by this function.
bool local_private = false;
if (var->AddressSpace() == builtin::AddressSpace::kPrivate) {
local_private = true;
for (auto* user : var->Users()) {
auto* stmt = user->Stmt();
if (!stmt || stmt->Function() != func_sem) {
local_private = false;
break;
}
}
}
if (local_private) {
// Redeclare the variable at function scope.
auto* disable_validation =
ctx.dst->Disable(ast::DisabledValidation::kIgnoreAddressSpace);
auto* initializer = ctx.Clone(var->Declaration()->initializer);
auto* local_var = ctx.dst->Var(new_var_symbol,
CreateASTTypeFor(ctx, var->Type()->UnwrapRef()),
builtin::AddressSpace::kPrivate, initializer,
utils::Vector{disable_validation});
ctx.InsertFront(func_ast->body->statements, ctx.dst->Decl(local_var));
local_private_vars_.insert(var);
// Process the variable to redeclare it as a parameter or local variable.
if (is_entry_point) {
ProcessVariableInEntryPoint(func_ast, var, new_var_symbol, workgroup_param,
workgroup_parameter_members, is_pointer,
is_wrapped);
} else {
// Process the variable to redeclare it as a parameter or local variable.
if (is_entry_point) {
ProcessVariableInEntryPoint(func_ast, var, new_var_symbol, workgroup_param,
workgroup_parameter_members, is_pointer,
is_wrapped);
} else {
ProcessVariableInUserFunction(func_ast, var, new_var_symbol, is_pointer);
if (var->AddressSpace() == builtin::AddressSpace::kWorkgroup) {
needs_pointer_aliasing = true;
}
ProcessVariableInUserFunction(func_ast, var, new_var_symbol, is_pointer);
if (var->AddressSpace() == builtin::AddressSpace::kWorkgroup) {
needs_pointer_aliasing = true;
}
// Record the replacement symbol.
var_to_newvar[var] = {new_var_symbol, is_pointer, is_wrapped};
}
// Record the replacement symbol.
var_to_newvar[var] = {new_var_symbol, is_pointer, is_wrapped};
// Replace all uses of the module-scope variable.
ReplaceUsesInFunction(func_ast, var, new_var_symbol, is_pointer, is_wrapped);
ReplaceUsesInFunction(
func_ast, var, new_var_symbol, is_pointer,
is_wrapped ? ctx.dst->Sym(kWrappedArrayMemberName) : Symbol());
}
// Allow pointer aliasing if needed.
@ -468,6 +506,15 @@ struct ModuleScopeVarToEntryPointParam::State {
auto* call_sem = ctx.src->Sem().Get(call)->Unwrap()->As<sem::Call>();
auto* target_sem = call_sem->Target()->As<sem::Function>();
// Pass the private variable struct pointer if needed.
if (uses_privates.count(target_sem->Declaration())) {
const ast::Expression* arg = ctx.dst->Expr(PrivateStructVariableName());
if (is_entry_point) {
arg = ctx.dst->AddressOf(arg);
}
ctx.InsertBack(call->args, arg);
}
// 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->TransitivelyReferencedGlobals()) {
@ -509,16 +556,35 @@ struct ModuleScopeVarToEntryPointParam::State {
}
}
/// @returns the name of the structure that contains all of the module-scope private variables
Symbol PrivateStructName() {
if (!private_struct_name.IsValid()) {
private_struct_name = ctx.dst->Symbols().New("tint_private_vars_struct");
}
return private_struct_name;
}
/// @returns the name of the variable that contains all of the module-scope private variables
Symbol PrivateStructVariableName() {
if (!private_struct_variable_name.IsValid()) {
private_struct_variable_name = ctx.dst->Symbols().New("tint_private_vars");
}
return private_struct_variable_name;
}
private:
// The structures that have already been cloned by this transform.
std::unordered_set<const sem::Struct*> cloned_structs_;
// Set of a private variables that are local to a single function.
std::unordered_set<const sem::Variable*> local_private_vars_;
// Map from identifier expression to the address-of expression that uses it.
std::unordered_map<const ast::IdentifierExpression*, const ast::UnaryOpExpression*>
ident_to_address_of_;
// The name of the structure that contains all the module-scope private variables.
Symbol private_struct_name;
// The name of the structure variable that contains all the module-scope private variables.
Symbol private_struct_variable_name;
};
ModuleScopeVarToEntryPointParam::ModuleScopeVarToEntryPointParam() = default;

245
src/tint/transform/module_scope_var_to_entry_point_param_test.cc
View File

@ -49,11 +49,17 @@ fn main() {
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
p : f32,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol : f32;
@internal(disable_validation__ignore_address_space) var<private> tint_symbol_1 : f32;
tint_symbol = tint_symbol_1;
tint_symbol = tint_private_vars.p;
}
)";
@ -74,11 +80,17 @@ var<private> p : f32;
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
p : f32,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol : f32;
@internal(disable_validation__ignore_address_space) var<private> tint_symbol_1 : f32;
tint_symbol = tint_symbol_1;
tint_symbol = tint_private_vars.p;
}
)";
@ -118,32 +130,38 @@ fn main() {
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
p : f32,
}
fn no_uses() {
}
fn zoo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<private, f32>) {
*(tint_symbol) = (*(tint_symbol) * 2.0);
fn zoo(tint_private_vars : ptr<private, tint_private_vars_struct>) {
(*(tint_private_vars)).p = ((*(tint_private_vars)).p * 2.0);
}
@internal(disable_validation__ignore_pointer_aliasing)
fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<private, f32>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<workgroup, f32>) {
*(tint_symbol_1) = a;
*(tint_symbol_2) = b;
zoo(tint_symbol_1);
fn bar(a : f32, b : f32, tint_private_vars : ptr<private, tint_private_vars_struct>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<workgroup, f32>) {
(*(tint_private_vars)).p = a;
*(tint_symbol) = b;
zoo(tint_private_vars);
}
@internal(disable_validation__ignore_pointer_aliasing)
fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<private, f32>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_4 : ptr<workgroup, f32>) {
fn foo(a : f32, tint_private_vars : ptr<private, tint_private_vars_struct>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<workgroup, f32>) {
let b : f32 = 2.0;
bar(a, b, tint_symbol_3, tint_symbol_4);
bar(a, b, tint_private_vars, tint_symbol_1);
no_uses();
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol_5 : f32;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol_6 : f32;
foo(1.0, &(tint_symbol_5), &(tint_symbol_6));
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol_2 : f32;
foo(1.0, &(tint_private_vars), &(tint_symbol_2));
}
)";
@ -183,17 +201,23 @@ var<workgroup> w : f32;
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
p : f32,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol_5 : f32;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol_6 : f32;
foo(1.0, &(tint_symbol_5), &(tint_symbol_6));
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol_2 : f32;
foo(1.0, &(tint_private_vars), &(tint_symbol_2));
}
@internal(disable_validation__ignore_pointer_aliasing)
fn foo(a : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_3 : ptr<private, f32>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_4 : ptr<workgroup, f32>) {
fn foo(a : f32, tint_private_vars : ptr<private, tint_private_vars_struct>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<workgroup, f32>) {
let b : f32 = 2.0;
bar(a, b, tint_symbol_3, tint_symbol_4);
bar(a, b, tint_private_vars, tint_symbol_1);
no_uses();
}
@ -201,14 +225,14 @@ fn no_uses() {
}
@internal(disable_validation__ignore_pointer_aliasing)
fn bar(a : f32, b : f32, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_1 : ptr<private, f32>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol_2 : ptr<workgroup, f32>) {
*(tint_symbol_1) = a;
*(tint_symbol_2) = b;
zoo(tint_symbol_1);
fn bar(a : f32, b : f32, tint_private_vars : ptr<private, tint_private_vars_struct>, @internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<workgroup, f32>) {
(*(tint_private_vars)).p = a;
*(tint_symbol) = b;
zoo(tint_private_vars);
}
fn zoo(@internal(disable_validation__ignore_address_space) @internal(disable_validation__ignore_invalid_pointer_argument) tint_symbol : ptr<private, f32>) {
*(tint_symbol) = (*(tint_symbol) * 2.0);
fn zoo(tint_private_vars : ptr<private, tint_private_vars_struct>) {
(*(tint_private_vars)).p = ((*(tint_private_vars)).p * 2.0);
}
)";
@ -229,11 +253,19 @@ fn main() {
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
a : f32,
b : f32,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol : f32 = 1.0;
@internal(disable_validation__ignore_address_space) var<private> tint_symbol_1 : f32 = f32();
let x : f32 = (tint_symbol + tint_symbol_1);
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
tint_private_vars.a = 1.0;
tint_private_vars.b = f32();
let x : f32 = (tint_private_vars.a + tint_private_vars.b);
}
)";
@ -254,11 +286,19 @@ var<private> a : f32 = 1.0;
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
a : f32,
b : f32,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol : f32 = 1.0;
@internal(disable_validation__ignore_address_space) var<private> tint_symbol_1 : f32 = f32();
let x : f32 = (tint_symbol + tint_symbol_1);
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
tint_private_vars.a = 1.0;
tint_private_vars.b = f32();
let x : f32 = (tint_private_vars.a + tint_private_vars.b);
}
)";
@ -282,12 +322,18 @@ fn main() {
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
p : f32,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol : f32;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol_1 : f32;
let p_ptr : ptr<private, f32> = &(tint_symbol);
let w_ptr : ptr<workgroup, f32> = &(tint_symbol_1);
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol : f32;
let p_ptr : ptr<private, f32> = &(tint_private_vars.p);
let w_ptr : ptr<workgroup, f32> = &(tint_symbol);
let x : f32 = (*(p_ptr) + *(w_ptr));
*(p_ptr) = x;
}
@ -313,12 +359,18 @@ var<private> p : f32;
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
p : f32,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol : f32;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol_1 : f32;
let p_ptr : ptr<private, f32> = &(tint_symbol);
let w_ptr : ptr<workgroup, f32> = &(tint_symbol_1);
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
@internal(disable_validation__ignore_address_space) var<workgroup> tint_symbol : f32;
let p_ptr : ptr<private, f32> = &(tint_private_vars.p);
let w_ptr : ptr<workgroup, f32> = &(tint_symbol);
let x : f32 = (*(p_ptr) + *(w_ptr));
*(p_ptr) = x;
}
@ -1151,6 +1203,7 @@ struct S {
var<private> p : f32;
var<workgroup> w : f32;
var<private> p_with_init : f32 = 42;
@group(0) @binding(0)
var<uniform> ub : S;
@ -1166,6 +1219,13 @@ fn main() {
)";
auto* expect = R"(
enable chromium_experimental_full_ptr_parameters;
struct tint_private_vars_struct {
p : f32,
p_with_init : f32,
}
struct S {
a : f32,
}
@ -1180,16 +1240,22 @@ fn main() {
EXPECT_EQ(expect, str(got));
}
// Test that a private variable that is only referenced by a single user-defined function is
// promoted to a function scope variable, rather than passed as a parameter.
TEST_F(ModuleScopeVarToEntryPointParamTest, PromotePrivateToFunctionScope) {
TEST_F(ModuleScopeVarToEntryPointParamTest, MultiplePrivateVariables) {
auto* src = R"(
var<private> p : f32;
struct S {
a : f32,
b : f32,
c : f32,
}
fn foo(a : f32) -> f32 {
let x = p;
p = x * a;
return p;
var<private> a : f32;
var<private> b : f32 = 42;
var<private> c : S = S(1, 2, 3);
var<private> d : S;
var<private> unused : f32;
fn foo(x : f32) -> f32 {
return (a + b + c.a + d.c) * x;
}
@compute @workgroup_size(1)
@ -1199,16 +1265,32 @@ fn main() {
)";
auto* expect = R"(
fn foo(a : f32) -> f32 {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol : f32;
let x = tint_symbol;
tint_symbol = (x * a);
return tint_symbol;
enable chromium_experimental_full_ptr_parameters;
struct S {
a : f32,
b : f32,
c : f32,
}
struct tint_private_vars_struct {
a : f32,
b : f32,
c : S,
d : S,
unused : f32,
}
fn foo(x : f32, tint_private_vars : ptr<private, tint_private_vars_struct>) -> f32 {
return (((((*(tint_private_vars)).a + (*(tint_private_vars)).b) + (*(tint_private_vars)).c.a) + (*(tint_private_vars)).d.c) * x);
}
@compute @workgroup_size(1)
fn main() {
_ = foo(1.0);
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
tint_private_vars.b = 42;
tint_private_vars.c = S(1, 2, 3);
_ = foo(1.0, &(tint_private_vars));
}
)";
@ -1217,36 +1299,59 @@ fn main() {
EXPECT_EQ(expect, str(got));
}
// Test that a private variable that is only referenced by a single user-defined function is
// promoted to a function scope variable, rather than passed as a parameter.
TEST_F(ModuleScopeVarToEntryPointParamTest, PromotePrivateToFunctionScope_OutOfOrder) {
TEST_F(ModuleScopeVarToEntryPointParamTest, MultiplePrivateVariables_OutOfOrder) {
auto* src = R"(
var<private> p : f32;
var<private> a : f32;
var<private> c : S = S(1, 2, 3);
var<private> unused : f32;
@compute @workgroup_size(1)
fn main() {
_ = foo(1.0);
}
fn foo(a : f32) -> f32 {
let x = p;
p = x * a;
return p;
fn foo(x : f32) -> f32 {
return (a + b + c.a + d.c) * x;
}
var<private> b : f32 = 42;
struct S {
a : f32,
b : f32,
c : f32,
}
var<private> d : S;
)";
auto* expect = R"(
@compute @workgroup_size(1)
fn main() {
_ = foo(1.0);
enable chromium_experimental_full_ptr_parameters;
struct S {
a : f32,
b : f32,
c : f32,
}
fn foo(a : f32) -> f32 {
@internal(disable_validation__ignore_address_space) var<private> tint_symbol : f32;
let x = tint_symbol;
tint_symbol = (x * a);
return tint_symbol;
struct tint_private_vars_struct {
a : f32,
c : S,
unused : f32,
b : f32,
d : S,
}
@compute @workgroup_size(1)
fn main() {
@internal(disable_validation__ignore_address_space) var<private> tint_private_vars : tint_private_vars_struct;
tint_private_vars.c = S(1, 2, 3);
tint_private_vars.b = 42;
_ = foo(1.0, &(tint_private_vars));
}
fn foo(x : f32, tint_private_vars : ptr<private, tint_private_vars_struct>) -> f32 {
return (((((*(tint_private_vars)).a + (*(tint_private_vars)).b) + (*(tint_private_vars)).c.a) + (*(tint_private_vars)).d.c) * x);
}
)";

8
src/tint/writer/msl/generator_impl_builtin_test.cc
View File

@ -1089,9 +1089,13 @@ template<typename T>
T tint_dot3(vec<T,3> a, vec<T,3> b) {
return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
}
struct tint_private_vars_struct {
int3 v;
};
kernel void test_function() {
thread int3 tint_symbol = 0;
int r = tint_dot3(tint_symbol, tint_symbol);
thread tint_private_vars_struct tint_private_vars = {};
int r = tint_dot3(tint_private_vars.v, tint_private_vars.v);
return;
}

5
src/tint/writer/msl/generator_impl_variable_decl_statement_test.cc
View File

@ -527,7 +527,10 @@ TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Private) {
gen.increment_indent();
ASSERT_TRUE(gen.Generate()) << gen.error();
EXPECT_THAT(gen.result(), HasSubstr("thread float tint_symbol_1 = 0.0f;\n"));
EXPECT_THAT(gen.result(), HasSubstr(R"(thread tint_private_vars_struct tint_private_vars = {};
float const tint_symbol = tint_private_vars.a;
return;
)"));
}
TEST_F(MslGeneratorImplTest, Emit_VariableDeclStatement_Workgroup) {

13
test/tint/array/assign_to_function_var.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
};
struct S {
/* 0x0000 */ tint_array<int4, 4> arr;
};
@ -28,8 +32,7 @@ S ret_struct_arr() {
return tint_symbol_2;
}
void foo(tint_array<int4, 4> src_param, threadgroup tint_array<int4, 4>* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7) {
thread tint_array<int4, 4> tint_symbol_4 = {};
void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) {
tint_array<int4, 4> src_function = {};
tint_array<int4, 4> dst = {};
tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
@ -39,12 +42,12 @@ void foo(tint_array<int4, 4> src_param, threadgroup tint_array<int4, 4>* const t
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
dst = src_let;
dst = src_function;
dst = tint_symbol_4;
dst = *(tint_symbol_5);
dst = (*(tint_private_vars)).src_private;
dst = *(tint_symbol_4);
S const tint_symbol = ret_struct_arr();
dst = tint_symbol.arr;
dst = (*(tint_symbol_5)).arr;
dst = (*(tint_symbol_6)).arr;
dst = (*(tint_symbol_7)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> dst_nested = {};
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
dst_nested = src_nested;

33
test/tint/array/assign_to_private_var.wgsl.expected.msl
View File

@ -14,6 +14,12 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
tint_array<int4, 4> dst;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> dst_nested;
};
struct S {
/* 0x0000 */ tint_array<int4, 4> arr;
};
@ -28,25 +34,22 @@ S ret_struct_arr() {
return tint_symbol_2;
}
void foo(tint_array<int4, 4> src_param, threadgroup tint_array<int4, 4>* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8) {
thread tint_array<int4, 4> tint_symbol_4 = {};
thread tint_array<int4, 4> tint_symbol_5 = {};
thread tint_array<tint_array<tint_array<int, 2>, 3>, 4> tint_symbol_9 = {};
void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) {
tint_array<int4, 4> src_function = {};
tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
tint_symbol_4 = tint_symbol_3;
tint_symbol_4 = src_param;
tint_symbol_4 = ret_arr();
(*(tint_private_vars)).dst = tint_symbol_3;
(*(tint_private_vars)).dst = src_param;
(*(tint_private_vars)).dst = ret_arr();
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
tint_symbol_4 = src_let;
tint_symbol_4 = src_function;
tint_symbol_4 = tint_symbol_5;
tint_symbol_4 = *(tint_symbol_6);
(*(tint_private_vars)).dst = src_let;
(*(tint_private_vars)).dst = src_function;
(*(tint_private_vars)).dst = (*(tint_private_vars)).src_private;
(*(tint_private_vars)).dst = *(tint_symbol_4);
S const tint_symbol = ret_struct_arr();
tint_symbol_4 = tint_symbol.arr;
tint_symbol_4 = (*(tint_symbol_7)).arr;
tint_symbol_4 = (*(tint_symbol_8)).arr;
(*(tint_private_vars)).dst = tint_symbol.arr;
(*(tint_private_vars)).dst = (*(tint_symbol_5)).arr;
(*(tint_private_vars)).dst = (*(tint_symbol_6)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
tint_symbol_9 = src_nested;
(*(tint_private_vars)).dst_nested = src_nested;
}

15
test/tint/array/assign_to_storage_var.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
};
struct S {
/* 0x0000 */ tint_array<int4, 4> arr;
};
@ -32,8 +36,7 @@ S ret_struct_arr() {
return tint_symbol_2;
}
void foo(tint_array<int4, 4> src_param, device S* const tint_symbol_4, threadgroup tint_array<int4, 4>* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8, device S_nested* const tint_symbol_9) {
thread tint_array<int4, 4> tint_symbol_5 = {};
void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, device S* const tint_symbol_4, threadgroup tint_array<int4, 4>* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, device S_nested* const tint_symbol_8) {
tint_array<int4, 4> src_function = {};
tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
(*(tint_symbol_4)).arr = tint_symbol_3;
@ -42,13 +45,13 @@ void foo(tint_array<int4, 4> src_param, device S* const tint_symbol_4, threadgro
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
(*(tint_symbol_4)).arr = src_let;
(*(tint_symbol_4)).arr = src_function;
(*(tint_symbol_4)).arr = tint_symbol_5;
(*(tint_symbol_4)).arr = *(tint_symbol_6);
(*(tint_symbol_4)).arr = (*(tint_private_vars)).src_private;
(*(tint_symbol_4)).arr = *(tint_symbol_5);
S const tint_symbol = ret_struct_arr();
(*(tint_symbol_4)).arr = tint_symbol.arr;
(*(tint_symbol_4)).arr = (*(tint_symbol_6)).arr;
(*(tint_symbol_4)).arr = (*(tint_symbol_7)).arr;
(*(tint_symbol_4)).arr = (*(tint_symbol_8)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
(*(tint_symbol_9)).arr = src_nested;
(*(tint_symbol_8)).arr = src_nested;
}

15
test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<int4, 4> src_private;
};
struct S {
/* 0x0000 */ tint_array<int4, 4> arr;
};
@ -28,8 +32,7 @@ S ret_struct_arr() {
return tint_symbol_2;
}
void foo(tint_array<int4, 4> src_param, threadgroup tint_array<int4, 4>* const tint_symbol_4, threadgroup tint_array<int4, 4>* const tint_symbol_6, const constant S* const tint_symbol_7, device S* const tint_symbol_8, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_9) {
thread tint_array<int4, 4> tint_symbol_5 = {};
void foo(tint_array<int4, 4> src_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_4, threadgroup tint_array<int4, 4>* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_8) {
tint_array<int4, 4> src_function = {};
tint_array<int4, 4> const tint_symbol_3 = tint_array<int4, 4>{int4(1), int4(2), int4(3), int4(3)};
*(tint_symbol_4) = tint_symbol_3;
@ -38,13 +41,13 @@ void foo(tint_array<int4, 4> src_param, threadgroup tint_array<int4, 4>* const t
tint_array<int4, 4> const src_let = tint_array<int4, 4>{};
*(tint_symbol_4) = src_let;
*(tint_symbol_4) = src_function;
*(tint_symbol_4) = tint_symbol_5;
*(tint_symbol_4) = *(tint_symbol_6);
*(tint_symbol_4) = (*(tint_private_vars)).src_private;
*(tint_symbol_4) = *(tint_symbol_5);
S const tint_symbol = ret_struct_arr();
*(tint_symbol_4) = tint_symbol.arr;
*(tint_symbol_4) = (*(tint_symbol_6)).arr;
*(tint_symbol_4) = (*(tint_symbol_7)).arr;
*(tint_symbol_4) = (*(tint_symbol_8)).arr;
tint_array<tint_array<tint_array<int, 2>, 3>, 4> src_nested = {};
*(tint_symbol_9) = src_nested;
*(tint_symbol_8) = src_nested;
}

25
test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float2x2, 4> const l_a = *(tint_symbol_3);
float2x2 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
float2 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<float2x2, 4> const l_a = *(tint_symbol_2);
float2x2 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
float2 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat2x2_f32/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float2x2, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].yx;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<float2x2, 4> p;
};
kernel void f(const constant tint_array<float2x2, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].yx;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -35,20 +39,21 @@ tint_array<half2x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<tin
return result;
}
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<half2x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_3));
half2x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_3))[p_a_i_save]);
half3 const l_a_i_i = half3((*(tint_symbol_3))[p_a_i_save][p_a_i_i_save].elements);
tint_array<half2x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_2));
half2x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_2))[p_a_i_save]);
half3 const l_a_i_i = half3((*(tint_symbol_2))[p_a_i_save][p_a_i_i_save].elements);
return;
}

16
test/tint/buffer/uniform/std140/array/mat2x3_f16/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<half2x3, 4> p;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -35,12 +39,12 @@ tint_array<half2x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<tin
return result;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<half2x3, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_1(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2]);
tint_symbol[1][0] = half3((*(tint_symbol_1))[0][1].elements).zxy;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1].elements[0];
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_1(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite((*(tint_symbol))[2]);
tint_private_vars.p[1][0] = half3((*(tint_symbol))[0][1].elements).zxy;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1].elements[0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -35,20 +39,21 @@ tint_array<float2x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<ti
return result;
}
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float2x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_3));
float2x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_3))[p_a_i_save]);
float3 const l_a_i_i = float3((*(tint_symbol_3))[p_a_i_save][p_a_i_i_save].elements);
tint_array<float2x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_2));
float2x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_2))[p_a_i_save]);
float3 const l_a_i_i = float3((*(tint_symbol_2))[p_a_i_save][p_a_i_i_save].elements);
return;
}

16
test/tint/buffer/uniform/std140/array/mat2x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<float2x3, 4> p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -35,12 +39,12 @@ tint_array<float2x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<ti
return result;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float2x3, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_1(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2]);
tint_symbol[1][0] = float3((*(tint_symbol_1))[0][1].elements).zxy;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1].elements[0];
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_1(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite((*(tint_symbol))[2]);
tint_private_vars.p[1][0] = float3((*(tint_symbol))[0][1].elements).zxy;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1].elements[0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<half2x4, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<half2x4, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<half2x4, 4> const l_a = *(tint_symbol_3);
half2x4 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
half4 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<half2x4, 4> const l_a = *(tint_symbol_2);
half2x4 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
half4 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat2x4_f16/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<half2x4, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<half2x4, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].ywxz;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<half2x4, 4> p;
};
kernel void f(const constant tint_array<half2x4, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].ywxz;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<float2x4, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<float2x4, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float2x4, 4> const l_a = *(tint_symbol_3);
float2x4 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
float4 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<float2x4, 4> const l_a = *(tint_symbol_2);
float2x4 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
float4 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat2x4_f32/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<float2x4, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float2x4, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].ywxz;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<float2x4, 4> p;
};
kernel void f(const constant tint_array<float2x4, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].ywxz;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -35,20 +39,21 @@ tint_array<float3x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<ti
return result;
}
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float3x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_3));
float3x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_3))[p_a_i_save]);
float3 const l_a_i_i = float3((*(tint_symbol_3))[p_a_i_save][p_a_i_i_save].elements);
tint_array<float3x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_2));
float3x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_2))[p_a_i_save]);
float3 const l_a_i_i = float3((*(tint_symbol_2))[p_a_i_save][p_a_i_i_save].elements);
return;
}

16
test/tint/buffer/uniform/std140/array/mat3x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<float3x3, 4> p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -35,12 +39,12 @@ tint_array<float3x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<ti
return result;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float3x3, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_1(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2]);
tint_symbol[1][0] = float3((*(tint_symbol_1))[0][1].elements).zxy;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1].elements[0];
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_1(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite((*(tint_symbol))[2]);
tint_private_vars.p[1][0] = float3((*(tint_symbol))[0][1].elements).zxy;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1].elements[0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<float3x4, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<float3x4, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float3x4, 4> const l_a = *(tint_symbol_3);
float3x4 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
float4 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<float3x4, 4> const l_a = *(tint_symbol_2);
float3x4 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
float4 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat3x4_f32/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<float3x4, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float3x4, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].ywxz;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<float3x4, 4> p;
};
kernel void f(const constant tint_array<float3x4, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].ywxz;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<half4x2, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<half4x2, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<half4x2, 4> const l_a = *(tint_symbol_3);
half4x2 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
half2 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<half4x2, 4> const l_a = *(tint_symbol_2);
half4x2 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
half2 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat4x2_f16/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<half4x2, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<half4x2, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].yx;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<half4x2, 4> p;
};
kernel void f(const constant tint_array<half4x2, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].yx;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float4x2, 4> const l_a = *(tint_symbol_3);
float4x2 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
float2 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<float4x2, 4> const l_a = *(tint_symbol_2);
float4x2 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
float2 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat4x2_f32/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float4x2, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].yx;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<float4x2, 4> p;
};
kernel void f(const constant tint_array<float4x2, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].yx;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -35,20 +39,21 @@ tint_array<half4x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<tin
return result;
}
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<half4x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_3));
half4x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_3))[p_a_i_save]);
half3 const l_a_i_i = half3((*(tint_symbol_3))[p_a_i_save][p_a_i_i_save].elements);
tint_array<half4x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_2));
half4x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_2))[p_a_i_save]);
half3 const l_a_i_i = half3((*(tint_symbol_2))[p_a_i_save][p_a_i_i_save].elements);
return;
}

16
test/tint/buffer/uniform/std140/array/mat4x3_f16/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<half4x3, 4> p;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -35,12 +39,12 @@ tint_array<half4x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<tin
return result;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<half4x3, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_1(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2]);
tint_symbol[1][0] = half3((*(tint_symbol_1))[0][1].elements).zxy;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1].elements[0];
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_1(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite((*(tint_symbol))[2]);
tint_private_vars.p[1][0] = half3((*(tint_symbol))[0][1].elements).zxy;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1].elements[0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -35,20 +39,21 @@ tint_array<float4x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<ti
return result;
}
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float4x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_3));
float4x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_3))[p_a_i_save]);
float3 const l_a_i_i = float3((*(tint_symbol_3))[p_a_i_save][p_a_i_i_save].elements);
tint_array<float4x3, 4> const l_a = tint_unpack_vec3_in_composite_1(*(tint_symbol_2));
float4x3 const l_a_i = tint_unpack_vec3_in_composite((*(tint_symbol_2))[p_a_i_save]);
float3 const l_a_i_i = float3((*(tint_symbol_2))[p_a_i_save][p_a_i_i_save].elements);
return;
}

16
test/tint/buffer/uniform/std140/array/mat4x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
tint_array<float4x3, 4> p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -35,12 +39,12 @@ tint_array<float4x3, 4> tint_unpack_vec3_in_composite_1(tint_array<tint_array<ti
return result;
}
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float4x3, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_1(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2]);
tint_symbol[1][0] = float3((*(tint_symbol_1))[0][1].elements).zxy;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1].elements[0];
kernel void f(const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_1(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite((*(tint_symbol))[2]);
tint_private_vars.p[1][0] = float3((*(tint_symbol))[0][1].elements).zxy;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1].elements[0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<half4x4, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<half4x4, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<half4x4, 4> const l_a = *(tint_symbol_3);
half4x4 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
half4 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<half4x4, 4> const l_a = *(tint_symbol_2);
half4x4 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
half4 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat4x4_f16/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<half4x4, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<half4x4, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].ywxz;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<half4x4, 4> p;
};
kernel void f(const constant tint_array<half4x4, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].ywxz;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

25
test/tint/buffer/uniform/std140/array/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,20 +14,25 @@ struct tint_array {
T elements[N];
};
int i() {
thread int tint_symbol_2 = 0;
tint_symbol_2 = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
return tint_symbol_2;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<float4x4, 4>* tint_symbol_3 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<float4x4, 4>* tint_symbol_2 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_i_save = tint_symbol_1;
tint_array<float4x4, 4> const l_a = *(tint_symbol_3);
float4x4 const l_a_i = (*(tint_symbol_3))[p_a_i_save];
float4 const l_a_i_i = (*(tint_symbol_3))[p_a_i_save][p_a_i_i_save];
tint_array<float4x4, 4> const l_a = *(tint_symbol_2);
float4x4 const l_a_i = (*(tint_symbol_2))[p_a_i_save];
float4 const l_a_i_i = (*(tint_symbol_2))[p_a_i_save][p_a_i_i_save];
return;
}

16
test/tint/buffer/uniform/std140/array/mat4x4_f32/to_private.wgsl.expected.msl
View File

@ -14,12 +14,16 @@ struct tint_array {
T elements[N];
};
kernel void f(const constant tint_array<float4x4, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<float4x4, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[1][0] = (*(tint_symbol_1))[0][1].ywxz;
tint_symbol[1][0][0] = (*(tint_symbol_1))[0][1][0];
struct tint_private_vars_struct {
tint_array<float4x4, 4> p;
};
kernel void f(const constant tint_array<float4x4, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[1][0] = (*(tint_symbol))[0][1].ywxz;
tint_private_vars.p[1][0][0] = (*(tint_symbol))[0][1][0];
return;
}

37
test/tint/buffer/uniform/std140/struct/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ half2x2 m;
/* 0x0008 */ tint_array<int8_t, 56> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
half2x2 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
half2 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
half2x2 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
half2 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_private.wgsl.expected.msl
View File

@ -22,12 +22,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].yx;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].yx;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ float2x2 m;
/* 0x0010 */ tint_array<int8_t, 48> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
float2x2 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
float2 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
float2x2 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
float2 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].yx;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].yx;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -72,27 +76,28 @@ tint_array<Outer, 4> tint_unpack_vec3_in_composite_4(tint_array<Outer_tint_packe
return result;
}
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_5));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_5))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_5))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save]);
half2x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m);
half3 const l_a_i_a_i_m_i = half3((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_4));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_4))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_4))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save]);
half2x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m);
half3 const l_a_i_a_i_m_i = half3((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
return;
}

28
test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_private.wgsl.expected.msl
View File

@ -14,6 +14,16 @@ struct tint_array {
T elements[N];
};
struct S {
int before;
half2x3 m;
int after;
};
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -36,12 +46,6 @@ half2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
struct S {
int before;
half2x3 m;
int after;
};
S tint_unpack_vec3_in_composite_1(S_tint_packed_vec3 in) {
S result = {};
result.before = in.before;
@ -58,12 +62,12 @@ tint_array<S, 4> tint_unpack_vec3_in_composite_2(tint_array<S_tint_packed_vec3,
return result;
}
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_2(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol_1))[2]);
tint_symbol[3].m = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2].m);
tint_symbol[1].m[0] = half3((*(tint_symbol_1))[0].m[1].elements).zxy;
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_2(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol))[2]);
tint_private_vars.p[3].m = tint_unpack_vec3_in_composite((*(tint_symbol))[2].m);
tint_private_vars.p[1].m[0] = half3((*(tint_symbol))[0].m[1].elements).zxy;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -72,27 +76,28 @@ tint_array<Outer, 4> tint_unpack_vec3_in_composite_4(tint_array<Outer_tint_packe
return result;
}
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_5));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_5))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_5))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save]);
float2x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m);
float3 const l_a_i_a_i_m_i = float3((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_4));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_4))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_4))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save]);
float2x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m);
float3 const l_a_i_a_i_m_i = float3((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
return;
}

28
test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,16 @@ struct tint_array {
T elements[N];
};
struct S {
int before;
float2x3 m;
int after;
};
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -36,12 +46,6 @@ float2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
struct S {
int before;
float2x3 m;
int after;
};
S tint_unpack_vec3_in_composite_1(S_tint_packed_vec3 in) {
S result = {};
result.before = in.before;
@ -58,12 +62,12 @@ tint_array<S, 4> tint_unpack_vec3_in_composite_2(tint_array<S_tint_packed_vec3,
return result;
}
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_2(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol_1))[2]);
tint_symbol[3].m = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2].m);
tint_symbol[1].m[0] = float3((*(tint_symbol_1))[0].m[1].elements).zxy;
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_2(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol))[2]);
tint_private_vars.p[3].m = tint_unpack_vec3_in_composite((*(tint_symbol))[2].m);
tint_private_vars.p[1].m[0] = float3((*(tint_symbol))[0].m[1].elements).zxy;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ half2x4 m;
/* 0x0010 */ tint_array<int8_t, 48> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
half2x4 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
half4 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
half2x4 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
half4 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].ywxz;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].ywxz;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ float2x4 m;
/* 0x0020 */ tint_array<int8_t, 32> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
float2x4 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
float4 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
float2x4 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
float4 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].ywxz;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].ywxz;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ half3x2 m;
/* 0x000c */ tint_array<int8_t, 52> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
half3x2 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
half2 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
half3x2 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
half2 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_private.wgsl.expected.msl
View File

@ -22,12 +22,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].yx;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].yx;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ float3x2 m;
/* 0x0018 */ tint_array<int8_t, 40> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
float3x2 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
float2 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
float3x2 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
float2 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].yx;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].yx;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -72,27 +76,28 @@ tint_array<Outer, 4> tint_unpack_vec3_in_composite_4(tint_array<Outer_tint_packe
return result;
}
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_5));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_5))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_5))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save]);
half3x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m);
half3 const l_a_i_a_i_m_i = half3((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_4));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_4))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_4))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save]);
half3x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m);
half3 const l_a_i_a_i_m_i = half3((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
return;
}

28
test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_private.wgsl.expected.msl
View File

@ -14,6 +14,16 @@ struct tint_array {
T elements[N];
};
struct S {
int before;
half3x3 m;
int after;
};
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -36,12 +46,6 @@ half3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
struct S {
int before;
half3x3 m;
int after;
};
S tint_unpack_vec3_in_composite_1(S_tint_packed_vec3 in) {
S result = {};
result.before = in.before;
@ -58,12 +62,12 @@ tint_array<S, 4> tint_unpack_vec3_in_composite_2(tint_array<S_tint_packed_vec3,
return result;
}
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_2(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol_1))[2]);
tint_symbol[3].m = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2].m);
tint_symbol[1].m[0] = half3((*(tint_symbol_1))[0].m[1].elements).zxy;
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_2(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol))[2]);
tint_private_vars.p[3].m = tint_unpack_vec3_in_composite((*(tint_symbol))[2].m);
tint_private_vars.p[1].m[0] = half3((*(tint_symbol))[0].m[1].elements).zxy;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -72,27 +76,28 @@ tint_array<Outer, 4> tint_unpack_vec3_in_composite_4(tint_array<Outer_tint_packe
return result;
}
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_5));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_5))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_5))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save]);
float3x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m);
float3 const l_a_i_a_i_m_i = float3((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_4));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_4))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_4))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save]);
float3x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m);
float3 const l_a_i_a_i_m_i = float3((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
return;
}

28
test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,16 @@ struct tint_array {
T elements[N];
};
struct S {
int before;
float3x3 m;
int after;
};
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -35,12 +45,6 @@ float3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
struct S {
int before;
float3x3 m;
int after;
};
S tint_unpack_vec3_in_composite_1(S_tint_packed_vec3 in) {
S result = {};
result.before = in.before;
@ -57,12 +61,12 @@ tint_array<S, 4> tint_unpack_vec3_in_composite_2(tint_array<S_tint_packed_vec3,
return result;
}
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_2(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol_1))[2]);
tint_symbol[3].m = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2].m);
tint_symbol[1].m[0] = float3((*(tint_symbol_1))[0].m[1].elements).zxy;
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_2(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol))[2]);
tint_private_vars.p[3].m = tint_unpack_vec3_in_composite((*(tint_symbol))[2].m);
tint_private_vars.p[1].m[0] = float3((*(tint_symbol))[0].m[1].elements).zxy;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ half3x4 m;
/* 0x0018 */ tint_array<int8_t, 40> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
half3x4 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
half4 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
half3x4 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
half4 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].ywxz;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].ywxz;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ float3x4 m;
/* 0x0030 */ tint_array<int8_t, 16> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
float3x4 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
float4 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
float3x4 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
float4 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_private.wgsl.expected.msl
View File

@ -22,12 +22,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].ywxz;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].ywxz;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ half4x2 m;
/* 0x0010 */ tint_array<int8_t, 48> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
half4x2 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
half2 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
half4x2 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
half2 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_private.wgsl.expected.msl
View File

@ -22,12 +22,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_1;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].yx;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].yx;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ float4x2 m;
/* 0x0020 */ tint_array<int8_t, 32> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
float4x2 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
float2 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
float4x2 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
float2 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].yx;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].yx;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -72,27 +76,28 @@ tint_array<Outer, 4> tint_unpack_vec3_in_composite_4(tint_array<Outer_tint_packe
return result;
}
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_5));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_5))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_5))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save]);
half4x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m);
half3 const l_a_i_a_i_m_i = half3((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_4));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_4))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_4))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save]);
half4x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m);
half3 const l_a_i_a_i_m_i = half3((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
return;
}

28
test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_private.wgsl.expected.msl
View File

@ -14,6 +14,16 @@ struct tint_array {
T elements[N];
};
struct S {
int before;
half4x3 m;
int after;
};
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -36,12 +46,6 @@ half4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
struct S {
int before;
half4x3 m;
int after;
};
S tint_unpack_vec3_in_composite_1(S_tint_packed_vec3 in) {
S result = {};
result.before = in.before;
@ -58,12 +62,12 @@ tint_array<S, 4> tint_unpack_vec3_in_composite_2(tint_array<S_tint_packed_vec3,
return result;
}
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_2(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol_1))[2]);
tint_symbol[3].m = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2].m);
tint_symbol[1].m[0] = half3((*(tint_symbol_1))[0].m[1].elements).zxy;
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_2(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol))[2]);
tint_private_vars.p[3].m = tint_unpack_vec3_in_composite((*(tint_symbol))[2].m);
tint_private_vars.p[1].m[0] = half3((*(tint_symbol))[0].m[1].elements).zxy;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -71,27 +75,28 @@ tint_array<Outer, 4> tint_unpack_vec3_in_composite_4(tint_array<Outer_tint_packe
return result;
}
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer_tint_packed_vec3, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_5));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_5))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_5))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save]);
float4x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m);
float3 const l_a_i_a_i_m_i = float3((*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
tint_array<Outer, 4> const l_a = tint_unpack_vec3_in_composite_4(*(tint_symbol_4));
Outer const l_a_i = tint_unpack_vec3_in_composite_3((*(tint_symbol_4))[p_a_i_save]);
tint_array<Inner, 4> const l_a_i_a = tint_unpack_vec3_in_composite_2((*(tint_symbol_4))[p_a_i_save].a);
Inner const l_a_i_a_i = tint_unpack_vec3_in_composite_1((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save]);
float4x3 const l_a_i_a_i_m = tint_unpack_vec3_in_composite((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m);
float3 const l_a_i_a_i_m_i = float3((*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements);
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save].elements[tint_symbol_3];
return;
}

28
test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,16 @@ struct tint_array {
T elements[N];
};
struct S {
int before;
float4x3 m;
int after;
};
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -36,12 +46,6 @@ float4x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
struct S {
int before;
float4x3 m;
int after;
};
S tint_unpack_vec3_in_composite_1(S_tint_packed_vec3 in) {
S result = {};
result.before = in.before;
@ -58,12 +62,12 @@ tint_array<S, 4> tint_unpack_vec3_in_composite_2(tint_array<S_tint_packed_vec3,
return result;
}
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = tint_unpack_vec3_in_composite_2(*(tint_symbol_1));
tint_symbol[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol_1))[2]);
tint_symbol[3].m = tint_unpack_vec3_in_composite((*(tint_symbol_1))[2].m);
tint_symbol[1].m[0] = float3((*(tint_symbol_1))[0].m[1].elements).zxy;
kernel void f(const constant tint_array<S_tint_packed_vec3, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite_2(*(tint_symbol));
tint_private_vars.p[1] = tint_unpack_vec3_in_composite_1((*(tint_symbol))[2]);
tint_private_vars.p[3].m = tint_unpack_vec3_in_composite((*(tint_symbol))[2].m);
tint_private_vars.p[1].m[0] = float3((*(tint_symbol))[0].m[1].elements).zxy;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ half4x4 m;
/* 0x0020 */ tint_array<int8_t, 32> tint_pad;
@ -23,27 +27,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
half4x4 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
half4 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
half const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
half4x4 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
half4 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
half const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0044 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].ywxz;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].ywxz;
return;
}

37
test/tint/buffer/uniform/std140/struct/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct Inner {
/* 0x0000 */ float4x4 m;
};
@ -22,27 +26,28 @@ struct Outer {
/* 0x0000 */ tint_array<Inner, 4> a;
};
int i() {
thread int tint_symbol_4 = 0;
tint_symbol_4 = as_type<int>((as_type<uint>(tint_symbol_4) + as_type<uint>(1)));
return tint_symbol_4;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_5 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<Outer, 4>* tint_symbol_4 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_a_i_save = tint_symbol;
int const tint_symbol_1 = i();
int const tint_symbol_1 = i(&(tint_private_vars));
int const p_a_i_a_i_save = tint_symbol_1;
int const tint_symbol_2 = i();
int const tint_symbol_2 = i(&(tint_private_vars));
int const p_a_i_a_i_m_i_save = tint_symbol_2;
tint_array<Outer, 4> const l_a = *(tint_symbol_5);
Outer const l_a_i = (*(tint_symbol_5))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_5))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save];
float4x4 const l_a_i_a_i_m = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m;
float4 const l_a_i_a_i_m_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i();
float const l_a_i_a_i_m_i_i = (*(tint_symbol_5))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
tint_array<Outer, 4> const l_a = *(tint_symbol_4);
Outer const l_a_i = (*(tint_symbol_4))[p_a_i_save];
tint_array<Inner, 4> const l_a_i_a = (*(tint_symbol_4))[p_a_i_save].a;
Inner const l_a_i_a_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save];
float4x4 const l_a_i_a_i_m = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m;
float4 const l_a_i_a_i_m_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save];
int const tint_symbol_3 = i(&(tint_private_vars));
float const l_a_i_a_i_m_i_i = (*(tint_symbol_4))[p_a_i_save].a[p_a_i_a_i_save].m[p_a_i_a_i_m_i_save][tint_symbol_3];
return;
}

16
test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_private.wgsl.expected.msl
View File

@ -23,12 +23,16 @@ struct S {
/* 0x0084 */ tint_array<int8_t, 60> tint_pad_2;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol_1 [[buffer(0)]]) {
thread tint_array<S, 4> tint_symbol = {};
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[2];
tint_symbol[3].m = (*(tint_symbol_1))[2].m;
tint_symbol[1].m[0] = (*(tint_symbol_1))[0].m[1].ywxz;
struct tint_private_vars_struct {
tint_array<S, 4> p;
};
kernel void f(const constant tint_array<S, 4>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[2];
tint_private_vars.p[3].m = (*(tint_symbol))[2].m;
tint_private_vars.p[1].m[0] = (*(tint_symbol))[0].m[1].ywxz;
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -1,17 +1,22 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant half2x2* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant half2x2* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
half2x2 const l_m = *(tint_symbol_2);
half2 const l_m_i = (*(tint_symbol_2))[p_m_i_save];
half2x2 const l_m = *(tint_symbol_1);
half2 const l_m_i = (*(tint_symbol_1))[p_m_i_save];
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat2x2_f16/static_index_via_ptr.wgsl.expected.msl
View File

@ -1,15 +1,18 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant half2x2* tint_symbol_1 [[buffer(0)]]) {
half2x2 const l_m = *(tint_symbol_1);
half2 const l_m_1 = (*(tint_symbol_1))[1];
kernel void f(const constant half2x2* tint_symbol [[buffer(0)]]) {
half2x2 const l_m = *(tint_symbol);
half2 const l_m_1 = (*(tint_symbol))[1];
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_private.wgsl.expected.msl
View File

@ -1,12 +1,16 @@
#include <metal_stdlib>
using namespace metal;
kernel void f(const constant half2x2* tint_symbol_1 [[buffer(0)]]) {
thread half2x2 tint_symbol = half2x2(0.0h);
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[0];
tint_symbol[1] = (*(tint_symbol_1))[0].yx;
tint_symbol[0][1] = (*(tint_symbol_1))[1][0];
struct tint_private_vars_struct {
half2x2 p;
};
kernel void f(const constant half2x2* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[0];
tint_private_vars.p[1] = (*(tint_symbol))[0].yx;
tint_private_vars.p[0][1] = (*(tint_symbol))[1][0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -1,17 +1,22 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant float2x2* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant float2x2* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
float2x2 const l_m = *(tint_symbol_2);
float2 const l_m_i = (*(tint_symbol_2))[p_m_i_save];
float2x2 const l_m = *(tint_symbol_1);
float2 const l_m_i = (*(tint_symbol_1))[p_m_i_save];
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat2x2_f32/static_index_via_ptr.wgsl.expected.msl
View File

@ -1,15 +1,18 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant float2x2* tint_symbol_1 [[buffer(0)]]) {
float2x2 const l_m = *(tint_symbol_1);
float2 const l_m_1 = (*(tint_symbol_1))[1];
kernel void f(const constant float2x2* tint_symbol [[buffer(0)]]) {
float2x2 const l_m = *(tint_symbol);
float2 const l_m_1 = (*(tint_symbol))[1];
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_private.wgsl.expected.msl
View File

@ -1,12 +1,16 @@
#include <metal_stdlib>
using namespace metal;
kernel void f(const constant float2x2* tint_symbol_1 [[buffer(0)]]) {
thread float2x2 tint_symbol = float2x2(0.0f);
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[0];
tint_symbol[1] = (*(tint_symbol_1))[0].yx;
tint_symbol[0][1] = (*(tint_symbol_1))[1][0];
struct tint_private_vars_struct {
float2x2 p;
};
kernel void f(const constant float2x2* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[0];
tint_private_vars.p[1] = (*(tint_symbol))[0].yx;
tint_private_vars.p[0][1] = (*(tint_symbol))[1][0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -27,17 +31,18 @@ half2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
half2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_2));
half3 const l_m_i = half3((*(tint_symbol_2))[p_m_i_save].elements);
half2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
half3 const l_m_i = half3((*(tint_symbol_1))[p_m_i_save].elements);
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat2x3_f16/static_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -27,15 +31,14 @@ half2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol_1 [[buffer(0)]]) {
half2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
half3 const l_m_1 = half3((*(tint_symbol_1))[1].elements);
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol [[buffer(0)]]) {
half2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol));
half3 const l_m_1 = half3((*(tint_symbol))[1].elements);
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
half2x3 p;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -27,12 +31,12 @@ half2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol_1 [[buffer(0)]]) {
thread half2x3 tint_symbol = half2x3(0.0h);
tint_symbol = tint_unpack_vec3_in_composite(*(tint_symbol_1));
tint_symbol[1] = half3((*(tint_symbol_1))[0].elements);
tint_symbol[1] = half3((*(tint_symbol_1))[0].elements).zxy;
tint_symbol[0][1] = (*(tint_symbol_1))[1].elements[0];
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 2>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite(*(tint_symbol));
tint_private_vars.p[1] = half3((*(tint_symbol))[0].elements);
tint_private_vars.p[1] = half3((*(tint_symbol))[0].elements).zxy;
tint_private_vars.p[0][1] = (*(tint_symbol))[1].elements[0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -27,17 +31,18 @@ float2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
float2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_2));
float3 const l_m_i = float3((*(tint_symbol_2))[p_m_i_save].elements);
float2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
float3 const l_m_i = float3((*(tint_symbol_1))[p_m_i_save].elements);
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat2x3_f32/static_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -27,15 +31,14 @@ float2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol_1 [[buffer(0)]]) {
float2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
float3 const l_m_1 = float3((*(tint_symbol_1))[1].elements);
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol [[buffer(0)]]) {
float2x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol));
float3 const l_m_1 = float3((*(tint_symbol))[1].elements);
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
float2x3 p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -27,12 +31,12 @@ float2x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol_1 [[buffer(0)]]) {
thread float2x3 tint_symbol = float2x3(0.0f);
tint_symbol = tint_unpack_vec3_in_composite(*(tint_symbol_1));
tint_symbol[1] = float3((*(tint_symbol_1))[0].elements);
tint_symbol[1] = float3((*(tint_symbol_1))[0].elements).zxy;
tint_symbol[0][1] = (*(tint_symbol_1))[1].elements[0];
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 2>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite(*(tint_symbol));
tint_private_vars.p[1] = float3((*(tint_symbol))[0].elements);
tint_private_vars.p[1] = float3((*(tint_symbol))[0].elements).zxy;
tint_private_vars.p[0][1] = (*(tint_symbol))[1].elements[0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -1,17 +1,22 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant half2x4* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant half2x4* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
half2x4 const l_m = *(tint_symbol_2);
half4 const l_m_i = (*(tint_symbol_2))[p_m_i_save];
half2x4 const l_m = *(tint_symbol_1);
half4 const l_m_i = (*(tint_symbol_1))[p_m_i_save];
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat2x4_f16/static_index_via_ptr.wgsl.expected.msl
View File

@ -1,15 +1,18 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant half2x4* tint_symbol_1 [[buffer(0)]]) {
half2x4 const l_m = *(tint_symbol_1);
half4 const l_m_1 = (*(tint_symbol_1))[1];
kernel void f(const constant half2x4* tint_symbol [[buffer(0)]]) {
half2x4 const l_m = *(tint_symbol);
half4 const l_m_1 = (*(tint_symbol))[1];
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_private.wgsl.expected.msl
View File

@ -1,12 +1,16 @@
#include <metal_stdlib>
using namespace metal;
kernel void f(const constant half2x4* tint_symbol_1 [[buffer(0)]]) {
thread half2x4 tint_symbol = half2x4(0.0h);
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[0];
tint_symbol[1] = (*(tint_symbol_1))[0].ywxz;
tint_symbol[0][1] = (*(tint_symbol_1))[1][0];
struct tint_private_vars_struct {
half2x4 p;
};
kernel void f(const constant half2x4* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[0];
tint_private_vars.p[1] = (*(tint_symbol))[0].ywxz;
tint_private_vars.p[0][1] = (*(tint_symbol))[1][0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -1,17 +1,22 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant float2x4* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant float2x4* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
float2x4 const l_m = *(tint_symbol_2);
float4 const l_m_i = (*(tint_symbol_2))[p_m_i_save];
float2x4 const l_m = *(tint_symbol_1);
float4 const l_m_i = (*(tint_symbol_1))[p_m_i_save];
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat2x4_f32/static_index_via_ptr.wgsl.expected.msl
View File

@ -1,15 +1,18 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant float2x4* tint_symbol_1 [[buffer(0)]]) {
float2x4 const l_m = *(tint_symbol_1);
float4 const l_m_1 = (*(tint_symbol_1))[1];
kernel void f(const constant float2x4* tint_symbol [[buffer(0)]]) {
float2x4 const l_m = *(tint_symbol);
float4 const l_m_1 = (*(tint_symbol))[1];
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_private.wgsl.expected.msl
View File

@ -1,12 +1,16 @@
#include <metal_stdlib>
using namespace metal;
kernel void f(const constant float2x4* tint_symbol_1 [[buffer(0)]]) {
thread float2x4 tint_symbol = float2x4(0.0f);
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[0];
tint_symbol[1] = (*(tint_symbol_1))[0].ywxz;
tint_symbol[0][1] = (*(tint_symbol_1))[1][0];
struct tint_private_vars_struct {
float2x4 p;
};
kernel void f(const constant float2x4* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[0];
tint_private_vars.p[1] = (*(tint_symbol))[0].ywxz;
tint_private_vars.p[0][1] = (*(tint_symbol))[1][0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -1,17 +1,22 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant half3x2* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant half3x2* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
half3x2 const l_m = *(tint_symbol_2);
half2 const l_m_i = (*(tint_symbol_2))[p_m_i_save];
half3x2 const l_m = *(tint_symbol_1);
half2 const l_m_i = (*(tint_symbol_1))[p_m_i_save];
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat3x2_f16/static_index_via_ptr.wgsl.expected.msl
View File

@ -1,15 +1,18 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant half3x2* tint_symbol_1 [[buffer(0)]]) {
half3x2 const l_m = *(tint_symbol_1);
half2 const l_m_1 = (*(tint_symbol_1))[1];
kernel void f(const constant half3x2* tint_symbol [[buffer(0)]]) {
half3x2 const l_m = *(tint_symbol);
half2 const l_m_1 = (*(tint_symbol))[1];
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_private.wgsl.expected.msl
View File

@ -1,12 +1,16 @@
#include <metal_stdlib>
using namespace metal;
kernel void f(const constant half3x2* tint_symbol_1 [[buffer(0)]]) {
thread half3x2 tint_symbol = half3x2(0.0h);
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[0];
tint_symbol[1] = (*(tint_symbol_1))[0].yx;
tint_symbol[0][1] = (*(tint_symbol_1))[1][0];
struct tint_private_vars_struct {
half3x2 p;
};
kernel void f(const constant half3x2* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[0];
tint_private_vars.p[1] = (*(tint_symbol))[0].yx;
tint_private_vars.p[0][1] = (*(tint_symbol))[1][0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -1,17 +1,22 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant float3x2* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant float3x2* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
float3x2 const l_m = *(tint_symbol_2);
float2 const l_m_i = (*(tint_symbol_2))[p_m_i_save];
float3x2 const l_m = *(tint_symbol_1);
float2 const l_m_i = (*(tint_symbol_1))[p_m_i_save];
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat3x2_f32/static_index_via_ptr.wgsl.expected.msl
View File

@ -1,15 +1,18 @@
#include <metal_stdlib>
using namespace metal;
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
struct tint_private_vars_struct {
int counter;
};
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant float3x2* tint_symbol_1 [[buffer(0)]]) {
float3x2 const l_m = *(tint_symbol_1);
float2 const l_m_1 = (*(tint_symbol_1))[1];
kernel void f(const constant float3x2* tint_symbol [[buffer(0)]]) {
float3x2 const l_m = *(tint_symbol);
float2 const l_m_1 = (*(tint_symbol))[1];
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_private.wgsl.expected.msl
View File

@ -1,12 +1,16 @@
#include <metal_stdlib>
using namespace metal;
kernel void f(const constant float3x2* tint_symbol_1 [[buffer(0)]]) {
thread float3x2 tint_symbol = float3x2(0.0f);
tint_symbol = *(tint_symbol_1);
tint_symbol[1] = (*(tint_symbol_1))[0];
tint_symbol[1] = (*(tint_symbol_1))[0].yx;
tint_symbol[0][1] = (*(tint_symbol_1))[1][0];
struct tint_private_vars_struct {
float3x2 p;
};
kernel void f(const constant float3x2* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = *(tint_symbol);
tint_private_vars.p[1] = (*(tint_symbol))[0];
tint_private_vars.p[1] = (*(tint_symbol))[0].yx;
tint_private_vars.p[0][1] = (*(tint_symbol))[1][0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -27,17 +31,18 @@ half3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
half3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_2));
half3 const l_m_i = half3((*(tint_symbol_2))[p_m_i_save].elements);
half3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
half3 const l_m_i = half3((*(tint_symbol_1))[p_m_i_save].elements);
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat3x3_f16/static_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -27,15 +31,14 @@ half3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol_1 [[buffer(0)]]) {
half3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
half3 const l_m_1 = half3((*(tint_symbol_1))[1].elements);
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol [[buffer(0)]]) {
half3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol));
half3 const l_m_1 = half3((*(tint_symbol))[1].elements);
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
half3x3 p;
};
struct tint_packed_vec3_f16_array_element {
/* 0x0000 */ packed_half3 elements;
/* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@ -27,12 +31,12 @@ half3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f16_array_elem
return result;
}
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol_1 [[buffer(0)]]) {
thread half3x3 tint_symbol = half3x3(0.0h);
tint_symbol = tint_unpack_vec3_in_composite(*(tint_symbol_1));
tint_symbol[1] = half3((*(tint_symbol_1))[0].elements);
tint_symbol[1] = half3((*(tint_symbol_1))[0].elements).zxy;
tint_symbol[0][1] = (*(tint_symbol_1))[1].elements[0];
kernel void f(const constant tint_array<tint_packed_vec3_f16_array_element, 3>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite(*(tint_symbol));
tint_private_vars.p[1] = half3((*(tint_symbol))[0].elements);
tint_private_vars.p[1] = half3((*(tint_symbol))[0].elements).zxy;
tint_private_vars.p[0][1] = (*(tint_symbol))[1].elements[0];
return;
}

21
test/tint/buffer/uniform/std140/unnested/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -27,17 +31,18 @@ float3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
int i() {
thread int tint_symbol_1 = 0;
tint_symbol_1 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
return tint_symbol_1;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_2 [[buffer(0)]]) {
int const tint_symbol = i();
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_1 [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.counter = 0;
int const tint_symbol = i(&(tint_private_vars));
int const p_m_i_save = tint_symbol;
float3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_2));
float3 const l_m_i = float3((*(tint_symbol_2))[p_m_i_save].elements);
float3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
float3 const l_m_i = float3((*(tint_symbol_1))[p_m_i_save].elements);
return;
}

17
test/tint/buffer/uniform/std140/unnested/mat3x3_f32/static_index_via_ptr.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
int counter;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -27,15 +31,14 @@ float3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
int i() {
thread int tint_symbol = 0;
tint_symbol = as_type<int>((as_type<uint>(tint_symbol) + as_type<uint>(1)));
return tint_symbol;
int i(thread tint_private_vars_struct* const tint_private_vars) {
(*(tint_private_vars)).counter = as_type<int>((as_type<uint>((*(tint_private_vars)).counter) + as_type<uint>(1)));
return (*(tint_private_vars)).counter;
}
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_1 [[buffer(0)]]) {
float3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol_1));
float3 const l_m_1 = float3((*(tint_symbol_1))[1].elements);
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol [[buffer(0)]]) {
float3x3 const l_m = tint_unpack_vec3_in_composite(*(tint_symbol));
float3 const l_m_1 = float3((*(tint_symbol))[1].elements);
return;
}

16
test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_private.wgsl.expected.msl
View File

@ -14,6 +14,10 @@ struct tint_array {
T elements[N];
};
struct tint_private_vars_struct {
float3x3 p;
};
struct tint_packed_vec3_f32_array_element {
/* 0x0000 */ packed_float3 elements;
/* 0x000c */ tint_array<int8_t, 4> tint_pad;
@ -27,12 +31,12 @@ float3x3 tint_unpack_vec3_in_composite(tint_array<tint_packed_vec3_f32_array_ele
return result;
}
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_1 [[buffer(0)]]) {
thread float3x3 tint_symbol = float3x3(0.0f);
tint_symbol = tint_unpack_vec3_in_composite(*(tint_symbol_1));
tint_symbol[1] = float3((*(tint_symbol_1))[0].elements);
tint_symbol[1] = float3((*(tint_symbol_1))[0].elements).zxy;
tint_symbol[0][1] = (*(tint_symbol_1))[1].elements[0];
kernel void f(const constant tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol [[buffer(0)]]) {
thread tint_private_vars_struct tint_private_vars = {};
tint_private_vars.p = tint_unpack_vec3_in_composite(*(tint_symbol));
tint_private_vars.p[1] = float3((*(tint_symbol))[0].elements);
tint_private_vars.p[1] = float3((*(tint_symbol))[0].elements).zxy;
tint_private_vars.p[0][1] = (*(tint_symbol))[1].elements[0];
return;
}

Some files were not shown because too many files have changed in this diff Show More